OpenCL ICD Loader updates for OpenCL 3.1

.github/dependabot.yml

@@ -0,0 +1,10 @@
+version: 2
+updates:
+  - package-ecosystem: "github-actions"
+    directory: "/"
+    schedule:
+      interval: "monthly"
+    groups:
+      github-actions:
+        patterns:
+          - "*"

.github/workflows/presubmit.yml

@@ -79,17 +79,17 @@ jobs:
 
     steps:
     - name: Install system CMake
-      if: ${{matrix.CMAKE}} == 'system'
+      if: matrix.CMAKE == 'system'
       run: apt-get update -qq && apt-get install -y cmake &&
         echo "CMAKE_EXE=cmake" >> "$GITHUB_ENV" &&
         echo "CPACK_EXE=cpack" >> "$GITHUB_ENV" &&
         echo "CTEST_EXE=ctest" >> "$GITHUB_ENV"
 
     - name: Checkout OpenCL-ICD-Loader
-      uses: actions/checkout@v4
+      uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
 
     - name: Checkout OpenCL-Headers
-      uses: actions/checkout@v4
+      uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
       with:
         path: external/OpenCL-Headers
         repository: KhronosGroup/OpenCL-Headers
@@ -267,7 +267,7 @@ jobs:
     - name: Cache Ninja install
       if: matrix.GEN == 'Ninja Multi-Config'
       id: ninja-install
-      uses: actions/cache@v4
+      uses: actions/cache@27d5ce7f107fe9357f9df03efb73ab90386fccae # v5.0.5
       with:
         path: |
           C:\Tools\Ninja
@@ -281,10 +281,10 @@ jobs:
         Remove-Item ~\Downloads\*
 
     - name: Checkout OpenCL-ICD-Loader
-      uses: actions/checkout@v4
+      uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
 
     - name: Checkout OpenCL-Headers
-      uses: actions/checkout@v4
+      uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
       with:
         repository: KhronosGroup/OpenCL-Headers
         path: external/OpenCL-Headers
@@ -623,10 +623,10 @@ jobs:
 
     steps:
     - name: Checkout OpenCL-ICD-Loader
-      uses: actions/checkout@v4
+      uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
 
     - name: Checkout OpenCL-Headers
-      uses: actions/checkout@v4
+      uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
       with:
         repository: KhronosGroup/OpenCL-Headers
         path: external/OpenCL-Headers
@@ -728,10 +728,10 @@ jobs:
       CFLAGS: -Wall -Wextra -pedantic -Werror
     steps:
     - name: Checkout OpenCL-ICD-Loader
-      uses: actions/checkout@v4
+      uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
 
     - name: Checkout OpenCL-Headers
-      uses: actions/checkout@v4
+      uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
       with:
         repository: KhronosGroup/OpenCL-Headers
         path: external/OpenCL-Headers

CMakeLists.txt

