--- /dev/null
+#define DECL_KERNEL_SUB(type)\
+__kernel void \
+compiler_sub_##type(__global type *src0, __global type *src1, __global type *dst) \
+{ \
+ int id = (int)get_global_id(0); \
+ dst[id] = src0[id] - src1[id]; \
+}
+
+#define DECL_KERNEL_ADD(type)\
+__kernel void \
+compiler_add_##type(__global type *src0, __global type *src1, __global type *dst) \
+{ \
+ int id = (int)get_global_id(0); \
+ dst[id] = src0[id] + src1[id]; \
+}
+
+#define DECL_KERNEL_MUL(type)\
+__kernel void \
+compiler_mul_##type(__global type *src0, __global type *src1, __global type *dst) \
+{ \
+ int id = (int)get_global_id(0); \
+ dst[id] = src0[id] * src1[id]; \
+}
+
+#define DECL_KERNEL_DIV(type)\
+__kernel void \
+compiler_div_##type(__global type *src0, __global type *src1, __global type *dst) \
+{ \
+ int id = (int)get_global_id(0); \
+ dst[id] = src0[id] / src1[id]; \
+}
+
+#define DECL_KERNEL_REM(type)\
+__kernel void \
+compiler_rem_##type(__global type *src0, __global type *src1, __global type *dst) \
+{ \
+ int id = (int)get_global_id(0); \
+ dst[id] = src0[id] % src1[id]; \
+}
+
+#define DECL_KERNEL_FOR_ALL_TYPE(op) \
+DECL_KERNEL_##op(char) \
+DECL_KERNEL_##op(uchar) \
+DECL_KERNEL_##op(short) \
+DECL_KERNEL_##op(ushort) \
+DECL_KERNEL_##op(int) \
+DECL_KERNEL_##op(uint)
+
+DECL_KERNEL_FOR_ALL_TYPE(SUB)
+DECL_KERNEL_FOR_ALL_TYPE(ADD)
+DECL_KERNEL_FOR_ALL_TYPE(MUL)
+DECL_KERNEL_FOR_ALL_TYPE(DIV)
+DECL_KERNEL_FOR_ALL_TYPE(REM)
+++ /dev/null
-__kernel void
-compiler_sub_bytes(__global char *src0, __global char *src1, __global char *dst)
-{
- int id = (int)get_global_id(0);
- dst[id] = src0[id] - src1[id];
-}
-
+++ /dev/null
-__kernel void
-compiler_sub_shorts(__global short *src0, __global short *src1, __global short *dst)
-{
- int id = (int)get_global_id(0);
- dst[id] = src0[id] - src1[id];
-}
-
set (utests_sources
cl_create_kernel.cpp
utest_error.c
+ compiler_basic_arithmetic.cpp
compiler_displacement_map_element.cpp
compiler_shader_toy.cpp
compiler_mandelbrot.cpp
compiler_saturate_sub.cpp
compiler_shift_right.cpp
compiler_short_scatter.cpp
- compiler_sub_bytes.cpp
- compiler_sub_shorts.cpp
compiler_uint2_copy.cpp
compiler_uint3_copy.cpp
compiler_uint8_copy.cpp
--- /dev/null
+#include "utest_helper.hpp"
+
+enum eTestOP {
+ TEST_OP_ADD =0,
+ TEST_OP_SUB,
+ TEST_OP_MUL,
+ TEST_OP_DIV,
+ TEST_OP_REM
+};
+
+template <typename T, eTestOP op>
+static void test_exec(const char* kernel_name)
+{
+ const size_t n = 160;
+
+ // Setup kernel and buffers
+ OCL_CREATE_KERNEL_FROM_FILE("compiler_basic_arithmetic", kernel_name);
+std::cout <<"kernel name: " << kernel_name << std::endl;
+ buf_data[0] = (T*) malloc(sizeof(T) * n);
+ buf_data[1] = (T*) malloc(sizeof(T) * n);
+ for (uint32_t i = 0; i < n; ++i) ((T*)buf_data[0])[i] = (T) rand();
+ for (uint32_t i = 0; i < n; ++i) ((T*)buf_data[1])[i] = (T) rand();
+ if(op == TEST_OP_DIV || op == TEST_OP_REM) {
+ for (uint32_t i = 0; i < n; ++i) {
+ if(((T*)buf_data[1])[i] == 0)
+ ((T*)buf_data[1])[i] = (T) 1;
+ }
+ }
+ OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(T), buf_data[0]);
+ OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(T), buf_data[1]);
+ OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(T), NULL);
+
+ // Run the kernel
+ 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]);
+ globals[0] = n;
+ locals[0] = 16;
+ OCL_NDRANGE(1);
+
+ // Check result
+ OCL_MAP_BUFFER(2);
+ if(op == TEST_OP_SUB) {
+ for (uint32_t i = 0; i < n; ++i)
+ OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] - ((T*)buf_data[1])[i]));
+ } else if(op == TEST_OP_ADD) {
+ for (uint32_t i = 0; i < n; ++i)
+ OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] + ((T*)buf_data[1])[i]));
+ } else if(op == TEST_OP_MUL) {
+ for (uint32_t i = 0; i < n; ++i)
+ OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] * ((T*)buf_data[1])[i]));
+ } else if(op == TEST_OP_DIV) {
+ for (uint32_t i = 0; i < n; ++i)
+ OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] / ((T*)buf_data[1])[i]));
+ } else {
+ for (uint32_t i = 0; i < n; ++i)
+ OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] % ((T*)buf_data[1])[i]));
+ }
+ free(buf_data[0]);
+ free(buf_data[1]);
+ buf_data[0] = buf_data[1] = NULL;
+}
+
+#define DECL_TEST_SUB(type, alias) \
+static void compiler_sub_ ##alias(void)\
+{\
+ test_exec<type, TEST_OP_SUB>("compiler_sub_" # alias);\
+}\
+MAKE_UTEST_FROM_FUNCTION(compiler_sub_ ## alias)
+
+#define DECL_TEST_ADD(type, alias) \
+static void compiler_add_ ##alias(void)\
+{\
+ test_exec<type, TEST_OP_ADD>("compiler_add_" # alias);\
+}\
+MAKE_UTEST_FROM_FUNCTION(compiler_add_ ## alias)
+
+#define DECL_TEST_MUL(type, alias) \
+static void compiler_mul_ ##alias(void)\
+{\
+ test_exec<type, TEST_OP_MUL>("compiler_mul_" # alias);\
+}\
+MAKE_UTEST_FROM_FUNCTION(compiler_mul_ ## alias)
+
+#define DECL_TEST_DIV(type, alias) \
+static void compiler_div_ ##alias(void)\
+{\
+ test_exec<type, TEST_OP_DIV>("compiler_div_" # alias);\
+}\
+MAKE_UTEST_FROM_FUNCTION(compiler_div_ ## alias)
+
+#define DECL_TEST_REM(type, alias) \
+static void compiler_rem_ ##alias(void)\
+{\
+ test_exec<type, TEST_OP_REM>("compiler_rem_" # alias);\
+}\
+MAKE_UTEST_FROM_FUNCTION(compiler_rem_ ## alias)
+
+#define DECL_TEST_FOR_ALL_TYPE(op)\
+DECL_TEST_##op(int8_t, char) \
+DECL_TEST_##op(uint8_t, uchar) \
+DECL_TEST_##op(int16_t, short) \
+DECL_TEST_##op(uint16_t, ushort) \
+DECL_TEST_##op(int32_t, int) \
+DECL_TEST_##op(uint32_t, uint)
+
+DECL_TEST_FOR_ALL_TYPE(SUB)
+DECL_TEST_FOR_ALL_TYPE(ADD)
+DECL_TEST_FOR_ALL_TYPE(MUL)
+DECL_TEST_FOR_ALL_TYPE(DIV)
+DECL_TEST_FOR_ALL_TYPE(REM)
+#undef DECL_TEST_FOR_ALL_TYPE
+++ /dev/null
-#include "utest_helper.hpp"
-
-static void compiler_sub_bytes(void)
-{
- const size_t n = 16;
-
- // Setup kernel and buffers
- OCL_CREATE_KERNEL("compiler_sub_bytes");
- buf_data[0] = (int8_t*) malloc(sizeof(int8_t) * n);
- buf_data[1] = (int8_t*) malloc(sizeof(int8_t) * n);
- for (uint32_t i = 0; i < n; ++i) ((int8_t*)buf_data[0])[i] = (int8_t) rand();
- for (uint32_t i = 0; i < n; ++i) ((int8_t*)buf_data[1])[i] = (int8_t) rand();
- OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(int8_t), buf_data[0]);
- OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(int8_t), buf_data[1]);
- OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int8_t), NULL);
-
- // Run the kernel
- 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]);
- globals[0] = n;
- locals[0] = 16;
- OCL_NDRANGE(1);
-
- // Check result
- OCL_MAP_BUFFER(2);
- for (uint32_t i = 0; i < n; ++i)
- OCL_ASSERT(((int8_t*)buf_data[2])[i] == (int8_t)(((int8_t*)buf_data[0])[i] - ((int8_t*)buf_data[1])[i]));
- free(buf_data[0]);
- free(buf_data[1]);
- buf_data[0] = buf_data[1] = NULL;
-}
-
-MAKE_UTEST_FROM_FUNCTION(compiler_sub_bytes);
-
+++ /dev/null
-#include "utest_helper.hpp"
-
-static void compiler_sub_shorts(void)
-{
- const size_t n = 16;
-
- // Setup kernel and buffers
- OCL_CREATE_KERNEL("compiler_sub_shorts");
- buf_data[0] = (int16_t*) malloc(sizeof(int16_t) * n);
- buf_data[1] = (int16_t*) malloc(sizeof(int16_t) * n);
- for (uint32_t i = 0; i < n; ++i) ((int16_t*)buf_data[0])[i] = (int16_t) rand();
- for (uint32_t i = 0; i < n; ++i) ((int16_t*)buf_data[1])[i] = (int16_t) rand();
- OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(int16_t), buf_data[0]);
- OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(int16_t), buf_data[1]);
- OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int16_t), NULL);
-
- // Run the kernel
- 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]);
- globals[0] = n;
- locals[0] = 16;
- OCL_NDRANGE(1);
-
- // Check result
- OCL_MAP_BUFFER(2);
- for (uint32_t i = 0; i < n; ++i)
- OCL_ASSERT(((int16_t*)buf_data[2])[i] == (int16_t)(((int16_t*)buf_data[0])[i] - ((int16_t*)buf_data[1])[i]));
- free(buf_data[0]);
- free(buf_data[1]);
- buf_data[0] = buf_data[1] = NULL;
-}
-
-MAKE_UTEST_FROM_FUNCTION(compiler_sub_shorts);
-
-