diff --git a/docs/install.md b/docs/install.md index e3143043..817a8214 100644 --- a/docs/install.md +++ b/docs/install.md @@ -16,6 +16,67 @@ OpenCL Applications or to revert back to normal operation. See the [cliloader](cliloader.md) documentation for more detail. +## As an ICD Loader Layer + +The Intercept Layer for OpenCL Applications can be loaded by the OpenCL ICD +loader as an ICD loader layer. +When used as an ICD loader layer, the OpenCL ICD loader is responsible for +loading and initializing the Intercept Layer for OpenCL applications, then +calling it instead of calling through a dispatch table to the target ICD. + +To setup the Intercept Layer for OpenCL Applications as an ICD loader layer, set +the `OPENCL_LAYERS` environment variable so it includes the Intercept Layer for +OpenCL Applications library. +This will usually be a full path to the library, for example: + +```sh +$ export OPENCL_LAYERS="/path/to/libOpenCL.so" +``` + +or: + +```sh +> set OPENCL_LAYERS="c:\path\to\OpenCL.dll" +``` + +The Intercept Layer for OpenCL Applications may also be configured as a system +ICD loader layer. +Please consult your OpenCL ICD loader documentation for information about +installing and using system layers. + +When the Intercept Layer for OpenCL Applications is used as an ICD loader layer, +it will first go through typical initialization when it is initially loaded: + +``` +-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-=-= +CLIntercept (64-bit) is loading... + +... loading complete. +``` + +Then, it will go through a separate layer initialization, initiated by the +OpenCL ICD loader: + +``` +CLIntercept layer is initializing... +... layer initialization complete. +``` + +After the Intercept Layer for OpenCL Applications has been initialized as an ICD +loader layer, it should operate the same as if it were installed and enabled +through other mechanisms. + +For more information about ICD loader layers, please refer to the +[cl_loader_layers](https://github.com/KhronosGroup/OpenCL-Docs/blob/main/extensions/cl_loader_layers.asciidoc) +documentation, or follow the [OpenCL Layers Tutorial](https://github.com/Kerilk/OpenCL-Layers-Tutorial). + +To troubleshoot issues using the Intercept Layer for OpenCL Applications as an +ICD loader layer, consider setting the `OCL_ICD_ENABLE_TRACE` environment +variable (for the Khronos +[OpenCL-ICD-Loader](https://github.com/KhronosGroup/OpenCL-ICD-Loader)) or the +`OCL_ICD_DEBUG` environment variable (for the +[ocl-icd](https://github.com/OCL-dev/ocl-icd) loader). + ## Windows ### Local Install diff --git a/intercept/CL/cl_function_types.h b/intercept/CL/cl_function_types.h index a7963373..991dbdc7 100644 --- a/intercept/CL/cl_function_types.h +++ b/intercept/CL/cl_function_types.h @@ -1191,7 +1191,7 @@ typedef cl_int CL_API_CALL clGetKernelSuggestedLocalWorkSize_t( 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; + size_t* suggested_local_work_size); typedef clGetKernelSuggestedLocalWorkSize_t * clGetKernelSuggestedLocalWorkSize_fn CL_API_SUFFIX__VERSION_3_1; diff --git a/intercept/CL/cl_layer.h b/intercept/CL/cl_layer.h new file mode 100644 index 00000000..2b272eb3 --- /dev/null +++ b/intercept/CL/cl_layer.h @@ -0,0 +1,165 @@ +/******************************************************************************* + * Copyright (c) 2008-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. + ******************************************************************************/ + +#ifndef OPENCL_CL_LAYER_H_ +#define OPENCL_CL_LAYER_H_ + +/* +** This header is generated from the Khronos OpenCL XML API Registry. +*/ + +// Unlike the Khronos header file, we want to unconditionally include the +// CLIntercept cl_icd.h and cl.h, system headers. +#include "CL/cl_icd.h" + +#include "CL/cl.h" + +/* CL_NO_PROTOTYPES implies CL_NO_EXTENSION_PROTOTYPES: */ +#if defined(CL_NO_PROTOTYPES) && !defined(CL_NO_EXTENSION_PROTOTYPES) +#define CL_NO_EXTENSION_PROTOTYPES +#endif + +/* CL_NO_EXTENSION_PROTOTYPES implies + CL_NO_ICD_DISPATCH_EXTENSION_PROTOTYPES and + CL_NO_NON_ICD_DISPATCH_EXTENSION_PROTOTYPES: */ +#if defined(CL_NO_EXTENSION_PROTOTYPES) && \ + !defined(CL_NO_ICD_DISPATCH_EXTENSION_PROTOTYPES) +#define CL_NO_ICD_DISPATCH_EXTENSION_PROTOTYPES +#endif +#if defined(CL_NO_EXTENSION_PROTOTYPES) && \ + !defined(CL_NO_NON_ICD_DISPATCH_EXTENSION_PROTOTYPES) +#define CL_NO_NON_ICD_DISPATCH_EXTENSION_PROTOTYPES +#endif + +#ifdef __cplusplus +extern "C" { +#endif + +/*************************************************************** +* cl_loader_layers +***************************************************************/ +#define cl_loader_layers 1 +#define CL_LOADER_LAYERS_EXTENSION_NAME \ + "cl_loader_layers" + + +#define CL_LOADER_LAYERS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 1) + +typedef cl_uint cl_layer_info; +typedef cl_uint cl_layer_api_version; +typedef cl_properties cl_layer_properties; + +/* cl_layer_info */ +#define CL_LAYER_API_VERSION 0x4240 +#define CL_LAYER_NAME 0x4241 + +/* Misc API enums */ +#define CL_LAYER_API_VERSION_100 100 +#define CL_LAYER_PROPERTIES_LIST_END ((cl_layer_properties)0) + + +typedef cl_int CL_API_CALL +clGetLayerInfo_t( + cl_layer_info param_name, + size_t param_value_size, + void* param_value, + size_t* param_value_size_ret); + +typedef clGetLayerInfo_t * +clGetLayerInfo_fn ; + +typedef cl_int CL_API_CALL +clInitLayer_t( + cl_uint num_entries, + const cl_icd_dispatch* target_dispatch, + cl_uint* num_entries_ret, + const cl_icd_dispatch** layer_dispatch_ret); + +typedef clInitLayer_t * +clInitLayer_fn ; + +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 * +clInitLayerWithProperties_fn ; + +typedef cl_int CL_API_CALL +clDeinitLayer_t( + void ); + +typedef clDeinitLayer_t * +clDeinitLayer_fn ; + +/* +** The function pointer typedefs prefixed with "pfn_" are provided for +** compatibility with earlier versions of the headers. New code is +** encouraged to use the function pointer typedefs that are suffixed with +** "_fn" instead, for consistency. +*/ + +typedef clGetLayerInfo_t * +pfn_clGetLayerInfo ; + +typedef clInitLayer_t * +pfn_clInitLayer ; + +typedef clInitLayerWithProperties_t * +pfn_clInitLayerWithProperties ; + +typedef clDeinitLayer_t * +pfn_clDeinitLayer ; + +#if !defined(CL_NO_NON_ICD_DISPATCH_EXTENSION_PROTOTYPES) + +extern CL_API_ENTRY cl_int CL_API_CALL +clGetLayerInfo( + cl_layer_info param_name, + size_t param_value_size, + void* param_value, + size_t* param_value_size_ret) ; + +extern CL_API_ENTRY cl_int CL_API_CALL +clInitLayer( + cl_uint num_entries, + const cl_icd_dispatch* target_dispatch, + cl_uint* num_entries_ret, + const cl_icd_dispatch** layer_dispatch_ret) ; + +extern 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) ; + +extern CL_API_ENTRY cl_int CL_API_CALL +clDeinitLayer( + void ) ; + +#endif /* !defined(CL_NO_NON_ICD_DISPATCH_EXTENSION_PROTOTYPES) */ + +#ifdef __cplusplus +} +#endif + +#endif /* OPENCL_CL_LAYER_H_ */ diff --git a/intercept/CL/cl_platform.h b/intercept/CL/cl_platform.h index a1c20775..95172c49 100644 --- a/intercept/CL/cl_platform.h +++ b/intercept/CL/cl_platform.h @@ -243,7 +243,7 @@ typedef double cl_double; #define CL_NAN (CL_INFINITY - CL_INFINITY) #define CL_HUGE_VALF ((cl_float) 1e50) -#define CL_HUGE_VAL ((cl_double) 1e500) +#define CL_HUGE_VAL ((cl_double) (1e300 * 1e300)) #define CL_MAXFLOAT CL_FLT_MAX #define CL_INFINITY CL_HUGE_VALF diff --git a/intercept/CMakeLists.txt b/intercept/CMakeLists.txt index 11e81e7e..fc163848 100644 --- a/intercept/CMakeLists.txt +++ b/intercept/CMakeLists.txt @@ -72,6 +72,7 @@ set(CLINTERCEPT_SOURCE_FILES src/instrumentation.h src/intercept.cpp src/intercept.h + src/layer.cpp src/main.cpp src/objtracker.cpp src/objtracker.h diff --git a/intercept/OS/OS_linux.h b/intercept/OS/OS_linux.h index dc30a1d7..74081835 100644 --- a/intercept/OS/OS_linux.h +++ b/intercept/OS/OS_linux.h @@ -38,12 +38,7 @@ #endif #endif -// Since we just need the symbol clGetPlatformIDs, we can use any OpenCL -// version. -#if !defined(CL_TARGET_OPENCL_VERSION) -#define CL_TARGET_OPENCL_VERSION 100 -#endif -#include "CL/cl.h" +void CLIntercept_Load(void); namespace OS { @@ -75,7 +70,8 @@ class Services const std::string& str ) const; void* LoadLibrary( - const std::string& libraryName ) const; + const std::string& libraryName, + const bool deepBind = false ) const; void UnloadLibrary( void*& pLibrary ) const; @@ -192,7 +188,8 @@ inline void Services::OutputDebugString( } inline void* Services::LoadLibrary( - const std::string& libraryName ) const + const std::string& libraryName, + const bool deepBind ) const { void* pLibrary = dlopen( libraryName.c_str(), RTLD_NOW | RTLD_GLOBAL ); //if( pLibrary == NULL ) @@ -298,9 +295,10 @@ inline bool Services::GetCLInterceptName( std::string& name ) const { Dl_info info; - if( dladdr( (void*)clGetPlatformIDs, &info ) ) + if( dladdr( (void*)CLIntercept_Load, &info ) ) { name = info.dli_fname; + return true; } return false; } diff --git a/intercept/OS/OS_mac.h b/intercept/OS/OS_mac.h index 59ca604d..993e5f1e 100644 --- a/intercept/OS/OS_mac.h +++ b/intercept/OS/OS_mac.h @@ -52,7 +52,8 @@ class Services const std::string& str ) const; void* LoadLibrary( - const std::string& libraryName ) const; + const std::string& libraryName, + const bool deepBind = false ) const; void UnloadLibrary( void*& pLibrary ) const; @@ -145,9 +146,11 @@ inline void Services::OutputDebugString( } inline void* Services::LoadLibrary( - const std::string& libraryName ) const + const std::string& libraryName, + const bool deepBind ) const { - void* pLibrary = dlopen( libraryName.c_str(), RTLD_NOW ); + const int flags = RTLD_NOW | (deepBind ? RTLD_FIRST : 0); + void* pLibrary = dlopen( libraryName.c_str(), flags ); return pLibrary; } @@ -241,6 +244,7 @@ inline bool Services::GetCLInterceptName( if( dladdr( (void*)CLIntercept_Load, &info ) ) { name = info.dli_fname; + return true; } return false; } diff --git a/intercept/OS/OS_windows.h b/intercept/OS/OS_windows.h index 3522655e..15324b97 100644 --- a/intercept/OS/OS_windows.h +++ b/intercept/OS/OS_windows.h @@ -41,7 +41,8 @@ class Services const std::string& str ) const; void* LoadLibrary( - const std::string& libraryName ) const; + const std::string& libraryName, + bool deepBind = false ) const; void UnloadLibrary( void*& pLibrary ) const; @@ -223,7 +224,8 @@ inline void Services::OutputDebugString( } inline void* Services::LoadLibrary( - const std::string& libraryName ) const + const std::string& libraryName, + const bool ) const { HMODULE hModule = ::LoadLibraryA( libraryName.c_str() ); return hModule; diff --git a/intercept/src/clIntercept.def b/intercept/src/clIntercept.def index 16742bee..73e00208 100644 --- a/intercept/src/clIntercept.def +++ b/intercept/src/clIntercept.def @@ -12,6 +12,7 @@ EXPORTS ; ========================================================== clGetLayerInfo clInitLayer + clInitLayerWithProperties ; ========================================================== ; Standard OpenCL APIs diff --git a/intercept/src/clIntercept.map b/intercept/src/clIntercept.map index a009bc2e..caf2b8dd 100644 --- a/intercept/src/clIntercept.map +++ b/intercept/src/clIntercept.map @@ -10,6 +10,9 @@ INTERNAL { _binary_kernels_builtin_kernels_cl_end; _binary_kernels_precompiled_kernels_cl_start; _binary_kernels_precompiled_kernels_cl_end; + clGetLayerInfo; + clInitLayer; + clInitLayerWithProperties; local: *; }; diff --git a/intercept/src/git_version.cpp.in b/intercept/src/git_version.cpp.in index fb741a59..d7fcb9e6 100644 --- a/intercept/src/git_version.cpp.in +++ b/intercept/src/git_version.cpp.in @@ -6,6 +6,6 @@ #include "intercept.h" -const char* CLIntercept::sc_GitDescribe = "@GIT_DESCRIBE@"; -const char* CLIntercept::sc_GitRefSpec = "@GIT_REFSPEC@"; -const char* CLIntercept::sc_GitHash = "@GIT_SHA1@"; +const char* const CLIntercept::sc_GitDescribe = "@GIT_DESCRIBE@"; +const char* const CLIntercept::sc_GitRefSpec = "@GIT_REFSPEC@"; +const char* const CLIntercept::sc_GitHash = "@GIT_SHA1@"; diff --git a/intercept/src/intercept.cpp b/intercept/src/intercept.cpp index 07c63023..d11918e3 100644 --- a/intercept/src/intercept.cpp +++ b/intercept/src/intercept.cpp @@ -87,12 +87,12 @@ static inline uint64_t Hash( } #undef HASH_JENKINS_MIX -const char* CLIntercept::sc_URL = "https://github.com/intel/opencl-intercept-layer"; -const char* CLIntercept::sc_DumpDirectoryName = "CLIntercept_Dump"; -const char* CLIntercept::sc_ReportFileName = "clintercept_report.txt"; -const char* CLIntercept::sc_LogFileName = "clintercept_log.txt"; -const char* CLIntercept::sc_PerfCountersFileNamePrefix = "clintercept_perfcounter"; -const char* CLIntercept::sc_TraceFileName = "clintercept_trace.json"; +const char* const CLIntercept::sc_URL = "https://github.com/intel/opencl-intercept-layer"; +const char* const CLIntercept::sc_DumpDirectoryName = "CLIntercept_Dump"; +const char* const CLIntercept::sc_ReportFileName = "clintercept_report.txt"; +const char* const CLIntercept::sc_LogFileName = "clintercept_log.txt"; +const char* const CLIntercept::sc_PerfCountersFileNamePrefix = "clintercept_perfcounter"; +const char* const CLIntercept::sc_TraceFileName = "clintercept_trace.json"; /////////////////////////////////////////////////////////////////////////////// // @@ -653,6 +653,241 @@ bool CLIntercept::init() return true; } +/////////////////////////////////////////////////////////////////////////////// +// +#define INIT_NATIVE_LAYER_FUNC(pLibrary, funcname) \ +{ \ + constexpr cl_uint entry = offsetof(cl_icd_dispatch, funcname) / \ + sizeof(layerDispatch . funcname); \ + if( entry < num_entries ) \ + { \ + void* func = OS().GetFunctionPointer( pLibrary, #funcname); \ + void** pfunc = (void**)( &layerDispatch . funcname ); \ + *pfunc = reinterpret_cast(func); \ + m_Dispatch . funcname = target_dispatch -> funcname; \ + } \ + else \ + { \ + logf("Not enough layer entries for function %s (entry %u)!\n", \ + #funcname, entry ); \ + } \ +} + +void CLIntercept::initLayer( + cl_uint num_entries, + const cl_icd_dispatch *target_dispatch, + cl_uint *num_entries_out, + const cl_icd_dispatch **layer_dispatch_ret ) +{ + // This should have been checked earlier, but check again just in case. + if( target_dispatch == nullptr || + num_entries_out == nullptr || + layer_dispatch_ret == nullptr ) + { + CLI_ASSERT( 0 ); + return; + } + + std::lock_guard lock(m_Mutex); + + log( "CLIntercept layer is initializing...\n" ); + + static cl_icd_dispatch layerDispatch; + + // In the layer path, the real OpenCL ICD loader has been loaded first, so + // (on some operating systems, at least) its symbols are used for the OpenCL + // APIs. Since we want to install the CLIntercept symbols into the layer + // dispatch instead, we need to explicitly load the CLIntercept library to + // get the right symbols. + + std::string name; + OS().GetCLInterceptName(name); + + void* pLibrary = OS().LoadLibrary( name.c_str(), true ); + if( pLibrary ) + { + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetPlatformIDs ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetPlatformInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetDeviceIDs ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetDeviceInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateContext ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateContextFromType ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clRetainContext ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clReleaseContext ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetContextInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateCommandQueue ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clRetainCommandQueue ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clReleaseCommandQueue ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetCommandQueueInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetCommandQueueProperty ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateBuffer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateImage2D ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateImage3D ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clRetainMemObject ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clReleaseMemObject ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetSupportedImageFormats ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetMemObjectInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetImageInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateSampler ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clRetainSampler ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clReleaseSampler ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetSamplerInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateProgramWithSource ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateProgramWithBinary ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clRetainProgram ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clReleaseProgram ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clBuildProgram ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clUnloadCompiler ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetProgramInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetProgramBuildInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateKernel ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateKernelsInProgram ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clRetainKernel ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clReleaseKernel ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetKernelArg ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetKernelInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetKernelWorkGroupInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clWaitForEvents ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetEventInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clRetainEvent ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clReleaseEvent ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetEventProfilingInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clFlush ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clFinish ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueReadBuffer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueWriteBuffer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueCopyBuffer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueReadImage ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueWriteImage ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueCopyImage ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueCopyImageToBuffer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueCopyBufferToImage ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueMapBuffer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueMapImage ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueUnmapMemObject ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueNDRangeKernel ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueTask ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueNativeKernel ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueMarker ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueWaitForEvents ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueBarrier ); + + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetExtensionFunctionAddress ); + + // cl_khr_gl_sharing + // The entry points for this extension are exported from the ICD + // loader even though they are extension APIs. + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetGLContextInfoKHR ); +#if defined (_WIN32) || defined (__linux__) || defined (__FreeBSD__) + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromGLBuffer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromGLTexture ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromGLTexture2D ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromGLTexture3D ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromGLRenderbuffer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetGLObjectInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetGLTextureInfo ); // OpenCL 1.2 + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueAcquireGLObjects ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueReleaseGLObjects ); +#endif + +#if defined(_WIN32) + // cl_khr_d3d10_sharing + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetDeviceIDsFromD3D10KHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromD3D10BufferKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromD3D10Texture2DKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromD3D10Texture3DKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueAcquireD3D10ObjectsKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueReleaseD3D10ObjectsKHR ); + + // cl_khr_d3d11_sharing + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetDeviceIDsFromD3D11KHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromD3D11BufferKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromD3D11Texture2DKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromD3D11Texture3DKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueAcquireD3D11ObjectsKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueReleaseD3D11ObjectsKHR ); + + // cl_khr_dx9_media_sharing + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetDeviceIDsFromDX9MediaAdapterKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateFromDX9MediaSurfaceKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueAcquireDX9MediaSurfacesKHR ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueReleaseDX9MediaSurfacesKHR ); +#endif + + // OpenCL 1.1 Entry Points (optional) + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetEventCallback ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateSubBuffer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetMemObjectDestructorCallback ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateUserEvent ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetUserEventStatus ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueReadBufferRect ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueWriteBufferRect ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueCopyBufferRect ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateEventFromGLsyncKHR ); + + // OpenCL 1.2 Entry Points (optional) + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateSubDevices ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clRetainDevice ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clReleaseDevice ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateImage ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateProgramWithBuiltInKernels ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCompileProgram ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clLinkProgram ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clUnloadPlatformCompiler ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetKernelArgInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueFillBuffer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueFillImage ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueMigrateMemObjects ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueMarkerWithWaitList ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueBarrierWithWaitList ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetExtensionFunctionAddressForPlatform ); + + // OpenCL 2.0 Entry Points + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateCommandQueueWithProperties ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreatePipe ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetPipeInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSVMAlloc ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSVMFree ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueSVMFree ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueSVMMemcpy ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueSVMMemFill ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueSVMMap ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueSVMUnmap ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateSamplerWithProperties ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetKernelArgSVMPointer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetKernelExecInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetKernelSubGroupInfoKHR ); + + // OpenCL 2.1 Entry Points + INIT_NATIVE_LAYER_FUNC( pLibrary, clCloneKernel ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateProgramWithIL ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clEnqueueSVMMigrateMem ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetDeviceAndHostTimer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetHostTimer ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetKernelSubGroupInfo ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetDefaultDeviceCommandQueue ); + + // OpenCL 2.2 Entry Points + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetProgramReleaseCallback ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetProgramSpecializationConstant ); + + // OpenCL 3.0 Entry Points + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateBufferWithProperties ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clCreateImageWithProperties ); + INIT_NATIVE_LAYER_FUNC( pLibrary, clSetContextDestructorCallback ); + + // OpenCL 3.1 Entry Points + INIT_NATIVE_LAYER_FUNC( pLibrary, clGetKernelSuggestedLocalWorkSize ); + + *num_entries_out = num_entries; + *layer_dispatch_ret = &layerDispatch; + } + + OS().UnloadLibrary( pLibrary ); + + log( "... layer initialization complete.\n" ); +} + /////////////////////////////////////////////////////////////////////////////// // void CLIntercept::report() @@ -10807,7 +11042,7 @@ cl_int CLIntercept::writeStringToMemory( size_t param_value_size, const std::string& param, size_t* param_value_size_ret, - char* pointer ) const + char* pointer ) { cl_int errorCode = CL_SUCCESS; @@ -13952,190 +14187,6 @@ bool CLIntercept::initDispatch( void ) #error Unknown OS! #endif -/////////////////////////////////////////////////////////////////////////////// -// -#define INIT_NATIVE_LAYER_FUNC(funcname) \ -{ \ - void** pfunc = (void**)( &layer_dispatch . funcname ); \ - *pfunc = reinterpret_cast(funcname); \ -} - -void CLIntercept::initLayer(cl_icd_dispatch& layer_dispatch, const cl_icd_dispatch* target_dispatch) -{ - m_Dispatch = *target_dispatch; - layer_dispatch = *target_dispatch; - - INIT_NATIVE_LAYER_FUNC( clGetPlatformIDs ); - INIT_NATIVE_LAYER_FUNC( clGetPlatformInfo ); - INIT_NATIVE_LAYER_FUNC( clGetDeviceIDs ); - INIT_NATIVE_LAYER_FUNC( clGetDeviceInfo ); - INIT_NATIVE_LAYER_FUNC( clCreateContext ); - INIT_NATIVE_LAYER_FUNC( clCreateContextFromType ); - INIT_NATIVE_LAYER_FUNC( clRetainContext ); - INIT_NATIVE_LAYER_FUNC( clReleaseContext ); - INIT_NATIVE_LAYER_FUNC( clGetContextInfo ); - INIT_NATIVE_LAYER_FUNC( clCreateCommandQueue ); - INIT_NATIVE_LAYER_FUNC( clRetainCommandQueue ); - INIT_NATIVE_LAYER_FUNC( clReleaseCommandQueue ); - INIT_NATIVE_LAYER_FUNC( clGetCommandQueueInfo ); - INIT_NATIVE_LAYER_FUNC( clSetCommandQueueProperty ); - INIT_NATIVE_LAYER_FUNC( clCreateBuffer ); - INIT_NATIVE_LAYER_FUNC( clCreateImage2D ); - INIT_NATIVE_LAYER_FUNC( clCreateImage3D ); - INIT_NATIVE_LAYER_FUNC( clRetainMemObject ); - INIT_NATIVE_LAYER_FUNC( clReleaseMemObject ); - INIT_NATIVE_LAYER_FUNC( clGetSupportedImageFormats ); - INIT_NATIVE_LAYER_FUNC( clGetMemObjectInfo ); - INIT_NATIVE_LAYER_FUNC( clGetImageInfo ); - INIT_NATIVE_LAYER_FUNC( clCreateSampler ); - INIT_NATIVE_LAYER_FUNC( clRetainSampler ); - INIT_NATIVE_LAYER_FUNC( clReleaseSampler ); - INIT_NATIVE_LAYER_FUNC( clGetSamplerInfo ); - INIT_NATIVE_LAYER_FUNC( clCreateProgramWithSource ); - INIT_NATIVE_LAYER_FUNC( clCreateProgramWithBinary ); - INIT_NATIVE_LAYER_FUNC( clRetainProgram ); - INIT_NATIVE_LAYER_FUNC( clReleaseProgram ); - INIT_NATIVE_LAYER_FUNC( clBuildProgram ); - INIT_NATIVE_LAYER_FUNC( clUnloadCompiler ); - INIT_NATIVE_LAYER_FUNC( clGetProgramInfo ); - INIT_NATIVE_LAYER_FUNC( clGetProgramBuildInfo ); - INIT_NATIVE_LAYER_FUNC( clCreateKernel ); - INIT_NATIVE_LAYER_FUNC( clCreateKernelsInProgram ); - INIT_NATIVE_LAYER_FUNC( clRetainKernel ); - INIT_NATIVE_LAYER_FUNC( clReleaseKernel ); - INIT_NATIVE_LAYER_FUNC( clSetKernelArg ); - INIT_NATIVE_LAYER_FUNC( clGetKernelInfo ); - INIT_NATIVE_LAYER_FUNC( clGetKernelWorkGroupInfo ); - INIT_NATIVE_LAYER_FUNC( clWaitForEvents ); - INIT_NATIVE_LAYER_FUNC( clGetEventInfo ); - INIT_NATIVE_LAYER_FUNC( clRetainEvent ); - INIT_NATIVE_LAYER_FUNC( clReleaseEvent ); - INIT_NATIVE_LAYER_FUNC( clGetEventProfilingInfo ); - INIT_NATIVE_LAYER_FUNC( clFlush ); - INIT_NATIVE_LAYER_FUNC( clFinish ); - INIT_NATIVE_LAYER_FUNC( clEnqueueReadBuffer ); - INIT_NATIVE_LAYER_FUNC( clEnqueueWriteBuffer ); - INIT_NATIVE_LAYER_FUNC( clEnqueueCopyBuffer ); - INIT_NATIVE_LAYER_FUNC( clEnqueueReadImage ); - INIT_NATIVE_LAYER_FUNC( clEnqueueWriteImage ); - INIT_NATIVE_LAYER_FUNC( clEnqueueCopyImage ); - INIT_NATIVE_LAYER_FUNC( clEnqueueCopyImageToBuffer ); - INIT_NATIVE_LAYER_FUNC( clEnqueueCopyBufferToImage ); - INIT_NATIVE_LAYER_FUNC( clEnqueueMapBuffer ); - INIT_NATIVE_LAYER_FUNC( clEnqueueMapImage ); - INIT_NATIVE_LAYER_FUNC( clEnqueueUnmapMemObject ); - INIT_NATIVE_LAYER_FUNC( clEnqueueNDRangeKernel ); - INIT_NATIVE_LAYER_FUNC( clEnqueueTask ); - INIT_NATIVE_LAYER_FUNC( clEnqueueNativeKernel ); - INIT_NATIVE_LAYER_FUNC( clEnqueueMarker ); - INIT_NATIVE_LAYER_FUNC( clEnqueueWaitForEvents ); - INIT_NATIVE_LAYER_FUNC( clEnqueueBarrier ); - - INIT_NATIVE_LAYER_FUNC( clGetExtensionFunctionAddress ); - - // cl_khr_gl_sharing (optional) - // The entry points for this extension are exported from the ICD - // loader even though they are extension APIs. - INIT_NATIVE_LAYER_FUNC( clGetGLContextInfoKHR ); - -#if defined (_WIN32) || defined (__linux__) || defined (__FreeBSD__) - INIT_NATIVE_LAYER_FUNC( clCreateFromGLBuffer ); - INIT_NATIVE_LAYER_FUNC( clCreateFromGLTexture2D ); - INIT_NATIVE_LAYER_FUNC( clCreateFromGLTexture3D ); - INIT_NATIVE_LAYER_FUNC( clCreateFromGLRenderbuffer ); - INIT_NATIVE_LAYER_FUNC( clGetGLObjectInfo ); - INIT_NATIVE_LAYER_FUNC( clGetGLTextureInfo ); // OpenCL 1.2 - INIT_NATIVE_LAYER_FUNC( clEnqueueAcquireGLObjects ); - INIT_NATIVE_LAYER_FUNC( clEnqueueReleaseGLObjects ); -#endif - -#if defined(_WIN32) - // cl_khr_d3d10_sharing - INIT_NATIVE_LAYER_FUNC( clGetDeviceIDsFromD3D10KHR ); - INIT_NATIVE_LAYER_FUNC( clCreateFromD3D10BufferKHR ); - INIT_NATIVE_LAYER_FUNC( clCreateFromD3D10Texture2DKHR ); - INIT_NATIVE_LAYER_FUNC( clCreateFromD3D10Texture3DKHR ); - INIT_NATIVE_LAYER_FUNC( clEnqueueAcquireD3D10ObjectsKHR ); - INIT_NATIVE_LAYER_FUNC( clEnqueueReleaseD3D10ObjectsKHR ); - - // cl_khr_d3d11_sharing - INIT_NATIVE_LAYER_FUNC( clGetDeviceIDsFromD3D11KHR ); - INIT_NATIVE_LAYER_FUNC( clCreateFromD3D11BufferKHR ); - INIT_NATIVE_LAYER_FUNC( clCreateFromD3D11Texture2DKHR ); - INIT_NATIVE_LAYER_FUNC( clCreateFromD3D11Texture3DKHR ); - INIT_NATIVE_LAYER_FUNC( clEnqueueAcquireD3D11ObjectsKHR ); - INIT_NATIVE_LAYER_FUNC( clEnqueueReleaseD3D11ObjectsKHR ); - - // cl_khr_dx9_media_sharing - INIT_NATIVE_LAYER_FUNC( clGetDeviceIDsFromDX9MediaAdapterKHR ); - INIT_NATIVE_LAYER_FUNC( clCreateFromDX9MediaSurfaceKHR ); - INIT_NATIVE_LAYER_FUNC( clEnqueueAcquireDX9MediaSurfacesKHR ); - INIT_NATIVE_LAYER_FUNC( clEnqueueReleaseDX9MediaSurfacesKHR ); -#endif - - INIT_NATIVE_LAYER_FUNC( clSetEventCallback ); - INIT_NATIVE_LAYER_FUNC( clCreateSubBuffer ); - INIT_NATIVE_LAYER_FUNC( clSetMemObjectDestructorCallback ); - INIT_NATIVE_LAYER_FUNC( clCreateUserEvent ); - INIT_NATIVE_LAYER_FUNC( clSetUserEventStatus ); - INIT_NATIVE_LAYER_FUNC( clEnqueueReadBufferRect ); - INIT_NATIVE_LAYER_FUNC( clEnqueueWriteBufferRect ); - INIT_NATIVE_LAYER_FUNC( clEnqueueCopyBufferRect ); - INIT_NATIVE_LAYER_FUNC( clCreateEventFromGLsyncKHR ); - INIT_NATIVE_LAYER_FUNC( clCreateSubDevices ); - INIT_NATIVE_LAYER_FUNC( clRetainDevice ); - INIT_NATIVE_LAYER_FUNC( clReleaseDevice ); - INIT_NATIVE_LAYER_FUNC( clCreateImage ); - INIT_NATIVE_LAYER_FUNC( clCreateProgramWithBuiltInKernels ); - INIT_NATIVE_LAYER_FUNC( clCompileProgram ); - INIT_NATIVE_LAYER_FUNC( clLinkProgram ); - INIT_NATIVE_LAYER_FUNC( clUnloadPlatformCompiler ); - INIT_NATIVE_LAYER_FUNC( clGetKernelArgInfo ); - INIT_NATIVE_LAYER_FUNC( clEnqueueFillBuffer ); - INIT_NATIVE_LAYER_FUNC( clEnqueueFillImage ); - INIT_NATIVE_LAYER_FUNC( clEnqueueMigrateMemObjects ); - INIT_NATIVE_LAYER_FUNC( clEnqueueMarkerWithWaitList ); - INIT_NATIVE_LAYER_FUNC( clEnqueueBarrierWithWaitList ); - INIT_NATIVE_LAYER_FUNC( clGetExtensionFunctionAddressForPlatform ); - INIT_NATIVE_LAYER_FUNC( clCreateFromGLTexture ); - -#if !defined(__APPLE__) - // OpenCL 2.0 Entry Points (optional) - INIT_NATIVE_LAYER_FUNC( clCreateCommandQueueWithProperties ); - INIT_NATIVE_LAYER_FUNC( clCreatePipe ); - INIT_NATIVE_LAYER_FUNC( clGetPipeInfo ); - INIT_NATIVE_LAYER_FUNC( clSVMAlloc ); - INIT_NATIVE_LAYER_FUNC( clSVMFree ); - INIT_NATIVE_LAYER_FUNC( clEnqueueSVMFree ); - INIT_NATIVE_LAYER_FUNC( clEnqueueSVMMemcpy ); - INIT_NATIVE_LAYER_FUNC( clEnqueueSVMMemFill ); - INIT_NATIVE_LAYER_FUNC( clEnqueueSVMMap ); - INIT_NATIVE_LAYER_FUNC( clEnqueueSVMUnmap ); - INIT_NATIVE_LAYER_FUNC( clCreateSamplerWithProperties ); - INIT_NATIVE_LAYER_FUNC( clSetKernelArgSVMPointer ); - INIT_NATIVE_LAYER_FUNC( clSetKernelExecInfo ); - INIT_NATIVE_LAYER_FUNC( clGetKernelSubGroupInfoKHR ); - - // OpenCL 2.1 Entry Points (optional) - INIT_NATIVE_LAYER_FUNC( clCloneKernel ); - INIT_NATIVE_LAYER_FUNC( clCreateProgramWithIL ); - INIT_NATIVE_LAYER_FUNC( clEnqueueSVMMigrateMem ); - INIT_NATIVE_LAYER_FUNC( clGetDeviceAndHostTimer ); - INIT_NATIVE_LAYER_FUNC( clGetHostTimer ); - INIT_NATIVE_LAYER_FUNC( clGetKernelSubGroupInfo ); - INIT_NATIVE_LAYER_FUNC( clSetDefaultDeviceCommandQueue ); - - // OpenCL 2.2 Entry Points (optional) - INIT_NATIVE_LAYER_FUNC( clSetProgramReleaseCallback ); - INIT_NATIVE_LAYER_FUNC( clSetProgramSpecializationConstant ); - - // OpenCL 3.0 Entry Points (optional) - INIT_NATIVE_LAYER_FUNC( clCreateBufferWithProperties ); - INIT_NATIVE_LAYER_FUNC( clCreateImageWithProperties ); - INIT_NATIVE_LAYER_FUNC( clSetContextDestructorCallback ); -#endif -} - /////////////////////////////////////////////////////////////////////////////// // void CLIntercept::dumpMemoryToFile( diff --git a/intercept/src/intercept.h b/intercept/src/intercept.h index a35e2a66..c53952de 100644 --- a/intercept/src/intercept.h +++ b/intercept/src/intercept.h @@ -61,7 +61,11 @@ class CLIntercept static bool Create( void* pGlobalData, CLIntercept*& pIntercept ); static void Delete( CLIntercept*& pIntercept ); - void initLayer(cl_icd_dispatch& layer_dispatch, const cl_icd_dispatch* target_dispatch); + void initLayer( + cl_uint num_entries, + const cl_icd_dispatch *target_dispatch, + cl_uint *num_entries_out, + const cl_icd_dispatch **layer_dispatch_ret ); void report(); @@ -684,23 +688,23 @@ class CLIntercept void initBuiltinKernelOverrides( const cl_context context ); - cl_int writeStringToMemory( - size_t param_value_size, - const std::string& param, - size_t* param_value_size_ret, - char* pointer ) const; + static cl_int writeStringToMemory( + size_t param_value_size, + const std::string& param, + size_t* param_value_size_ret, + char* pointer ); template< class T > - cl_int writeVectorToMemory( - size_t param_value_size, - const std::vector& param, - size_t* param_value_size_ret, - T* pointer ) const; + static cl_int writeVectorToMemory( + size_t param_value_size, + const std::vector& param, + size_t* param_value_size_ret, + T* pointer ); template< class T > - cl_int writeParamToMemory( - size_t param_value_size, - T param, - size_t* param_value_size_ret, - T* pointer ) const; + static cl_int writeParamToMemory( + size_t param_value_size, + T param, + size_t* param_value_size_ret, + T* pointer ); bool overrideGetPlatformInfo( cl_platform_id platform, @@ -1007,20 +1011,20 @@ class CLIntercept cl_int finishAll( cl_context conetxt ); -private: - static const char* sc_URL; - static const char* sc_DumpDirectoryName; - static const char* sc_ReportFileName; - static const char* sc_LogFileName; - static const char* sc_TraceFileName; - static const char* sc_PerfCountersFileNamePrefix; - #if defined(CLINTERCEPT_CMAKE) - static const char* sc_GitDescribe; - static const char* sc_GitRefSpec; - static const char* sc_GitHash; + static const char* const sc_GitDescribe; + static const char* const sc_GitRefSpec; + static const char* const sc_GitHash; #endif +private: + static const char* const sc_URL; + static const char* const sc_DumpDirectoryName; + static const char* const sc_ReportFileName; + static const char* const sc_LogFileName; + static const char* const sc_TraceFileName; + static const char* const sc_PerfCountersFileNamePrefix; + CLIntercept( void* pGlobalData ); ~CLIntercept(); CLIntercept( const CLIntercept& ) = delete; @@ -1532,7 +1536,7 @@ cl_int CLIntercept::writeVectorToMemory( size_t param_value_size, const std::vector& param, size_t *param_value_size_ret, - T* pointer ) const + T* pointer ) { cl_int errorCode = CL_SUCCESS; @@ -1565,7 +1569,7 @@ cl_int CLIntercept::writeParamToMemory( size_t param_value_size, T param, size_t *param_value_size_ret, - T* pointer ) const + T* pointer ) { cl_int errorCode = CL_SUCCESS; diff --git a/intercept/src/layer.cpp b/intercept/src/layer.cpp new file mode 100644 index 00000000..f4eb1d4a --- /dev/null +++ b/intercept/src/layer.cpp @@ -0,0 +1,80 @@ +/* +// Copyright (c) 2018-2026 Intel Corporation +// +// SPDX-License-Identifier: MIT +*/ + +#include "intercept.h" + +#include "CL/cl_layer.h" + +extern "C" CL_API_ENTRY cl_int CL_API_CALL clGetLayerInfo( + cl_layer_info param_name, + size_t param_value_size, + void *param_value, + size_t *param_value_size_ret) +{ + switch (param_name) { + case CL_LAYER_API_VERSION: { + const cl_layer_api_version version = CL_LAYER_API_VERSION_100; + auto* ptr = (cl_layer_api_version*)param_value; + return CLIntercept::writeParamToMemory( + param_value_size, + version, + param_value_size_ret, + ptr ); + } + case CL_LAYER_NAME: { + std::string name("CLIntercept"); +#if defined(CLINTERCEPT_CMAKE) + name += " (Version: "; + name += CLIntercept::sc_GitDescribe; + name += ")"; +#endif + char* ptr = (char*)param_value; + return CLIntercept::writeStringToMemory( + param_value_size, + name, + param_value_size_ret, + ptr ); + } + + default: + return CL_INVALID_VALUE; + } +} + +extern "C" CL_API_ENTRY cl_int CL_API_CALL clInitLayerWithProperties( + cl_uint num_entries, + const cl_icd_dispatch *target_dispatch, + cl_uint *num_entries_out, + const cl_icd_dispatch **layer_dispatch_ret, + const cl_layer_properties* properties) +{ + if ( target_dispatch == NULL || num_entries_out == NULL || layer_dispatch_ret == NULL ) { + return CL_INVALID_VALUE; + } + + CLIntercept* pIntercept = GetIntercept(); + pIntercept->initLayer( + num_entries, + target_dispatch, + num_entries_out, + layer_dispatch_ret ); + + return CL_SUCCESS; +} + +extern "C" CL_API_ENTRY cl_int CL_API_CALL clInitLayer( + cl_uint num_entries, + const cl_icd_dispatch *target_dispatch, + cl_uint *num_entries_out, + const cl_icd_dispatch **layer_dispatch_ret) +{ + return clInitLayerWithProperties( + num_entries, + target_dispatch, + num_entries_out, + layer_dispatch_ret, + nullptr); +} diff --git a/intercept/src/main.cpp b/intercept/src/main.cpp index 00215759..24aeed1b 100644 --- a/intercept/src/main.cpp +++ b/intercept/src/main.cpp @@ -68,67 +68,3 @@ void CLIntercept_Unload(void) #else #error Unknown OS! #endif - -typedef cl_uint cl_layer_info; -typedef cl_uint cl_layer_api_version; - -#define CL_LAYER_API_VERSION 0x4240 -#define CL_LAYER_NAME 0x4241 -#define CL_LAYER_API_VERSION_100 100 - -static cl_icd_dispatch layer_dispatch; - -extern "C" CL_API_ENTRY cl_int CL_API_CALL clGetLayerInfo( - cl_layer_info param_name, - size_t param_value_size, - void *param_value, - size_t *param_value_size_ret) -{ - switch (param_name) { - case CL_LAYER_API_VERSION: { - const cl_layer_api_version version = CL_LAYER_API_VERSION_100; - if (param_value_size_ret) - *param_value_size_ret = sizeof(version); - - if (param_value) { - if (param_value_size < sizeof(version)) - return CL_INVALID_VALUE; - - CLI_MEMCPY(param_value, param_value_size, &version, sizeof(version)); - } - return CL_SUCCESS; - } - case CL_LAYER_NAME: { - const char name[] = "CLIntercept"; - if (param_value_size_ret) - *param_value_size_ret = sizeof(name); - - if (param_value) { - if (param_value_size < sizeof(name)) - return CL_INVALID_VALUE; - - CLI_MEMCPY(param_value, param_value_size, name, sizeof(name)); - } - return CL_SUCCESS; - } - - default: - return CL_INVALID_VALUE; - } -} - -extern "C" CL_API_ENTRY cl_int CL_API_CALL clInitLayer( - cl_uint num_entries, - const cl_icd_dispatch *target_dispatch, - cl_uint *num_entries_out, - const cl_icd_dispatch **layer_dispatch_ret) -{ - if ( target_dispatch == NULL || num_entries_out == NULL || layer_dispatch_ret == NULL ) { - return CL_INVALID_VALUE; - } - - g_pIntercept->initLayer( layer_dispatch, target_dispatch ); - *layer_dispatch_ret = &layer_dispatch; - *num_entries_out = num_entries; - return CL_SUCCESS; -}