diff options
| -rw-r--r-- | kernels/compiler_vector_load_store.cl | 8 | ||||
| -rw-r--r-- | utests/compiler_vector_load_store.cpp | 32 |
2 files changed, 31 insertions, 9 deletions
diff --git a/kernels/compiler_vector_load_store.cl b/kernels/compiler_vector_load_store.cl index aec38b19..d423174f 100644 --- a/kernels/compiler_vector_load_store.cl +++ b/kernels/compiler_vector_load_store.cl @@ -17,6 +17,7 @@ __kernel void test_##type ##n(__global type *pin, \ vstore ##n(value, x, pout); \ } +#ifndef HALF #define TEST_ALL_TYPE(n) \ TEST_TYPE(char,n) \ TEST_TYPE(uchar,n) \ @@ -27,11 +28,12 @@ __kernel void test_##type ##n(__global type *pin, \ TEST_TYPE(float,n) \ TEST_TYPE(long,n) \ TEST_TYPE(ulong,n) -// TEST_TYPE(double,n) - -#if 0 +#else +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +#define TEST_ALL_TYPE(n) \ TEST_TYPE(half,n) #endif +// TEST_TYPE(double,n) TEST_ALL_TYPE(2) TEST_ALL_TYPE(3) diff --git a/utests/compiler_vector_load_store.cpp b/utests/compiler_vector_load_store.cpp index 5a1a8d12..80e72a98 100644 --- a/utests/compiler_vector_load_store.cpp +++ b/utests/compiler_vector_load_store.cpp @@ -1,15 +1,27 @@ #include "utest_helper.hpp" #include <string.h> +#include <math.h> template<typename T> static void compiler_vector_load_store(int elemNum, const char *kernelName) { const size_t n = elemNum * 256; + if (strstr(kernelName, "half") != NULL) + if (!cl_check_half()) + return; // Setup kernel and buffers - OCL_CREATE_KERNEL_FROM_FILE("compiler_vector_load_store", kernelName); + if (strstr(kernelName, "half") != NULL) + OCL_CALL(cl_kernel_init, "compiler_vector_load_store.cl", kernelName, + SOURCE, "-DHALF"); + else + 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; + for (uint32_t i = 0; i < n; ++i) { + if (strstr(kernelName, "half") != NULL) + ((T*)buf_data[0])[i] = __float_to_half(as_uint((float)i/(float)n)); + else + ((T*)buf_data[0])[i] = i; + } OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(T), buf_data[0]); OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(T), NULL); free(buf_data[0]); @@ -28,10 +40,17 @@ static void compiler_vector_load_store(int elemNum, const char *kernelName) for (uint32_t i = 0; i < n; ++i) { int shift = ((i % elemNum) + 1); - if (strstr(kernelName, "double") == NULL) - OCL_ASSERT(((T*)buf_data[1])[i] == (T)(((T*)buf_data[0])[i] + shift)); - else + if (strstr(kernelName, "double") != NULL) OCL_ASSERT((((T*)buf_data[1])[i] - ((T)((T*)buf_data[0])[i] + shift)) < 1e-5); + else if (strstr(kernelName, "half") != NULL) { + float fdst = as_float(__half_to_float(((T*)buf_data[1])[i])); + float fsrc = as_float(__half_to_float((T)(((T*)buf_data[0])[i]))); + fsrc += shift; + //printf("%d (%f, %f)\n",i, fdst, fsrc); + OCL_ASSERT((fabs(fsrc - fdst) <= 0.03 * fabs(fdst))); + } + else + OCL_ASSERT(((T*)buf_data[1])[i] == (T)(((T*)buf_data[0])[i] + shift)); } OCL_UNMAP_BUFFER(0); OCL_UNMAP_BUFFER(1); @@ -61,3 +80,4 @@ test_all_vector(float, float, true) //test_all_vector(double, double, true) test_all_vector(int64_t, long, true) test_all_vector(uint64_t, ulong, false) +test_all_vector(uint16_t, half, false) |