From 1fbecb13fa3f842acb867e12fffb9e71da74c937 Mon Sep 17 00:00:00 2001 From: Brice Videau Date: Wed, 22 Oct 2025 14:05:51 -0500 Subject: [PATCH 1/3] Deinitialization support. Co-authored-by: Ben Ashbaugh --- CMakeLists.txt | 2 +- README.md | 2 + include/cl_khr_icd2.h | 25 +++ loader/cllayerinfo.c | 17 ++- loader/icd.c | 161 +++++++++++++++---- loader/icd.h | 21 ++- loader/icd_dispatch_generated.c | 212 +++++++++++++++++++++++++- loader/linux/icd_linux.c | 9 +- loader/windows/icd_windows.c | 13 +- scripts/icd_dispatch_generated.c.mako | 8 +- test/driver_stub/cl.c | 77 ++++++++-- test/layer/icd_print_layer.c | 58 ++++++- test/layer/icd_print_layer.def | 2 + test/layer/icd_print_layer.map | 2 + 14 files changed, 544 insertions(+), 65 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 1ce37cff..63c2d857 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -161,7 +161,7 @@ set (OPENCL_COMPILE_DEFINITIONS CL_NO_NON_ICD_DISPATCH_EXTENSION_PROTOTYPES OPENCL_ICD_LOADER_VERSION_MAJOR=3 OPENCL_ICD_LOADER_VERSION_MINOR=0 - OPENCL_ICD_LOADER_VERSION_REV=7 + OPENCL_ICD_LOADER_VERSION_REV=8 $<$:CL_ENABLE_LAYERS> $<$:CL_ENABLE_LOADER_MANAGED_DISPATCH> ) diff --git a/README.md b/README.md index 0fee18cf..c2a56531 100644 --- a/README.md +++ b/README.md @@ -145,3 +145,5 @@ The following debug environment variables are available for use with the OpenCL | OPENCL_LAYERS | Specifies a list of layers to load. | `export OPENCL_LAYERS=libLayerA.so:libLayerB.so`

