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
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/beignet
