Added kernel corrections

KhronosGroup · Oct 8, 2024 · 2da26b0 · 2da26b0
1 parent bb4645c
commit 2da26b0
Show file tree

Hide file tree

Showing 5 changed files with 266 additions and 131 deletions.
diff --git a/test_conformance/api/CMakeLists.txt b/test_conformance/api/CMakeLists.txt
@@ -11,6 +11,7 @@ set(${MODULE_NAME}_SOURCES
          test_queries.cpp
          test_create_kernels.cpp
          test_kernels.cpp
+         test_kernel_local_memory_size.cpp
          test_kernel_private_memory_size.cpp
          test_api_min_max.cpp
          test_kernel_arg_changes.cpp

diff --git a/test_conformance/api/main.cpp b/test_conformance/api/main.cpp
@@ -163,7 +163,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),
+    ADD_TEST(kernel_local_memory_size),
 };
 
 const int test_num = ARRAY_SIZE(test_list);

diff --git a/test_conformance/api/procs.h b/test_conformance/api/procs.h
@@ -214,8 +214,10 @@ extern int test_consistency_requirements_fp16(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_kernel_local_memory_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
@@ -18,41 +18,14 @@
 #include "testBase.h"
 #include "harness/testHarness.h"
 #include "harness/deviceInfo.h"
+#include <memory>
 
 static const char* test_kernel = R"CLC(
 __kernel void test(__global int* dst) {
     dst[0] = 0;
 }
 )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)
 {
@@ -1180,106 +1153,6 @@ 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;
-}
-
 int test_consistency_requirements_fp64(cl_device_id deviceID,
                                        cl_context context,
                                        cl_command_queue queue, int num_elements)