From 3d58b8b333b552003d4418b2d767ee4a37ffd1fd Mon Sep 17 00:00:00 2001 From: Vikas Katariya Date: Thu, 15 Apr 2021 17:03:55 +0100 Subject: [PATCH 1/7] Add test coverage for `clCloneKernel` with args Use `clSetKernelArg` to set args after kernel is cloned. Enqueue and read the buffer to validate. The test uses `buf_write_kernel` kernel program with 2 arguments. Signed-off-by: Vikas Katariya --- test_conformance/api/test_clone_kernel.cpp | 200 +++++++++++++++++---- 1 file changed, 166 insertions(+), 34 deletions(-) diff --git a/test_conformance/api/test_clone_kernel.cpp b/test_conformance/api/test_clone_kernel.cpp index cc95c9b055..9a44b2107a 100644 --- a/test_conformance/api/test_clone_kernel.cpp +++ b/test_conformance/api/test_clone_kernel.cpp @@ -1,6 +1,6 @@ // -// Copyright (c) 2017 The Khronos Group Inc. -// +// Copyright (c) 2017-2021 The Khronos Group Inc. +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at @@ -54,37 +54,38 @@ const char *clone_kernel_test_double[] = "}\n" }; -const char *clone_kernel_test_kernel[] = { -"typedef struct\n" -"{\n" -" int i;\n" -" float f;\n" -"} structArg;\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" - - }; +const char* clone_kernel_test_kernel[] = { + "typedef struct\n" + "{\n" + " int i;\n" + " float f;\n" + "} structArg;\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" + +}; const int BUF_SIZE = 128; @@ -230,7 +231,122 @@ int test_double_arg_clone(cl_device_id deviceID, cl_context context, cl_command_ return 0; } -int test_clone_kernel(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_args_enqueue_helper(cl_context context, cl_command_queue queue, + 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, + 0, NULL, NULL); + test_error(error, "clEnqueueNDRangeKernel failed"); + error = clEnqueueReadBuffer(queue, *bufOut, CL_TRUE, 0, sizeof(cl_int), + &bufRes, 0, NULL, NULL); + test_error(error, "clEnqueueReadBuffer failed"); + + test_assert_error(bufRes == *value, + "clCloneKernel test failed to verify integer value.\n"); + + return TEST_PASS; +} + +int test_cloned_kernel_args(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) +{ + cl_int error; + clProgramWrapper program; + clKernelWrapper srcKernel; + cl_int intargs[] = { 1, 2, 3, 4 }; + clMemWrapper bufOut; + + // Create srcKernel to test with + error = create_single_kernel_helper(context, &program, &srcKernel, 1, + clone_kernel_test_kernel, + "buf_write_kernel"); + test_error(error, "Unable to create srcKernel for test_cloned_kernel_args"); + + bufOut = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL, + &error); + test_error(error, "clCreateBuffer failed."); + + // srcKernel, set args + error = clSetKernelArg(srcKernel, 1, sizeof(cl_int), &intargs[0]); + error |= clSetKernelArg(srcKernel, 0, sizeof(cl_mem), &bufOut); + test_error(error, "clSetKernelArg failed for srcKernel"); + + // clone the srcKernel and set different arg + clKernelWrapper cloneKernel_1 = clCloneKernel(srcKernel, &error); + test_error(error, "clCloneKernel failed for cloneKernel_1"); + error = clSetKernelArg(cloneKernel_1, 1, sizeof(cl_int), &intargs[1]); + test_error(error, "clSetKernelArg failed for cloneKernel_1"); + + // clone the cloneKernel_1 and set different arg + clKernelWrapper cloneKernel_2 = clCloneKernel(cloneKernel_1, &error); + test_error(error, "clCloneKernel failed for cloneKernel_2"); + error = clSetKernelArg(cloneKernel_2, 1, sizeof(cl_int), &intargs[2]); + 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"); + } + + // 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"); + } + + // 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"); + } + + // 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"); + } + + // 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"); + } + + // 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"); + } + + return TEST_PASS; +} + +int test_buff_image_multiple_args(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { int error; clProgramWrapper program; @@ -395,3 +511,19 @@ int test_clone_kernel(cl_device_id deviceID, cl_context context, cl_command_queu return 0; } +int test_clone_kernel(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) +{ + if (test_buff_image_multiple_args(deviceID, context, queue, num_elements) + != 0) + { + test_fail("clCloneKernel test_buff_image_multiple_args failed.\n"); + } + + if (test_cloned_kernel_args(deviceID, context, queue, num_elements) != 0) + { + test_fail("clCloneKernel test_cloned_kernel_args failed.\n"); + } + + return TEST_PASS; +} From e422edc18d5384cedeaeabd8dbc84731c9227eb3 Mon Sep 17 00:00:00 2001 From: Vikas Katariya Date: Fri, 16 Apr 2021 09:43:10 +0100 Subject: [PATCH 2/7] Fix: Replace `test_error` with `test_assert_error` In test_clone_kernel, if `clEnqueueReadBuffer` was a success then the error code would be `CL_SUCCESS`, which will not print the error message when buffer validation fails, therefore replace with `test_assert_error` to print the error message. Signed-off-by: Vikas Katariya --- test_conformance/api/test_clone_kernel.cpp | 97 +++++++++------------- 1 file changed, 38 insertions(+), 59 deletions(-) diff --git a/test_conformance/api/test_clone_kernel.cpp b/test_conformance/api/test_clone_kernel.cpp index 9a44b2107a..e936e1b0fa 100644 --- a/test_conformance/api/test_clone_kernel.cpp +++ b/test_conformance/api/test_clone_kernel.cpp @@ -167,31 +167,19 @@ int test_image_arg_shallow_clone(cl_device_id deviceID, cl_context context, cl_c error = clEnqueueReadBuffer(queue, bufOut, CL_TRUE, 0, 128, pbufRes, 0, NULL, NULL); test_error( error, "clEnqueueReadBuffer failed." ); - if (((cl_uint*)pbufRes)[7] != color[0]) - { - test_error( error, "clCloneKernel test failed." ); - return -1; - } + test_assert_error(((cl_uint*)pbufRes)[7] == color[0], + "clCloneKernel test failed."); - if (((cl_uint*)pbufRes)[8] != color[1]) - { - test_error( error, "clCloneKernel test failed." ); - return -1; - } + test_assert_error(((cl_uint*)pbufRes)[8] == color[1], + "clCloneKernel test failed."); - if (((cl_uint*)pbufRes)[9] != color[2]) - { - test_error( error, "clCloneKernel test failed." ); - return -1; - } + test_assert_error(((cl_uint*)pbufRes)[9] == color[2], + "clCloneKernel test failed."); - if (((cl_uint*)pbufRes)[10] != color[3]) - { - test_error( error, "clCloneKernel test failed." ); - return -1; - } + test_assert_error(((cl_uint*)pbufRes)[10] == color[3], + "clCloneKernel test failed."); - return 0; + return TEST_PASS; } int test_double_arg_clone(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, void* pbufRes, clMemWrapper& bufOut) @@ -222,13 +210,10 @@ int test_double_arg_clone(cl_device_id deviceID, cl_context context, cl_command_ error = clEnqueueReadBuffer(queue, bufOut, CL_TRUE, 0, BUF_SIZE, pbufRes, 0, NULL, NULL); test_error( error, "clEnqueueReadBuffer failed." ); - if (abs(((cl_double*)pbufRes)[2] - d) > 0.0000001) - { - test_error( error, "clCloneKernel test failed." ); - return -1; - } + test_assert_error(abs(((cl_double*)pbufRes)[2] - d) <= 0.0000001, + "clCloneKernel test failed."); - return 0; + return TEST_PASS; } int test_args_enqueue_helper(cl_context context, cl_command_queue queue, @@ -463,52 +448,46 @@ int test_buff_image_multiple_args(cl_device_id deviceID, cl_context context, test_error( error, "clEnqueueReadBuffer failed." ); // Compare the results - if (((int*)pbufRes)[0] != intarg) - { - test_error( error, "clCloneKernel test failed. Failed to clone integer type argument." ); - return -1; - } + test_assert_error(((int*)pbufRes)[0] == intarg, + "clCloneKernel test failed. Failed to clone integer type " + "argument."); - if (((int*)pbufRes)[1] != sa.i) - { - test_error( error, "clCloneKernel test failed. Failed to clone structure type argument." ); - return -1; - } + test_assert_error( + ((int*)pbufRes)[1] == sa.i, + "clCloneKernel test failed. Failed to clone structure type " + "argument."); - if (((float*)pbufRes)[2] != farg) - { - test_error( error, "clCloneKernel test failed. Failed to clone structure type argument." ); - return -1; - } + test_assert_error( + ((float*)pbufRes)[2] == farg, + "clCloneKernel test failed. Failed to clone float type argument."); - if (((float*)pbufRes)[3] != sa.f) - { - test_error( error, "clCloneKernel test failed. Failed to clone float type argument." ); - return -1; - } + test_assert_error( + ((float*)pbufRes)[3] == sa.f, + "clCloneKernel test failed. Failed to clone structure type " + "argument."); - if (((int*)pbufRes)[6] != write_val) - { - test_error( error, "clCloneKernel test failed. Failed to clone cl_mem argument." ); - return -1; - } + test_assert_error( + ((int*)pbufRes)[6] == write_val, + "clCloneKernel test failed. Failed to clone cl_mem argument."); if (bimg) { - error = test_image_arg_shallow_clone(deviceID, context, queue, num_elements, pbufRes, bufOut); - test_error( error, "image arg shallow clone test failed." ); + error = test_image_arg_shallow_clone(deviceID, context, queue, + num_elements, pbufRes, bufOut); + test_error(error, "image arg shallow clone test failed."); } if (bdouble) { - error = test_double_arg_clone(deviceID, context, queue, num_elements, pbufRes, bufOut); - test_error( error, "double arg clone test failed." ); + error = test_double_arg_clone(deviceID, context, queue, num_elements, + pbufRes, bufOut); + test_error(error, "double arg clone test failed."); } - delete [] pbuf; - delete [] pbufRes; + delete[] pbuf; + delete[] pbufRes; - return 0; + return TEST_PASS; } int test_clone_kernel(cl_device_id deviceID, cl_context context, From 771b05a780bee6da98a8be2c547b31da3a7f1b63 Mon Sep 17 00:00:00 2001 From: Vikas Katariya Date: Tue, 20 Apr 2021 14:54:19 +0100 Subject: [PATCH 3/7] 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 e936e1b0fa..c3ca30de95 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 @@ -490,6 +503,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) { @@ -504,5 +692,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; } From 2bdcdf738ca95b2752b4fe830488f6ea53b55a00 Mon Sep 17 00:00:00 2001 From: Vikas Katariya Date: Mon, 10 May 2021 16:22:39 +0100 Subject: [PATCH 4/7] Add test covergae for `clCloneKernel` with no args Clone a kernel with no args and enqueue. The test uses `test_kernel_empty` kernel program with no arguments. Signed-off-by: Vikas Katariya --- test_conformance/api/test_clone_kernel.cpp | 61 ++++++++++++++++++++++ 1 file changed, 61 insertions(+) diff --git a/test_conformance/api/test_clone_kernel.cpp b/test_conformance/api/test_clone_kernel.cpp index c3ca30de95..4f8560dc40 100644 --- a/test_conformance/api/test_clone_kernel.cpp +++ b/test_conformance/api/test_clone_kernel.cpp @@ -93,6 +93,7 @@ const char* clone_kernel_test_kernel[] = { "{\n" " buffer->store[0] = iarg;\n" "}\n" + "__kernel void test_kernel_empty(){}\n" }; typedef struct @@ -678,6 +679,60 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, return TEST_PASS; } +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, + 0, NULL, NULL); + test_error(error, "clEnqueueNDRangeKernel failed"); + + error = clFinish(queue); + test_error(error, "clFinish failed"); + + return TEST_PASS; +} + +int test_cloned_kernel_empty_args(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) +{ + cl_int error; + clProgramWrapper program; + clKernelWrapper srcKernel; + + // Create srcKernel to test with + error = create_single_kernel_helper(context, &program, &srcKernel, 1, + clone_kernel_test_kernel, + "test_kernel_empty"); + test_error(error, + "Unable to create srcKernel for test_cloned_kernel_empty_args"); + + // enqueue - srcKernel + if (test_empty_enqueue_helper(queue, &srcKernel) != 0) + { + test_fail("test_empty_enqueue_helper failed for srcKernel.\n"); + } + + // clone the srcKernel + 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; +} + int test_clone_kernel(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { @@ -698,5 +753,11 @@ int test_clone_kernel(cl_device_id deviceID, cl_context context, test_fail("clCloneKernel test_cloned_kernel_exec_info failed.\n"); } + if (test_cloned_kernel_empty_args(deviceID, context, queue, num_elements) + != 0) + { + test_fail("clCloneKernel test_cloned_kernel_empty_args failed.\n"); + } + return TEST_PASS; } From 791ed1bc78be311fa8c753f48493829b78a7de09 Mon Sep 17 00:00:00 2001 From: Vikas Katariya Date: Wed, 12 May 2021 11:20:38 +0100 Subject: [PATCH 5/7] Add test coverge for `clCloneKernel` with SVM pointer Use `clSetKernelArgSVMPointer` to set args after kernel is cloned. Enqueue and read the buffer to validate. The test uses `buf_write_kernel` kernel program with 2 arguments. Signed-off-by: Vikas Katariya --- test_conformance/api/test_clone_kernel.cpp | 157 +++++++++++++++++++++ 1 file changed, 157 insertions(+) diff --git a/test_conformance/api/test_clone_kernel.cpp b/test_conformance/api/test_clone_kernel.cpp index 4f8560dc40..dc29578798 100644 --- a/test_conformance/api/test_clone_kernel.cpp +++ b/test_conformance/api/test_clone_kernel.cpp @@ -733,6 +733,158 @@ int test_cloned_kernel_empty_args(cl_device_id deviceID, cl_context context, return TEST_PASS; } +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, + 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_ptr_helper(cl_context context, cl_command_queue queue, + cl_int* svmPtr_Kernel, cl_kernel* srcKernel, + cl_int* value) +{ + cl_int error; + + error = clSetKernelArgSVMPointer(*srcKernel, 0, svmPtr_Kernel); + test_error(error, "clSetKernelArgSVMPointer failed"); + 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"); + } + + return TEST_PASS; +} + +int test_cloned_kernel_svm_ptr(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, + "buf_write_kernel"); + test_error(error, "Unable to create srcKernel"); + + 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( + svmPtr_srcKernel != NULL || svmPtr_cloneKernel_1 != NULL + || svmPtr_srcKernel_1 != NULL || svmPtr_cloneKernel_2 != NULL, + "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"); + } + 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_ptr_helper(context, queue, svmPtr_cloneKernel_1, + &cloneKernel_1, &intargs[1]) + != 0) + { + test_fail("test_svm_ptr_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_ptr_helper(context, queue, svmPtr_cloneKernel_2, + &cloneKernel_2, &intargs[2]) + != 0) + { + test_fail("test_svm_ptr_helper failed for cloneKernel_2.\n"); + } + + // 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"); + } + clSVMFree(context, svmPtr_srcKernel_1); + + // 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"); + } + clSVMFree(context, svmPtr_cloneKernel_1); + + // 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"); + } + clSVMFree(context, svmPtr_cloneKernel_2); + } + + return TEST_PASS; +} + int test_clone_kernel(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { @@ -759,5 +911,10 @@ int test_clone_kernel(cl_device_id deviceID, cl_context context, test_fail("clCloneKernel test_cloned_kernel_empty_args failed.\n"); } + if (test_cloned_kernel_svm_ptr(deviceID, context, queue, num_elements) != 0) + { + test_fail("clCloneKernel test_cloned_kernel_svm_ptr failed.\n"); + } + return TEST_PASS; } From 0c2fb2c55e97b6ade5e6163263d28a731ec6ed85 Mon Sep 17 00:00:00 2001 From: Vikas Katariya Date: Thu, 19 Aug 2021 15:08:19 +0100 Subject: [PATCH 6/7] fixup! Add test coverage for `clCloneKernel` with Execinfo Review comments --- test_conformance/api/test_clone_kernel.cpp | 20 +++++++++++--------- 1 file changed, 11 insertions(+), 9 deletions(-) diff --git a/test_conformance/api/test_clone_kernel.cpp b/test_conformance/api/test_clone_kernel.cpp index dc29578798..e3a57d0adb 100644 --- a/test_conformance/api/test_clone_kernel.cpp +++ b/test_conformance/api/test_clone_kernel.cpp @@ -542,8 +542,6 @@ int test_svm_exec_info_helper(cl_context context, cl_command_queue queue, 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; @@ -619,7 +617,6 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, { 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); @@ -651,7 +648,6 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, 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, @@ -661,7 +657,6 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, 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, @@ -671,12 +666,19 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, test_fail("test_exec_enqueue_helper failed for cloneKernel_2 on " "retry.\n"); } - clSVMFree(context, svmPtr_cloneKernel_2); + clSVMFree(context, svmPtr_srcKernel); + clSVMFree(context, svmPtr_srcKernel_1); + clSVMFree(context, svmPtr_cloneKernel_1); + clSVMFree(context, svmPtr_cloneKernel_2); clSVMFree(context, pBuf); - } - return TEST_PASS; + return TEST_PASS; + } + else + { + return TEST_SKIPPED_ITSELF; + } } int test_empty_enqueue_helper(cl_command_queue queue, cl_kernel* srcKernel) @@ -900,7 +902,7 @@ int test_clone_kernel(cl_device_id deviceID, cl_context context, } if (test_cloned_kernel_exec_info(deviceID, context, queue, num_elements) - != 0) + == TEST_FAIL) { test_fail("clCloneKernel test_cloned_kernel_exec_info failed.\n"); } From 938be73ec49f86a2bee4ea7f0c4353e0111babeb Mon Sep 17 00:00:00 2001 From: Vikas Katariya Date: Thu, 19 Aug 2021 15:11:42 +0100 Subject: [PATCH 7/7] fixup! Add test coverge for `clCloneKernel` with SVM pointer Review comments --- test_conformance/api/test_clone_kernel.cpp | 228 +++++++++------------ 1 file changed, 99 insertions(+), 129 deletions(-) diff --git a/test_conformance/api/test_clone_kernel.cpp b/test_conformance/api/test_clone_kernel.cpp index e3a57d0adb..03228c16d5 100644 --- a/test_conformance/api/test_clone_kernel.cpp +++ b/test_conformance/api/test_clone_kernel.cpp @@ -55,51 +55,52 @@ const char *clone_kernel_test_double[] = }; 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" + 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 +struct BufPtr { cl_int* store; -} BufPtr; +}; const int BUF_SIZE = 128; @@ -231,22 +232,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; @@ -289,24 +289,23 @@ 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) + if (test_args_enqueue_helper(context, queue, srcKernel, intargs[0], bufOut) != 0) { test_fail("test_args_enqueue_helper failed for srcKernel.\n"); } // enqueue - cloneKernel_1 - if (test_args_enqueue_helper(context, queue, &cloneKernel_1, &intargs[1], - &bufOut) + if (test_args_enqueue_helper(context, queue, cloneKernel_1, intargs[1], + bufOut) != 0) { test_fail("test_args_enqueue_helper failed for cloneKernel_1.\n"); } // enqueue - cloneKernel_2 - if (test_args_enqueue_helper(context, queue, &cloneKernel_2, &intargs[2], - &bufOut) + if (test_args_enqueue_helper(context, queue, cloneKernel_2, intargs[2], + bufOut) != 0) { test_fail("test_args_enqueue_helper failed for cloneKernel_2.\n"); @@ -316,16 +315,15 @@ int test_cloned_kernel_args(cl_device_id deviceID, cl_context context, 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) + if (test_args_enqueue_helper(context, queue, srcKernel, intargs[3], bufOut) != 0) { test_fail("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) + if (test_args_enqueue_helper(context, queue, cloneKernel_1, intargs[1], + bufOut) != 0) { test_fail( @@ -333,8 +331,8 @@ int test_cloned_kernel_args(cl_device_id deviceID, cl_context context, } // enqueue - cloneKernel_2 again, to check if the args were not modified - if (test_args_enqueue_helper(context, queue, &cloneKernel_2, &intargs[2], - &bufOut) + if (test_args_enqueue_helper(context, queue, cloneKernel_2, intargs[2], + bufOut) != 0) { test_fail( @@ -504,15 +502,15 @@ 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); @@ -521,7 +519,7 @@ int test_exec_enqueue_helper(cl_context context, cl_command_queue queue, 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, + test_assert_error(svmPtr_Kernel[0] == value, "clCloneKernel test failed, Failed to verify " "integer value from SVM pointer. "); @@ -535,7 +533,7 @@ int test_exec_enqueue_helper(cl_context context, cl_command_queue queue, 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; @@ -550,19 +548,15 @@ 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); - 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) + if (test_svm_enqueue_helper(context, queue, svmPtr_Kernel, srcKernel, value) != 0) { test_fail("test_exec_enqueue_helper failed.\n"); @@ -612,7 +606,7 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, // srcKernel, set args if (test_svm_exec_info_helper(context, queue, pBuf, svmPtr_srcKernel, - &srcKernel, &intargs[0]) + srcKernel, intargs[0]) != 0) { test_fail("test_svm_exec_info_helper failed for srcKernel.\n"); @@ -622,8 +616,8 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, 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]) + svmPtr_cloneKernel_1, cloneKernel_1, + intargs[1]) != 0) { test_fail("test_svm_exec_info_helper failed for cloneKernel_1.\n"); @@ -633,8 +627,8 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, 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]) + svmPtr_cloneKernel_2, cloneKernel_2, + intargs[2]) != 0) { test_fail("test_svm_exec_info_helper failed for cloneKernel_2.\n"); @@ -642,7 +636,7 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, // enqueue - srcKernel again with different svm_ptr and args if (test_svm_exec_info_helper(context, queue, pBuf, svmPtr_srcKernel_1, - &srcKernel, &intargs[3]) + srcKernel, intargs[3]) != 0) { test_fail("test_svm_exec_info_helper failed for srcKernel with " @@ -650,8 +644,8 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, } // 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]) + if (test_svm_enqueue_helper(context, queue, svmPtr_cloneKernel_1, + cloneKernel_1, intargs[1]) != 0) { test_fail("test_exec_enqueue_helper failed for cloneKernel_1 on " @@ -659,8 +653,8 @@ int test_cloned_kernel_exec_info(cl_device_id deviceID, cl_context context, } // 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]) + if (test_svm_enqueue_helper(context, queue, svmPtr_cloneKernel_2, + cloneKernel_2, intargs[2]) != 0) { test_fail("test_exec_enqueue_helper failed for cloneKernel_2 on " @@ -735,45 +729,15 @@ int test_cloned_kernel_empty_args(cl_device_id deviceID, cl_context context, return TEST_PASS; } -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, - 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_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); @@ -825,19 +789,18 @@ 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]) + if (test_svm_ptr_helper(context, queue, svmPtr_srcKernel, srcKernel, + intargs[0]) != 0) { test_fail("test_svm_ptr_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_ptr_helper(context, queue, svmPtr_cloneKernel_1, - &cloneKernel_1, &intargs[1]) + cloneKernel_1, intargs[1]) != 0) { test_fail("test_svm_ptr_helper failed for cloneKernel_1.\n"); @@ -847,44 +810,50 @@ int test_cloned_kernel_svm_ptr(cl_device_id deviceID, cl_context context, 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]) + cloneKernel_2, intargs[2]) != 0) { test_fail("test_svm_ptr_helper failed for cloneKernel_2.\n"); } // enqueue - srcKernel again with different svm_ptr and args - if (test_svm_ptr_helper(context, queue, svmPtr_srcKernel_1, &srcKernel, - &intargs[3]) + 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"); } - clSVMFree(context, svmPtr_srcKernel_1); // 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]) + cloneKernel_1, intargs[1]) != 0) { test_fail( "test_svm_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_svm_enqueue_helper(context, queue, svmPtr_cloneKernel_2, - &cloneKernel_2, &intargs[2]) + cloneKernel_2, intargs[2]) != 0) { test_fail("test_svm_enqueue_helper failed for cloneKernel_2 on " "retry.\n"); } + + clSVMFree(context, svmPtr_srcKernel); + clSVMFree(context, svmPtr_srcKernel_1); + clSVMFree(context, svmPtr_cloneKernel_1); clSVMFree(context, svmPtr_cloneKernel_2); - } - return TEST_PASS; + return TEST_PASS; + } + else + { + return TEST_SKIPPED_ITSELF; + } } int test_clone_kernel(cl_device_id deviceID, cl_context context, @@ -913,7 +882,8 @@ int test_clone_kernel(cl_device_id deviceID, cl_context context, test_fail("clCloneKernel test_cloned_kernel_empty_args failed.\n"); } - if (test_cloned_kernel_svm_ptr(deviceID, context, queue, num_elements) != 0) + if (test_cloned_kernel_svm_ptr(deviceID, context, queue, num_elements) + == TEST_FAIL) { test_fail("clCloneKernel test_cloned_kernel_svm_ptr failed.\n"); }