code review

Author: Alexander Khokhlov

CLW/CL/CLW.cl

@@ -1513,7 +1513,8 @@ __kernel void segmented_distribute_part_sum_int_nocut(
 #define DEFINE_ATOMIC(operation)\
     __attribute__((always_inline)) void atomic_##operation##_float(volatile __global float* addr, float value)\
     {\
-        union{\
+        union\
+        {\
         unsigned int u32;\
         float        f32;\
         } next, expected, current;\
@@ -1549,7 +1550,7 @@ __attribute__((always_inline)) void atomic_min_int(volatile __global int* addr,
 // --------------------- REDUCTION ------------------------
 
 #define DEFINE_REDUCTION(bin_op, type)\
-__kernel void reduction_##bin_op##_##type(__global type* buffer,\
+__kernel void reduction_##bin_op##_##type(const __global type* buffer,\
                                           int count,\
                                           __local type* shared_mem,\
                                           __global type* out,\
@@ -1577,10 +1578,10 @@ __kernel void reduction_##bin_op##_##type(__global type* buffer,\
 // --------------------- NORMALIZATION ------------------------
 
 #define DEFINE_BUFFER_NORMALIZATION(type)\
-__kernel void buffer_normalization_##type(__global type* input,\
+__kernel void buffer_normalization_##type(const __global type* input,\
                                           __global type* output,\
                                           int count,\
-                                          __global type* storage)\
+                                          const __global type* storage)\
 {\
     type norm_coef = storage[0] - storage[1];\
     int global_id = get_global_id(0);\

CLW/CLWParallelPrimitives.cpp

@@ -1047,39 +1047,39 @@ CLWEvent CLWParallelPrimitives::Normalize(const char* normalizeKernelName,
     return context_.Launch1D(deviceIdx, NUM_BLOCKS * WG_SIZE, WG_SIZE, normalizeKernel);
 }
 
-//CLWEvent CLWParallelPrimitives::Normalize(unsigned int deviceIdx, CLWBuffer<cl_int> input, CLWBuffer<cl_int> output, int numElems)
-//{
-//    CLWBuffer<cl_int> cache = GetTempIntBuffer(2);
-//
-//    CLWEvent event = Normalize("buffer_normalization_int",
-//                               "reduction_min_int",
-//                               "reduction_max_int",
-//                               deviceIdx,
-//                               input,
-//                               output,
-//                               numElems,
-//                               cache);
-//
-//    ReclaimTempIntBuffer(cache);
-//    return event;
-//}
-//
-//CLWEvent CLWParallelPrimitives::Normalize(unsigned int deviceIdx, CLWBuffer<cl_float> input, CLWBuffer<cl_float> output, int numElems)
-//{
-//    CLWBuffer<cl_float> cache = GetTempFloatBuffer(2);
-//
-//    CLWEvent event = Normalize("buffer_normalization_float",
-//                               "reduction_min_float",
-//                               "reduction_max_float",
-//                               deviceIdx,
-//                               input,
-//                               output,
-//                               numElems,
-//                               cache);
-//
-//    ReclaimTempFloatBuffer(cache);
-//    return event;
-//}
+CLWEvent CLWParallelPrimitives::Normalize(unsigned int deviceIdx, CLWBuffer<cl_int> input, CLWBuffer<cl_int> output, int numElems)
+{
+    CLWBuffer<cl_int> cache = GetTempIntBuffer(2);
+
+    CLWEvent event = Normalize("buffer_normalization_int",
+                               "reduction_min_int",
+                               "reduction_max_int",
+                               deviceIdx,
+                               input,
+                               output,
+                               numElems,
+                               cache);
+
+    ReclaimTempIntBuffer(cache);
+    return event;
+}
+
+CLWEvent CLWParallelPrimitives::Normalize(unsigned int deviceIdx, CLWBuffer<cl_float> input, CLWBuffer<cl_float> output, int numElems)
+{
+    CLWBuffer<cl_float> cache = GetTempFloatBuffer(2);
+
+    CLWEvent event = Normalize("buffer_normalization_float",
+                               "reduction_min_float",
+                               "reduction_max_float",
+                               deviceIdx,
+                               input,
+                               output,
+                               numElems,
+                               cache);
+
+    ReclaimTempFloatBuffer(cache);
+    return event;
+}
 
 CLWEvent CLWParallelPrimitives::Normalize(unsigned int deviceIdx, CLWBuffer<cl_float3> input, CLWBuffer<cl_float3> output, int numElems)
 {

RadeonRays/src/async/thread_pool.h

@@ -29,6 +29,7 @@ THE SOFTWARE.
 #include <atomic>
 #include <future>
 #include <iostream>
+#include <algorithm>
 
 namespace RadeonRays
 {
@@ -66,7 +67,7 @@ namespace RadeonRays
         {
             std::unique_lock<std::mutex> lock(mutex_);
             cv_.wait(lock, [this](){return !queue_.empty();});
-            t = queue_.front();
+            t = std::move(queue_.front());
             queue_.pop();
         }