`set OPENCL_LAYERS=libLayerA.dll;libLayerB.dll` | | OPENCL_LAYER_PATH | On Linux and Android, specifies a directory to scan for layers to enumerate in place of the default `/etc/OpenCL/layers'. | `export OPENCL_LAYER_PATH=/my/local/layers/search/path` | | OCL_ICD_ENABLE_TRACE | Enable the trace mechanism | `export OCL_ICD_ENABLE_TRACE=True`

`set OCL_ICD_ENABLE_TRACE=True`
`true, T, 1 can also be used here.` | +| OCL_ICD_FORCE_LEGACY_TERMINATION | Forces using the legacy termination scheme. Legacy termination supports older layers, but does not support layer de-initialization and does not close library handles for layers or ICDs. | `export OCL_ICD_FORCE_LEGACY_TERMINATION=True`

`set OCL_ICD_FORCE_LEGACY_TERMINATION=True`
`true, T, 1 can also be used here.` | +| OCL_ICD_DISABLE_DYNAMIC_LIBRARY_UNLOADING | Prevents the ICD loader from closing dynamic library handles for layers and ICDs. This can be used for debugging purposes or to allow leak sanitizers to have full stack traces. | `export OCL_ICD_DISABLE_DYNAMIC_LIBRARY_UNLOADING=True`

`set OCL_ICD_DISABLE_DYNAMIC_LIBRARY_UNLOADING=True`
`true, T, 1 can also be used here.` | diff --git a/include/cl_khr_icd2.h b/include/cl_khr_icd2.h index 3672401a..5bf05cec 100644 --- a/include/cl_khr_icd2.h +++ b/include/cl_khr_icd2.h @@ -18,6 +18,31 @@ #include +#if !defined(CL_PLATFORM_UNLOADABLE_KHR) +#define CL_PLATFORM_UNLOADABLE_KHR 0x0921 +#endif + +#if defined(CL_ENABLE_LAYERS) && !defined(CL_LAYER_PROPERTIES_LIST_END) +typedef cl_properties cl_layer_properties; + +#define CL_LAYER_PROPERTIES_LIST_END ((cl_layer_properties)0) + +typedef cl_int CL_API_CALL +clInitLayerWithProperties_t( + cl_uint num_entries, + const cl_icd_dispatch *target_dispatch, + cl_uint *num_entries_ret, + const cl_icd_dispatch **layer_dispatch_ret, + const cl_layer_properties *properties); + +typedef clInitLayerWithProperties_t *pfn_clInitLayerWithProperties; + +typedef cl_int CL_API_CALL +clDeinitLayer_t(void); + +typedef clDeinitLayer_t *pfn_clDeinitLayer; +#endif //defined(CL_ENABLE_LAYERS) && !defined(CL_LAYER_PROPERTIES_LIST_END) + #if !defined(CL_ICD2_TAG_KHR) #if INTPTR_MAX == INT32_MAX #define CL_ICD2_TAG_KHR ((intptr_t)0x434C3331) diff --git a/loader/cllayerinfo.c b/loader/cllayerinfo.c index 5a85ba82..cbdbb43c 100644 --- a/loader/cllayerinfo.c +++ b/loader/cllayerinfo.c @@ -19,7 +19,6 @@ #include "icd.h" #include #include -#include #if defined(_WIN32) #include #include @@ -90,7 +89,7 @@ static void restore_outputs(void) void printLayerInfo(const struct KHRLayer *layer) { cl_layer_api_version api_version = 0; - pfn_clGetLayerInfo p_clGetLayerInfo = (pfn_clGetLayerInfo)(size_t)layer->p_clGetLayerInfo; + pfn_clGetLayerInfo p_clGetLayerInfo = layer->p_clGetLayerInfo; cl_int result = CL_SUCCESS; size_t sz; @@ -113,20 +112,26 @@ void printLayerInfo(const struct KHRLayer *layer) } } -int main (int argc, char *argv[]) +static void run_silently(void (*pfn)(void)) { - (void)argc; - (void)argv; silence_layers(); atexit(restore_outputs); - khrIcdInitialize(); + pfn(); restore_outputs(); atexit(silence_layers); +} + +int main (int argc, char *argv[]) +{ + (void)argc; + (void)argv; + run_silently(&khrIcdInitialize); const struct KHRLayer *layer = khrFirstLayer; while (layer) { printLayerInfo(layer); layer = layer->next; } + run_silently(&khrIcdDeinitialize); return 0; } diff --git a/loader/icd.c b/loader/icd.c index 6cb2b993..f90eba26 100644 --- a/loader/icd.c +++ b/loader/icd.c @@ -19,32 +19,48 @@ #include "icd.h" #include "icd_dispatch.h" #include "icd_envvars.h" -#if defined(CL_ENABLE_LAYERS) -#include -#endif // defined(CL_ENABLE_LAYERS) #include #include KHRicdVendor *khrIcdVendors = NULL; +static KHRicdVendor *lastVendor = NULL; int khrEnableTrace = 0; +static int khrDisableLibraryUnloading = 0; +static int khrForceLegacyTermination = 0; #if defined(CL_ENABLE_LAYERS) struct KHRLayer *khrFirstLayer = NULL; #endif // defined(CL_ENABLE_LAYERS) -// entrypoint to check and initialize trace. -void khrIcdInitializeTrace(void) +static inline int khrIcdCheckEnvTrue(const char *variable) +{ + return (variable && (strcmp(variable, "True") == 0 || + strcmp(variable, "true") == 0 || + strcmp(variable, "T") == 0 || + strcmp(variable, "1") == 0)); +} + +// Set a given flag if the given environement variable is true +static void khrInitializeFlagWithEnv(int *flag, const char *variable) { - char *enableTrace = khrIcd_getenv("OCL_ICD_ENABLE_TRACE"); - if (enableTrace && (strcmp(enableTrace, "True") == 0 || - strcmp(enableTrace, "true") == 0 || - strcmp(enableTrace, "T") == 0 || - strcmp(enableTrace, "1") == 0)) + char *variableStr = khrIcd_getenv(variable); + if (khrIcdCheckEnvTrue(variableStr)) + { + *flag = 1; + } + if (variableStr) { - khrEnableTrace = 1; + khrIcd_free_getenv(variableStr); } } +void khrIcdInitializeEnvOptions(void) +{ + khrInitializeFlagWithEnv(&khrEnableTrace, "OCL_ICD_ENABLE_TRACE"); + khrInitializeFlagWithEnv(&khrDisableLibraryUnloading, "OCL_ICD_DISABLE_DYNAMIC_LIBRARY_UNLOADING"); + khrInitializeFlagWithEnv(&khrForceLegacyTermination, "OCL_ICD_FORCE_LEGACY_TERMINATION"); +} + // entrypoint to initialize the ICD and add all vendors void khrIcdInitialize(void) { @@ -186,6 +202,14 @@ void khrIcdVendorAdd(const char *libraryName) #endif // call clGetPlatformInfo on the returned platform to get the suffix + + KHR_ICD2_DISPATCH(platforms[i])->clGetPlatformInfo( + platforms[i], + CL_PLATFORM_UNLOADABLE_KHR, + sizeof(vendor->unloadable), + &vendor->unloadable, + NULL); + result = KHR_ICD2_DISPATCH(platforms[i])->clGetPlatformInfo( platforms[i], CL_PLATFORM_ICD_SUFFIX_KHR, @@ -230,11 +254,13 @@ void khrIcdVendorAdd(const char *libraryName) vendor->suffix = suffix; // add this vendor to the list of vendors at the tail - { - KHRicdVendor **prevNextPointer = NULL; - for (prevNextPointer = &khrIcdVendors; *prevNextPointer; prevNextPointer = &( (*prevNextPointer)->next) ); - *prevNextPointer = vendor; + if (lastVendor) { + lastVendor->next = vendor; + vendor->prev = lastVendor; + } else { + khrIcdVendors = vendor; } + lastVendor = vendor; KHR_ICD_TRACE("successfully added vendor %s with suffix %s\n", libraryName, suffix); @@ -259,6 +285,8 @@ void khrIcdLayerAdd(const char *libraryName) cl_int result = CL_SUCCESS; pfn_clGetLayerInfo p_clGetLayerInfo = NULL; pfn_clInitLayer p_clInitLayer = NULL; + pfn_clInitLayerWithProperties p_clInitLayerWithProperties = NULL; + pfn_clDeinitLayer p_clDeinitLayer = NULL; struct KHRLayer *layerIterator = NULL; struct KHRLayer *layer = NULL; cl_layer_api_version api_version = 0; @@ -300,14 +328,6 @@ void khrIcdLayerAdd(const char *libraryName) goto Done; } - // use that function to get the clInitLayer function pointer - p_clInitLayer = (pfn_clInitLayer)(size_t)khrIcdOsLibraryGetFunctionAddress(library, "clInitLayer"); - if (!p_clInitLayer) - { - KHR_ICD_TRACE("failed to get function address clInitLayer\n"); - goto Done; - } - result = p_clGetLayerInfo(CL_LAYER_API_VERSION, sizeof(api_version), &api_version, NULL); if (CL_SUCCESS != result) { @@ -321,6 +341,31 @@ void khrIcdLayerAdd(const char *libraryName) goto Done; } + // Support old version of layers, which should rely on at_exit for termination. + // In this case use clInitLayer to initialize layers + if (khrForceLegacyTermination) + { + p_clInitLayer = (pfn_clInitLayer)(size_t)khrIcdOsLibraryGetFunctionAddress(library, "clInitLayer"); + if (!p_clInitLayer) + { + KHR_ICD_TRACE("failed to get function address clInitLayer\n"); + goto Done; + } + } else { // New scheme, relies on clInitLayerWithProperties and the optional clDeinitLayer for termination + p_clInitLayerWithProperties = (pfn_clInitLayerWithProperties)(size_t)khrIcdOsLibraryGetFunctionAddress(library, "clInitLayerWithProperties"); + if (!p_clInitLayerWithProperties) + { + KHR_ICD_TRACE("failed to get function address clInitLayerWithProperties\n"); + goto Done; + } + + p_clDeinitLayer = (pfn_clDeinitLayer)(size_t)khrIcdOsLibraryGetFunctionAddress(library, "clDeinitLayer"); + if (!p_clDeinitLayer) + { + KHR_ICD_TRACE("layer does not support clDeinitLayer\n"); + } + } + layer = (struct KHRLayer*)calloc(sizeof(struct KHRLayer), 1); if (!layer) { @@ -338,9 +383,10 @@ void khrIcdLayerAdd(const char *libraryName) goto Done; } memcpy(layer->libraryName, libraryName, sz_name); - layer->p_clGetLayerInfo = (void *)(size_t)p_clGetLayerInfo; + layer->p_clGetLayerInfo = p_clGetLayerInfo; } #endif + layer->p_clDeinitLayer = p_clDeinitLayer; if (khrFirstLayer) { targetDispatch = &(khrFirstLayer->dispatch); @@ -349,11 +395,21 @@ void khrIcdLayerAdd(const char *libraryName) } loaderDispatchNumEntries = sizeof(khrMainDispatch)/sizeof(void*); - result = p_clInitLayer( - loaderDispatchNumEntries, - targetDispatch, - &layerDispatchNumEntries, - &layerDispatch); + if (khrForceLegacyTermination) + { + result = p_clInitLayer( + loaderDispatchNumEntries, + targetDispatch, + &layerDispatchNumEntries, + &layerDispatch); + } else { + result = p_clInitLayerWithProperties( + loaderDispatchNumEntries, + targetDispatch, + &layerDispatchNumEntries, + &layerDispatch, + NULL); + } if (CL_SUCCESS != result) { KHR_ICD_TRACE("failed to initialize layer\n"); @@ -472,3 +528,50 @@ void khrIcdContextPropertiesGetPlatform(const cl_context_properties *properties, } } +#if defined(CL_ENABLE_LAYERS) +static struct KHRLayer deinitLayer = {0}; +#endif + +void khrIcdDeinitialize(void) { + if (khrForceLegacyTermination) + { + KHR_ICD_TRACE("ICD Loader deinitialization disabled\n"); + return; + } + + KHR_ICD_TRACE("ICD Loader deinitialization\n"); + +#if defined(CL_ENABLE_LAYERS) + // free layers first in reverse order of their creation (front to back) + // they may still need to use vendors while terminating + KHR_ICD_TRACE("Finalizing and unloading layers\n"); + struct KHRLayer *head = khrFirstLayer; + deinitLayer.dispatch = khrDeinitDispatch; + khrFirstLayer = &deinitLayer; + + while(head) { + struct KHRLayer *cur = head; +#ifdef CL_LAYER_INFO + free(cur->libraryName); +#endif + if (cur->p_clDeinitLayer) + cur->p_clDeinitLayer(); + if (!khrDisableLibraryUnloading) + khrIcdOsLibraryUnload(cur->library); + head = cur->next; + free(cur); + } +#endif // defined(CL_ENABLE_LAYERS) + + // free vendor in reverse order of their creation (back to front) + KHR_ICD_TRACE("Finalizing and unloading vendors\n"); + while (lastVendor) { + KHRicdVendor *cur = lastVendor; + free(cur->suffix); + if (cur->unloadable && !khrDisableLibraryUnloading) + khrIcdOsLibraryUnload(cur->library); + lastVendor = cur->prev; + free(cur); + } + khrIcdVendors = NULL; +} diff --git a/loader/icd.h b/loader/icd.h index 326cf82c..59b370f9 100644 --- a/loader/icd.h +++ b/loader/icd.h @@ -49,6 +49,9 @@ #include #include #include +#if defined(CL_ENABLE_LAYERS) +#include +#endif // defined(CL_ENABLE_LAYERS) #include /* @@ -85,6 +88,9 @@ struct KHRicdVendorRec // the extension suffix for this platform char *suffix; + // can this vendor library be unloaded? + cl_bool unloadable; + // function pointer to the ICD platform IDs extracted from the library pfn_clGetExtensionFunctionAddress clGetExtensionFunctionAddress; @@ -98,6 +104,7 @@ struct KHRicdVendorRec // next vendor in the list vendors KHRicdVendor *next; + KHRicdVendor *prev; }; // the global state @@ -123,14 +130,17 @@ struct KHRLayer #ifdef CL_LAYER_INFO // The layer library name char *libraryName; - // the pointer to the clGetLayerInfo funciton - void *p_clGetLayerInfo; + // the pointer to the clGetLayerInfo function + pfn_clGetLayerInfo p_clGetLayerInfo; #endif + // the pointer to the clDeinitLayer function + pfn_clDeinitLayer p_clDeinitLayer; }; // the global layer state extern struct KHRLayer * khrFirstLayer; extern const struct _cl_icd_dispatch khrMainDispatch; +extern const struct _cl_icd_dispatch khrDeinitDispatch; #endif // defined(CL_ENABLE_LAYERS) /* @@ -144,8 +154,11 @@ extern const struct _cl_icd_dispatch khrMainDispatch; // API (e.g, getPlatformIDs, etc). void khrIcdInitialize(void); -// entrypoint to check and initialize trace. -void khrIcdInitializeTrace(void); +// entrypoint to check and initialize env options. +void khrIcdInitializeEnvOptions(void); + +// entrypoint to release icd resources +void khrIcdDeinitialize(void); // go through the list of vendors (in /etc/OpenCL.conf or through // the registry) and call khrIcdVendorAdd for each vendor encountered diff --git a/loader/icd_dispatch_generated.c b/loader/icd_dispatch_generated.c index 02e2af0a..5a914985 100644 --- a/loader/icd_dispatch_generated.c +++ b/loader/icd_dispatch_generated.c @@ -7031,7 +7031,7 @@ const struct _cl_icd_dispatch khrMainDispatch = { ; #endif // defined(CL_ENABLE_LAYERS) -#if defined(CL_ENABLE_LOADER_MANAGED_DISPATCH) +#if defined(CL_ENABLE_LOADER_MANAGED_DISPATCH) || defined(CL_ENABLE_LAYERS) /////////////////////////////////////////////////////////////////////////////// // Core APIs: static cl_int CL_API_CALL clGetPlatformIDs_unsupp( @@ -9168,7 +9168,217 @@ static cl_int CL_API_CALL clGetKernelSubGroupInfoKHR_unsupp( } /////////////////////////////////////////////////////////////////////////////// +#endif // defined(CL_ENABLE_LOADER_MANAGED_DISPATCH) || defined(CL_ENABLE_LAYERS) + +#if defined(CL_ENABLE_LAYERS) +const struct _cl_icd_dispatch khrDeinitDispatch = { + ICD_ANON_UNION_INIT_MEMBER(&clGetPlatformIDs_unsupp), + &clGetPlatformInfo_unsupp, + &clGetDeviceIDs_unsupp, + &clGetDeviceInfo_unsupp, + &clCreateContext_unsupp, + &clCreateContextFromType_unsupp, + &clRetainContext_unsupp, + &clReleaseContext_unsupp, + &clGetContextInfo_unsupp, + &clCreateCommandQueue_unsupp, + &clRetainCommandQueue_unsupp, + &clReleaseCommandQueue_unsupp, + &clGetCommandQueueInfo_unsupp, + &clSetCommandQueueProperty_unsupp, + &clCreateBuffer_unsupp, + &clCreateImage2D_unsupp, + &clCreateImage3D_unsupp, + &clRetainMemObject_unsupp, + &clReleaseMemObject_unsupp, + &clGetSupportedImageFormats_unsupp, + &clGetMemObjectInfo_unsupp, + &clGetImageInfo_unsupp, + &clCreateSampler_unsupp, + &clRetainSampler_unsupp, + &clReleaseSampler_unsupp, + &clGetSamplerInfo_unsupp, + &clCreateProgramWithSource_unsupp, + &clCreateProgramWithBinary_unsupp, + &clRetainProgram_unsupp, + &clReleaseProgram_unsupp, + &clBuildProgram_unsupp, + ICD_ANON_UNION_INIT_MEMBER(&clUnloadCompiler_unsupp), + &clGetProgramInfo_unsupp, + &clGetProgramBuildInfo_unsupp, + &clCreateKernel_unsupp, + &clCreateKernelsInProgram_unsupp, + &clRetainKernel_unsupp, + &clReleaseKernel_unsupp, + &clSetKernelArg_unsupp, + &clGetKernelInfo_unsupp, + &clGetKernelWorkGroupInfo_unsupp, + &clWaitForEvents_unsupp, + &clGetEventInfo_unsupp, + &clRetainEvent_unsupp, + &clReleaseEvent_unsupp, + &clGetEventProfilingInfo_unsupp, + &clFlush_unsupp, + &clFinish_unsupp, + &clEnqueueReadBuffer_unsupp, + &clEnqueueWriteBuffer_unsupp, + &clEnqueueCopyBuffer_unsupp, + &clEnqueueReadImage_unsupp, + &clEnqueueWriteImage_unsupp, + &clEnqueueCopyImage_unsupp, + &clEnqueueCopyImageToBuffer_unsupp, + &clEnqueueCopyBufferToImage_unsupp, + &clEnqueueMapBuffer_unsupp, + &clEnqueueMapImage_unsupp, + &clEnqueueUnmapMemObject_unsupp, + &clEnqueueNDRangeKernel_unsupp, + &clEnqueueTask_unsupp, + &clEnqueueNativeKernel_unsupp, + &clEnqueueMarker_unsupp, + &clEnqueueWaitForEvents_unsupp, + &clEnqueueBarrier_unsupp, + &clGetExtensionFunctionAddress_unsupp, + &clCreateFromGLBuffer_unsupp, + &clCreateFromGLTexture2D_unsupp, + &clCreateFromGLTexture3D_unsupp, + &clCreateFromGLRenderbuffer_unsupp, + &clGetGLObjectInfo_unsupp, + &clGetGLTextureInfo_unsupp, + &clEnqueueAcquireGLObjects_unsupp, + &clEnqueueReleaseGLObjects_unsupp, + &clGetGLContextInfoKHR_unsupp, + + /* cl_khr_d3d10_sharing */ +#if defined(_WIN32) + &clGetDeviceIDsFromD3D10KHR_unsupp, + &clCreateFromD3D10BufferKHR_unsupp, + &clCreateFromD3D10Texture2DKHR_unsupp, + &clCreateFromD3D10Texture3DKHR_unsupp, + &clEnqueueAcquireD3D10ObjectsKHR_unsupp, + &clEnqueueReleaseD3D10ObjectsKHR_unsupp, +#else + NULL, + NULL, + NULL, + NULL, + NULL, + NULL, +#endif + + /* OpenCL 1.1 */ + &clSetEventCallback_unsupp, + &clCreateSubBuffer_unsupp, + &clSetMemObjectDestructorCallback_unsupp, + &clCreateUserEvent_unsupp, + &clSetUserEventStatus_unsupp, + &clEnqueueReadBufferRect_unsupp, + &clEnqueueWriteBufferRect_unsupp, + &clEnqueueCopyBufferRect_unsupp, + + /* cl_ext_device_fission */ + &clCreateSubDevicesEXT_unsupp, + &clRetainDeviceEXT_unsupp, + &clReleaseDeviceEXT_unsupp, + + /* cl_khr_gl_event */ + &clCreateEventFromGLsyncKHR_unsupp, + + /* OpenCL 1.2 */ + &clCreateSubDevices_unsupp, + &clRetainDevice_unsupp, + &clReleaseDevice_unsupp, + &clCreateImage_unsupp, + &clCreateProgramWithBuiltInKernels_unsupp, + &clCompileProgram_unsupp, + &clLinkProgram_unsupp, + &clUnloadPlatformCompiler_unsupp, + &clGetKernelArgInfo_unsupp, + &clEnqueueFillBuffer_unsupp, + &clEnqueueFillImage_unsupp, + &clEnqueueMigrateMemObjects_unsupp, + &clEnqueueMarkerWithWaitList_unsupp, + &clEnqueueBarrierWithWaitList_unsupp, + &clGetExtensionFunctionAddressForPlatform_unsupp, + &clCreateFromGLTexture_unsupp, + + /* cl_khr_d3d11_sharing */ +#if defined(_WIN32) + &clGetDeviceIDsFromD3D11KHR_unsupp, + &clCreateFromD3D11BufferKHR_unsupp, + &clCreateFromD3D11Texture2DKHR_unsupp, + &clCreateFromD3D11Texture3DKHR_unsupp, + &clCreateFromDX9MediaSurfaceKHR_unsupp, + &clEnqueueAcquireD3D11ObjectsKHR_unsupp, + &clEnqueueReleaseD3D11ObjectsKHR_unsupp, +#else + NULL, + NULL, + NULL, + NULL, + NULL, + NULL, + NULL, +#endif + + /* cl_khr_dx9_media_sharing */ +#if defined(_WIN32) + &clGetDeviceIDsFromDX9MediaAdapterKHR_unsupp, + &clEnqueueAcquireDX9MediaSurfacesKHR_unsupp, + &clEnqueueReleaseDX9MediaSurfacesKHR_unsupp, +#else + NULL, + NULL, + NULL, +#endif + + /* cl_khr_egl_image */ + &clCreateFromEGLImageKHR_unsupp, + &clEnqueueAcquireEGLObjectsKHR_unsupp, + &clEnqueueReleaseEGLObjectsKHR_unsupp, + + /* cl_khr_egl_event */ + &clCreateEventFromEGLSyncKHR_unsupp, + /* OpenCL 2.0 */ + &clCreateCommandQueueWithProperties_unsupp, + &clCreatePipe_unsupp, + &clGetPipeInfo_unsupp, + &clSVMAlloc_unsupp, + &clSVMFree_unsupp, + &clEnqueueSVMFree_unsupp, + &clEnqueueSVMMemcpy_unsupp, + &clEnqueueSVMMemFill_unsupp, + &clEnqueueSVMMap_unsupp, + &clEnqueueSVMUnmap_unsupp, + &clCreateSamplerWithProperties_unsupp, + &clSetKernelArgSVMPointer_unsupp, + &clSetKernelExecInfo_unsupp, + + /* cl_khr_sub_groups */ + &clGetKernelSubGroupInfoKHR_unsupp, + + /* OpenCL 2.1 */ + &clCloneKernel_unsupp, + &clCreateProgramWithIL_unsupp, + &clEnqueueSVMMigrateMem_unsupp, + &clGetDeviceAndHostTimer_unsupp, + &clGetHostTimer_unsupp, + &clGetKernelSubGroupInfo_unsupp, + &clSetDefaultDeviceCommandQueue_unsupp, + + /* OpenCL 2.2 */ + &clSetProgramReleaseCallback_unsupp, + &clSetProgramSpecializationConstant_unsupp, + + /* OpenCL 3.0 */ + &clCreateBufferWithProperties_unsupp, + &clCreateImageWithProperties_unsupp, + &clSetContextDestructorCallback_unsupp +} +; +#endif // defined(CL_ENABLE_LAYERS) + +#if defined(CL_ENABLE_LOADER_MANAGED_DISPATCH) void khrIcd2PopulateDispatchTable( cl_platform_id platform, clIcdGetFunctionAddressForPlatformKHR_fn p_clIcdGetFunctionAddressForPlatform, diff --git a/loader/linux/icd_linux.c b/loader/linux/icd_linux.c index c6f6b542..d3f0649e 100644 --- a/loader/linux/icd_linux.c +++ b/loader/linux/icd_linux.c @@ -218,7 +218,7 @@ static inline void khrIcdOsDirEnumerate(const char *path, const char *env, // go through the list of vendors in the two configuration files void khrIcdOsVendorsEnumerate(void) { - khrIcdInitializeTrace(); + khrIcdInitializeEnvOptions(); khrIcdVendorsEnumerateEnv(); khrIcdOsDirEnumerate(ICD_VENDOR_PATH, "OCL_ICD_VENDORS", ".icd", khrIcdVendorAdd, 0); @@ -265,3 +265,10 @@ void khrIcdOsLibraryUnload(void *library) { dlclose(library); } + +#ifndef CL_LAYER_INFO +static +void __attribute__((destructor)) khrIcdDestructor(void) { + khrIcdDeinitialize(); +} +#endif diff --git a/loader/windows/icd_windows.c b/loader/windows/icd_windows.c index a92a999d..dbe336a9 100644 --- a/loader/windows/icd_windows.c +++ b/loader/windows/icd_windows.c @@ -195,7 +195,7 @@ BOOL CALLBACK khrIcdOsVendorsEnumerate(PINIT_ONCE InitOnce, PVOID Parameter, PVO HKEY platformsKey = NULL; DWORD dwIndex; - khrIcdInitializeTrace(); + khrIcdInitializeEnvOptions(); khrIcdVendorsEnumerateEnv(); currentStatus = khrIcdOsVendorsEnumerateDXGK(); @@ -447,3 +447,14 @@ void khrIcdOsLibraryUnload(void *library) { FreeLibrary( (HMODULE)library); } + +#ifndef CL_LAYER_INFO +BOOL APIENTRY DllMain(HINSTANCE hinst, DWORD reason, LPVOID reserved) { + (void)hinst; + (void)reserved; + if (reason == DLL_PROCESS_DETACH) { + khrIcdDeinitialize(); + } + return TRUE; +} +#endif diff --git a/scripts/icd_dispatch_generated.c.mako b/scripts/icd_dispatch_generated.c.mako index 276ab8ac..9dba6de7 100644 --- a/scripts/icd_dispatch_generated.c.mako +++ b/scripts/icd_dispatch_generated.c.mako @@ -285,7 +285,7 @@ ${("CL_API_ENTRY", "static")[disp]} ${api.RetType} CL_API_CALL ${api.Name + ("", const struct _cl_icd_dispatch khrMainDispatch = ${table_template.render(suffix = 'disp')}; #endif // defined(CL_ENABLE_LAYERS) -#if defined(CL_ENABLE_LOADER_MANAGED_DISPATCH) +#if defined(CL_ENABLE_LOADER_MANAGED_DISPATCH) || defined(CL_ENABLE_LAYERS) /////////////////////////////////////////////////////////////////////////////// // Core APIs: %for apis in coreapis.values(): @@ -355,7 +355,13 @@ static ${api.RetType} CL_API_CALL ${api.Name}_unsupp( /////////////////////////////////////////////////////////////////////////////// %endfor +#endif // defined(CL_ENABLE_LOADER_MANAGED_DISPATCH) || defined(CL_ENABLE_LAYERS) +#if defined(CL_ENABLE_LAYERS) +const struct _cl_icd_dispatch khrDeinitDispatch = ${table_template.render(suffix = 'unsupp')}; +#endif // defined(CL_ENABLE_LAYERS) + +#if defined(CL_ENABLE_LOADER_MANAGED_DISPATCH) void khrIcd2PopulateDispatchTable( cl_platform_id platform, clIcdGetFunctionAddressForPlatformKHR_fn p_clIcdGetFunctionAddressForPlatform, diff --git a/test/driver_stub/cl.c b/test/driver_stub/cl.c index 7bb7205a..c3286924 100644 --- a/test/driver_stub/cl.c +++ b/test/driver_stub/cl.c @@ -17,6 +17,7 @@ #include #include #include "icd_structs.h" +#include "cl_khr_icd2.h" #define CL_PLATFORM_ICD_SUFFIX_KHR 0x0920 CL_API_ENTRY cl_int CL_API_CALL @@ -31,6 +32,7 @@ struct _cl_platform_id const char *vendor; const char *extensions; const char *suffix; + cl_device_id device; }; struct _cl_device_id @@ -74,7 +76,7 @@ struct _cl_sampler }; static CLIicdDispatchTable* dispatchTable = NULL; -static cl_platform_id platform = NULL; +static cl_platform_id stub_platform = NULL; static cl_bool initialized = CL_FALSE; CL_API_ENTRY cl_int CL_API_CALL @@ -124,6 +126,18 @@ clGetPlatformInfo(cl_platform_id platform_id, cl_platform_info param_name, case CL_PLATFORM_ICD_SUFFIX_KHR: returnString = platform_id->suffix; break; + case CL_PLATFORM_UNLOADABLE_KHR: + if (param_value_size && param_value_size < sizeof(cl_bool)) { + ret = CL_INVALID_VALUE; + goto done; + } + if (param_value) { + *(cl_bool *)param_value = CL_TRUE; + } + if (param_value_size_ret) { + *param_value_size_ret = sizeof(cl_bool); + } + goto done; default: ret = CL_INVALID_VALUE; goto done; @@ -164,9 +178,11 @@ CL_API_ENTRY cl_int CL_API_CALL clGetDeviceIDs( } if (devices != NULL) { - cl_device_id obj = (cl_device_id) malloc(sizeof(struct _cl_device_id)); - CL_INIT_OBJECT(obj, platform); - devices[0] = obj; + if (!platform_id->device) { + platform_id->device = (cl_device_id) malloc(sizeof(struct _cl_device_id)); + CL_INIT_OBJECT(platform_id->device, stub_platform); + } + devices[0] = platform_id->device; } if (num_devices) { *num_devices = 1; @@ -280,7 +296,7 @@ clCreateContextFromType(const cl_context_properties * properties, cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0 { cl_context obj = (cl_context) malloc(sizeof(struct _cl_context)); - cl_platform_id plt = platform; + cl_platform_id plt = stub_platform; for (const cl_context_properties * property = properties; *property; property += 2) if (*property == (cl_context_properties)CL_CONTEXT_PLATFORM) plt = (cl_platform_id)property[1]; @@ -1930,6 +1946,7 @@ clEnqueueBarrier(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0 } extern cl_int cliIcdDispatchTableCreate(CLIicdDispatchTable **outDispatchTable); +extern void cliIcdDispatchTableDestroy(CLIicdDispatchTable *dispatchTable); CL_API_ENTRY cl_int CL_API_CALL clIcdGetPlatformIDsKHR(cl_uint num_entries, @@ -1939,20 +1956,20 @@ clIcdGetPlatformIDsKHR(cl_uint num_entries, cl_int result = CL_SUCCESS; if (!initialized) { result = cliIcdDispatchTableCreate(&dispatchTable); - platform = (cl_platform_id) malloc(sizeof(struct _cl_platform_id)); - memset(platform, 0, sizeof(struct _cl_platform_id)); + stub_platform = (cl_platform_id) malloc(sizeof(struct _cl_platform_id)); + memset(stub_platform, 0, sizeof(struct _cl_platform_id)); - CL_INIT_PLATFORM(platform, dispatchTable); - platform->version = "OpenCL 1.2 Stub"; - platform->vendor = "stubvendorxxx"; - platform->profile = "stubprofilexxx"; + CL_INIT_PLATFORM(stub_platform, dispatchTable); + stub_platform->version = "OpenCL 1.2 Stub"; + stub_platform->vendor = "stubvendorxxx"; + stub_platform->profile = "stubprofilexxx"; #if defined(CL_ENABLE_ICD2) - platform->name = "ICD_LOADER_TEST_OPENCL_STUB_ICD2"; + stub_platform->name = "ICD_LOADER_TEST_OPENCL_STUB_ICD2"; #else - platform->name = "ICD_LOADER_TEST_OPENCL_STUB"; + stub_platform->name = "ICD_LOADER_TEST_OPENCL_STUB"; #endif - platform->extensions = "cl_khr_icd cl_khr_gl cl_khr_d3d10"; - platform->suffix = "ilts"; + stub_platform->extensions = "cl_khr_icd cl_khr_gl cl_khr_d3d10"; + stub_platform->suffix = "ilts"; initialized = CL_TRUE; } @@ -1964,7 +1981,7 @@ clIcdGetPlatformIDsKHR(cl_uint num_entries, } if (platforms && num_entries == 1) { - platforms[0] = platform; + platforms[0] = stub_platform; } Done: @@ -1975,3 +1992,31 @@ clIcdGetPlatformIDsKHR(cl_uint num_entries, return result; } +static void deinit(void) { + if (initialized) { + free(stub_platform->device); + stub_platform->device = NULL; + free(stub_platform); + stub_platform = NULL; + cliIcdDispatchTableDestroy(dispatchTable); + dispatchTable = NULL; + initialized = CL_FALSE; + } +} + +#if defined(_WIN32) +#include +BOOL APIENTRY DllMain(HINSTANCE hinst, DWORD reason, LPVOID reserved) { + (void)hinst; + (void)reserved; + if (reason == DLL_PROCESS_DETACH) { + deinit(); + } + return TRUE; +} +#else +static +void __attribute__((destructor)) khrIcdDestructor(void) { + deinit(); +} +#endif diff --git a/test/layer/icd_print_layer.c b/test/layer/icd_print_layer.c index d3231ad2..f50ccb23 100644 --- a/test/layer/icd_print_layer.c +++ b/test/layer/icd_print_layer.c @@ -21,10 +21,16 @@ #include #include +#if !defined(CL_LAYER_PROPERTIES_LIST_END) +typedef cl_properties cl_layer_properties; + +#define CL_LAYER_PROPERTIES_LIST_END ((cl_layer_properties)0) +#endif //!defined(CL_LAYER_PROPERTIES_LIST_END) const struct _cl_icd_dispatch *tdispatch; static cl_layer_api_version api_version = CL_LAYER_API_VERSION_100; + static const char name[] = "print_layer"; static inline cl_int @@ -70,20 +76,62 @@ clGetLayerInfo( return set_param_value(param_value_size, param_value, param_value_size_ret, sz, src); } -CL_API_ENTRY cl_int CL_API_CALL -clInitLayer( +static void deinitLayerAtExit(void) { + printf("Cleaning up %s, API version: %d\n", name, api_version); + printf("atexit() deinitialization\n"); +} + +static void deinitLayerFunction(void) { + printf("Cleaning up %s, API version: %d\n", name, api_version); + printf("Function deinitialization\n"); +} + +static cl_int +checkAndInitDispatch( cl_uint num_entries, const struct _cl_icd_dispatch *target_dispatch, - cl_uint *num_entries_out, + cl_uint *num_entries_ret, const struct _cl_icd_dispatch **layer_dispatch_ret) { - if (!target_dispatch || !layer_dispatch_ret || !num_entries_out || num_entries < sizeof(dispatch)/sizeof(dispatch.clGetPlatformIDs)) + if (!target_dispatch || !layer_dispatch_ret || !num_entries_ret || num_entries < sizeof(dispatch)/sizeof(dispatch.clGetPlatformIDs)) return CL_INVALID_VALUE; tdispatch = target_dispatch; *layer_dispatch_ret = &dispatch; - *num_entries_out = sizeof(dispatch)/sizeof(dispatch.clGetPlatformIDs); + *num_entries_ret = sizeof(dispatch)/sizeof(dispatch.clGetPlatformIDs); return CL_SUCCESS; } +CL_API_ENTRY cl_int CL_API_CALL +clInitLayer( + cl_uint num_entries, + const struct _cl_icd_dispatch *target_dispatch, + cl_uint *num_entries_ret, + const struct _cl_icd_dispatch **layer_dispatch_ret) +{ + cl_int result = CL_SUCCESS; + result = checkAndInitDispatch(num_entries, target_dispatch, num_entries_ret, layer_dispatch_ret); + if (CL_SUCCESS != result) + return result; + + atexit(deinitLayerAtExit); + return CL_SUCCESS; +} + +CL_API_ENTRY cl_int CL_API_CALL +clInitLayerWithProperties( + cl_uint num_entries, + const cl_icd_dispatch *target_dispatch, + cl_uint *num_entries_ret, + const cl_icd_dispatch **layer_dispatch_ret, + const cl_layer_properties *properties) +{ + (void)properties; + return checkAndInitDispatch(num_entries, target_dispatch, num_entries_ret, layer_dispatch_ret); +} +CL_API_ENTRY cl_int CL_API_CALL +clDeinitLayer(void) { + deinitLayerFunction(); + return CL_SUCCESS; +} diff --git a/test/layer/icd_print_layer.def b/test/layer/icd_print_layer.def index c33a80b7..b9f7ac01 100644 --- a/test/layer/icd_print_layer.def +++ b/test/layer/icd_print_layer.def @@ -1,3 +1,5 @@ EXPORTS clGetLayerInfo clInitLayer +clInitLayerWithProperties +clDeinitLayer diff --git a/test/layer/icd_print_layer.map b/test/layer/icd_print_layer.map index b32d582a..60c57abd 100644 --- a/test/layer/icd_print_layer.map +++ b/test/layer/icd_print_layer.map @@ -2,6 +2,8 @@ global: clGetLayerInfo; clInitLayer; +clInitLayerWithProperties; +clDeinitLayer; local: *; From 5c9be3b6237411241495675ecca2aaf877f1f955 Mon Sep 17 00:00:00 2001 From: Brice Videau Date: Wed, 26 Nov 2025 10:11:41 -0600 Subject: [PATCH 2/3] Apply suggestions from code review --- loader/icd.c | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/loader/icd.c b/loader/icd.c index f90eba26..28b47d55 100644 --- a/loader/icd.c +++ b/loader/icd.c @@ -544,7 +544,7 @@ void khrIcdDeinitialize(void) { #if defined(CL_ENABLE_LAYERS) // free layers first in reverse order of their creation (front to back) // they may still need to use vendors while terminating - KHR_ICD_TRACE("Finalizing and unloading layers\n"); + KHR_ICD_TRACE("finalizing and unloading layers\n"); struct KHRLayer *head = khrFirstLayer; deinitLayer.dispatch = khrDeinitDispatch; khrFirstLayer = &deinitLayer; @@ -555,7 +555,13 @@ void khrIcdDeinitialize(void) { free(cur->libraryName); #endif if (cur->p_clDeinitLayer) - cur->p_clDeinitLayer(); + { + cl_int res = cur->p_clDeinitLayer(); + if (CL_SUCCESS != res) + { + KHR_ICD_TRACE("error reported in layer deinitialization\n"); + } + } if (!khrDisableLibraryUnloading) khrIcdOsLibraryUnload(cur->library); head = cur->next; @@ -564,7 +570,7 @@ void khrIcdDeinitialize(void) { #endif // defined(CL_ENABLE_LAYERS) // free vendor in reverse order of their creation (back to front) - KHR_ICD_TRACE("Finalizing and unloading vendors\n"); + KHR_ICD_TRACE("finalizing and unloading vendors\n"); while (lastVendor) { KHRicdVendor *cur = lastVendor; free(cur->suffix); From 41a93d583d7969eb6a8f9f41f5f604b17c6ad985 Mon Sep 17 00:00:00 2001 From: Brice Videau Date: Thu, 11 Dec 2025 11:54:54 -0600 Subject: [PATCH 3/3] Implement cl_khr_icd_unloadable proposal. --- loader/icd.c | 71 +++++++++++++++++++++++++++++++++++++------ test/driver_stub/cl.c | 2 +- 2 files changed, 63 insertions(+), 10 deletions(-) diff --git a/loader/icd.c b/loader/icd.c index 28b47d55..a2a64dbc 100644 --- a/loader/icd.c +++ b/loader/icd.c @@ -84,7 +84,7 @@ void khrIcdVendorAdd(const char *libraryName) KHRicdVendor *vendorIterator = NULL; // require that the library name be valid - if (!libraryName) + if (!libraryName) { goto Done; } @@ -155,6 +155,8 @@ void khrIcdVendorAdd(const char *libraryName) for (i = 0; i < platformCount; ++i) { KHRicdVendor* vendor = NULL; + char *extensions; + size_t extensionsSize; char *suffix; size_t suffixSize; @@ -201,15 +203,63 @@ void khrIcdVendorAdd(const char *libraryName) } #endif - // call clGetPlatformInfo on the returned platform to get the suffix - - KHR_ICD2_DISPATCH(platforms[i])->clGetPlatformInfo( + // call clGetPlatformInfo on the returned platform to get the supported extensions + result = KHR_ICD2_DISPATCH(platforms[i])->clGetPlatformInfo( platforms[i], - CL_PLATFORM_UNLOADABLE_KHR, - sizeof(vendor->unloadable), - &vendor->unloadable, + CL_PLATFORM_EXTENSIONS, + 0, + NULL, + &extensionsSize); + if (CL_SUCCESS != result) + { + KHR_ICD_TRACE("failed query platform extensions\n"); + free(vendor); + continue; + } + extensions = (char *)malloc(extensionsSize); + if (!extensions) + { + KHR_ICD_TRACE("failed to allocate memory\n"); + free(vendor); + continue; + } + result = KHR_ICD2_DISPATCH(platforms[i])->clGetPlatformInfo( + platforms[i], + CL_PLATFORM_EXTENSIONS, + extensionsSize, + extensions, NULL); + if (CL_SUCCESS != result) + { + KHR_ICD_TRACE("failed query platform extensions\n"); + free(extensions); + free(vendor); + continue; + } + + if (strstr(extensions, "cl_khr_icd_unloadable")) + { + KHR_ICD_TRACE("found cl_khr_icd_unloadable extension support\n"); + free(extensions); + result = KHR_ICD2_DISPATCH(platforms[i])->clGetPlatformInfo( + platforms[i], + CL_PLATFORM_UNLOADABLE_KHR, + sizeof(vendor->unloadable), + &vendor->unloadable, + NULL); + if (vendor->unloadable) + { + KHR_ICD_TRACE("platform is unloadable\n"); + } + if (CL_SUCCESS != result) + { + KHR_ICD_TRACE("found cl_khr_icd_unloadable but clGetPlatformInfo CL_PLATFORM_UNLOADABLE_KHR query failed\n"); + free(vendor); + continue; + } + } + // call clGetPlatformInfo on the returned platform to get the suffix result = KHR_ICD2_DISPATCH(platforms[i])->clGetPlatformInfo( platforms[i], CL_PLATFORM_ICD_SUFFIX_KHR, @@ -218,12 +268,14 @@ void khrIcdVendorAdd(const char *libraryName) &suffixSize); if (CL_SUCCESS != result) { + KHR_ICD_TRACE("failed query platform ICD suffix\n"); free(vendor); continue; } suffix = (char *)malloc(suffixSize); if (!suffix) { + KHR_ICD_TRACE("failed to allocate memory\n"); free(vendor); continue; } @@ -232,9 +284,10 @@ void khrIcdVendorAdd(const char *libraryName) CL_PLATFORM_ICD_SUFFIX_KHR, suffixSize, suffix, - NULL); + NULL); if (CL_SUCCESS != result) { + KHR_ICD_TRACE("failed query platform ICD suffix\n"); free(suffix); free(vendor); continue; @@ -242,7 +295,7 @@ void khrIcdVendorAdd(const char *libraryName) // populate vendor data vendor->library = khrIcdOsLibraryLoad(libraryName); - if (!vendor->library) + if (!vendor->library) { free(suffix); free(vendor); diff --git a/test/driver_stub/cl.c b/test/driver_stub/cl.c index c3286924..3e5d134d 100644 --- a/test/driver_stub/cl.c +++ b/test/driver_stub/cl.c @@ -1968,7 +1968,7 @@ clIcdGetPlatformIDsKHR(cl_uint num_entries, #else stub_platform->name = "ICD_LOADER_TEST_OPENCL_STUB"; #endif - stub_platform->extensions = "cl_khr_icd cl_khr_gl cl_khr_d3d10"; + stub_platform->extensions = "cl_khr_icd cl_khr_icd_unloadable cl_khr_gl cl_khr_d3d10"; stub_platform->suffix = "ilts"; initialized = CL_TRUE; }