From 68a5b7477a11630041451636f68ee9a6bb44fd31 Mon Sep 17 00:00:00 2001 From: Yang Rong Date: Mon, 22 Apr 2013 13:11:52 +0800 Subject: [PATCH] Add constant ptr argument test case. Signed-off-by: Yang Rong Reviewed-by: Zhigang Gong --- kernels/compiler_function_constant.cl | 6 +++++ kernels/compiler_function_constant0.cl | 6 +++++ utests/CMakeLists.txt | 3 +++ utests/compiler_function_constant.cpp | 34 ++++++++++++++++++++++++ utests/compiler_function_constant0.cpp | 42 ++++++++++++++++++++++++++++++ utests/compiler_function_constant1.cpp | 47 ++++++++++++++++++++++++++++++++++ 6 files changed, 138 insertions(+) create mode 100644 kernels/compiler_function_constant.cl create mode 100644 kernels/compiler_function_constant0.cl create mode 100644 utests/compiler_function_constant.cpp create mode 100644 utests/compiler_function_constant0.cpp create mode 100644 utests/compiler_function_constant1.cpp diff --git a/kernels/compiler_function_constant.cl b/kernels/compiler_function_constant.cl new file mode 100644 index 0000000..ca7e874 --- /dev/null +++ b/kernels/compiler_function_constant.cl @@ -0,0 +1,6 @@ +__kernel void +compiler_function_constant(__constant short *c, __global int *dst, int value) +{ + int id = (int)get_global_id(0); + dst[id] = value + c[id%69]; +} diff --git a/kernels/compiler_function_constant0.cl b/kernels/compiler_function_constant0.cl new file mode 100644 index 0000000..f6efcef --- /dev/null +++ b/kernels/compiler_function_constant0.cl @@ -0,0 +1,6 @@ +__kernel void +compiler_function_constant0(__constant short *c0, __constant char *c1, __global int *dst, int value) +{ + int id = (int)get_global_id(0); + dst[id] = value + c0[id%69] + c1[15]; +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index bed0159..b2e3c97 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -29,6 +29,9 @@ set (utests_sources compiler_function_argument0.cpp compiler_function_argument1.cpp compiler_function_argument.cpp + compiler_function_constant0.cpp + compiler_function_constant1.cpp + compiler_function_constant.cpp compiler_if_else.cpp compiler_integer_division.cpp compiler_integer_remainder.cpp diff --git a/utests/compiler_function_constant.cpp b/utests/compiler_function_constant.cpp new file mode 100644 index 0000000..20f0ece --- /dev/null +++ b/utests/compiler_function_constant.cpp @@ -0,0 +1,34 @@ +#include "utest_helper.hpp" + +void compiler_function_constant(void) +{ + const size_t n = 2048; + const uint32_t value = 34; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_function_constant"); + OCL_CREATE_BUFFER(buf[0], 0, 75 * sizeof(short), 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]); + OCL_SET_ARG(2, sizeof(uint32_t), &value); + + OCL_MAP_BUFFER(0); + for(uint32_t i = 0; i < 69; ++i) + ((short *)buf_data[0])[i] = i; + OCL_UNMAP_BUFFER(0); + + // Run the kernel + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + OCL_MAP_BUFFER(1); + + // Check results + for (uint32_t i = 0; i < n; ++i) + OCL_ASSERT(((uint32_t *)buf_data[1])[i] == (value + i%69)); + + OCL_UNMAP_BUFFER(1); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_function_constant); diff --git a/utests/compiler_function_constant0.cpp b/utests/compiler_function_constant0.cpp new file mode 100644 index 0000000..de564f3 --- /dev/null +++ b/utests/compiler_function_constant0.cpp @@ -0,0 +1,42 @@ +#include "utest_helper.hpp" + +void compiler_function_constant0(void) +{ + const size_t n = 2048; + const uint32_t value = 34; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_function_constant0"); + OCL_CREATE_BUFFER(buf[0], 0, 75 * sizeof(short), NULL); + OCL_CREATE_BUFFER(buf[1], 0, 256 * sizeof(char), NULL); + OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(uint32_t), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]); + OCL_SET_ARG(3, sizeof(uint32_t), &value); + + OCL_MAP_BUFFER(0); + for(uint32_t i = 0; i < 69; ++i) + ((short *)buf_data[0])[i] = i; + OCL_UNMAP_BUFFER(0); + + OCL_MAP_BUFFER(1); + for(uint32_t i = 0; i < 256; ++i) + ((char *)buf_data[1])[i] = 10; + ((char *)buf_data[1])[15] = 15; + OCL_UNMAP_BUFFER(1); + + // Run the kernel + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + OCL_MAP_BUFFER(2); + + // Check results + for (uint32_t i = 0; i < n; ++i) + OCL_ASSERT(((uint32_t *)buf_data[2])[i] == (value + 15 + i%69)); + + OCL_UNMAP_BUFFER(2); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_function_constant0); diff --git a/utests/compiler_function_constant1.cpp b/utests/compiler_function_constant1.cpp new file mode 100644 index 0000000..b92e6ca --- /dev/null +++ b/utests/compiler_function_constant1.cpp @@ -0,0 +1,47 @@ +#include "utest_helper.hpp" + +void compiler_function_constant1(void) +{ + const size_t n = 2048; + const uint32_t value = 34; + + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_function_constant"); + OCL_CREATE_BUFFER(buf[0], 0, 75 * sizeof(short), 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]); + OCL_SET_ARG(2, sizeof(uint32_t), &value); + + OCL_MAP_BUFFER(0); + for(uint32_t i = 0; i < 69; ++i) + ((short *)buf_data[0])[i] = i; + OCL_UNMAP_BUFFER(0); + + // Run the kernel + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + + OCL_CREATE_BUFFER(buf[2], 0, 101 * sizeof(short), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[2]); + OCL_MAP_BUFFER(2); + for(uint32_t i = 0; i < 69; ++i) + ((short *)buf_data[2])[i] = 2*i; + OCL_UNMAP_BUFFER(2); + + // Run the kernel + globals[0] = n; + locals[0] = 16; + OCL_NDRANGE(1); + + OCL_MAP_BUFFER(1); + + // Check results + for (uint32_t i = 0; i < n; ++i) + OCL_ASSERT(((uint32_t *)buf_data[1])[i] == (value + (i%69)*2)); + + OCL_UNMAP_BUFFER(1); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_function_constant1); -- 2.7.4