diff --git a/test_common/CMakeLists.txt b/test_common/CMakeLists.txt index 3acc742c74..ea9c4d4648 100644 --- a/test_common/CMakeLists.txt +++ b/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) diff --git a/test_conformance/extensions/CMakeLists.txt b/test_conformance/extensions/CMakeLists.txt index a2af536e73..1057911c4e 100644 --- a/test_conformance/extensions/CMakeLists.txt +++ b/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 ) diff --git a/test_conformance/extensions/cl_khr_icd_unloadable/CMakeLists.txt b/test_conformance/extensions/cl_khr_icd_unloadable/CMakeLists.txt new file mode 100644 index 0000000000..6d9b0794d9 --- /dev/null +++ b/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}/$) +install(TARGETS ${${PLUGIN_NAME}_OUT} + LIBRARY DESTINATION ${CMAKE_INSTALL_BINDIR}/$) diff --git a/test_conformance/extensions/cl_khr_icd_unloadable/load_unload_plugin.cpp b/test_conformance/extensions/cl_khr_icd_unloadable/load_unload_plugin.cpp new file mode 100644 index 0000000000..ec576361f5 --- /dev/null +++ b/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 + +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 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(i)) + { + test_fail("Data mismatch at index %zu: expected %d, got %d", i, + static_cast(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 diff --git a/test_conformance/extensions/cl_khr_icd_unloadable/main.cpp b/test_conformance/extensions/cl_khr_icd_unloadable/main.cpp new file mode 100644 index 0000000000..7aeaedeb36 --- /dev/null +++ b/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 +#else +#include +#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( + 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; +}