From 6bf4c0feddd971b3cce6bfcec5d8f832f1d32b9a Mon Sep 17 00:00:00 2001 From: Vikas Katariya Date: Tue, 20 Apr 2021 14:54:19 +0100 Subject: [PATCH] Add test coverage for `clCloneKernel` with Execinfo Use `clSetKernelExecInfo` after kernel is cloned and read the buffer to validate. The test uses the `set_kernel_exec_info_kernel` kernel program with 2 arguments. Signed-off-by: Vikas Katariya --- test_conformance/api/test_clone_kernel.cpp | 196 ++++++++++++++++++++- 1 file changed, 195 insertions(+), 1 deletion(-) diff --git a/test_conformance/api/test_clone_kernel.cpp b/test_conformance/api/test_clone_kernel.cpp index 7e98bec76..1a94ff2e7 100644 --- a/test_conformance/api/test_clone_kernel.cpp +++ b/test_conformance/api/test_clone_kernel.cpp @@ -61,6 +61,10 @@ const char* clone_kernel_test_kernel[] = { " 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" @@ -84,9 +88,18 @@ const char* clone_kernel_test_kernel[] = { "{\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" }; +typedef struct +{ + cl_int* store; +} BufPtr; + const int BUF_SIZE = 128; struct structArg @@ -494,6 +507,181 @@ 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) +{ + cl_int error; + size_t ndrange1 = 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"); + test_assert_error(svmPtr_Kernel[0] == *value, + "clCloneKernel test failed, Failed to verify " + "integer value from SVM pointer. "); + + 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_exec_info_helper(cl_context context, cl_command_queue queue, + BufPtr* pBuf, cl_int* svmPtr_Kernel, + cl_kernel* srcKernel, cl_int* value) +{ + cl_int error; + + error = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, pBuf, + sizeof(BufPtr), 0, NULL, NULL); + test_error(error, "clEnqueueSVMMap failed"); + error = clFinish(queue); + test_error(error, "clFinish failed"); + + pBuf->store = svmPtr_Kernel; + + error = clEnqueueSVMUnmap(queue, pBuf, 0, NULL, NULL); + test_error(error, "clEnqueueSVMUnmap failed"); + error = clFinish(queue); + test_error(error, "clFinish failed"); + + error = clSetKernelArg(*srcKernel, 0, sizeof(cl_int), value); + test_error(error, "clSetKernelArg failed"); + error = clSetKernelArgSVMPointer(*srcKernel, 1, pBuf); + test_error(error, "clSetKernelArgSVMPointer failed"); + error = clSetKernelExecInfo(*srcKernel, CL_KERNEL_EXEC_INFO_SVM_PTRS, + sizeof(pBuf), pBuf); + 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"); + } + + return TEST_PASS; +} + +int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) +{ + cl_int error; + + clMemWrapper bufOut; + clProgramWrapper program; + clKernelWrapper srcKernel; + + cl_int intargs[] = { 1, 2, 3, 4 }; + cl_device_svm_capabilities svmCaps = 0; + + error = clGetDeviceInfo(deviceID, CL_DEVICE_SVM_CAPABILITIES, + sizeof(svmCaps), &svmCaps, NULL); + test_error(error, "Unable to query CL_DEVICE_SVM_CAPABILITIES"); + + if (svmCaps != 0) + { + error = create_single_kernel_helper(context, &program, &srcKernel, 1, + clone_kernel_test_kernel, + "set_kernel_exec_info_kernel"); + test_error(error, "Unable to create srcKernel"); + + BufPtr* pBuf = + (BufPtr*)clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(BufPtr), 0); + cl_int* svmPtr_srcKernel = + (cl_int*)clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(cl_int), 0); + cl_int* svmPtr_srcKernel_1 = + (cl_int*)clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(cl_int), 0); + cl_int* svmPtr_cloneKernel_1 = + (cl_int*)clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(cl_int), 0); + cl_int* svmPtr_cloneKernel_2 = + (cl_int*)clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(cl_int), 0); + test_assert_error(pBuf != NULL || svmPtr_srcKernel != NULL + || svmPtr_cloneKernel_1 != NULL + || svmPtr_srcKernel_1 != NULL + || svmPtr_cloneKernel_2 != NULL, + "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"); + } + clSVMFree(context, svmPtr_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"); + } + + // 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"); + } + + // 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"); + } + clSVMFree(context, svmPtr_srcKernel_1); + + // 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"); + } + clSVMFree(context, svmPtr_cloneKernel_1); + + // 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"); + } + clSVMFree(context, svmPtr_cloneKernel_2); + + clSVMFree(context, pBuf); + } + + return TEST_PASS; +} + int test_clone_kernel(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { @@ -508,5 +696,11 @@ int test_clone_kernel(cl_device_id deviceID, cl_context context, test_fail("clCloneKernel test_cloned_kernel_args failed.\n"); } + if (test_cloned_kernel_exec_info(deviceID, context, queue, num_elements) + != 0) + { + test_fail("clCloneKernel test_cloned_kernel_exec_info failed.\n"); + } + return TEST_PASS; }