diff --git a/test_conformance/api/main.cpp b/test_conformance/api/main.cpp index 97dcc84df..fd8b4668d 100644 --- 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 index ee39c9823..b67781d4c 100644 --- 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 index 56c831815..1fa5b37e2 100644 --- 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; +}