From 3afae0e4caae5aa592a29bd9717370ea3242d3cd Mon Sep 17 00:00:00 2001 From: Zhigang Gong Date: Thu, 24 Jul 2014 13:13:32 +0800 Subject: [PATCH] utest: add new test for constant expression processing. If we use 3-component vector in a union, it may introduce some complex constant expression as below: float bitcast (i32 trunc (i128 bitcast (<4 x i32> to i128) to i32) to float). To test the constant expression processing function. Signed-off-by: Zhigang Gong Reviewed-by: "Song, Ruiling" --- kernels/compiler_constant_expr.cl | 23 +++++++++++++++++++++++ utests/CMakeLists.txt | 1 + utests/compiler_constant_expr.cpp | 35 +++++++++++++++++++++++++++++++++++ 3 files changed, 59 insertions(+) create mode 100644 kernels/compiler_constant_expr.cl create mode 100644 utests/compiler_constant_expr.cpp diff --git a/kernels/compiler_constant_expr.cl b/kernels/compiler_constant_expr.cl new file mode 100644 index 0000000..d40cead --- /dev/null +++ b/kernels/compiler_constant_expr.cl @@ -0,0 +1,23 @@ +float3 foo_pow3(float3 src0, float3 src1) +{ + union { + float3 f3; + float farray[4]; + } s0, s1, dst; + s0.f3 = src0; + s1.f3 = src1; + int i; + for(i = 0; i < 3; i++) + dst.farray[i] = pow(s0.farray[i], s1.farray[i]); + return dst.f3; +} + +__kernel void +compiler_constant_expr(__global float* src, __global float *dst) +{ + int gid = get_global_id(0); + float3 f3 = vload3(gid, src); + float3 cf3 = (float3)(1.f, 2.f, 3.f); + float3 result = foo_pow3(f3, cf3); + vstore3(result, gid, dst); +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index c6dea8c..721e6f7 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -183,6 +183,7 @@ set (utests_sources enqueue_built_in_kernels.cpp image_1D_buffer.cpp compare_image_2d_and_1d_array.cpp + compiler_constant_expr.cpp utest_assert.cpp utest.cpp utest_file_map.cpp diff --git a/utests/compiler_constant_expr.cpp b/utests/compiler_constant_expr.cpp new file mode 100644 index 0000000..8bed03b --- /dev/null +++ b/utests/compiler_constant_expr.cpp @@ -0,0 +1,35 @@ +#include "utest_helper.hpp" +#include + +static void compiler_constant_expr(void) +{ + const size_t n = 48; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_constant_expr"); + buf_data[0] = (uint32_t*) malloc(sizeof(float) * n); + for (uint32_t i = 0; i < n; ++i) ((float*)buf_data[0])[i] = i; + OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(float), buf_data[0]); + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL); + free(buf_data[0]); + buf_data[0] = NULL; + + // Run the kernel + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = 16; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < n; ++i) { + float expect = pow(((float*)buf_data[0])[i], (i % 3) + 1); + float err = fabs(((float*)buf_data[1])[i] - expect); + OCL_ASSERT(err <= 100 * cl_FLT_ULP(expect)); + } +} + +MAKE_UTEST_FROM_FUNCTION(compiler_constant_expr); + -- 2.7.4