Only use private memory

histogram/histogram.c

@@ -14,6 +14,7 @@ __kernel void histogram(__global unsigned char* bytes, __global unsigned int* pi
   unsigned int basePixelIdx = lid * BIN_SIZE + gid * gsize * BIN_SIZE;
   unsigned int baseResultIdx = globalId * RESULT_SIZE;
   unsigned int maxPixel = *pixelCount;
+  unsigned int privateBuffer[RESULT_SIZE];
 
   for (i = 0; i < RESULT_SIZE; i++) {
     result[baseResultIdx + i] = 0;
@@ -24,13 +25,17 @@ __kernel void histogram(__global unsigned char* bytes, __global unsigned int* pi
     // data partition of bytes is RGBRGBRGB....
     bValue = bytes[basePixelIdx * 3 + processIndex * 3];
     // result partition is RR..RRGG..GGBB..BB.
-    result[baseResultIdx + bValue]++;
+    privateBuffer[bValue]++;
     // G
     bValue = bytes[basePixelIdx * 3 + processIndex * 3 + 1];
-    result[baseResultIdx + 256 + bValue]++;
+    privateBuffer[256 + bValue]++;
     // B
     bValue = bytes[basePixelIdx * 3 + processIndex * 3 + 2];
-    result[baseResultIdx + 512 + bValue]++;
+    privateBuffer[512 + bValue]++;
     processIndex++;
   }
+
+  for (i = 0; i < RESULT_SIZE; i++) {
+    result[baseResultIdx + i] = privateBuffer[i];
+  }
 }

histogram/histogram.py

@@ -12,6 +12,7 @@ def cpu_histogram(img):
 def opencl_histogram(pixels):
   # format of pixels is RGBRGBRGB each of character in a byte
   # calculate buffer size
+  start_time = time()
   groupSize = 4
   binSize = 1024
   pixelSize = len(pixels) / 3 
@@ -36,11 +37,14 @@ def opencl_histogram(pixels):
   semiFinal = numpy.zeros(outputBufSize, dtype=numpy.uint32)
   evt = cl.enqueue_read_buffer(clQueue, bufOutput, semiFinal)
   evt.wait()
-
+  end_time = time()
+  print ('time elapsed with open cl: {0}s'.format(end_time - start_time))
+  start_time = time()
   finalResult = [0] * 768
   for i in range(outputBufSize):
     finalResult[i % 768] += semiFinal[i]
-
+  end_time = time()
+  print ('time elapsed with merging: {0}s'.format(end_time - start_time))
   return finalResult
 
 parser = argparse.ArgumentParser(description='Dump histogram data.')