add test for loading and unloading ICDs (cl_khr_icd_unloadable)

test_common/CMakeLists.txt

@@ -22,3 +22,4 @@ set(HARNESS_SOURCES
 )
 
 add_library(harness STATIC ${HARNESS_SOURCES})
+set_property(TARGET harness PROPERTY POSITION_INDEPENDENT_CODE ON)

test_conformance/extensions/CMakeLists.txt

@@ -9,8 +9,9 @@ if(ANDROID_PLATFORM GREATER 28)
     add_subdirectory( cl_khr_external_memory_ahb )
 endif ()
 add_subdirectory( cl_khr_external_memory_dma_buf )
-add_subdirectory( cl_khr_semaphore )
+add_subdirectory( cl_khr_icd_unloadable )
 add_subdirectory( cl_khr_kernel_clock )
+add_subdirectory( cl_khr_semaphore )
 add_subdirectory( cl_ext_buffer_device_address )
 if(VULKAN_IS_SUPPORTED)
     add_subdirectory( cl_khr_external_semaphore )

test_conformance/extensions/cl_khr_icd_unloadable/CMakeLists.txt

@@ -0,0 +1,29 @@
+# Test Executable:
+
+set(MODULE_NAME cl_khr_icd_unloadable)
+set(${MODULE_NAME}_OUT ${CONFORMANCE_PREFIX}${MODULE_NAME}${CONFORMANCE_SUFFIX})
+
+add_executable(${${MODULE_NAME}_OUT} main.cpp)
+set_property(TARGET ${${MODULE_NAME}_OUT} PROPERTY FOLDER "CONFORMANCE${CONFORMANCE_SUFFIX}")
+target_link_libraries(${${MODULE_NAME}_OUT} ${CMAKE_DL_LIBS})
+
+# Test Plugin (dynamically loaded and unloaded):
+
+set(PLUGIN_NAME LoadUnloadPlugin)
+set(${PLUGIN_NAME}_OUT ${PLUGIN_NAME}${CONFORMANCE_SUFFIX})
+
+add_library(${${PLUGIN_NAME}_OUT} MODULE load_unload_plugin.cpp)
+set_property(TARGET ${${PLUGIN_NAME}_OUT} PROPERTY FOLDER "CONFORMANCE${CONFORMANCE_SUFFIX}")
+target_link_libraries(${${PLUGIN_NAME}_OUT} ${HARNESS_LIB} ${CLConform_LIBRARIES})
+
+# Dependency
+add_dependencies(${${MODULE_NAME}_OUT} ${${PLUGIN_NAME}_OUT})
+
+# Install Targets:
+
+include(GNUInstallDirs)
+
+install(TARGETS ${${MODULE_NAME}_OUT}
+        RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}/$<CONFIG>)
+install(TARGETS ${${PLUGIN_NAME}_OUT}
+        LIBRARY DESTINATION ${CMAKE_INSTALL_BINDIR}/$<CONFIG>)

test_conformance/extensions/cl_khr_icd_unloadable/load_unload_plugin.cpp

