Add testing for CL_KERNEL_LOCAL_MEM_SIZE #1235

KhronosGroup · Sep 25, 2024 · 301a77f · 301a77f
1 parent 97045f8
commit 301a77f
Show file tree

Hide file tree

Showing 3 changed files with 131 additions and 0 deletions.
diff --git a/test_conformance/api/main.cpp b/test_conformance/api/main.cpp
@@ -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),
 };
 
 const int test_num = ARRAY_SIZE(test_list);

diff --git a/test_conformance/api/procs.h b/test_conformance/api/procs.h
@@ -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,

diff --git a/test_conformance/api/test_api_consistency.cpp b/test_conformance/api/test_api_consistency.cpp
@@ -25,6 +25,34 @@ __kernel void test(__global int* dst) {
 }
 )CLC";
 
+const char* empty_kernel[] = { "__kernel void empty_kernel()\n"
+                               "{\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"
+};
+
+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)
 {
@@ -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;
+    }
+
+    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");
+
+    // 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);
+    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;
+}