[Beignet] [PATCH 1/3] Benchmark: Fix benchmark kernel of workgroup, function input

[grigore.lupescu at intel.com]

From: Grigore Lupescu <grigore.lupescu at intel.com>

Signed-off-by: Grigore Lupescu <grigore.lupescu at intel.com>
---
 kernels/bench_workgroup.cl | 78 ++++++++++++++++++++++++++++++----------------
 1 file changed, 51 insertions(+), 27 deletions(-)

diff --git a/kernels/bench_workgroup.cl b/kernels/bench_workgroup.cl
index 596e936..f26537b 100644
--- a/kernels/bench_workgroup.cl
+++ b/kernels/bench_workgroup.cl
@@ -6,11 +6,13 @@ kernel void bench_workgroup_reduce_add_short(
   global short *dst,
   int reduce_loop)
 {
-  short val = src[get_global_id(0)];
+  short val;
   short result;
 
-  for(int i = 0; i < reduce_loop; i++)
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
     result = work_group_reduce_add(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -20,11 +22,13 @@ kernel void bench_workgroup_reduce_add_int(
   global int *dst,
   int reduce_loop)
 {
-  int val = src[get_global_id(0)];
+  int val;
   int result;
 
-  for(int i = 0; i < reduce_loop; i++)
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
     result = work_group_reduce_add(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -34,11 +38,13 @@ kernel void bench_workgroup_reduce_add_long(
   global long *dst,
   int reduce_loop)
 {
-  long val = src[get_global_id(0)];
+  long val;
   long result;
 
-  for(int i = 0; i < reduce_loop; i++)
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
     result = work_group_reduce_add(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -51,11 +57,13 @@ kernel void bench_workgroup_reduce_min_short(
   global short *dst,
   int reduce_loop)
 {
-  short val = src[get_global_id(0)];
+  short val;
   short result;
 
-  for(int i = 0; i < reduce_loop; i++)
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
     result = work_group_reduce_min(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -65,11 +73,13 @@ kernel void bench_workgroup_reduce_min_int(
   global int *dst,
   int reduce_loop)
 {
-  int val = src[get_global_id(0)];
+  int val;
   int result;
 
-  for(int i = 0; i < reduce_loop; i++)
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
     result = work_group_reduce_min(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -79,11 +89,13 @@ kernel void bench_workgroup_reduce_min_long(
   global long *dst,
   int reduce_loop)
 {
-  long val = src[get_global_id(0)];
+  long val;
   long result;
 
-  for(int i = 0; i < reduce_loop; i++)
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
     result = work_group_reduce_min(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -96,11 +108,13 @@ kernel void bench_workgroup_scan_inclusive_add_short(
   global short *dst,
   int reduce_loop)
 {
-  short val = src[get_global_id(0)];
+  short val;
   short result;
 
-  for(int i = 0; i < reduce_loop; i++)
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
     result = work_group_scan_inclusive_add(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -110,11 +124,13 @@ kernel void bench_workgroup_scan_inclusive_add_int(
   global int *dst,
   int reduce_loop)
 {
-  int val = src[get_global_id(0)];
+  int val;
   int result;
 
-  for(int i = 0; i < reduce_loop; i++)
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
     result = work_group_scan_inclusive_add(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -124,11 +140,13 @@ kernel void bench_workgroup_scan_inclusive_add_long(
   global long *dst,
   int reduce_loop)
 {
-  long val = src[get_global_id(0)];
+  long val;
   long result;
 
-  for(int i = 0; i < reduce_loop; i++)
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
     result = work_group_scan_inclusive_add(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -141,11 +159,13 @@ kernel void bench_workgroup_scan_inclusive_min_short(
   global short *dst,
   int reduce_loop)
 {
-  short val = src[get_global_id(0)];
+  short val;
   short result;
 
-  for(int i = 0; i < reduce_loop; i++)
-    result = work_group_scan_inclusive_min(val);
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
+    result = work_group_scan_inclusive_add(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -155,11 +175,13 @@ kernel void bench_workgroup_scan_inclusive_min_int(
   global int *dst,
   int reduce_loop)
 {
-  int val = src[get_global_id(0)];
+  int val;
   int result;
 
-  for(int i = 0; i < reduce_loop; i++)
-    result = work_group_scan_inclusive_min(val);
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
+    result = work_group_scan_inclusive_add(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
@@ -169,11 +191,13 @@ kernel void bench_workgroup_scan_inclusive_min_long(
   global long *dst,
   int reduce_loop)
 {
-  long val = src[get_global_id(0)];
+  long val;
   long result;
 
-  for(int i = 0; i < reduce_loop; i++)
-    result = work_group_scan_inclusive_min(val);
+  for(; reduce_loop > 0; reduce_loop--){
+    val = src[get_global_id(0)];
+    result = work_group_scan_inclusive_add(val);
+  }
 
   dst[get_global_id(0)] = result;
 }
-- 
2.5.0