Updated benchmarks for workgroup reduce

Submitted by Grigore Lupescu on Feb. 13, 2016, 11:26 p.m.

Details

Message ID 1455405987-8036-1-git-send-email-grigore.lupescu@intel.com
State New
Headers show
Series "Series without cover letter" ( rev: 4 ) in Beignet

Not browsing as part of any series.

Commit Message

Grigore Lupescu Feb. 13, 2016, 11:26 p.m.
Signed-off-by: Grigore Lupescu <grigore.lupescu@intel.com>
---
 benchmark/benchmark_workgroup_reduce.cpp | 157 +++++++++++++++++++++++++------
 kernels/bench_workgroup_reduce.cl        |  32 ++++++-
 2 files changed, 160 insertions(+), 29 deletions(-)

Patch hide | download patch | download mbox

diff --git a/benchmark/benchmark_workgroup_reduce.cpp b/benchmark/benchmark_workgroup_reduce.cpp
index 815b6b5..c93ef26 100644
--- a/benchmark/benchmark_workgroup_reduce.cpp
+++ b/benchmark/benchmark_workgroup_reduce.cpp
@@ -9,30 +9,30 @@  double benchmark_workgroup_add_uint(void)
 {
   double elapsed = 0;
   struct timeval start,stop;
-  const size_t set_size = 512 * 256;
-  const size_t set_local_size = 64;
+  const size_t global_size = 512 * 256;
+  const size_t local_size = 128;
   const uint32_t reduce_loop = 10000;
 
   /* Input set will be generated */
-  uint32_t* src = (uint32_t*)calloc(sizeof(uint32_t), set_size);
+  uint32_t* src = (uint32_t*)calloc(sizeof(uint32_t), global_size);
   OCL_ASSERT(src != NULL);
-  for(uint32_t i = 0; i < set_size; i++){
-    src[i] = 1;
+  for(uint32_t i = 0; i < global_size; i++){
+    src[i] = (i / local_size);
   }
 
   /* Setup kernel and buffers */
   OCL_CREATE_KERNEL_FROM_FILE("bench_workgroup_reduce",
                   "bench_workgroup_reduce_add_uint");
 
-  OCL_CREATE_BUFFER(buf[0], 0, (set_size) * sizeof(uint32_t), NULL);
-  OCL_CREATE_BUFFER(buf[1], 0, (set_size) * sizeof(uint32_t), NULL);
+  OCL_CREATE_BUFFER(buf[0], 0, (global_size) * sizeof(uint32_t), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, (global_size) * sizeof(uint32_t), NULL);
 
   OCL_MAP_BUFFER(0);
-  memcpy(buf_data[0], src, set_size * sizeof(uint32_t));
+  memcpy(buf_data[0], src, global_size * sizeof(uint32_t));
   OCL_UNMAP_BUFFER(0);
 
-  globals[0] = set_size;
-  locals[0] = set_local_size;
+  globals[0] = global_size;
+  locals[0] = local_size;
 
   OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
   OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
@@ -47,43 +47,95 @@  double benchmark_workgroup_add_uint(void)
 
   /* Check results */
   OCL_MAP_BUFFER(1);
-  for(uint32_t i = 1; i < set_size; i += set_size){
+  for(uint32_t i = 0; i < global_size; i += local_size){
     //printf(" %u", ((uint32_t*)buf_data[1])[i]);
-    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == set_local_size);
+    OCL_ASSERT( ((uint32_t*)buf_data[1])[i] == i );
   }
   OCL_UNMAP_BUFFER(1);
 
-  return BANDWIDTH(set_size * reduce_loop, elapsed);
+  return BANDWIDTH(global_size * reduce_loop, elapsed);
 }
 MAKE_BENCHMARK_FROM_FUNCTION(benchmark_workgroup_add_uint, "M/sec");
 
+double benchmark_workgroup_min_uint(void)
+{
+  double elapsed = 0;
+  struct timeval start,stop;
+  const size_t global_size = 512 * 256;
+  const size_t local_size = 128;
+  const uint32_t reduce_loop = 10000;
+
+  /* Input set will be generated */
+  uint32_t* src = (uint32_t*)calloc(sizeof(uint32_t), global_size);
+  OCL_ASSERT(src != NULL);
+  for(uint32_t i = 0; i < global_size; i++){
+    src[i] = i;
+  }
+
+  /* Setup kernel and buffers */
+  OCL_CREATE_KERNEL_FROM_FILE("bench_workgroup_reduce",
+                  "bench_workgroup_reduce_min_uint");
+
+  OCL_CREATE_BUFFER(buf[0], 0, (global_size) * sizeof(uint32_t), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, (global_size) * sizeof(uint32_t), NULL);
+
+  OCL_MAP_BUFFER(0);
+  memcpy(buf_data[0], src, global_size * sizeof(uint32_t));
+  OCL_UNMAP_BUFFER(0);
+
+  globals[0] = global_size;
+  locals[0] = local_size;
+
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  OCL_SET_ARG(2, sizeof(cl_uint), &reduce_loop);
+
+  /* Measure performance */
+  gettimeofday(&start,0);
+  OCL_NDRANGE(1);
+  clFinish(queue);
+  gettimeofday(&stop,0);
+  elapsed = time_subtract(&stop, &start, 0);
+
+  /* Check results */
+  OCL_MAP_BUFFER(1);
+  for(uint32_t i = local_size/2; i < global_size; i += local_size){
+    //printf(" %u", ((uint32_t*)buf_data[1])[i]);
+      OCL_ASSERT( ((uint32_t*)buf_data[1])[i] == (src[i] - (local_size / 2)) );
+  }
+  OCL_UNMAP_BUFFER(1);
+
+  return BANDWIDTH(global_size * reduce_loop, elapsed);
+}
+MAKE_BENCHMARK_FROM_FUNCTION(benchmark_workgroup_min_uint, "M/sec");
+
 double benchmark_workgroup_add_float(void)
 {
   double elapsed = 0;
   struct timeval start,stop;
-  const size_t set_size = 512 * 256;
-  const size_t set_local_size = 64;
+  const size_t global_size = 512 * 256;
+  const size_t local_size = 128;
   const uint32_t reduce_loop = 10000;
 
   /* Input set will be generated */
-  float* src = (float*)calloc(sizeof(float), set_size);
+  float* src = (float*)calloc(sizeof(float), global_size);
   OCL_ASSERT(src != NULL);
-  for(uint32_t i = 0; i < set_size; i++)
-    src[i] = 1.0f;
+  for(uint32_t i = 0; i < global_size; i++)
+    src[i] = (i / local_size);
 
   /* Setup kernel and buffers */
   OCL_CREATE_KERNEL_FROM_FILE("bench_workgroup_reduce",
     "bench_workgroup_reduce_add_float");
 
-  OCL_CREATE_BUFFER(buf[0], 0, (set_size) * sizeof(float), NULL);
-  OCL_CREATE_BUFFER(buf[1], 0, (set_size) * sizeof(float), NULL);
+  OCL_CREATE_BUFFER(buf[0], 0, (global_size) * sizeof(float), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, (global_size) * sizeof(float), NULL);
 
   OCL_MAP_BUFFER(0);
-  memcpy(buf_data[0], src, set_size * sizeof(float));
+  memcpy(buf_data[0], src, global_size * sizeof(float));
   OCL_UNMAP_BUFFER(0);
 
-  globals[0] = set_size;
-  locals[0] = set_local_size;
+  globals[0] = global_size;
+  locals[0] = local_size;
 
   OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
   OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
@@ -98,12 +150,63 @@  double benchmark_workgroup_add_float(void)
 
   /* Check results */
   OCL_MAP_BUFFER(1);
-  for(uint32_t i = 1; i < set_size; i += set_size){
-    //printf("%f ", ((float*)buf_data[1])[i]);
-    OCL_ASSERT(((float*)buf_data[1])[i] == set_local_size);
+  for(uint32_t i = 0; i < global_size; i += local_size){
+    //printf(" %f", ((float*)buf_data[1])[i]);
+    OCL_ASSERT( ((float*)buf_data[1])[i] == (float)i );
   }
   OCL_UNMAP_BUFFER(1);
 
-  return BANDWIDTH(set_size * reduce_loop, elapsed);
+  return BANDWIDTH(global_size * reduce_loop, elapsed);
 }
 MAKE_BENCHMARK_FROM_FUNCTION(benchmark_workgroup_add_float, "M/sec");
+
+double benchmark_workgroup_min_float(void)
+{
+  double elapsed = 0;
+  struct timeval start,stop;
+  const size_t global_size = 512 * 256;
+  const size_t local_size = 128;
+  const uint32_t reduce_loop = 10000;
+
+  /* Input set will be generated */
+  float* src = (float*)calloc(sizeof(float), global_size);
+  OCL_ASSERT(src != NULL);
+  for(uint32_t i = 0; i < global_size; i++)
+    src[i] = 1.0f * i + 1;
+
+  /* Setup kernel and buffers */
+  OCL_CREATE_KERNEL_FROM_FILE("bench_workgroup_reduce",
+    "bench_workgroup_reduce_min_float");
+
+  OCL_CREATE_BUFFER(buf[0], 0, (global_size) * sizeof(float), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, (global_size) * sizeof(float), NULL);
+
+  OCL_MAP_BUFFER(0);
+  memcpy(buf_data[0], src, global_size * sizeof(float));
+  OCL_UNMAP_BUFFER(0);
+
+  globals[0] = global_size;
+  locals[0] = local_size;
+
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  OCL_SET_ARG(2, sizeof(cl_uint), &reduce_loop);
+
+  /* Measure performance */
+  gettimeofday(&start,0);
+  OCL_NDRANGE(1);
+  clFinish(queue);
+  gettimeofday(&stop,0);
+  elapsed = time_subtract(&stop, &start, 0);
+
+  /* Check results */
+  OCL_MAP_BUFFER(1);
+  for(uint32_t i = local_size/2; i < global_size; i += local_size){
+    //printf(" %f", ((float*)buf_data[1])[i]);
+    OCL_ASSERT( ((float*)buf_data[1])[i] == (src[i] - (local_size / 2)) );
+  }
+  OCL_UNMAP_BUFFER(1);
+
+  return BANDWIDTH(global_size * reduce_loop, elapsed);
+}
+MAKE_BENCHMARK_FROM_FUNCTION(benchmark_workgroup_min_float, "M/sec");
diff --git a/kernels/bench_workgroup_reduce.cl b/kernels/bench_workgroup_reduce.cl
index 9e2f848..ba1c709 100644
--- a/kernels/bench_workgroup_reduce.cl
+++ b/kernels/bench_workgroup_reduce.cl
@@ -3,7 +3,7 @@  kernel void bench_workgroup_reduce_add_uint(
   global uint *dst,
   uint reduce_loop)
 {
-  uint val = src[get_local_id(0)];
+  uint val = src[get_global_id(0)];
   uint sum = work_group_reduce_add(val);
 
   for(; reduce_loop > 0; reduce_loop--)
@@ -12,12 +12,26 @@  kernel void bench_workgroup_reduce_add_uint(
   dst[get_global_id(0)] = sum;
 }
 
+kernel void bench_workgroup_reduce_min_uint(
+  global uint *src,
+  global uint *dst,
+  uint reduce_loop)
+{
+  uint val = src[get_global_id(0)];
+  uint min = work_group_reduce_min(val);
+
+  for(; reduce_loop > 0; reduce_loop--)
+    min = work_group_reduce_min(val);
+
+  dst[get_global_id(0)] = min;
+}
+
 kernel void bench_workgroup_reduce_add_float(
   global float *src,
   global float *dst,
   uint reduce_loop)
 {
-  float val = src[get_local_id(0)];
+  float val = src[get_global_id(0)];
   float sum = work_group_reduce_add(val);
 
   for(; reduce_loop > 0; reduce_loop--)
@@ -25,3 +39,17 @@  kernel void bench_workgroup_reduce_add_float(
 
   dst[get_global_id(0)] = sum;
 }
+
+kernel void bench_workgroup_reduce_min_float(
+  global float *src,
+  global float *dst,
+  uint reduce_loop)
+{
+  float val = src[get_global_id(0)];
+  float min = work_group_reduce_min(val);
+
+  for(; reduce_loop > 0; reduce_loop--)
+    min = work_group_reduce_min(val);
+
+  dst[get_global_id(0)] = min;
+}