diff --git a/test_conformance/api/test_clone_kernel.cpp b/test_conformance/api/test_clone_kernel.cpp index b559b49efc..0bcc298564 100644 --- a/test_conformance/api/test_clone_kernel.cpp +++ b/test_conformance/api/test_clone_kernel.cpp @@ -22,79 +22,78 @@ using namespace std; -const char *clone_kernel_test_img[] = +const char* clone_kernel_test_img[] = { R"( +__kernel void img_read_kernel(read_only image2d_t img, sampler_t sampler, __global int* outbuf) { - "__kernel void img_read_kernel(read_only image2d_t img, sampler_t sampler, __global int* outbuf)\n" - "{\n" - " uint4 color;\n" - "\n" - " color = read_imageui(img, sampler, (int2)(0,0));\n" - " \n" - " // 7, 8, 9, 10th DWORD\n" - " outbuf[7] = color.x;\n" - " outbuf[8] = color.y;\n" - " outbuf[9] = color.z;\n" - " outbuf[10] = color.w;\n" - "}\n" - "\n" - "__kernel void img_write_kernel(write_only image2d_t img, uint4 color)\n" - "{\n" - " write_imageui (img, (int2)(0, 0), color);\n" - "}\n" + uint4 color; + + color = read_imageui(img, sampler, (int2)(0,0)); + + // 7, 8, 9, 10th DWORD + outbuf[7] = color.x; + outbuf[8] = color.y; + outbuf[9] = color.z; + outbuf[10] = color.w; +} -}; +__kernel void img_write_kernel(write_only image2d_t img, uint4 color) +{ + write_imageui (img, (int2)(0, 0), color); +} +)" }; -const char *clone_kernel_test_double[] = +const char* clone_kernel_test_double[] = { R"( +#pragma OPENCL EXTENSION cl_khr_fp64 : enable +__kernel void clone_kernel_test1(double d, __global double* outbuf) { - "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" - "__kernel void clone_kernel_test1(double d, __global double* outbuf)\n" - "{\n" - " // use the same outbuf as rest of the tests\n" - " outbuf[2] = d;\n" - "}\n" -}; + // use the same outbuf as rest of the tests + outbuf[2] = d; +} +)" }; -const char* clone_kernel_test_kernel[] = { - "typedef struct\n" - "{\n" - " int i;\n" - " float f;\n" - "} structArg;\n" - "\n" - "typedef struct {\n" - " __global int *store;\n" - "} BufPtr;\n" - "\n" - "// value type test\n" - "__kernel void clone_kernel_test0(int iarg, float farg, structArg sarg, " - "__local int* localbuf, __global int* outbuf)\n" - "{\n" - " int tid = get_global_id(0);\n" - "\n" - " outbuf[0] = iarg;\n" - " outbuf[1] = sarg.i;\n" - " \n" - " ((__global float*)outbuf)[2] = farg;\n" - " ((__global float*)outbuf)[3] = sarg.f;\n" - "}\n" - "\n" - "__kernel void buf_read_kernel(__global int* buf, __global int* outbuf)\n" - "{\n" - " // 6th DWORD\n" - " outbuf[6] = buf[0];\n" - "}\n" - "\n" - "__kernel void buf_write_kernel(__global int* buf, int write_val)\n" - "{\n" - " buf[0] = write_val;\n" - "}\n" - "__kernel void set_kernel_exec_info_kernel(int iarg, __global BufPtr* " - "buffer)\n" - "{\n" - " buffer->store[0] = iarg;\n" - "}\n" - "__kernel void test_kernel_empty(){}\n" -}; +const char* clone_kernel_test_kernel[] = { R"( +typedef struct +{ + int i; + float f; +} structArg; + +typedef struct { + __global int *store; +} BufPtr; + +// value type test +__kernel void clone_kernel_test0(int iarg, float farg, structArg sarg, +__local int* localbuf, __global int* outbuf) +{ + int tid = get_global_id(0); + + outbuf[0] = iarg; + outbuf[1] = sarg.i; + + ((__global float*)outbuf)[2] = farg; + ((__global float*)outbuf)[3] = sarg.f; +} + +__kernel void buf_read_kernel(__global int* buf, __global int* outbuf) +{ + // 6th DWORD + outbuf[6] = buf[0]; +} + +__kernel void buf_write_kernel(__global int* buf, int write_val) +{ + buf[0] = write_val; +} + +__kernel void set_kernel_exec_info_kernel(int iarg, __global BufPtr* buffer) +{ + buffer->store[0] = iarg; +} + +__kernel void test_kernel_empty(){} + +)" }; typedef struct { @@ -243,22 +242,21 @@ int test_double_arg_clone(cl_device_id deviceID, cl_context context, cl_command_ } int test_args_enqueue_helper(cl_context context, cl_command_queue queue, - cl_kernel* srcKernel, cl_int* value, - cl_mem* bufOut) + cl_kernel srcKernel, cl_int value, cl_mem bufOut) { cl_int error; size_t ndrange1 = 1; cl_int bufRes; // enqueue - srcKernel - error = clEnqueueNDRangeKernel(queue, *srcKernel, 1, NULL, &ndrange1, NULL, + error = clEnqueueNDRangeKernel(queue, srcKernel, 1, NULL, &ndrange1, NULL, 0, NULL, NULL); test_error(error, "clEnqueueNDRangeKernel failed"); - error = clEnqueueReadBuffer(queue, *bufOut, CL_TRUE, 0, sizeof(cl_int), + error = clEnqueueReadBuffer(queue, bufOut, CL_TRUE, 0, sizeof(cl_int), &bufRes, 0, NULL, NULL); test_error(error, "clEnqueueReadBuffer failed"); - test_assert_error(bufRes == *value, + test_assert_error(bufRes == value, "clCloneKernel test failed to verify integer value.\n"); return TEST_PASS; @@ -301,57 +299,41 @@ int test_cloned_kernel_args(cl_device_id deviceID, cl_context context, test_error(error, "clSetKernelArg failed for cloneKernel_2"); // enqueue - srcKernel - if (test_args_enqueue_helper(context, queue, &srcKernel, &intargs[0], - &bufOut) - != 0) - { - test_fail("test_args_enqueue_helper failed for srcKernel.\n"); - } + error = + test_args_enqueue_helper(context, queue, srcKernel, intargs[0], bufOut); + test_error(error, "test_args_enqueue_helper failed for srcKernel.\n"); // enqueue - cloneKernel_1 - if (test_args_enqueue_helper(context, queue, &cloneKernel_1, &intargs[1], - &bufOut) - != 0) - { - test_fail("test_args_enqueue_helper failed for cloneKernel_1.\n"); - } + error = test_args_enqueue_helper(context, queue, cloneKernel_1, intargs[1], + bufOut); + test_error(error, "test_args_enqueue_helper failed for cloneKernel_1.\n"); // enqueue - cloneKernel_2 - if (test_args_enqueue_helper(context, queue, &cloneKernel_2, &intargs[2], - &bufOut) - != 0) - { - test_fail("test_args_enqueue_helper failed for cloneKernel_2.\n"); - } + error = test_args_enqueue_helper(context, queue, cloneKernel_2, intargs[2], + bufOut); + test_error(error, "test_args_enqueue_helper failed for cloneKernel_2.\n"); // srcKernel, set different arg and enqueue error = clSetKernelArg(srcKernel, 1, sizeof(cl_int), &intargs[3]); test_error(error, "clSetKernelArg failed for srcKernel with different value"); - if (test_args_enqueue_helper(context, queue, &srcKernel, &intargs[3], - &bufOut) - != 0) - { - test_fail("test_args_enqueue_helper failed for srcKernel on retry.\n"); - } + error = + test_args_enqueue_helper(context, queue, srcKernel, intargs[3], bufOut); + test_error(error, + "test_args_enqueue_helper failed for srcKernel on retry.\n"); // enqueue - cloneKernel_1 again, to check if the args were not modified - if (test_args_enqueue_helper(context, queue, &cloneKernel_1, &intargs[1], - &bufOut) - != 0) - { - test_fail( - "test_args_enqueue_helper failed for cloneKernel_1 on retry.\n"); - } + error = test_args_enqueue_helper(context, queue, cloneKernel_1, intargs[1], + bufOut); + test_error(error, + "test_args_enqueue_helper failed for cloneKernel_1 on retry.\n"); // enqueue - cloneKernel_2 again, to check if the args were not modified - if (test_args_enqueue_helper(context, queue, &cloneKernel_2, &intargs[2], - &bufOut) - != 0) - { - test_fail( - "test_args_enqueue_helper failed for cloneKernel_2 on retry.\n"); - } + error = test_args_enqueue_helper(context, queue, cloneKernel_2, intargs[2], + bufOut); + test_error(error, + "test_args_enqueue_helper failed for cloneKernel_2 on retry.\n"); + return TEST_PASS; } @@ -508,24 +490,23 @@ int test_buff_image_multiple_args(cl_device_id deviceID, cl_context context, return TEST_PASS; } -int test_exec_enqueue_helper(cl_context context, cl_command_queue queue, - BufPtr* pBuf, cl_int* svmPtr_Kernel, - cl_kernel* srcKernel, cl_int* value) +int test_svm_enqueue_helper(cl_context context, cl_command_queue queue, + cl_int* svmPtr_Kernel, cl_kernel srcKernel, + cl_int value) { cl_int error; size_t ndrange1 = 1; // enqueue - srcKernel - error = clEnqueueNDRangeKernel(queue, *srcKernel, 1, NULL, &ndrange1, NULL, + error = clEnqueueNDRangeKernel(queue, srcKernel, 1, NULL, &ndrange1, NULL, 0, NULL, NULL); test_error(error, "clEnqueueNDRangeKernel failed"); - error = clFinish(queue); - test_error(error, "clFinish failed"); - error = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, - svmPtr_Kernel, sizeof(cl_int), 0, NULL, NULL); + error = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ, svmPtr_Kernel, + sizeof(cl_int), 0, NULL, NULL); test_error(error, "clEnqueueSVMMap failed"); - test_assert_error(svmPtr_Kernel[0] == *value, + + test_assert_error(svmPtr_Kernel[0] == value, "clCloneKernel test failed, Failed to verify " "integer value from SVM pointer. "); @@ -537,9 +518,10 @@ int test_exec_enqueue_helper(cl_context context, cl_command_queue queue, return TEST_PASS; } + int test_svm_exec_info_helper(cl_context context, cl_command_queue queue, BufPtr* pBuf, cl_int* svmPtr_Kernel, - cl_kernel* srcKernel, cl_int* value) + cl_kernel srcKernel, cl_int value) { cl_int error; @@ -554,23 +536,22 @@ int test_svm_exec_info_helper(cl_context context, cl_command_queue queue, error = clFinish(queue); test_error(error, "clFinish failed"); - error = clSetKernelArg(*srcKernel, 0, sizeof(cl_int), value); + error = clSetKernelArg(srcKernel, 0, sizeof(cl_int), &value); test_error(error, "clSetKernelArg failed"); - error = clSetKernelArgSVMPointer(*srcKernel, 1, pBuf); + error = clSetKernelArgSVMPointer(srcKernel, 1, pBuf); test_error(error, "clSetKernelArgSVMPointer failed"); - error = clSetKernelExecInfo(*srcKernel, CL_KERNEL_EXEC_INFO_SVM_PTRS, - sizeof(pBuf), pBuf); + + void* exec_info_ptrs[] = { svmPtr_Kernel }; + error = clSetKernelExecInfo(srcKernel, CL_KERNEL_EXEC_INFO_SVM_PTRS, + sizeof(exec_info_ptrs), exec_info_ptrs); test_error(error, "clSetKernelExecInfo failed"); error = clFinish(queue); test_error(error, "clFinish failed"); - if (test_exec_enqueue_helper(context, queue, pBuf, svmPtr_Kernel, srcKernel, - value) - != 0) - { - test_fail("test_exec_enqueue_helper failed.\n"); - } + error = test_svm_enqueue_helper(context, queue, svmPtr_Kernel, srcKernel, + value); + test_error(error, "test_svm_enqueue_helper failed."); return TEST_PASS; } @@ -615,61 +596,54 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, "clSVMAlloc returned NULL"); // srcKernel, set args - if (test_svm_exec_info_helper(context, queue, pBuf, svmPtr_srcKernel, - &srcKernel, &intargs[0]) - != 0) - { - test_fail("test_svm_exec_info_helper failed for srcKernel.\n"); - } + error = test_svm_exec_info_helper( + context, queue, pBuf, svmPtr_srcKernel, srcKernel, intargs[0]); + test_error(error, "test_svm_exec_info_helper failed for srcKernel."); // clone the srcKernel and set args clKernelWrapper cloneKernel_1 = clCloneKernel(srcKernel, &error); test_error(error, "clCloneKernel failed for cloneKernel_1"); - if (test_svm_exec_info_helper(context, queue, pBuf, - svmPtr_cloneKernel_1, &cloneKernel_1, - &intargs[1]) - != 0) - { - test_fail("test_svm_exec_info_helper failed for cloneKernel_1.\n"); - } + error = test_svm_exec_info_helper(context, queue, pBuf, + svmPtr_cloneKernel_1, cloneKernel_1, + intargs[1]); + test_error(error, + "test_svm_exec_info_helper failed for cloneKernel_1."); + // clone the cloneKernel_1 and set args clKernelWrapper cloneKernel_2 = clCloneKernel(cloneKernel_1, &error); test_error(error, "clCloneKernel failed for cloneKernel_2"); - if (test_svm_exec_info_helper(context, queue, pBuf, - svmPtr_cloneKernel_2, &cloneKernel_2, - &intargs[2]) - != 0) - { - test_fail("test_svm_exec_info_helper failed for cloneKernel_2.\n"); - } + + error = test_svm_exec_info_helper(context, queue, pBuf, + svmPtr_cloneKernel_2, cloneKernel_2, + intargs[2]); + test_error(error, + "test_svm_exec_info_helper failed for cloneKernel_2."); + // enqueue - srcKernel again with different svm_ptr and args - if (test_svm_exec_info_helper(context, queue, pBuf, svmPtr_srcKernel_1, - &srcKernel, &intargs[3]) - != 0) - { - test_fail("test_svm_exec_info_helper failed for srcKernel with " - "different values.\n"); - } + error = test_svm_exec_info_helper( + context, queue, pBuf, svmPtr_srcKernel_1, srcKernel, intargs[3]); + test_error(error, + "test_svm_exec_info_helper failed for srcKernel with " + "different values."); + // enqueue - cloneKernel_1 again, to check if the args were not modified - if (test_exec_enqueue_helper(context, queue, pBuf, svmPtr_cloneKernel_1, - &cloneKernel_1, &intargs[1]) - != 0) - { - test_fail("test_exec_enqueue_helper failed for cloneKernel_1 on " - "retry.\n"); - } + error = test_svm_enqueue_helper(context, queue, svmPtr_cloneKernel_1, + cloneKernel_1, intargs[1]); + test_error( + error, + "test_svm_enqueue_helper failed for cloneKernel_1 on retry."); + // enqueue - cloneKernel_2 again, to check if the args were not modified - if (test_exec_enqueue_helper(context, queue, pBuf, svmPtr_cloneKernel_2, - &cloneKernel_2, &intargs[2]) - != 0) - { - test_fail("test_exec_enqueue_helper failed for cloneKernel_2 on " - "retry.\n"); - } + error = test_svm_enqueue_helper(context, queue, svmPtr_cloneKernel_2, + cloneKernel_2, intargs[2]); + test_error( + error, + "test_svm_enqueue_helper failed for cloneKernel_2 on retry."); + clSVMFree(context, svmPtr_srcKernel); clSVMFree(context, svmPtr_srcKernel_1); @@ -685,13 +659,13 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, } } -int test_empty_enqueue_helper(cl_command_queue queue, cl_kernel* srcKernel) +int test_empty_enqueue_helper(cl_command_queue queue, cl_kernel srcKernel) { cl_int error; size_t ndrange1 = 1; // enqueue - srcKernel - error = clEnqueueNDRangeKernel(queue, *srcKernel, 1, NULL, &ndrange1, NULL, + error = clEnqueueNDRangeKernel(queue, srcKernel, 1, NULL, &ndrange1, NULL, 0, NULL, NULL); test_error(error, "clEnqueueNDRangeKernel failed"); @@ -716,7 +690,7 @@ int test_cloned_kernel_empty_args(cl_device_id deviceID, cl_context context, "Unable to create srcKernel for test_cloned_kernel_empty_args"); // enqueue - srcKernel - if (test_empty_enqueue_helper(queue, &srcKernel) != 0) + if (test_empty_enqueue_helper(queue, srcKernel) != 0) { test_fail("test_empty_enqueue_helper failed for srcKernel.\n"); } @@ -725,69 +699,39 @@ int test_cloned_kernel_empty_args(cl_device_id deviceID, cl_context context, clKernelWrapper cloneKernel_1 = clCloneKernel(srcKernel, &error); test_error(error, "clCloneKernel failed for cloneKernel_1"); - if (test_empty_enqueue_helper(queue, &cloneKernel_1) != 0) - { - test_fail("test_empty_enqueue_helper failed for cloneKernel_1.\n"); - } - - // enqueue - srcKernel again - if (test_empty_enqueue_helper(queue, &srcKernel) != 0) - { - test_fail("test_empty_enqueue_helper failed for srcKernel on retry.\n"); - } - - return TEST_PASS; -} + error = test_empty_enqueue_helper(queue, cloneKernel_1); -int test_svm_enqueue_helper(cl_context context, cl_command_queue queue, - cl_int* svmPtr_Kernel, cl_kernel* srcKernel, - cl_int* value) -{ - cl_int error; - size_t ndrange1 = 1; + test_error(error, "test_empty_enqueue_helper failed for cloneKernel_1."); - // enqueue - srcKernel - error = clEnqueueNDRangeKernel(queue, *srcKernel, 1, NULL, &ndrange1, NULL, - 0, NULL, NULL); - test_error(error, "clEnqueueNDRangeKernel failed"); - error = clFinish(queue); - test_error(error, "clFinish failed"); - error = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, - svmPtr_Kernel, sizeof(cl_int), 0, NULL, NULL); - test_error(error, "clEnqueueSVMMap failed"); + // enqueue - srcKernel again + error = test_empty_enqueue_helper(queue, srcKernel); - test_assert_error(svmPtr_Kernel[0] == *value, - "clCloneKernel test failed, Failed to verify " - "integer value from SVM pointer. "); + test_error(error, + "test_empty_enqueue_helper failed for srcKernel on retry."); - error = clEnqueueSVMUnmap(queue, svmPtr_Kernel, 0, NULL, NULL); - test_error(error, "clEnqueueSVMUnmap failed"); - error = clFinish(queue); - test_error(error, "clFinish failed"); return TEST_PASS; } + int test_svm_ptr_helper(cl_context context, cl_command_queue queue, - cl_int* svmPtr_Kernel, cl_kernel* srcKernel, - cl_int* value) + cl_int* svmPtr_Kernel, cl_kernel srcKernel, + cl_int value) { cl_int error; - error = clSetKernelArgSVMPointer(*srcKernel, 0, svmPtr_Kernel); + error = clSetKernelArgSVMPointer(srcKernel, 0, svmPtr_Kernel); test_error(error, "clSetKernelArgSVMPointer failed"); - error = clSetKernelArg(*srcKernel, 1, sizeof(cl_int), value); + error = clSetKernelArg(srcKernel, 1, sizeof(cl_int), &value); test_error(error, "clSetKernelArg failed"); error = clFinish(queue); test_error(error, "clFinish failed"); - if (test_svm_enqueue_helper(context, queue, svmPtr_Kernel, srcKernel, value) - != 0) - { - test_fail("test_svm_enqueue_helper failed.\n"); - } + error = test_svm_enqueue_helper(context, queue, svmPtr_Kernel, srcKernel, + value); + test_error(error, "test_svm_enqueue_helper failed."); return TEST_PASS; } @@ -829,59 +773,49 @@ int test_cloned_kernel_svm_ptr(cl_device_id deviceID, cl_context context, "clSVMAlloc returned NULL"); // srcKernel, set args - if (test_svm_ptr_helper(context, queue, svmPtr_srcKernel, &srcKernel, - &intargs[0]) - != 0) - { - test_fail("test_svm_ptr_helper failed for srcKernel.\n"); - } + error = test_svm_ptr_helper(context, queue, svmPtr_srcKernel, srcKernel, + intargs[0]); + + test_error(error, "test_svm_ptr_helper failed for srcKernel."); + // clone the srcKernel and set args clKernelWrapper cloneKernel_1 = clCloneKernel(srcKernel, &error); test_error(error, "clCloneKernel failed for cloneKernel_1"); - if (test_svm_ptr_helper(context, queue, svmPtr_cloneKernel_1, - &cloneKernel_1, &intargs[1]) - != 0) - { - test_fail("test_svm_ptr_helper failed for cloneKernel_1.\n"); - } + error = test_svm_ptr_helper(context, queue, svmPtr_cloneKernel_1, + cloneKernel_1, intargs[1]); + test_error(error, "test_svm_ptr_helper failed for cloneKernel_1."); + // clone the cloneKernel_1 and set args clKernelWrapper cloneKernel_2 = clCloneKernel(cloneKernel_1, &error); test_error(error, "clCloneKernel failed for cloneKernel_2"); - if (test_svm_ptr_helper(context, queue, svmPtr_cloneKernel_2, - &cloneKernel_2, &intargs[2]) - != 0) - { - test_fail("test_svm_ptr_helper failed for cloneKernel_2.\n"); - } + error = test_svm_ptr_helper(context, queue, svmPtr_cloneKernel_2, + cloneKernel_2, intargs[2]); + test_error(error, "test_svm_ptr_helper failed for cloneKernel_2."); // enqueue - srcKernel again with different svm_ptr and args - if (test_svm_ptr_helper(context, queue, svmPtr_srcKernel_1, &srcKernel, - &intargs[3]) - != 0) - { - test_fail("test_svm_ptr_helper failed for srcKernel with " - "different values.\n"); - } + error = test_svm_ptr_helper(context, queue, svmPtr_srcKernel_1, + srcKernel, intargs[3]); + test_error( + error, + "test_svm_ptr_helper failed for srcKernel with different values."); + // enqueue - cloneKernel_1 again, to check if the args were not modified - if (test_svm_enqueue_helper(context, queue, svmPtr_cloneKernel_1, - &cloneKernel_1, &intargs[1]) - != 0) - { - test_fail( - "test_svm_enqueue_helper failed for cloneKernel_1 on retry.\n"); - } + error = test_svm_enqueue_helper(context, queue, svmPtr_cloneKernel_1, + cloneKernel_1, intargs[1]); + test_error( + error, + "test_svm_enqueue_helper failed for cloneKernel_1 on retry."); // enqueue - cloneKernel_2 again, to check if the args were not modified - if (test_svm_enqueue_helper(context, queue, svmPtr_cloneKernel_2, - &cloneKernel_2, &intargs[2]) - != 0) - { - test_fail("test_svm_enqueue_helper failed for cloneKernel_2 on " - "retry.\n"); - } + error = test_svm_enqueue_helper(context, queue, svmPtr_cloneKernel_2, + cloneKernel_2, intargs[2]); + test_error( + error, + "test_svm_enqueue_helper failed for cloneKernel_2 on retry."); + clSVMFree(context, svmPtr_srcKernel); clSVMFree(context, svmPtr_srcKernel_1);