diff options
| -rw-r--r-- | kernels/compiler_vector_load_store.cl | 52 | ||||
| -rw-r--r-- | utests/CMakeLists.txt | 1 | ||||
| -rw-r--r-- | utests/compiler_vector_load_store.cpp | 57 |
3 files changed, 91 insertions, 19 deletions
diff --git a/kernels/compiler_vector_load_store.cl b/kernels/compiler_vector_load_store.cl index 28fd93a8..b362412c 100644 --- a/kernels/compiler_vector_load_store.cl +++ b/kernels/compiler_vector_load_store.cl @@ -1,18 +1,40 @@ /* test OpenCL 1.1 Vector Data Load/Store Functions (section 6.11.7) */ -kernel void compiler_vector_load_store() { - float p[16], f; - float4 f4; - f4 = vload4(0, p); - vstore4(f4, 0, p); - - long x[16], l; - long16 l16; - l = vload16(0, x); - vstore16(l16, 0, x); - half h[16]; - half4 h4; - f = vload_half(0, h); - f4 = vload_half4(0, h); - vstore_half(f, 0, h); +#define OFFSET2(type) (type ##2) {(type)1, (type)2} +#define OFFSET3(type) (type ##3) {(type)1, (type)2, (type)3} +#define OFFSET4(type) (type ##4) {(type)1, (type)2, (type)3, (type)4} +#define OFFSET8(type) (type ##8) {(type)1, (type)2, (type)3, (type)4, (type)5, (type)6, (type)7, (type)8} +#define OFFSET16(type) (type ##16) {(type)1, (type)2, (type)3, (type)4, (type)5, (type)6, (type)7, (type)8, (type)9, (type)10, (type)11, (type)12, (type)13, (type)14, (type)15, (type)16} + +#define TEST_TYPE(type, n) \ +__kernel void test_##type ##n(__global type *pin, \ + __global type *pout) \ +{\ + int x = get_global_id(0); \ + type ##n value; \ + value = vload ##n(x, pin); \ + value += OFFSET ##n(type); \ + vstore ##n(value, x, pout); \ } + +#define TEST_ALL_TYPE(n) \ + TEST_TYPE(char,n) \ + TEST_TYPE(uchar,n) \ + TEST_TYPE(short,n) \ + TEST_TYPE(ushort,n) \ + TEST_TYPE(int,n) \ + TEST_TYPE(uint,n) \ + TEST_TYPE(float,n) + +#if 0 + TEST_TYPE(double,n) + TEST_TYPE(long,n) + TEST_TYPE(ulong,n) + TEST_TYPE(half,n) +#endif + +TEST_ALL_TYPE(2) +//TEST_ALL_TYPE(3) +TEST_ALL_TYPE(4) +TEST_ALL_TYPE(8) +TEST_ALL_TYPE(16) diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 63c873d6..268cc51a 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -76,6 +76,7 @@ set (utests_sources compiler_volatile.cpp compiler_copy_image1.cpp compiler_get_image_info.cpp + compiler_vector_load_store.cpp runtime_createcontext.cpp utest_assert.cpp utest.cpp diff --git a/utests/compiler_vector_load_store.cpp b/utests/compiler_vector_load_store.cpp index 96fcfa97..76c12a11 100644 --- a/utests/compiler_vector_load_store.cpp +++ b/utests/compiler_vector_load_store.cpp @@ -1,10 +1,59 @@ #include "utest_helper.hpp" - -void compiler_vector_load_store(void) +template<typename T> +static void compiler_vector_load_store(int elemNum, const char *kernelName) { - OCL_CREATE_KERNEL("compiler_vector_load_store"); + const size_t n = elemNum * 256; + + // Setup kernel and buffers + OCL_CREATE_KERNEL_FROM_FILE("compiler_vector_load_store", kernelName); + buf_data[0] = (T*) malloc(sizeof(T) * n); + for (uint32_t i = 0; i < n; ++i) + ((T*)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] = n / elemNum; + locals[0] = 16; + OCL_NDRANGE(1); + + // Check result + OCL_MAP_BUFFER(0); + OCL_MAP_BUFFER(1); + for (uint32_t i = 0; i < n; ++i) + { + int shift = ((i % elemNum) + 1); + OCL_ASSERT(((T*)buf_data[1])[i] == (T)(((T*)buf_data[0])[i] + shift)); + } + OCL_UNMAP_BUFFER(0); + OCL_UNMAP_BUFFER(1); } -MAKE_UTEST_FROM_FUNCTION(compiler_vector_load_store); +#define compiler_vector_load_store(type, n, kernel_type) \ +static void compiler_vector_ ##kernel_type ##n ##_load_store(void)\ +{\ + compiler_vector_load_store<type>(n, "test_" #kernel_type #n);\ +}\ +MAKE_UTEST_FROM_FUNCTION(compiler_vector_ ## kernel_type ##n ##_load_store); +#define test_all_vector(type, kernel_type) \ + compiler_vector_load_store(type, 2, kernel_type) \ + /*compiler_vector_load_store(type, 3, kernel_type)*/ \ + compiler_vector_load_store(type, 4, kernel_type) \ + compiler_vector_load_store(type, 8, kernel_type) \ + compiler_vector_load_store(type, 16, kernel_type) +test_all_vector(int8_t, char) +test_all_vector(uint8_t, uchar) +test_all_vector(int16_t, short) +test_all_vector(uint16_t, ushort) +test_all_vector(int32_t, int) +test_all_vector(uint32_t, uint) +test_all_vector(float, float) +//test_all_vector(double, double) +//test_all_vector(int64_t, long) +//test_all_vector(uint64_t, ulong) |