Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Add testing for CL_KERNEL_LOCAL_MEM_SIZE #1235 #2089

Open
wants to merge 8 commits into
base: main
Choose a base branch
from
Open

Add testing for CL_KERNEL_LOCAL_MEM_SIZE #1235 #2089

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
301a77f
Add testing for CL_KERNEL_LOCAL_MEM_SIZE #1235
kamil-goras-mobica Sep 25, 2024
bb4645c
Merge branch 'main' into kernel_local_mem_size
kamil-goras-mobica Oct 3, 2024
7bd8dc0
Added kernel corrections
kamil-goras-mobica Oct 8, 2024
1cace37
Added verification of kernel results
kamil-goras-mobica Oct 9, 2024
95b817a
Added checking for param_value_size_ret
kamil-goras-mobica Oct 15, 2024
0a6751f
Added major and minor fixes
kamil-goras-mobica Oct 21, 2024
7d9575b
Merge branch 'main' into kernel_local_mem_size
kamil-goras-mobica Oct 21, 2024
132cd70
Removed commented code
kamil-goras-mobica Oct 21, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions test_conformance/api/main.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -161,6 +161,7 @@ test_definition test_list[] = {
ADD_TEST_VERSION(negative_create_command_queue_with_properties,
Version(2, 0)),
ADD_TEST(negative_create_command_queue_with_properties_khr),
ADD_TEST(kernel_local_mem_size),
bashbaug marked this conversation as resolved.
Show resolved Hide resolved
};

const int test_num = ARRAY_SIZE(test_list);
Expand Down
2 changes: 2 additions & 0 deletions test_conformance/api/procs.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -206,6 +206,8 @@ extern int test_consistency_3d_image_writes(cl_device_id deviceID,

extern int test_min_image_formats(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_kernel_local_mem_size(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_negative_get_platform_info(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
Expand Down
128 changes: 128 additions & 0 deletions test_conformance/api/test_api_consistency.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,34 @@ __kernel void test(__global int* dst) {
}
)CLC";

const char* empty_kernel[] = { "__kernel void empty_kernel()\n"
bashbaug marked this conversation as resolved.
Show resolved Hide resolved
"{\n"
"}\n" };

const char* local_memory_kernel[] = {
"__kernel void local_memory_kernel(__local int* ptr)\n"
"{\n"
"__local float array[10000];\n"
"for(int i = 0; i<10000; i++)\n"
" array[i]*=2;\n"
"}\n"
bashbaug marked this conversation as resolved.
Show resolved Hide resolved
};

const char* local_param_kernel[] = {
"__kernel void local_param_kernel(__local int* ptr)\n"
"{\n"
"}\n"
};

const char* local_param_local_memory_kernel[] = {
"__kernel void local_param_local_memory_kernel(__local int* ptr)\n"
"{\n"
"__local float array[10000];\n"
"for(int i = 0; i<10000; i++)\n"
" array[i]*=2;\n"
"}\n"
};

int test_consistency_svm(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
{
Expand Down Expand Up @@ -1151,3 +1179,103 @@ int test_consistency_3d_image_writes(cl_device_id deviceID, cl_context context,

return TEST_PASS;
}

int test_kernel_local_mem_size(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
{
int error;
clProgramWrapper program;
clKernelWrapper kernel;

// Check memory needed to execute empty kernel
if (create_single_kernel_helper(context, &program, &kernel, 1, empty_kernel,
"empty_kernel")
!= 0)
{
return -1;
}
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We really ought to return TEST_FAIL here rather than the magic number -1. Or, even better, something like:

Suggested change
if (create_single_kernel_helper(context, &program, &kernel, 1, empty_kernel,
"empty_kernel")
!= 0)
{
return -1;
}
error = create_single_kernel_helper(context, &program, &kernel, 1, empty_kernel,
"empty_kernel");
test_error(error, "Unable to create empty_kernel");


cl_ulong kernelLocalUsage = 0;
error = clGetKernelWorkGroupInfo(kernel, deviceID, CL_KERNEL_LOCAL_MEM_SIZE,
sizeof(kernelLocalUsage),
&kernelLocalUsage, NULL);
test_error(error,
"clGetKernelWorkGroupInfo for CL_KERNEL_LOCAL_MEM_SIZE failed");

test_assert_error(kernelLocalUsage > 0, "kernel local mem size failed");
bashbaug marked this conversation as resolved.
Show resolved Hide resolved

// Check memory needed to execute empty kernel with __local variable
if (create_single_kernel_helper(context, &program, &kernel, 1,
local_memory_kernel, "local_memory_kernel")
!= 0)
{
return -1;
}

kernelLocalUsage = 0;
error = clGetKernelWorkGroupInfo(kernel, deviceID, CL_KERNEL_LOCAL_MEM_SIZE,
sizeof(kernelLocalUsage),
&kernelLocalUsage, NULL);
test_error(error,
"clGetKernelWorkGroupInfo for CL_KERNEL_LOCAL_MEM_SIZE failed");

test_assert_error(kernelLocalUsage >= 10000 * sizeof(cl_float),
"kernel local mem size failed");

// Check memory needed to execute empty kernel with __local parameter with
// setKernelArg
if (create_single_kernel_helper(context, &program, &kernel, 1,
local_param_kernel, "local_param_kernel")
!= 0)
{
return -1;
}

size_t elements = 100;
size_t sizeToAllocate = elements * sizeof(cl_int);
int* localData = (cl_int*)malloc(sizeToAllocate);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can we use a smart pointer here rather than malloc (e.g. make_unique) so that if the function returns early before the free(localData) we don't leak.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

localData is unused. We may as well remove it entirely.

for (size_t i = 0; i < elements; i++)
{
localData[i] = i;
}
error = clSetKernelArg(kernel, 0, sizeToAllocate, NULL);
test_error(error, "Unable to set indexed kernel arguments");

kernelLocalUsage = 0;
error = clGetKernelWorkGroupInfo(kernel, deviceID, CL_KERNEL_LOCAL_MEM_SIZE,
sizeof(kernelLocalUsage),
&kernelLocalUsage, NULL);
test_error(error,
"clGetKernelWorkGroupInfo for CL_KERNEL_LOCAL_MEM_SIZE failed");

test_assert_error(kernelLocalUsage >= sizeToAllocate,
"kernel local mem size failed");

// Check memory needed to execute kernel with __local variable and __local
// parameter with setKernelArg
if (create_single_kernel_helper(context, &program, &kernel, 1,
local_param_local_memory_kernel,
"local_param_local_memory_kernel")
!= 0)
{
return -1;
}

error = clSetKernelArg(kernel, 0, sizeToAllocate, NULL);
test_error(error, "Unable to set indexed kernel arguments");

kernelLocalUsage = 0;
error = clGetKernelWorkGroupInfo(kernel, deviceID, CL_KERNEL_LOCAL_MEM_SIZE,
sizeof(kernelLocalUsage),
&kernelLocalUsage, NULL);
test_error(error,
"clGetKernelWorkGroupInfo for CL_KERNEL_LOCAL_MEM_SIZE failed");

test_assert_error(kernelLocalUsage
>= sizeToAllocate + 10000 * sizeof(cl_float),
"kernel local mem size failed");

free(localData);

return CL_SUCCESS;
}
Loading