diff --git a/intercept/src/clIntercept.def b/intercept/src/clIntercept.def index 5fbc4e6e..4b8c8eab 100644 --- a/intercept/src/clIntercept.def +++ b/intercept/src/clIntercept.def @@ -4,6 +4,18 @@ LIBRARY opencl.dll EXPORTS +; ========================================================== +; OpenCL Native Layer Entry Points +; +; Required exclusively by the Khronos ICD Loader to +; negotiate and initialize the layer (via OPENCL_LAYERS). +; ========================================================== + clGetLayerInfo + clInitLayer + +; ========================================================== +; Standard OpenCL APIs +; ========================================================== clBuildProgram clCloneKernel clCompileProgram diff --git a/intercept/src/intercept.cpp b/intercept/src/intercept.cpp index 82b9a144..b1b46800 100644 --- a/intercept/src/intercept.cpp +++ b/intercept/src/intercept.cpp @@ -14148,6 +14148,190 @@ 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 a54f1160..288dce5b 100644 --- a/intercept/src/intercept.h +++ b/intercept/src/intercept.h @@ -61,6 +61,8 @@ 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 report(); void callLoggingEnter( diff --git a/intercept/src/main.cpp b/intercept/src/main.cpp index 24aeed1b..00215759 100644 --- a/intercept/src/main.cpp +++ b/intercept/src/main.cpp @@ -68,3 +68,67 @@ 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; +}