More code review coments fixed

Author: Alexander Khokhlov

CLW/CL/CLW.cl

@@ -1511,14 +1511,15 @@ __kernel void segmented_distribute_part_sum_int_nocut(
 // --------------------- ATOMIC OPERTIONS ------------------------
 
 #define DEFINE_ATOMIC(operation)\
-    inline void atomic_##operation##_float(volatile __global float* addr, float value)\
+    __attribute__((always_inline)) void atomic_##operation##_float(volatile __global float* addr, float value)\
     {\
         union{\
         unsigned int u32;\
         float        f32;\
         } next, expected, current;\
         current.f32 = *addr;\
-        do{\
+        do\
+        {\
             expected.f32 = current.f32;\
             next.f32 = operation(expected.f32, value);\
             current.u32 = atomic_cmpxchg((volatile __global unsigned int *)addr,\
@@ -1527,20 +1528,20 @@ __kernel void segmented_distribute_part_sum_int_nocut(
     }
 
 #define DEFINE_ATOMIC_FLOAT3(operation)\
-    inline void atomic_##operation##_float3(volatile __global float3* addr, float3 value)\
+    __attribute__((always_inline)) void atomic_##operation##_float3(volatile __global float3* addr, float3 value)\
     {\
         volatile __global float* p = (volatile __global float*)addr;\
         atomic_##operation##_float(p, value.x);\
         atomic_##operation##_float(p + 1, value.y);\
         atomic_##operation##_float(p + 2, value.z);\
     }
 
-inline void atomic_max_int(volatile __global int* addr, int value)
+__attribute__((always_inline)) void atomic_max_int(volatile __global int* addr, int value)
 {
     atomic_max(addr, value);
 }
 
-inline void atomic_min_int(volatile __global int* addr, int value)
+__attribute__((always_inline)) void atomic_min_int(volatile __global int* addr, int value)
 {
     atomic_min(addr, value);
 }
@@ -1551,7 +1552,8 @@ inline void atomic_min_int(volatile __global int* addr, int value)
 __kernel void reduction_##bin_op##_##type(__global type* buffer,\
                                           int count,\
                                           __local type* shared_mem,\
-                                          __global type* out)\
+                                          __global type* out,\
+                                          int /* in elements */ out_offset)\
 {\
     int global_id = get_global_id(0);\
     int group_id = get_group_id(0);\
@@ -1569,7 +1571,7 @@ __kernel void reduction_##bin_op##_##type(__global type* buffer,\
         barrier(CLK_LOCAL_MEM_FENCE);\
     }\
     if (local_id == 0)\
-        atomic_##bin_op##_##type(out, shared_mem[0]);\
+        atomic_##bin_op##_##type(out + out_offset, shared_mem[0]);\
 }
 
 // --------------------- NORMALIZATION ------------------------
@@ -1578,12 +1580,12 @@ __kernel void reduction_##bin_op##_##type(__global type* buffer,\
 __kernel void buffer_normalization_##type(__global type* input,\
                                           __global type* output,\
                                           int count,\
-                                          type max,\
-                                          type min)\
+                                          __global type* storage)\
 {\
+    type norm_coef = storage[0] - storage[1];\
     int global_id = get_global_id(0);\
     if (global_id < count)\
-        output[global_id] = input[global_id] / (max - min);\
+        output[global_id] = input[global_id] / norm_coef;\
 }
 
 // Do not change the order

CLW/CLWParallelPrimitives.cpp

@@ -976,10 +976,11 @@ cl_float3 CLWParallelPrimitives::GetMinNum<cl_float3>()
 
 template <class T>
 CLWEvent CLWParallelPrimitives::Reduction(const char* kernelName,
-    unsigned int deviceIdx,
-    CLWBuffer<T> input,
-    int numElems,
-    CLWBuffer<T> out)
+                                          unsigned int deviceIdx,
+                                          CLWBuffer<T> input,
+                                          int numElems,
+                                          CLWBuffer<T> out,
+                                          int out_offset)
 {
     assert(input.GetElementCount() >= numElems);
 
@@ -993,6 +994,7 @@ CLWEvent CLWParallelPrimitives::Reduction(const char* kernelName,
     reductionKernel.SetArg(argc++, numElems);
     reductionKernel.SetArg(argc++, SharedMemory(sizeof(T) * WG_SIZE));
     reductionKernel.SetArg(argc++, out);
+    reductionKernel.SetArg(argc++, out_offset);
 
     return context_.Launch1D(deviceIdx, NUM_BLOCKS * WG_SIZE, WG_SIZE, reductionKernel);
 }
@@ -1015,24 +1017,22 @@ CLWEvent CLWParallelPrimitives::Normalize(const char* normalizeKernelName,
     T min = GetMaxNum<T>();
     T max = GetMinNum<T>();
 
-    context_.WriteBuffer<T>(deviceIdx, cache, &min, 1);
-
-   Reduction(minReductionKernelName,
-             0,
-             input,
-             numElems,
-             cache).Wait();
+    context_.WriteBuffer<T>(deviceIdx, cache, &max, 1);
+    context_.WriteBuffer<T>(deviceIdx, cache, &min, 1, 1);
 
-    context_.ReadBuffer<T>(deviceIdx, cache, &min, 1).Wait();
-    context_.WriteBuffer<T>(deviceIdx, cache, &max, 1).Wait();
+    Reduction(minReductionKernelName,
+              0,
+              input,
+              numElems,
+              cache,
+              1);
 
     Reduction(maxReductionKernelName,
               0,
               input,
               numElems,
-              cache).Wait();
-
-    context_.ReadBuffer<T>(deviceIdx, cache, &max, 1).Wait();
+              cache,
+              0);
 
     // launch normalization kernel
     CLWKernel normalizeKernel = program_.GetKernel(normalizeKernelName);
@@ -1042,49 +1042,48 @@ CLWEvent CLWParallelPrimitives::Normalize(const char* normalizeKernelName,
     normalizeKernel.SetArg(argc++, input);
     normalizeKernel.SetArg(argc++, output);
     normalizeKernel.SetArg(argc++, numElems);
-    normalizeKernel.SetArg(argc++, max);
-    normalizeKernel.SetArg(argc++, min);
+    normalizeKernel.SetArg(argc++, cache);
 
     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(1);
-
-    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(1);
-
-    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)
 {
-    CLWBuffer<cl_float3> cache = GetTempBuffer<cl_float3>(float3_BufferCache_, 1);
+    CLWBuffer<cl_float3> cache = GetTempBuffer<cl_float3>(float3_BufferCache_, 2);
 
     CLWEvent event = Normalize("buffer_normalization_float3",
                                "reduction_min_float3",

CLW/CLWParallelPrimitives.h

@@ -95,7 +95,8 @@ class CLWParallelPrimitives
                        unsigned int deviceIdx,
                        CLWBuffer<T> input,
                        int numElems,
-                       CLWBuffer<T> out);
+                       CLWBuffer<T> out,
+                       int /* in elements */out_offset = 0);
 
     template <class T>
     T GetMaxNum();