Skip to content

Commit

Permalink
Add test coverage for clCloneKernel with Execinfo
Browse files Browse the repository at this point in the history 
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 <vikas.katariya@arm.com>
  • Loading branch information
@jainvikas8
jainvikas8 committed Aug 5, 2021
1 parent 7b27969 commit 6bf4c0f
Showing 1 changed file with 195 additions and 1 deletion.
196 changes: 195 additions & 1 deletion test_conformance/api/test_clone_kernel.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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
Expand Down Expand Up @@ -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)
{
Expand All @@ -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;
}

0 comments on commit 6bf4c0f

Please sign in to comment.