@@ -0,0 +1,96 @@
+// Copyright (c) 2026 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "harness/testHarness.h"
+#include "harness/typeWrappers.h"
+
+#include <array>
+
+REGISTER_TEST(execute_kernel)
+{
+    REQUIRE_EXTENSION("cl_khr_icd_unloadable");
+
+    cl_int error = CL_SUCCESS;
+
+    const char* source = R"CLC(
+        __kernel void test_kernel(__global int* dst)
+        {
+            size_t id = get_global_id(0);
+            dst[id] = id;
+        }
+    )CLC";
+
+    clProgramWrapper program;
+    clKernelWrapper kernel;
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &source,
+                                        "test_kernel");
+    test_error(error, "Unable to create test kernel");
+
+    std::array<cl_int, 4> data = { -1, -1, -1, -1 };
+    clMemWrapper dst =
+        clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR,
+                       sizeof(data), data.data(), &error);
+    test_error(error, "clCreateBuffer failed");
+
+    error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &dst);
+    test_error(error, "clSetKernelArg failed");
+
+    size_t global_work_size = data.size();
+    error = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global_work_size,
+                                   nullptr, 0, nullptr, nullptr);
+    test_error(error, "clEnqueueNDRangeKernel failed");
+
+    error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, sizeof(data),
+                                data.data(), 0, nullptr, nullptr);
+    test_error(error, "clEnqueueReadBuffer failed");
+
+    for (size_t i = 0; i < data.size(); i++)
+    {
+        if (data[i] != static_cast<cl_int>(i))
+        {
+            test_fail("Data mismatch at index %zu: expected %d, got %d", i,
+                      static_cast<cl_int>(i), data[i]);
+        }
+    }
+
+    return TEST_PASS;
+}
+
+#if defined _WIN32 || defined __CYGWIN__
+#ifdef __GNUC__
+#define PLUGIN_API __attribute__((dllexport))
+#else
+#define PLUGIN_API __declspec(dllexport)
+#endif
+#else
+#if __GNUC__ >= 4
+#define PLUGIN_API __attribute__((visibility("default")))
+#else
+#define PLUGIN_API
+#endif
+#endif
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+PLUGIN_API int do_test(int argc, const char* argv[])
+{
+    return runTestHarness(argc, argv, test_registry::getInstance().num_tests(),
+                          test_registry::getInstance().definitions(), false, 0);
+}
+
+#ifdef __cplusplus
+}
+#endif

test_conformance/extensions/cl_khr_icd_unloadable/main.cpp

@@ -0,0 +1,77 @@
+// Copyright (c) 2026 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "harness/errorHelpers.h"
+
+#ifdef _WIN32
+#include <windows.h>
+#else
+#include <dlfcn.h>
+#endif
+
+#ifdef _WIN32
+using PluginHandle = HMODULE;
+#define LoadPlugin() ::LoadLibraryA("LoadUnloadPlugin.dll")
+#define ClosePlugin(_handle) ::FreeLibrary(_handle)
+#define GetFunctionAddress(_handle, _name) ::GetProcAddress(_handle, _name)
+#else
+using PluginHandle = void *;
+#define LoadPlugin() ::dlopen("./libLoadUnloadPlugin.so", RTLD_NOW)
+#define ClosePlugin(_handle) ::dlclose(_handle)
+#define GetFunctionAddress(_handle, _name) ::dlsym(_handle, _name)
+#endif
+
+typedef int TestFunction_t(int argc, const char *argv[]);
+typedef TestFunction_t *TestFunction_ptr;
+
+int main(int argc, const char *argv[])
+{
+    constexpr int iterations = 5;
+    int result = EXIT_SUCCESS;
+    for (int i = 0; i < iterations && result == EXIT_SUCCESS; i++)
+    {
+        log_info("Iteration %d of %d...\n", i + 1, iterations);
+
+        log_info("Loading plugin...\n");
+        PluginHandle plugin = LoadPlugin();
+        if (!plugin)
+        {
+            log_error("Failed to load plugin!\n");
+            return EXIT_FAILURE;
+        }
+
+        log_info("Getting test pointer...\n");
+        TestFunction_ptr testFunction = reinterpret_cast<TestFunction_ptr>(
+            GetFunctionAddress(plugin, "do_test"));
+        if (!testFunction)
+        {
+            log_error("Failed to get test function address!\n");
+            ClosePlugin(plugin);
+            return EXIT_FAILURE;
+        }
+
+        log_info("Running test...\n");
+        result = testFunction(argc, argv);
+        if (result != EXIT_SUCCESS)
+        {
+            log_error("Test function failed!\n");
+        }
+
+        log_info("Closing plugin...\n");
+        ClosePlugin(plugin);
+    }
+
+    log_info("All testing complete.\n");
+    return result;
+}