diff options
author | Junyan He <junyan.he@linux.intel.com> | 2015-12-01 16:10:40 +0800 |
---|---|---|
committer | Yang Rong <rong.r.yang@intel.com> | 2015-12-14 15:12:26 +0800 |
commit | 70a0bfdcf23b311c6edd5779a839aaea708419dd (patch) | |
tree | b914597e1c385c86df905902eafedb0acaea32a8 | |
parent | 81fbb83890db4eb676f126d16b572472d7d63e6a (diff) |
Utests: Add test cases for reduce add.
Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
-rw-r--r-- | kernels/compiler_workgroup_reduce.cl | 12 | ||||
-rw-r--r-- | utests/compiler_workgroup_reduce.cpp | 69 |
2 files changed, 81 insertions, 0 deletions
diff --git a/kernels/compiler_workgroup_reduce.cl b/kernels/compiler_workgroup_reduce.cl index 27d306bc..1fc57b51 100644 --- a/kernels/compiler_workgroup_reduce.cl +++ b/kernels/compiler_workgroup_reduce.cl @@ -26,3 +26,15 @@ kernel void compiler_workgroup_reduce_max_float(global float *src, global float float max_val = work_group_reduce_max(val); dst[get_local_id(0)] = max_val; } + +kernel void compiler_workgroup_reduce_add_uint(global uint *src, global uint *dst) { + uint val = src[get_local_id(0)]; + uint sum = work_group_reduce_add(val); + dst[get_local_id(0)] = sum; +} + +kernel void compiler_workgroup_reduce_add_float(global float *src, global float *dst) { + float val = src[get_local_id(0)]; + float sum = work_group_reduce_add(val); + dst[get_local_id(0)] = sum; +} diff --git a/utests/compiler_workgroup_reduce.cpp b/utests/compiler_workgroup_reduce.cpp index 6340cb2a..40978435 100644 --- a/utests/compiler_workgroup_reduce.cpp +++ b/utests/compiler_workgroup_reduce.cpp @@ -97,6 +97,41 @@ void compiler_workgroup_reduce_max_uint(void) } MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_reduce_max_uint); +void compiler_workgroup_reduce_add_uint(void) +{ + const size_t n = 50; + uint32_t* src = test_array_uint; + + // Setup kernel and buffers + OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce", "compiler_workgroup_reduce_add_uint"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = n; + + uint32_t cpu_res = 0; + for (size_t i = 0; i < n; i++) + cpu_res += src[i]; + + OCL_MAP_BUFFER(0); + memcpy(buf_data[0], src, n * sizeof(uint32_t)); + OCL_UNMAP_BUFFER(0); + + // Run the kernel on GPU + OCL_NDRANGE(1); + + // Compare + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < (int32_t) n; ++i) { + //printf("%u ", ((uint32_t *)buf_data[1])[i]); + OCL_ASSERT(((uint32_t *)buf_data[1])[i] == cpu_res); + } + OCL_UNMAP_BUFFER(1); +} +MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_reduce_add_uint); + static float test_array_float[64] = {1.0234f, 0.34e32f, -13441.4334f, 1893.21f, -9999.0f, -88.00f, 1.3f, 1.0f, 2.33f, 134.44f, 263.0f, 1.0f, 0.0f, 344.900043f, 0.1e30f, 1.0e10f, @@ -172,3 +207,37 @@ void compiler_workgroup_reduce_max_float(void) } MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_reduce_max_float); +void compiler_workgroup_reduce_add_float(void) +{ + const size_t n = 42; + float* src = test_array_float; + + // Setup kernel and buffers + OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce", "compiler_workgroup_reduce_add_float"); + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = n; + + float cpu_res = 0; + for (size_t i = 0; i < n; i++) + cpu_res += src[i]; + + OCL_MAP_BUFFER(0); + memcpy(buf_data[0], src, n * sizeof(float)); + OCL_UNMAP_BUFFER(0); + + // Run the kernel on GPU + OCL_NDRANGE(1); + + // Compare + OCL_MAP_BUFFER(1); + for (int32_t i = 0; i < (int32_t) n; ++i) { + //printf("%f ", ((float *)buf_data[1])[i]); + OCL_ASSERT(((float *)buf_data[1])[i] == cpu_res); + } + OCL_UNMAP_BUFFER(1); +} +MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_reduce_add_float); |