From f5e3493ace74494ceb68076bfe479344d3d87244 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Tue, 16 Jan 2024 14:35:11 +0000 Subject: [PATCH] [SPEC] Add urProgramGetGlobalVariablePointer entrypoint --- include/ur_api.h | 49 +++++++++++++++ include/ur_ddi.h | 10 +++ include/ur_print.hpp | 44 +++++++++++++ scripts/core/program.yml | 37 +++++++++++ scripts/core/registry.yml | 3 + source/adapters/cuda/program.cpp | 31 ++++++++++ source/adapters/cuda/ur_interface_loader.cpp | 1 + source/adapters/hip/program.cpp | 5 ++ source/adapters/hip/ur_interface_loader.cpp | 1 + source/adapters/level_zero/program.cpp | 5 ++ .../level_zero/ur_interface_loader.cpp | 1 + source/adapters/null/ur_nullddi.cpp | 36 +++++++++++ source/adapters/opencl/common.hpp | 10 ++- source/adapters/opencl/program.cpp | 39 ++++++++++++ .../adapters/opencl/ur_interface_loader.cpp | 1 + source/loader/layers/tracing/ur_trcddi.cpp | 45 ++++++++++++++ source/loader/layers/validation/ur_valddi.cpp | 52 ++++++++++++++++ source/loader/ur_ldrddi.cpp | 41 ++++++++++++ source/loader/ur_libapi.cpp | 52 ++++++++++++++++ source/ur_api.cpp | 43 +++++++++++++ .../conformance/device_code/device_global.cpp | 5 +- .../enqueue/enqueue_adapter_cuda.match | 3 - .../enqueue/enqueue_adapter_opencl.match | 1 - .../urEnqueueDeviceGlobalVariableRead.cpp | 2 +- test/conformance/program/CMakeLists.txt | 1 + .../program/program_adapter_hip.match | 2 + .../program/program_adapter_level_zero.match | 2 + .../program/program_adapter_opencl.match | 2 +- .../urProgramGetGlobalVariablePointer.cpp | 62 +++++++++++++++++++ .../testing/include/uur/fixtures.h | 25 +++++++- 30 files changed, 601 insertions(+), 10 deletions(-) create mode 100644 test/conformance/program/urProgramGetGlobalVariablePointer.cpp diff --git a/include/ur_api.h b/include/ur_api.h index 5c9c7af5da..cc49027599 100644 --- a/include/ur_api.h +++ b/include/ur_api.h @@ -215,6 +215,7 @@ typedef enum ur_function_t { UR_FUNCTION_COMMAND_BUFFER_APPEND_USM_ADVISE_EXP = 213, ///< Enumerator for ::urCommandBufferAppendUSMAdviseExp UR_FUNCTION_ENQUEUE_COOPERATIVE_KERNEL_LAUNCH_EXP = 214, ///< Enumerator for ::urEnqueueCooperativeKernelLaunchExp UR_FUNCTION_KERNEL_SUGGEST_MAX_COOPERATIVE_GROUP_COUNT_EXP = 215, ///< Enumerator for ::urKernelSuggestMaxCooperativeGroupCountExp + UR_FUNCTION_PROGRAM_GET_GLOBAL_VARIABLE_POINTER = 216, ///< Enumerator for ::urProgramGetGlobalVariablePointer /// @cond UR_FUNCTION_FORCE_UINT32 = 0x7fffffff /// @endcond @@ -4262,6 +4263,42 @@ urProgramGetFunctionPointer( void **ppFunctionPointer ///< [out] Returns the pointer to the function if it is found in the program. ); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Retrieves a pointer to a device global variable. +/// +/// @details +/// - Retrieves a pointer to a device global variable. +/// - The application may call this function from simultaneous threads for +/// the same device. +/// - The implementation of this function should be thread-safe. +/// +/// @remarks +/// _Analogues_ +/// - **clGetDeviceGlobalVariablePointerINTEL** +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hDevice` +/// + `NULL == hProgram` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pGlobalVariableName` +/// + `NULL == ppGlobalVariablePointerRet` +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// + `name` is not a valid variable in the program. +UR_APIEXPORT ur_result_t UR_APICALL +urProgramGetGlobalVariablePointer( + ur_device_handle_t hDevice, ///< [in] handle of the device to retrieve the pointer for. + ur_program_handle_t hProgram, ///< [in] handle of the program where the global variable is. + const char *pGlobalVariableName, ///< [in] mangled name of the global variable to retrieve the pointer for. + size_t *pGlobalVariableSizeRet, ///< [out][optional] Returns the size of the global variable if it is found + ///< in the program. + void **ppGlobalVariablePointerRet ///< [out] Returns the pointer to the global variable if it is found in the program. +); + /////////////////////////////////////////////////////////////////////////////// /// @brief Get Program object information typedef enum ur_program_info_t { @@ -9144,6 +9181,18 @@ typedef struct ur_program_get_function_pointer_params_t { void ***pppFunctionPointer; } ur_program_get_function_pointer_params_t; +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urProgramGetGlobalVariablePointer +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct ur_program_get_global_variable_pointer_params_t { + ur_device_handle_t *phDevice; + ur_program_handle_t *phProgram; + const char **ppGlobalVariableName; + size_t **ppGlobalVariableSizeRet; + void ***pppGlobalVariablePointerRet; +} ur_program_get_global_variable_pointer_params_t; + /////////////////////////////////////////////////////////////////////////////// /// @brief Function parameters for urProgramGetInfo /// @details Each entry is a pointer to the parameter passed to the function; diff --git a/include/ur_ddi.h b/include/ur_ddi.h index 92fc742f72..e84ef47e8c 100644 --- a/include/ur_ddi.h +++ b/include/ur_ddi.h @@ -329,6 +329,15 @@ typedef ur_result_t(UR_APICALL *ur_pfnProgramGetFunctionPointer_t)( const char *, void **); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urProgramGetGlobalVariablePointer +typedef ur_result_t(UR_APICALL *ur_pfnProgramGetGlobalVariablePointer_t)( + ur_device_handle_t, + ur_program_handle_t, + const char *, + size_t *, + void **); + /////////////////////////////////////////////////////////////////////////////// /// @brief Function-pointer for urProgramGetInfo typedef ur_result_t(UR_APICALL *ur_pfnProgramGetInfo_t)( @@ -380,6 +389,7 @@ typedef struct ur_program_dditable_t { ur_pfnProgramRetain_t pfnRetain; ur_pfnProgramRelease_t pfnRelease; ur_pfnProgramGetFunctionPointer_t pfnGetFunctionPointer; + ur_pfnProgramGetGlobalVariablePointer_t pfnGetGlobalVariablePointer; ur_pfnProgramGetInfo_t pfnGetInfo; ur_pfnProgramGetBuildInfo_t pfnGetBuildInfo; ur_pfnProgramSetSpecializationConstants_t pfnSetSpecializationConstants; diff --git a/include/ur_print.hpp b/include/ur_print.hpp index 70e5b9886d..051549b407 100644 --- a/include/ur_print.hpp +++ b/include/ur_print.hpp @@ -879,6 +879,9 @@ inline std::ostream &operator<<(std::ostream &os, ur_function_t value) { case UR_FUNCTION_KERNEL_SUGGEST_MAX_COOPERATIVE_GROUP_COUNT_EXP: os << "UR_FUNCTION_KERNEL_SUGGEST_MAX_COOPERATIVE_GROUP_COUNT_EXP"; break; + case UR_FUNCTION_PROGRAM_GET_GLOBAL_VARIABLE_POINTER: + os << "UR_FUNCTION_PROGRAM_GET_GLOBAL_VARIABLE_POINTER"; + break; default: os << "unknown enumerator"; break; @@ -10257,6 +10260,44 @@ inline std::ostream &operator<<(std::ostream &os, [[maybe_unused]] const struct return os; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_program_get_global_variable_pointer_params_t type +/// @returns +/// std::ostream & +inline std::ostream &operator<<(std::ostream &os, [[maybe_unused]] const struct ur_program_get_global_variable_pointer_params_t *params) { + + os << ".hDevice = "; + + ur::details::printPtr(os, + *(params->phDevice)); + + os << ", "; + os << ".hProgram = "; + + ur::details::printPtr(os, + *(params->phProgram)); + + os << ", "; + os << ".pGlobalVariableName = "; + + ur::details::printPtr(os, + *(params->ppGlobalVariableName)); + + os << ", "; + os << ".pGlobalVariableSizeRet = "; + + ur::details::printPtr(os, + *(params->ppGlobalVariableSizeRet)); + + os << ", "; + os << ".ppGlobalVariablePointerRet = "; + + ur::details::printPtr(os, + *(params->pppGlobalVariablePointerRet)); + + return os; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Print operator for the ur_program_get_info_params_t type /// @returns @@ -15992,6 +16033,9 @@ inline ur_result_t UR_APICALL printFunctionParams(std::ostream &os, ur_function_ case UR_FUNCTION_PROGRAM_GET_FUNCTION_POINTER: { os << (const struct ur_program_get_function_pointer_params_t *)params; } break; + case UR_FUNCTION_PROGRAM_GET_GLOBAL_VARIABLE_POINTER: { + os << (const struct ur_program_get_global_variable_pointer_params_t *)params; + } break; case UR_FUNCTION_PROGRAM_GET_INFO: { os << (const struct ur_program_get_info_params_t *)params; } break; diff --git a/scripts/core/program.yml b/scripts/core/program.yml index 88b652210b..52cfbf65d2 100644 --- a/scripts/core/program.yml +++ b/scripts/core/program.yml @@ -311,6 +311,43 @@ params: desc: | [out] Returns the pointer to the function if it is found in the program. --- #-------------------------------------------------------------------------- +type: function +desc: "Retrieves a pointer to a device global variable." +class: $xProgram +name: GetGlobalVariablePointer +decl: static +ordinal: "7" +analogue: + - "**clGetDeviceGlobalVariablePointerINTEL**" +details: + - "Retrieves a pointer to a device global variable." + - "The application may call this function from simultaneous threads for the same device." + - "The implementation of this function should be thread-safe." +params: + - type: "$x_device_handle_t" + name: hDevice + desc: | + [in] handle of the device to retrieve the pointer for. + - type: "$x_program_handle_t" + name: hProgram + desc: | + [in] handle of the program where the global variable is. + - type: "const char*" + name: pGlobalVariableName + desc: | + [in] mangled name of the global variable to retrieve the pointer for. + - type: "size_t*" + name: pGlobalVariableSizeRet + desc: | + [out][optional] Returns the size of the global variable if it is found in the program. + - type: "void**" + name: ppGlobalVariablePointerRet + desc: | + [out] Returns the pointer to the global variable if it is found in the program. +returns: + - $X_RESULT_ERROR_INVALID_SIZE: + - "`name` is not a valid variable in the program." +--- #-------------------------------------------------------------------------- type: enum desc: "Get Program object information" class: $xProgram diff --git a/scripts/core/registry.yml b/scripts/core/registry.yml index 6195cd4980..beb67354c6 100644 --- a/scripts/core/registry.yml +++ b/scripts/core/registry.yml @@ -559,6 +559,9 @@ etors: - name: KERNEL_SUGGEST_MAX_COOPERATIVE_GROUP_COUNT_EXP desc: Enumerator for $xKernelSuggestMaxCooperativeGroupCountExp value: '215' +- name: PROGRAM_GET_GLOBAL_VARIABLE_POINTER + desc: Enumerator for $xProgramGetGlobalVariablePointer + value: '216' --- type: enum desc: Defines structure types diff --git a/source/adapters/cuda/program.cpp b/source/adapters/cuda/program.cpp index 022fd258f7..f5a29abc43 100644 --- a/source/adapters/cuda/program.cpp +++ b/source/adapters/cuda/program.cpp @@ -489,3 +489,34 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramGetFunctionPointer( return Result; } + +UR_APIEXPORT ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( + ur_device_handle_t, ur_program_handle_t hProgram, + const char *pGlobalVariableName, size_t *pGlobalVariableSizeRet, + void **ppGlobalVariablePointerRet) { + + /* Since CUDA requires a global variable to be referenced by name, we use + * metadata to find the correct name to access it by. */ + auto DeviceGlobalNameIt = hProgram->GlobalIDMD.find(pGlobalVariableName); + if (DeviceGlobalNameIt == hProgram->GlobalIDMD.end()) + return UR_RESULT_ERROR_INVALID_VALUE; + std::string DeviceGlobalName = DeviceGlobalNameIt->second; + + ur_result_t Result = UR_RESULT_SUCCESS; + try { + CUdeviceptr DeviceGlobal = 0; + size_t DeviceGlobalSize = 0; + UR_CHECK_ERROR(cuModuleGetGlobal(&DeviceGlobal, &DeviceGlobalSize, + hProgram->get(), + DeviceGlobalName.c_str())); + + if (pGlobalVariableSizeRet) { + *pGlobalVariableSizeRet = DeviceGlobalSize; + } + *ppGlobalVariablePointerRet = reinterpret_cast(DeviceGlobal); + + } catch (ur_result_t Err) { + Result = Err; + } + return Result; +} diff --git a/source/adapters/cuda/ur_interface_loader.cpp b/source/adapters/cuda/ur_interface_loader.cpp index f31ffe6d87..856cab6bab 100644 --- a/source/adapters/cuda/ur_interface_loader.cpp +++ b/source/adapters/cuda/ur_interface_loader.cpp @@ -93,6 +93,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramProcAddrTable( pDdiTable->pfnCreateWithNativeHandle = urProgramCreateWithNativeHandle; pDdiTable->pfnGetBuildInfo = urProgramGetBuildInfo; pDdiTable->pfnGetFunctionPointer = urProgramGetFunctionPointer; + pDdiTable->pfnGetGlobalVariablePointer = urProgramGetGlobalVariablePointer; pDdiTable->pfnGetInfo = urProgramGetInfo; pDdiTable->pfnGetNativeHandle = urProgramGetNativeHandle; pDdiTable->pfnLink = urProgramLink; diff --git a/source/adapters/hip/program.cpp b/source/adapters/hip/program.cpp index 81f1be1194..23e7063e97 100644 --- a/source/adapters/hip/program.cpp +++ b/source/adapters/hip/program.cpp @@ -495,3 +495,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramGetFunctionPointer( return Result; } + +UR_APIEXPORT ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( + ur_device_handle_t, ur_program_handle_t, const char *, size_t *, void **) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} diff --git a/source/adapters/hip/ur_interface_loader.cpp b/source/adapters/hip/ur_interface_loader.cpp index 7707e78425..f5d3adb0cf 100644 --- a/source/adapters/hip/ur_interface_loader.cpp +++ b/source/adapters/hip/ur_interface_loader.cpp @@ -93,6 +93,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramProcAddrTable( pDdiTable->pfnCreateWithNativeHandle = urProgramCreateWithNativeHandle; pDdiTable->pfnGetBuildInfo = urProgramGetBuildInfo; pDdiTable->pfnGetFunctionPointer = urProgramGetFunctionPointer; + pDdiTable->pfnGetGlobalVariablePointer = urProgramGetGlobalVariablePointer; pDdiTable->pfnGetInfo = urProgramGetInfo; pDdiTable->pfnGetNativeHandle = urProgramGetNativeHandle; pDdiTable->pfnLink = urProgramLink; diff --git a/source/adapters/level_zero/program.cpp b/source/adapters/level_zero/program.cpp index f118a5b9dd..6c34b46d92 100644 --- a/source/adapters/level_zero/program.cpp +++ b/source/adapters/level_zero/program.cpp @@ -558,6 +558,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramGetFunctionPointer( return ze2urResult(ZeResult); } +UR_APIEXPORT ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( + ur_device_handle_t, ur_program_handle_t, const char *, size_t *, void **) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + UR_APIEXPORT ur_result_t UR_APICALL urProgramGetInfo( ur_program_handle_t Program, ///< [in] handle of the Program object ur_program_info_t PropName, ///< [in] name of the Program property to query diff --git a/source/adapters/level_zero/ur_interface_loader.cpp b/source/adapters/level_zero/ur_interface_loader.cpp index 74d0706b31..937416464b 100644 --- a/source/adapters/level_zero/ur_interface_loader.cpp +++ b/source/adapters/level_zero/ur_interface_loader.cpp @@ -215,6 +215,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramProcAddrTable( pDdiTable->pfnRetain = urProgramRetain; pDdiTable->pfnRelease = urProgramRelease; pDdiTable->pfnGetFunctionPointer = urProgramGetFunctionPointer; + pDdiTable->pfnGetGlobalVariablePointer = urProgramGetGlobalVariablePointer; pDdiTable->pfnGetInfo = urProgramGetInfo; pDdiTable->pfnGetBuildInfo = urProgramGetBuildInfo; pDdiTable->pfnSetSpecializationConstants = diff --git a/source/adapters/null/ur_nullddi.cpp b/source/adapters/null/ur_nullddi.cpp index f016830d11..b442c3c92d 100644 --- a/source/adapters/null/ur_nullddi.cpp +++ b/source/adapters/null/ur_nullddi.cpp @@ -1872,6 +1872,39 @@ __urdlllocal ur_result_t UR_APICALL urProgramGetFunctionPointer( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urProgramGetGlobalVariablePointer +__urdlllocal ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( + ur_device_handle_t + hDevice, ///< [in] handle of the device to retrieve the pointer for. + ur_program_handle_t + hProgram, ///< [in] handle of the program where the global variable is. + const char * + pGlobalVariableName, ///< [in] mangled name of the global variable to retrieve the pointer for. + size_t * + pGlobalVariableSizeRet, ///< [out][optional] Returns the size of the global variable if it is found + ///< in the program. + void ** + ppGlobalVariablePointerRet ///< [out] Returns the pointer to the global variable if it is found in the program. + ) try { + ur_result_t result = UR_RESULT_SUCCESS; + + // if the driver has created a custom function, then call it instead of using the generic path + auto pfnGetGlobalVariablePointer = + d_context.urDdiTable.Program.pfnGetGlobalVariablePointer; + if (nullptr != pfnGetGlobalVariablePointer) { + result = pfnGetGlobalVariablePointer( + hDevice, hProgram, pGlobalVariableName, pGlobalVariableSizeRet, + ppGlobalVariablePointerRet); + } else { + // generic implementation + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urProgramGetInfo __urdlllocal ur_result_t UR_APICALL urProgramGetInfo( @@ -5941,6 +5974,9 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramProcAddrTable( pDdiTable->pfnGetFunctionPointer = driver::urProgramGetFunctionPointer; + pDdiTable->pfnGetGlobalVariablePointer = + driver::urProgramGetGlobalVariablePointer; + pDdiTable->pfnGetInfo = driver::urProgramGetInfo; pDdiTable->pfnGetBuildInfo = driver::urProgramGetBuildInfo; diff --git a/source/adapters/opencl/common.hpp b/source/adapters/opencl/common.hpp index 0cb19694a6..55cdfea1cf 100644 --- a/source/adapters/opencl/common.hpp +++ b/source/adapters/opencl/common.hpp @@ -39,7 +39,7 @@ * error is mapped to UR */ #define CL_RETURN_ON_FAILURE_AND_SET_NULL(clCall, outPtr) \ - if (const cl_int cl_result_macro = clCall != CL_SUCCESS) { \ + if (const cl_int cl_result_macro = clCall; cl_result_macro != CL_SUCCESS) { \ if (outPtr != nullptr) { \ *outPtr = nullptr; \ } \ @@ -197,6 +197,8 @@ CONSTFIX char SetProgramSpecializationConstantName[] = "clSetProgramSpecializationConstant"; CONSTFIX char GetDeviceFunctionPointerName[] = "clGetDeviceFunctionPointerINTEL"; +CONSTFIX char GetDeviceGlobalVariablePointerName[] = + "clGetDeviceGlobalVariablePointerINTEL"; CONSTFIX char EnqueueWriteGlobalVariableName[] = "clEnqueueWriteGlobalVariableINTEL"; CONSTFIX char EnqueueReadGlobalVariableName[] = @@ -221,6 +223,10 @@ using clGetDeviceFunctionPointer_fn = CL_API_ENTRY cl_int(CL_API_CALL *)(cl_device_id device, cl_program program, const char *FuncName, cl_ulong *ret_ptr); +using clGetDeviceGlobalVariablePointer_fn = CL_API_ENTRY cl_int(CL_API_CALL *)( + cl_device_id device, cl_program program, const char *globalVariableName, + size_t *globalVariableSizeRet, void **globalVariablePointerRet); + using clEnqueueWriteGlobalVariable_fn = CL_API_ENTRY cl_int(CL_API_CALL *)(cl_command_queue, cl_program, const char *, cl_bool, size_t, size_t, const void *, cl_uint, const cl_event *, @@ -314,6 +320,8 @@ struct ExtFuncPtrCacheT { FuncPtrCache clDeviceMemAllocINTELCache; FuncPtrCache clSharedMemAllocINTELCache; FuncPtrCache clGetDeviceFunctionPointerCache; + FuncPtrCache + clGetDeviceGlobalVariablePointerCache; FuncPtrCache clCreateBufferWithPropertiesINTELCache; FuncPtrCache clMemBlockingFreeINTELCache; diff --git a/source/adapters/opencl/program.cpp b/source/adapters/opencl/program.cpp index f628c8152b..d973eba0f3 100644 --- a/source/adapters/opencl/program.cpp +++ b/source/adapters/opencl/program.cpp @@ -488,3 +488,42 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramGetFunctionPointer( return UR_RESULT_SUCCESS; } + +UR_APIEXPORT ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( + ur_device_handle_t hDevice, ur_program_handle_t hProgram, + const char *pGlobalVariableName, size_t *pGlobalVariableSizeRet, + void **ppGlobalVariablePointerRet) { + + cl_context CLContext = nullptr; + CL_RETURN_ON_FAILURE(clGetProgramInfo(cl_adapter::cast(hProgram), + CL_PROGRAM_CONTEXT, sizeof(CLContext), + &CLContext, nullptr)); + + cl_ext::clGetDeviceGlobalVariablePointer_fn FuncT = nullptr; + + UR_RETURN_ON_FAILURE(cl_ext::getExtFuncFromContext< + cl_ext::clGetDeviceGlobalVariablePointer_fn>( + CLContext, cl_ext::ExtFuncPtrCache->clGetDeviceGlobalVariablePointerCache, + cl_ext::GetDeviceGlobalVariablePointerName, &FuncT)); + + if (!FuncT) { + return UR_RESULT_ERROR_INVALID_FUNCTION_NAME; + } + + const cl_int CLResult = + FuncT(cl_adapter::cast(hDevice), + cl_adapter::cast(hProgram), pGlobalVariableName, + pGlobalVariableSizeRet, ppGlobalVariablePointerRet); + + if (CLResult != CL_SUCCESS) { + *ppGlobalVariablePointerRet = nullptr; + + if (CLResult == CL_INVALID_ARG_VALUE) { + return UR_RESULT_ERROR_INVALID_VALUE; + } + + CL_RETURN_ON_FAILURE(CLResult); + } + + return UR_RESULT_SUCCESS; +} diff --git a/source/adapters/opencl/ur_interface_loader.cpp b/source/adapters/opencl/ur_interface_loader.cpp index ac2c33475b..bea2cd9cf4 100644 --- a/source/adapters/opencl/ur_interface_loader.cpp +++ b/source/adapters/opencl/ur_interface_loader.cpp @@ -93,6 +93,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramProcAddrTable( pDdiTable->pfnCreateWithNativeHandle = urProgramCreateWithNativeHandle; pDdiTable->pfnGetBuildInfo = urProgramGetBuildInfo; pDdiTable->pfnGetFunctionPointer = urProgramGetFunctionPointer; + pDdiTable->pfnGetGlobalVariablePointer = urProgramGetGlobalVariablePointer; pDdiTable->pfnGetInfo = urProgramGetInfo; pDdiTable->pfnGetNativeHandle = urProgramGetNativeHandle; pDdiTable->pfnLink = urProgramLink; diff --git a/source/loader/layers/tracing/ur_trcddi.cpp b/source/loader/layers/tracing/ur_trcddi.cpp index 402b64d638..b4bbe2338d 100644 --- a/source/loader/layers/tracing/ur_trcddi.cpp +++ b/source/loader/layers/tracing/ur_trcddi.cpp @@ -2124,6 +2124,46 @@ __urdlllocal ur_result_t UR_APICALL urProgramGetFunctionPointer( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urProgramGetGlobalVariablePointer +__urdlllocal ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( + ur_device_handle_t + hDevice, ///< [in] handle of the device to retrieve the pointer for. + ur_program_handle_t + hProgram, ///< [in] handle of the program where the global variable is. + const char * + pGlobalVariableName, ///< [in] mangled name of the global variable to retrieve the pointer for. + size_t * + pGlobalVariableSizeRet, ///< [out][optional] Returns the size of the global variable if it is found + ///< in the program. + void ** + ppGlobalVariablePointerRet ///< [out] Returns the pointer to the global variable if it is found in the program. +) { + auto pfnGetGlobalVariablePointer = + context.urDdiTable.Program.pfnGetGlobalVariablePointer; + + if (nullptr == pfnGetGlobalVariablePointer) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + + ur_program_get_global_variable_pointer_params_t params = { + &hDevice, &hProgram, &pGlobalVariableName, &pGlobalVariableSizeRet, + &ppGlobalVariablePointerRet}; + uint64_t instance = + context.notify_begin(UR_FUNCTION_PROGRAM_GET_GLOBAL_VARIABLE_POINTER, + "urProgramGetGlobalVariablePointer", ¶ms); + + ur_result_t result = pfnGetGlobalVariablePointer( + hDevice, hProgram, pGlobalVariableName, pGlobalVariableSizeRet, + ppGlobalVariablePointerRet); + + context.notify_end(UR_FUNCTION_PROGRAM_GET_GLOBAL_VARIABLE_POINTER, + "urProgramGetGlobalVariablePointer", ¶ms, &result, + instance); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urProgramGetInfo __urdlllocal ur_result_t UR_APICALL urProgramGetInfo( @@ -6965,6 +7005,11 @@ __urdlllocal ur_result_t UR_APICALL urGetProgramProcAddrTable( pDdiTable->pfnGetFunctionPointer = ur_tracing_layer::urProgramGetFunctionPointer; + dditable.pfnGetGlobalVariablePointer = + pDdiTable->pfnGetGlobalVariablePointer; + pDdiTable->pfnGetGlobalVariablePointer = + ur_tracing_layer::urProgramGetGlobalVariablePointer; + dditable.pfnGetInfo = pDdiTable->pfnGetInfo; pDdiTable->pfnGetInfo = ur_tracing_layer::urProgramGetInfo; diff --git a/source/loader/layers/validation/ur_valddi.cpp b/source/loader/layers/validation/ur_valddi.cpp index 72e225028c..0bbe2e3ded 100644 --- a/source/loader/layers/validation/ur_valddi.cpp +++ b/source/loader/layers/validation/ur_valddi.cpp @@ -2645,6 +2645,53 @@ __urdlllocal ur_result_t UR_APICALL urProgramGetFunctionPointer( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urProgramGetGlobalVariablePointer +__urdlllocal ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( + ur_device_handle_t + hDevice, ///< [in] handle of the device to retrieve the pointer for. + ur_program_handle_t + hProgram, ///< [in] handle of the program where the global variable is. + const char * + pGlobalVariableName, ///< [in] mangled name of the global variable to retrieve the pointer for. + size_t * + pGlobalVariableSizeRet, ///< [out][optional] Returns the size of the global variable if it is found + ///< in the program. + void ** + ppGlobalVariablePointerRet ///< [out] Returns the pointer to the global variable if it is found in the program. +) { + auto pfnGetGlobalVariablePointer = + context.urDdiTable.Program.pfnGetGlobalVariablePointer; + + if (nullptr == pfnGetGlobalVariablePointer) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (context.enableParameterValidation) { + if (NULL == hDevice) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (NULL == hProgram) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (NULL == pGlobalVariableName) { + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + } + + if (NULL == ppGlobalVariablePointerRet) { + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + } + } + + ur_result_t result = pfnGetGlobalVariablePointer( + hDevice, hProgram, pGlobalVariableName, pGlobalVariableSizeRet, + ppGlobalVariablePointerRet); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urProgramGetInfo __urdlllocal ur_result_t UR_APICALL urProgramGetInfo( @@ -8851,6 +8898,11 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramProcAddrTable( pDdiTable->pfnGetFunctionPointer = ur_validation_layer::urProgramGetFunctionPointer; + dditable.pfnGetGlobalVariablePointer = + pDdiTable->pfnGetGlobalVariablePointer; + pDdiTable->pfnGetGlobalVariablePointer = + ur_validation_layer::urProgramGetGlobalVariablePointer; + dditable.pfnGetInfo = pDdiTable->pfnGetInfo; pDdiTable->pfnGetInfo = ur_validation_layer::urProgramGetInfo; diff --git a/source/loader/ur_ldrddi.cpp b/source/loader/ur_ldrddi.cpp index 6d3dda30f0..f824644e2c 100644 --- a/source/loader/ur_ldrddi.cpp +++ b/source/loader/ur_ldrddi.cpp @@ -2415,6 +2415,45 @@ __urdlllocal ur_result_t UR_APICALL urProgramGetFunctionPointer( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urProgramGetGlobalVariablePointer +__urdlllocal ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( + ur_device_handle_t + hDevice, ///< [in] handle of the device to retrieve the pointer for. + ur_program_handle_t + hProgram, ///< [in] handle of the program where the global variable is. + const char * + pGlobalVariableName, ///< [in] mangled name of the global variable to retrieve the pointer for. + size_t * + pGlobalVariableSizeRet, ///< [out][optional] Returns the size of the global variable if it is found + ///< in the program. + void ** + ppGlobalVariablePointerRet ///< [out] Returns the pointer to the global variable if it is found in the program. +) { + ur_result_t result = UR_RESULT_SUCCESS; + + // extract platform's function pointer table + auto dditable = reinterpret_cast(hDevice)->dditable; + auto pfnGetGlobalVariablePointer = + dditable->ur.Program.pfnGetGlobalVariablePointer; + if (nullptr == pfnGetGlobalVariablePointer) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + // convert loader handle to platform handle + hDevice = reinterpret_cast(hDevice)->handle; + + // convert loader handle to platform handle + hProgram = reinterpret_cast(hProgram)->handle; + + // forward to device-platform + result = pfnGetGlobalVariablePointer(hDevice, hProgram, pGlobalVariableName, + pGlobalVariableSizeRet, + ppGlobalVariablePointerRet); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urProgramGetInfo __urdlllocal ur_result_t UR_APICALL urProgramGetInfo( @@ -8083,6 +8122,8 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramProcAddrTable( pDdiTable->pfnRelease = ur_loader::urProgramRelease; pDdiTable->pfnGetFunctionPointer = ur_loader::urProgramGetFunctionPointer; + pDdiTable->pfnGetGlobalVariablePointer = + ur_loader::urProgramGetGlobalVariablePointer; pDdiTable->pfnGetInfo = ur_loader::urProgramGetInfo; pDdiTable->pfnGetBuildInfo = ur_loader::urProgramGetBuildInfo; pDdiTable->pfnSetSpecializationConstants = diff --git a/source/loader/ur_libapi.cpp b/source/loader/ur_libapi.cpp index 0a69fcd1e2..27d0afadb6 100644 --- a/source/loader/ur_libapi.cpp +++ b/source/loader/ur_libapi.cpp @@ -3183,6 +3183,58 @@ ur_result_t UR_APICALL urProgramGetFunctionPointer( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Retrieves a pointer to a device global variable. +/// +/// @details +/// - Retrieves a pointer to a device global variable. +/// - The application may call this function from simultaneous threads for +/// the same device. +/// - The implementation of this function should be thread-safe. +/// +/// @remarks +/// _Analogues_ +/// - **clGetDeviceGlobalVariablePointerINTEL** +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hDevice` +/// + `NULL == hProgram` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pGlobalVariableName` +/// + `NULL == ppGlobalVariablePointerRet` +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// + `name` is not a valid variable in the program. +ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( + ur_device_handle_t + hDevice, ///< [in] handle of the device to retrieve the pointer for. + ur_program_handle_t + hProgram, ///< [in] handle of the program where the global variable is. + const char * + pGlobalVariableName, ///< [in] mangled name of the global variable to retrieve the pointer for. + size_t * + pGlobalVariableSizeRet, ///< [out][optional] Returns the size of the global variable if it is found + ///< in the program. + void ** + ppGlobalVariablePointerRet ///< [out] Returns the pointer to the global variable if it is found in the program. + ) try { + auto pfnGetGlobalVariablePointer = + ur_lib::context->urDdiTable.Program.pfnGetGlobalVariablePointer; + if (nullptr == pfnGetGlobalVariablePointer) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + return pfnGetGlobalVariablePointer(hDevice, hProgram, pGlobalVariableName, + pGlobalVariableSizeRet, + ppGlobalVariablePointerRet); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Query information about a Program object /// diff --git a/source/ur_api.cpp b/source/ur_api.cpp index 2bcc229f29..8fa77e69a0 100644 --- a/source/ur_api.cpp +++ b/source/ur_api.cpp @@ -2702,6 +2702,49 @@ ur_result_t UR_APICALL urProgramGetFunctionPointer( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Retrieves a pointer to a device global variable. +/// +/// @details +/// - Retrieves a pointer to a device global variable. +/// - The application may call this function from simultaneous threads for +/// the same device. +/// - The implementation of this function should be thread-safe. +/// +/// @remarks +/// _Analogues_ +/// - **clGetDeviceGlobalVariablePointerINTEL** +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hDevice` +/// + `NULL == hProgram` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pGlobalVariableName` +/// + `NULL == ppGlobalVariablePointerRet` +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// + `name` is not a valid variable in the program. +ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( + ur_device_handle_t + hDevice, ///< [in] handle of the device to retrieve the pointer for. + ur_program_handle_t + hProgram, ///< [in] handle of the program where the global variable is. + const char * + pGlobalVariableName, ///< [in] mangled name of the global variable to retrieve the pointer for. + size_t * + pGlobalVariableSizeRet, ///< [out][optional] Returns the size of the global variable if it is found + ///< in the program. + void ** + ppGlobalVariablePointerRet ///< [out] Returns the pointer to the global variable if it is found in the program. +) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Query information about a Program object /// diff --git a/test/conformance/device_code/device_global.cpp b/test/conformance/device_code/device_global.cpp index 3ead86bbf5..37c08419ce 100644 --- a/test/conformance/device_code/device_global.cpp +++ b/test/conformance/device_code/device_global.cpp @@ -5,7 +5,10 @@ #include -sycl::ext::oneapi::experimental::device_global dev_var; +sycl::ext::oneapi::experimental::device_global< + int, decltype(sycl::ext::oneapi::experimental::properties( + sycl::ext::oneapi::experimental::device_image_scope))> + dev_var; int main() { diff --git a/test/conformance/enqueue/enqueue_adapter_cuda.match b/test/conformance/enqueue/enqueue_adapter_cuda.match index 2392247314..8fe2045d2c 100644 --- a/test/conformance/enqueue/enqueue_adapter_cuda.match +++ b/test/conformance/enqueue/enqueue_adapter_cuda.match @@ -1,6 +1,3 @@ -{{OPT}}urEnqueueDeviceGetGlobalVariableReadTest.Success/NVIDIA_CUDA_BACKEND___{{.*}}_ -{{OPT}}urEnqueueDeviceGetGlobalVariableReadTest.InvalidEventWaitInvalidEvent/NVIDIA_CUDA_BACKEND___{{.*}}_ -{{OPT}}urEnqueueDeviceGetGlobalVariableWriteTest.InvalidEventWaitInvalidEvent/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urEnqueueKernelLaunchWithVirtualMemory.Success/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urEnqueueMemBufferCopyRectTest.InvalidSize/NVIDIA_CUDA_BACKEND___{{.*}}_ {{OPT}}urEnqueueMemBufferFillTest.Success/NVIDIA_CUDA_BACKEND___{{.*}}___size__256__patternSize__256 diff --git a/test/conformance/enqueue/enqueue_adapter_opencl.match b/test/conformance/enqueue/enqueue_adapter_opencl.match index a034083c87..903d105b08 100644 --- a/test/conformance/enqueue/enqueue_adapter_opencl.match +++ b/test/conformance/enqueue/enqueue_adapter_opencl.match @@ -1,4 +1,3 @@ -{{OPT}}urEnqueueDeviceGetGlobalVariableReadTest.Success/Intel_R__OpenCL___{{.*}} {{OPT}}urEnqueueMemBufferCopyRectTest.InvalidSize/Intel_R__OpenCL___{{.*}} {{OPT}}urEnqueueMemBufferReadRectTest.InvalidSize/Intel_R__OpenCL___{{.*}} {{OPT}}urEnqueueMemBufferWriteRectTest.InvalidSize/Intel_R__OpenCL___{{.*}} diff --git a/test/conformance/enqueue/urEnqueueDeviceGlobalVariableRead.cpp b/test/conformance/enqueue/urEnqueueDeviceGlobalVariableRead.cpp index df1e70b2b3..799ce7a67d 100644 --- a/test/conformance/enqueue/urEnqueueDeviceGlobalVariableRead.cpp +++ b/test/conformance/enqueue/urEnqueueDeviceGlobalVariableRead.cpp @@ -20,7 +20,7 @@ TEST_P(urEnqueueDeviceGetGlobalVariableReadTest, Success) { // execute the kernel ASSERT_SUCCESS(urEnqueueKernelLaunch(queue, kernel, n_dimensions, &global_offset, &global_size, nullptr, - 1, nullptr, nullptr)); + 0, nullptr, nullptr)); ASSERT_SUCCESS(urQueueFinish(queue)); // read global var back to host diff --git a/test/conformance/program/CMakeLists.txt b/test/conformance/program/CMakeLists.txt index 8f58488c73..317e3df946 100644 --- a/test/conformance/program/CMakeLists.txt +++ b/test/conformance/program/CMakeLists.txt @@ -11,6 +11,7 @@ add_conformance_test_with_kernels_environment(program urProgramCreateWithNativeHandle.cpp urProgramGetBuildInfo.cpp urProgramGetFunctionPointer.cpp + urProgramGetGlobalVariablePointer.cpp urProgramGetInfo.cpp urProgramGetNativeHandle.cpp urProgramLink.cpp diff --git a/test/conformance/program/program_adapter_hip.match b/test/conformance/program/program_adapter_hip.match index 1f95931e09..a0414f4dee 100644 --- a/test/conformance/program/program_adapter_hip.match +++ b/test/conformance/program/program_adapter_hip.match @@ -23,3 +23,5 @@ {{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_KERNEL_NAMES {{OPT}}urProgramLinkTest.Success/AMD_HIP_BACKEND___{{.*}}_ {{OPT}}urProgramSetSpecializationConstantsTest.Success/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urProgramGetGlobalVariablePointerTest.Success/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urProgramGetGlobalVariablePointerTest.InvalidVariableName/AMD_HIP_BACKEND___{{.*}}_ diff --git a/test/conformance/program/program_adapter_level_zero.match b/test/conformance/program/program_adapter_level_zero.match index 5bbdfd554c..1598cf1224 100644 --- a/test/conformance/program/program_adapter_level_zero.match +++ b/test/conformance/program/program_adapter_level_zero.match @@ -3,4 +3,6 @@ urProgramCreateWithNativeHandleTest.InvalidNullHandleContext/Intel_R__oneAPI_Uni urProgramCreateWithNativeHandleTest.InvalidNullPointerProgram/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ urProgramGetBuildInfoTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_UR_PROGRAM_BUILD_INFO_STATUS urProgramGetFunctionPointerTest.InvalidFunctionName/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urProgramGetGlobalVariablePointerTest.Success/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ +urProgramGetGlobalVariablePointerTest.InvalidVariableName/Intel_R__oneAPI_Unified_Runtime_over_Level_Zero___{{.*}}_ Aborted diff --git a/test/conformance/program/program_adapter_opencl.match b/test/conformance/program/program_adapter_opencl.match index 0d429016ee..98e1940c19 100644 --- a/test/conformance/program/program_adapter_opencl.match +++ b/test/conformance/program/program_adapter_opencl.match @@ -1,3 +1,3 @@ -urProgramGetFunctionPointerTest.InvalidFunctionName/Intel_R__OpenCL___{{.*}}_ +urProgramGetFunctionPointerTest.InvalidFunctionName/Intel_R__OpenCL___{{.*}} urProgramGetInfoTest.Success/Intel_R__OpenCL___{{.*}}___UR_PROGRAM_INFO_SOURCE urProgramGetInfoTest.Success/Intel_R__OpenCL___{{.*}}___UR_PROGRAM_INFO_BINARIES diff --git a/test/conformance/program/urProgramGetGlobalVariablePointer.cpp b/test/conformance/program/urProgramGetGlobalVariablePointer.cpp new file mode 100644 index 0000000000..d1f66e7a75 --- /dev/null +++ b/test/conformance/program/urProgramGetGlobalVariablePointer.cpp @@ -0,0 +1,62 @@ +// Copyright (C) 2024 Intel Corporation +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions. +// See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include + +using urProgramGetGlobalVariablePointerTest = uur::urGlobalVariableTest; + +UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urProgramGetGlobalVariablePointerTest); + +TEST_P(urProgramGetGlobalVariablePointerTest, Success) { + + size_t global_variable_size = 0; + void *global_variable_pointer; + ASSERT_SUCCESS(urProgramGetGlobalVariablePointer( + device, program, global_var.name.c_str(), &global_variable_size, + &global_variable_pointer)); + ASSERT_GT(global_variable_size, 0); + ASSERT_NE(global_variable_pointer, nullptr); +} + +TEST_P(urProgramGetGlobalVariablePointerTest, InvalidNullHandleDevice) { + void *global_variable_pointer; + ASSERT_EQ_RESULT(urProgramGetGlobalVariablePointer( + nullptr, program, global_var.name.c_str(), nullptr, + &global_variable_pointer), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); +} + +TEST_P(urProgramGetGlobalVariablePointerTest, InvalidNullHandleProgram) { + void *global_variable_pointer; + ASSERT_EQ_RESULT(urProgramGetGlobalVariablePointer( + device, nullptr, global_var.name.c_str(), nullptr, + &global_variable_pointer), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); +} + +TEST_P(urProgramGetGlobalVariablePointerTest, InvalidVariableName) { + void *global_variable_pointer; + ASSERT_EQ_RESULT( + urProgramGetGlobalVariablePointer(device, program, "foo", nullptr, + &global_variable_pointer), + UR_RESULT_ERROR_INVALID_VALUE); +} + +TEST_P(urProgramGetGlobalVariablePointerTest, InvalidNullPointerVariableName) { + void *global_variable_pointer; + ASSERT_EQ_RESULT( + urProgramGetGlobalVariablePointer(device, program, nullptr, nullptr, + &global_variable_pointer), + UR_RESULT_ERROR_INVALID_NULL_POINTER); +} + +TEST_P(urProgramGetGlobalVariablePointerTest, + InvalidNullPointerVariablePointer) { + size_t global_variable_size = 0; + ASSERT_EQ_RESULT(urProgramGetGlobalVariablePointer( + device, program, global_var.name.c_str(), + &global_variable_size, nullptr), + UR_RESULT_ERROR_INVALID_NULL_POINTER); +} diff --git a/test/conformance/testing/include/uur/fixtures.h b/test/conformance/testing/include/uur/fixtures.h index 2ede84d135..ada9bdba6b 100644 --- a/test/conformance/testing/include/uur/fixtures.h +++ b/test/conformance/testing/include/uur/fixtures.h @@ -1052,8 +1052,14 @@ struct urProgramTest : urQueueTest { UUR_RETURN_ON_FATAL_FAILURE(urQueueTest::SetUp()); uur::KernelsEnvironment::instance->LoadSource(program_name, 0, il_binary); - ASSERT_SUCCESS(urProgramCreateWithIL( - context, il_binary->data(), il_binary->size(), nullptr, &program)); + + const ur_program_properties_t properties = { + UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES, nullptr, + static_cast(metadatas.size()), + metadatas.empty() ? nullptr : metadatas.data()}; + ASSERT_SUCCESS(urProgramCreateWithIL(context, il_binary->data(), + il_binary->size(), &properties, + &program)); } void TearDown() override { @@ -1066,6 +1072,7 @@ struct urProgramTest : urQueueTest { std::shared_ptr> il_binary; std::string program_name = "foo"; ur_program_handle_t program = nullptr; + std::vector metadatas{}; }; template struct urProgramTestWithParam : urContextTestWithParam { @@ -1266,13 +1273,27 @@ template struct GlobalVar { T value; }; +using namespace std::string_literals; struct urGlobalVariableTest : uur::urKernelExecutionTest { void SetUp() override { + program_name = "device_global"; global_var = {"_Z7dev_var", 0}; + + /* Some adapters cannot use the mangled variable name directly. + * Instead, in order to map the mangled variable to the internal name, + * they rely on metadata set when creating the program */ + const std::string metadata_name = "_Z7dev_var@global_id_mapping"; + ur_program_metadata_value_t metadata_value; + metadata_value.pData = (void *)metadataData.c_str(); + metadatas.push_back({metadata_name.c_str(), + UR_PROGRAM_METADATA_TYPE_BYTE_ARRAY, + metadataData.size(), metadata_value}); UUR_RETURN_ON_FATAL_FAILURE(uur::urKernelExecutionTest::SetUp()); } + /* We pad the first 8 bytes of the metadata since they are ignored */ + std::string metadataData = "\0\0\0\0\0\0\0\0dev_var"s; GlobalVar global_var; };