From 8d96ef5c5fe6b9500abe00a049fbe2b62963b83b Mon Sep 17 00:00:00 2001 From: Junyan He Date: Mon, 19 Jan 2015 15:35:04 +0800 Subject: Add test case for long bitcast. Signed-off-by: Junyan He Reviewed-by: Zhigang Gong --- kernels/compiler_long_bitcast.cl | 47 ++++++++ utests/CMakeLists.txt | 1 + utests/compiler_long_bitcast.cpp | 227 +++++++++++++++++++++++++++++++++++++++ 3 files changed, 275 insertions(+) create mode 100644 kernels/compiler_long_bitcast.cl create mode 100644 utests/compiler_long_bitcast.cpp diff --git a/kernels/compiler_long_bitcast.cl b/kernels/compiler_long_bitcast.cl new file mode 100644 index 00000000..4b008ef1 --- /dev/null +++ b/kernels/compiler_long_bitcast.cl @@ -0,0 +1,47 @@ +__kernel void compiler_bitcast_char8_to_long(__global char8 *src, __global ulong *dst) +{ + int tid = get_global_id(0); + char8 v = src[tid]; + ulong dl = as_ulong(v); + dst[tid] = dl; +} + +__kernel void compiler_bitcast_long_to_char8(__global ulong *src, __global uchar8 *dst) +{ + int tid = get_global_id(0); + ulong v = src[tid]; + uchar8 dl = as_uchar8(v); + dst[tid] = dl; +} + +__kernel void compiler_bitcast_int2_to_long(__global int2 *src, __global ulong *dst) +{ + int tid = get_global_id(0); + int2 v = src[tid]; + ulong dl = as_ulong(v); + dst[tid] = dl; +} + +__kernel void compiler_bitcast_long_to_int2(__global ulong *src, __global uint2 *dst) +{ + int tid = get_global_id(0); + ulong v = src[tid]; + uint2 dl = as_uint2(v); + dst[tid] = dl; +} + +__kernel void compiler_bitcast_short4_to_long(__global short4 *src, __global ulong *dst) +{ + int tid = get_global_id(0); + short4 v = src[tid]; + ulong dl = as_ulong(v); + dst[tid] = dl; +} + +__kernel void compiler_bitcast_long_to_short4(__global ulong *src, __global ushort4 *dst) +{ + int tid = get_global_id(0); + ulong v = src[tid]; + ushort4 dl = as_ushort4(v); + dst[tid] = dl; +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 0f56c6c9..f8fb9c6b 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -182,6 +182,7 @@ set (utests_sources compiler_long_asr.cpp compiler_long_mult.cpp compiler_long_cmp.cpp + compiler_long_bitcast.cpp compiler_function_argument3.cpp compiler_function_qualifiers.cpp compiler_bool_cross_basic_block.cpp diff --git a/utests/compiler_long_bitcast.cpp b/utests/compiler_long_bitcast.cpp new file mode 100644 index 00000000..5bd962de --- /dev/null +++ b/utests/compiler_long_bitcast.cpp @@ -0,0 +1,227 @@ +#include +#include +#include +#include "utest_helper.hpp" + +void compiler_bitcast_char8_to_long(void) +{ + const size_t n = 64; + const int v = 8; + char src[n * v]; + uint64_t *dst = (uint64_t *)src; + + // Setup kernel and buffers + OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_char8_to_long"); + OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL); + OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + + for (int32_t i = 0; i < (int32_t) n*v; ++i) { + src[i] = (char)rand(); + } + + OCL_MAP_BUFFER(0); + memcpy(buf_data[0], src, sizeof(src)); + 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) { + OCL_ASSERT(((uint64_t *)(buf_data[1]))[i] == dst[i]); + //printf("ref is 0x%lx, result is 0x%lx\n", dst[i], ((int64_t *)(buf_data[1]))[i]); + } + OCL_UNMAP_BUFFER(1); +} + +void compiler_bitcast_long_to_char8(void) +{ + const size_t n = 64; + const int v = 8; + ulong src[n]; + char *dst = (char *)src; + + // Setup kernel and buffers + OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_char8_to_long"); + OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL); + OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + + for (int32_t i = 0; i < (int32_t) n; ++i) { + src[i] = ((int64_t)rand() << 32) + rand(); + } + + OCL_MAP_BUFFER(0); + memcpy(buf_data[0], src, sizeof(src)); + 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*v; ++i) { + OCL_ASSERT(((char *)(buf_data[1]))[i] == dst[i]); +// printf("ref is 0x%2x, result is 0x%2x\n", dst[i], ((char *)(buf_data[1]))[i]); + } + OCL_UNMAP_BUFFER(1); +} + +void compiler_bitcast_int2_to_long(void) +{ + const size_t n = 64; + const int v = 2; + int src[n * v]; + uint64_t *dst = (uint64_t *)src; + + // Setup kernel and buffers + OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_int2_to_long"); + OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL); + OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + + for (int32_t i = 0; i < (int32_t) n*v; ++i) { + src[i] = (char)rand(); + } + + OCL_MAP_BUFFER(0); + memcpy(buf_data[0], src, sizeof(src)); + 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) { + OCL_ASSERT(((uint64_t *)(buf_data[1]))[i] == dst[i]); + //printf("ref is 0x%lx, result is 0x%lx\n", dst[i], ((int64_t *)(buf_data[1]))[i]); + } + OCL_UNMAP_BUFFER(1); +} + +void compiler_bitcast_long_to_int2(void) +{ + const size_t n = 64; + const int v = 2; + ulong src[n]; + uint32_t *dst = (uint32_t *)src; + + // Setup kernel and buffers + OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_long_to_int2"); + OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL); + OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + + for (int32_t i = 0; i < (int32_t) n; ++i) { + src[i] = ((int64_t)i << 32) + i; + } + + OCL_MAP_BUFFER(0); + memcpy(buf_data[0], src, sizeof(src)); + 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*v; ++i) { + OCL_ASSERT(((uint32_t *)(buf_data[1]))[i] == dst[i]); + //printf("ref is 0x%2x, result is 0x%2x\n", dst[i], ((uint32_t *)(buf_data[1]))[i]); + } + OCL_UNMAP_BUFFER(1); +} + +void compiler_bitcast_short4_to_long(void) +{ + const size_t n = 64; + const int v = 4; + short src[n * v]; + uint64_t *dst = (uint64_t *)src; + + // Setup kernel and buffers + OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_short4_to_long"); + OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL); + OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + + for (int32_t i = 0; i < (int32_t) n*v; ++i) { + src[i] = (char)rand(); + } + + OCL_MAP_BUFFER(0); + memcpy(buf_data[0], src, sizeof(src)); + 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) { + OCL_ASSERT(((uint64_t *)(buf_data[1]))[i] == dst[i]); + //printf("ref is 0x%lx, result is 0x%lx\n", dst[i], ((int64_t *)(buf_data[1]))[i]); + } + OCL_UNMAP_BUFFER(1); +} + +void compiler_bitcast_long_to_short4(void) +{ + const size_t n = 64; + const int v = 4; + ulong src[n]; + uint16_t *dst = (uint16_t *)src; + + // Setup kernel and buffers + OCL_CREATE_KERNEL_FROM_FILE("compiler_long_bitcast", "compiler_bitcast_long_to_short4"); + OCL_CREATE_BUFFER(buf[0], 0, sizeof(src), NULL); + OCL_CREATE_BUFFER(buf[1], 0, sizeof(src), NULL); + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); + globals[0] = n; + locals[0] = 16; + + for (int32_t i = 0; i < (int32_t) n; ++i) { + src[i] = ((int64_t)rand() << 32) + rand(); + } + + OCL_MAP_BUFFER(0); + memcpy(buf_data[0], src, sizeof(src)); + 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*v; ++i) { + OCL_ASSERT(((uint16_t *)(buf_data[1]))[i] == dst[i]); + //printf("ref is 0x%2x, result is 0x%2x\n", dst[i], ((uint16_t *)(buf_data[1]))[i]); + } + OCL_UNMAP_BUFFER(1); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_char8_to_long); +MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_long_to_char8); +MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_int2_to_long); +MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_long_to_int2); +MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_short4_to_long); +MAKE_UTEST_FROM_FUNCTION(compiler_bitcast_long_to_short4); -- cgit v1.2.3