@@ -165,7 +165,7 @@ else ()
 endif ()
 
 set (OPENCL_COMPILE_DEFINITIONS
-    CL_TARGET_OPENCL_VERSION=300
+    CL_TARGET_OPENCL_VERSION=310
     CL_NO_NON_ICD_DISPATCH_EXTENSION_PROTOTYPES
     OPENCL_ICD_LOADER_VERSION_MAJOR=3
     OPENCL_ICD_LOADER_VERSION_MINOR=0

loader/icd_dispatch_generated.c

@@ -5191,6 +5191,57 @@ static cl_mem CL_API_CALL clCreateImageWithProperties_disp(
 
 ///////////////////////////////////////////////////////////////////////////////
 
+CL_API_ENTRY cl_int CL_API_CALL clGetKernelSuggestedLocalWorkSize(
+    cl_command_queue command_queue,
+    cl_kernel kernel,
+    cl_uint work_dim,
+    const size_t* global_work_offset,
+    const size_t* global_work_size,
+    size_t* suggested_local_work_size)
+{
+#if defined(CL_ENABLE_LAYERS)
+    if (khrFirstLayer)
+        return khrFirstLayer->dispatch.clGetKernelSuggestedLocalWorkSize(
+            command_queue,
+            kernel,
+            work_dim,
+            global_work_offset,
+            global_work_size,
+            suggested_local_work_size);
+#endif // defined(CL_ENABLE_LAYERS)
+    KHR_ICD_VALIDATE_HANDLE_RETURN_ERROR(command_queue, CL_INVALID_COMMAND_QUEUE);
+    return KHR_ICD2_DISPATCH(command_queue)->clGetKernelSuggestedLocalWorkSize(
+        command_queue,
+        kernel,
+        work_dim,
+        global_work_offset,
+        global_work_size,
+        suggested_local_work_size);
+}
+
+///////////////////////////////////////////////////////////////////////////////
+#if defined(CL_ENABLE_LAYERS)
+static cl_int CL_API_CALL clGetKernelSuggestedLocalWorkSize_disp(
+    cl_command_queue command_queue,
+    cl_kernel kernel,
+    cl_uint work_dim,
+    const size_t* global_work_offset,
+    const size_t* global_work_size,
+    size_t* suggested_local_work_size)
+{
+    KHR_ICD_VALIDATE_HANDLE_RETURN_ERROR(command_queue, CL_INVALID_COMMAND_QUEUE);
+    return KHR_ICD2_DISPATCH(command_queue)->clGetKernelSuggestedLocalWorkSize(
+        command_queue,
+        kernel,
+        work_dim,
+        global_work_offset,
+        global_work_size,
+        suggested_local_work_size);
+}
+#endif // defined(CL_ENABLE_LAYERS)
+
+///////////////////////////////////////////////////////////////////////////////
+
 // cl_ext_device_fission
 
 CL_API_ENTRY cl_int CL_API_CALL clReleaseDeviceEXT(
@@ -7026,7 +7077,10 @@ const struct _cl_icd_dispatch khrMainDispatch = {
   /* OpenCL 3.0 */
     &clCreateBufferWithProperties_disp,
     &clCreateImageWithProperties_disp,
-    &clSetContextDestructorCallback_disp
+    &clSetContextDestructorCallback_disp,
+
+  /* OpenCL 3.1 */
+    &clGetKernelSuggestedLocalWorkSize_disp,
 }
 ;
 #endif // defined(CL_ENABLE_LAYERS)
@@ -8621,6 +8675,22 @@ static cl_mem CL_API_CALL clCreateImageWithProperties_unsupp(
     (void)errcode_ret;
     KHR_ICD_ERROR_RETURN_HANDLE(CL_INVALID_OPERATION);
 }
+static cl_int CL_API_CALL clGetKernelSuggestedLocalWorkSize_unsupp(
+    cl_command_queue command_queue,
+    cl_kernel kernel,
+    cl_uint work_dim,
+    const size_t* global_work_offset,
+    const size_t* global_work_size,
+    size_t* suggested_local_work_size)
+{
+    (void)command_queue;
+    (void)kernel;
+    (void)work_dim;
+    (void)global_work_offset;
+    (void)global_work_size;
+    (void)suggested_local_work_size;
+    KHR_ICD_ERROR_RETURN_ERROR(CL_INVALID_OPERATION);
+}
 
 ///////////////////////////////////////////////////////////////////////////////
 // cl_ext_device_fission
@@ -9373,7 +9443,10 @@ const struct _cl_icd_dispatch khrDeinitDispatch = {
   /* OpenCL 3.0 */
     &clCreateBufferWithProperties_unsupp,
     &clCreateImageWithProperties_unsupp,
-    &clSetContextDestructorCallback_unsupp
+    &clSetContextDestructorCallback_unsupp,
+
+  /* OpenCL 3.1 */
+    &clGetKernelSuggestedLocalWorkSize_unsupp,
 }
 ;
 #endif // defined(CL_ENABLE_LAYERS)
@@ -9728,6 +9801,9 @@ void khrIcd2PopulateDispatchTable(
    dispatch->clCreateImageWithProperties = (clCreateImageWithProperties_t *)(size_t)p_clIcdGetFunctionAddressForPlatform(platform, "clCreateImageWithProperties");
    if (!dispatch->clCreateImageWithProperties)
        dispatch->clCreateImageWithProperties = &clCreateImageWithProperties_unsupp;
+   dispatch->clGetKernelSuggestedLocalWorkSize = (clGetKernelSuggestedLocalWorkSize_t *)(size_t)p_clIcdGetFunctionAddressForPlatform(platform, "clGetKernelSuggestedLocalWorkSize");
+   if (!dispatch->clGetKernelSuggestedLocalWorkSize)
+       dispatch->clGetKernelSuggestedLocalWorkSize = &clGetKernelSuggestedLocalWorkSize_unsupp;
 
 ///////////////////////////////////////////////////////////////////////////////
 // cl_ext_device_fission

loader/icd_version.h

@@ -50,6 +50,9 @@
 #if CL_TARGET_OPENCL_VERSION == 300
 #define OPENCL_ICD_LOADER_OCL_VERSION_NUMBER "3.0"
 #endif
+#if CL_TARGET_OPENCL_VERSION == 310
+#define OPENCL_ICD_LOADER_OCL_VERSION_NUMBER "3.1"
+#endif
 
 #define OPENCL_ICD_LOADER_OCL_VERSION_STRING \
     "OpenCL " OPENCL_ICD_LOADER_OCL_VERSION_NUMBER

loader/linux/icd_exports.map

@@ -170,3 +170,8 @@ OPENCL_3.0 {
         clCreateImageWithProperties;
         clSetContextDestructorCallback;
 } OPENCL_2.2;
+
+OPENCL_3.1 {
+    global:
+        clGetKernelSuggestedLocalWorkSize;
+} OPENCL_3.0;

loader/windows/OpenCL-mingw-i686.def

@@ -160,3 +160,6 @@ clSetProgramSpecializationConstant@16 == clSetProgramSpecializationConstant
 clCreateBufferWithProperties@28 == clCreateBufferWithProperties
 clCreateImageWithProperties@32 == clCreateImageWithProperties
 clSetContextDestructorCallback@12 == clSetContextDestructorCallback
+
+; OpenCL 3.1 API
+clGetKernelSuggestedLocalWorkSize@24 == clGetKernelSuggestedLocalWorkSize

loader/windows/OpenCL.def

@@ -164,3 +164,6 @@ clSetProgramSpecializationConstant
 clCreateBufferWithProperties
 clCreateImageWithProperties
 clSetContextDestructorCallback
+
+; OpenCL 3.1 API
+clGetKernelSuggestedLocalWorkSize

scripts/dispatch_table.mako

@@ -200,5 +200,8 @@
   /* OpenCL 3.0 */
     &clCreateBufferWithProperties_${suffix},
     &clCreateImageWithProperties_${suffix},
-    &clSetContextDestructorCallback_${suffix}
+    &clSetContextDestructorCallback_${suffix},
+
+  /* OpenCL 3.1 */
+    &clGetKernelSuggestedLocalWorkSize_${suffix},
 }

test/driver_stub/CMakeLists.txt

@@ -9,7 +9,7 @@ add_library (OpenCLDriverStub SHARED ${OPENCL_DRIVER_STUB_SOURCES})
 
 target_link_libraries (OpenCLDriverStub IcdLog OpenCL::Headers)
 
-target_compile_definitions (OpenCLDriverStub PRIVATE CL_TARGET_OPENCL_VERSION=300)
+target_compile_definitions (OpenCLDriverStub PRIVATE CL_TARGET_OPENCL_VERSION=310)
 
 set (OPENCL_DRIVER_STUB_ICD2_SOURCES cl.c cl_ext.c cl_gl.c icd.c)
 
@@ -21,4 +21,4 @@ add_library (OpenCLDriverStubICD2 SHARED ${OPENCL_DRIVER_STUB_ICD2_SOURCES})
 
 target_link_libraries (OpenCLDriverStubICD2 IcdLog OpenCL::Headers)
 
-target_compile_definitions (OpenCLDriverStubICD2 PRIVATE CL_TARGET_OPENCL_VERSION=300 CL_ENABLE_ICD2=1)
+target_compile_definitions (OpenCLDriverStubICD2 PRIVATE CL_TARGET_OPENCL_VERSION=310 CL_ENABLE_ICD2=1)

test/driver_stub/cl.c

@@ -1090,6 +1090,27 @@ clGetKernelInfo(cl_kernel        kernel ,
     return return_value;
 }
 
+CL_API_ENTRY cl_int CL_API_CALL
+clGetKernelSuggestedLocalWorkSize(cl_command_queue command_queue,
+                                  cl_kernel        kernel,
+                                  cl_uint          work_dim,
+                                  const size_t *   global_work_offset,
+                                  const size_t *   global_work_size,
+                                  size_t *         suggested_local_work_size) CL_API_SUFFIX__VERSION_3_1
+{
+    cl_int return_value = CL_OUT_OF_RESOURCES;
+    test_icd_stub_log("clGetKernelSuggestedLocalWorkSize(%p, %p, %u, %p, %p, %p)\n",
+                      command_queue,
+                      kernel,
+                      work_dim,
+                      global_work_offset,
+                      global_work_size,
+                      suggested_local_work_size);
+
+    test_icd_stub_log("Value returned: %d\n", return_value);
+    return return_value;
+}
+
 CL_API_ENTRY cl_int CL_API_CALL
 clGetKernelArgInfo(cl_kernel        kernel ,
                    cl_uint          arg_indx ,

test/driver_stub/icd.c

@@ -248,6 +248,9 @@ cl_int cliIcdDispatchTableCreate(CLIicdDispatchTable **outDispatchTable)
     ICD_DISPATCH_TABLE_ENTRY ( clCreateImageWithProperties );
     ICD_DISPATCH_TABLE_ENTRY ( clSetContextDestructorCallback );
 
+    /* OpenCL 3.1 */
+    ICD_DISPATCH_TABLE_ENTRY ( clGetKernelSuggestedLocalWorkSize );
+
     // return success
     *outDispatchTable = dispatchTable;
     return CL_SUCCESS;
@@ -471,6 +474,9 @@ void * CL_API_CALL clIcdGetFunctionAddressForPlatformKHR(
     ICD_GET_FUNCTION_ADDRESS ( clCreateImageWithProperties );
     ICD_GET_FUNCTION_ADDRESS ( clSetContextDestructorCallback );
 
+    /* OpenCL 3.1 */
+    ICD_GET_FUNCTION_ADDRESS ( clGetKernelSuggestedLocalWorkSize );
+
     return NULL;
 }
 

test/driver_stub/rename_api.h

@@ -105,5 +105,6 @@
 #define clCreateBufferWithProperties             ___clCreateBufferWithProperties
 #define clCreateImageWithProperties              ___clCreateImageWithProperties
 #define clSetContextDestructorCallback           ___clSetContextDestructorCallback
+#define clGetKernelSuggestedLocalWorkSize        ___clGetKernelSuggestedLocalWorkSize
 
 #endif /* __RENAME_API_H__ */

test/layer/CMakeLists.txt

@@ -15,7 +15,7 @@ add_library (PrintLayer SHARED ${OPENCL_PRINT_LAYER_SOURCES})
 
 target_include_directories(PrintLayer PRIVATE ${PARENT_DIR}/include)
 target_link_libraries(PrintLayer PUBLIC OpenCL::Headers)
-target_compile_definitions (PrintLayer PRIVATE CL_TARGET_OPENCL_VERSION=300)
+target_compile_definitions (PrintLayer PRIVATE CL_TARGET_OPENCL_VERSION=310)
 
 if (NOT WIN32 AND NOT APPLE)
   set_target_properties (PrintLayer PROPERTIES LINK_FLAGS "-Wl,--version-script -Wl,${CMAKE_CURRENT_SOURCE_DIR}/icd_print_layer.map")

test/layer/icd_print_layer_generated.c

@@ -1949,6 +1949,25 @@ return tdispatch->clCreateImageWithProperties(
             errcode_ret);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+static cl_int CL_API_CALL clGetKernelSuggestedLocalWorkSize_wrap(
+    cl_command_queue command_queue,
+    cl_kernel kernel,
+    cl_uint work_dim,
+    const size_t* global_work_offset,
+    const size_t* global_work_size,
+    size_t* suggested_local_work_size) CL_API_SUFFIX__VERSION_3_1
+{
+printf("clGetKernelSuggestedLocalWorkSize\n");
+return tdispatch->clGetKernelSuggestedLocalWorkSize(
+            command_queue,
+            kernel,
+            work_dim,
+            global_work_offset,
+            global_work_size,
+            suggested_local_work_size);
+}
+
 ///////////////////////////////////////////////////////////////////////////////
 
 // cl_ext_device_fission
@@ -2742,6 +2761,9 @@ struct _cl_icd_dispatch dispatch = {
   /* OpenCL 3.0 */
     &clCreateBufferWithProperties_wrap,
     &clCreateImageWithProperties_wrap,
-    &clSetContextDestructorCallback_wrap
+    &clSetContextDestructorCallback_wrap,
+
+  /* OpenCL 3.1 */
+    &clGetKernelSuggestedLocalWorkSize_wrap,
 }
 ;

test/loader_test/CMakeLists.txt

@@ -14,7 +14,7 @@ add_executable (icd_loader_test
 
 target_compile_definitions(icd_loader_test
     PRIVATE
-        CL_TARGET_OPENCL_VERSION=300
+        CL_TARGET_OPENCL_VERSION=310
 )
 
 target_link_libraries (icd_loader_test OpenCL IcdLog OpenCL::Headers)

test/loader_test/param_struct.h

@@ -745,6 +745,7 @@ struct clReleaseKernel_st
 #define NUM_ITEMS_clGetKernelInfo 1
 #define NUM_ITEMS_clGetKernelArgInfo 1
 #define NUM_ITEMS_clGetKernelWorkGroupInfo 1
+#define NUM_ITEMS_clGetKernelSuggestedLocalWorkSize 1
 
 struct clSetKernelArg_st 
 {
@@ -783,6 +784,16 @@ struct clGetKernelWorkGroupInfo_st
     size_t *param_value_size_ret;
 };
 
+struct clGetKernelSuggestedLocalWorkSize_st
+{
+    cl_command_queue command_queue;
+    cl_kernel kernel;
+    cl_uint work_dim;
+    const size_t *global_work_offset;
+    const size_t *global_work_size;
+    size_t *suggested_local_work_size;
+};
+
 #define NUM_ITEMS_clEnqueueMigrateMemObjects 1
 #define NUM_ITEMS_clEnqueueNDRangeKernel 1
 #define NUM_ITEMS_clEnqueueTask 1