diff --git a/source/adapters/cuda/common.cpp b/source/adapters/cuda/common.cpp index 9f2a330262..f24df17beb 100644 --- a/source/adapters/cuda/common.cpp +++ b/source/adapters/cuda/common.cpp @@ -100,16 +100,6 @@ std::string getCudaVersionString() { return stream.str(); } -void detail::ur::die(const char *Message) { - std::cerr << "ur_die: " << Message << std::endl; - std::terminate(); -} - -void detail::ur::assertion(bool Condition, const char *Message) { - if (!Condition) - die(Message); -} - void detail::ur::cuPrint(const char *Message) { std::cerr << "ur_print: " << Message << std::endl; } diff --git a/source/adapters/cuda/common.hpp b/source/adapters/cuda/common.hpp index 67223c45bc..da77c6eeb9 100644 --- a/source/adapters/cuda/common.hpp +++ b/source/adapters/cuda/common.hpp @@ -12,6 +12,16 @@ #include #include +/** + * Call a UR API and, if the result is not UR_RESULT_SUCCESS, automatically + * return from the current function. + */ +#define UR_RETURN_ON_FAILURE(urCall) \ + if (const ur_result_t ur_result_macro = urCall; \ + ur_result_macro != UR_RESULT_SUCCESS) { \ + return ur_result_macro; \ + } + ur_result_t mapErrorUR(CUresult Result); /// Converts CUDA error into UR error codes, and outputs error information @@ -46,16 +56,8 @@ void setPluginSpecificMessage(CUresult cu_res); namespace detail { namespace ur { -// Report error and no return (keeps compiler from printing warnings). -// TODO: Probably change that to throw a catchable exception, -// but for now it is useful to see every failure. -// -[[noreturn]] void die(const char *Message); - // Reports error messages void cuPrint(const char *Message); -void assertion(bool Condition, const char *Message = nullptr); - } // namespace ur } // namespace detail diff --git a/source/adapters/cuda/device.cpp b/source/adapters/cuda/device.cpp index 69665fb456..b14d60ae59 100644 --- a/source/adapters/cuda/device.cpp +++ b/source/adapters/cuda/device.cpp @@ -57,10 +57,14 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_MAX_COMPUTE_UNITS: { int ComputeUnits = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &ComputeUnits, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, - hDevice->get())); - detail::ur::assertion(ComputeUnits >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &ComputeUnits, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(ComputeUnits >= 0); return ReturnValue(static_cast(ComputeUnits)); } case UR_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS: { @@ -72,17 +76,29 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } ReturnSizes; int MaxX = 0, MaxY = 0, MaxZ = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &MaxX, CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X, hDevice->get())); - detail::ur::assertion(MaxX >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MaxX, CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxX > 0); - UR_CHECK_ERROR(cuDeviceGetAttribute( - &MaxY, CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y, hDevice->get())); - detail::ur::assertion(MaxY >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MaxY, CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxY > 0); - UR_CHECK_ERROR(cuDeviceGetAttribute( - &MaxZ, CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z, hDevice->get())); - detail::ur::assertion(MaxZ >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MaxZ, CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxZ > 0); ReturnSizes.Sizes[0] = size_t(MaxX); ReturnSizes.Sizes[1] = size_t(MaxY); @@ -95,17 +111,29 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, size_t Sizes[MaxWorkItemDimensions]; } ReturnSizes; int MaxX = 0, MaxY = 0, MaxZ = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &MaxX, CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X, hDevice->get())); - detail::ur::assertion(MaxX >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MaxX, CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxX > 0); - UR_CHECK_ERROR(cuDeviceGetAttribute( - &MaxY, CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y, hDevice->get())); - detail::ur::assertion(MaxY >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MaxY, CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxY > 0); - UR_CHECK_ERROR(cuDeviceGetAttribute( - &MaxZ, CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z, hDevice->get())); - detail::ur::assertion(MaxZ >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MaxZ, CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxZ > 0); ReturnSizes.Sizes[0] = size_t(MaxX); ReturnSizes.Sizes[1] = size_t(MaxY); @@ -115,11 +143,14 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_MAX_WORK_GROUP_SIZE: { int MaxWorkGroupSize = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &MaxWorkGroupSize, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, - hDevice->get())); - - detail::ur::assertion(MaxWorkGroupSize >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MaxWorkGroupSize, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxWorkGroupSize > 0); return ReturnValue(size_t(MaxWorkGroupSize)); } @@ -168,12 +199,20 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_MAX_NUM_SUB_GROUPS: { // Number of sub-groups = max block size / warp size + possible remainder int MaxThreads = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &MaxThreads, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, - hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MaxThreads, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } int WarpSize = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } int MaxWarps = (MaxThreads + WarpSize - 1) / WarpSize; return ReturnValue(MaxWarps); } @@ -181,17 +220,26 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, // Volta provides independent thread scheduling // TODO: Revisit for previous generation GPUs int Major = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } bool IFP = (Major >= 7); return ReturnValue(IFP); } case UR_DEVICE_INFO_ATOMIC_64: { int Major = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hDevice->get())); - + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } bool Atomic64 = (Major >= 6) ? true : false; return ReturnValue(Atomic64); } @@ -212,8 +260,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES: { int Major = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } uint64_t Capabilities = (Major >= 7) ? UR_MEMORY_SCOPE_CAPABILITY_FLAG_WORK_ITEM | UR_MEMORY_SCOPE_CAPABILITY_FLAG_SUB_GROUP | @@ -258,25 +311,37 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_BFLOAT16: { int Major = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hDevice->get())); - + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } bool BFloat16 = (Major >= 8) ? true : false; return ReturnValue(BFloat16); } case UR_DEVICE_INFO_SUB_GROUP_SIZES_INTEL: { // NVIDIA devices only support one sub-group size (the warp size) int WarpSize = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } size_t Sizes[1] = {static_cast(WarpSize)}; return ReturnValue(Sizes, 1); } case UR_DEVICE_INFO_MAX_CLOCK_FREQUENCY: { int ClockFreq = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &ClockFreq, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, hDevice->get())); - detail::ur::assertion(ClockFreq >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &ClockFreq, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(ClockFreq > 0); return ReturnValue(static_cast(ClockFreq) / 1000u); } case UR_DEVICE_INFO_ADDRESS_BITS: { @@ -317,15 +382,23 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE2D_MAX_HEIGHT: { // Take the smaller of maximum surface and maximum texture height. int TexHeight = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &TexHeight, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT, - hDevice->get())); - detail::ur::assertion(TexHeight >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &TexHeight, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexHeight > 0); int SurfHeight = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &SurfHeight, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_HEIGHT, - hDevice->get())); - detail::ur::assertion(SurfHeight >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &SurfHeight, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_HEIGHT, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfHeight > 0); int Min = std::min(TexHeight, SurfHeight); @@ -334,15 +407,23 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE2D_MAX_WIDTH: { // Take the smaller of maximum surface and maximum texture width. int TexWidth = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &TexWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH, - hDevice->get())); - detail::ur::assertion(TexWidth >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &TexWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexWidth > 0); int SurfWidth = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &SurfWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_WIDTH, - hDevice->get())); - detail::ur::assertion(SurfWidth >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &SurfWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_WIDTH, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfWidth > 0); int Min = std::min(TexWidth, SurfWidth); @@ -351,15 +432,23 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE3D_MAX_HEIGHT: { // Take the smaller of maximum surface and maximum texture height. int TexHeight = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &TexHeight, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT, - hDevice->get())); - detail::ur::assertion(TexHeight >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &TexHeight, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexHeight > 0); int SurfHeight = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &SurfHeight, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_HEIGHT, - hDevice->get())); - detail::ur::assertion(SurfHeight >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &SurfHeight, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_HEIGHT, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfHeight > 0); int Min = std::min(TexHeight, SurfHeight); @@ -368,15 +457,23 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE3D_MAX_WIDTH: { // Take the smaller of maximum surface and maximum texture width. int TexWidth = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &TexWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH, - hDevice->get())); - detail::ur::assertion(TexWidth >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &TexWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexWidth > 0); int SurfWidth = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &SurfWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_WIDTH, - hDevice->get())); - detail::ur::assertion(SurfWidth >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &SurfWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_WIDTH, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfWidth > 0); int Min = std::min(TexWidth, SurfWidth); @@ -385,15 +482,23 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE3D_MAX_DEPTH: { // Take the smaller of maximum surface and maximum texture depth. int TexDepth = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &TexDepth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH, - hDevice->get())); - detail::ur::assertion(TexDepth >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &TexDepth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexDepth > 0); int SurfDepth = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &SurfDepth, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_DEPTH, - hDevice->get())); - detail::ur::assertion(SurfDepth >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &SurfDepth, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_DEPTH, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfDepth > 0); int Min = std::min(TexDepth, SurfDepth); @@ -402,15 +507,23 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE: { // Take the smaller of maximum surface and maximum texture width. int TexWidth = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &TexWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH, - hDevice->get())); - detail::ur::assertion(TexWidth >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &TexWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexWidth > 0); int SurfWidth = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &SurfWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_WIDTH, - hDevice->get())); - detail::ur::assertion(SurfWidth >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &SurfWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_WIDTH, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfWidth > 0); int Min = std::min(TexWidth, SurfWidth); @@ -432,9 +545,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_MEM_BASE_ADDR_ALIGN: { int MemBaseAddrAlign = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute(&MemBaseAddrAlign, - CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT, - hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute(&MemBaseAddrAlign, + CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } // Multiply by 8 as clGetDeviceInfo returns this value in bits MemBaseAddrAlign *= 8; return ReturnValue(MemBaseAddrAlign); @@ -477,25 +594,36 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_GLOBAL_MEM_CACHE_SIZE: { int CacheSize = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &CacheSize, CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE, hDevice->get())); - detail::ur::assertion(CacheSize >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &CacheSize, CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(CacheSize > 0); // The L2 cache is global to the GPU. return ReturnValue(static_cast(CacheSize)); } case UR_DEVICE_INFO_GLOBAL_MEM_SIZE: { size_t Bytes = 0; // Runtime API has easy access to this value, driver API info is scarse. - detail::ur::assertion(cuDeviceTotalMem(&Bytes, hDevice->get()) == - CUDA_SUCCESS); + try { + UR_CHECK_ERROR(cuDeviceTotalMem(&Bytes, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(uint64_t{Bytes}); } case UR_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE: { int ConstantMemory = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &ConstantMemory, CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY, - hDevice->get())); - detail::ur::assertion(ConstantMemory >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &ConstantMemory, CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(ConstantMemory > 0); return ReturnValue(static_cast(ConstantMemory)); } @@ -522,19 +650,27 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_ERROR_CORRECTION_SUPPORT: { int ECCEnabled = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &ECCEnabled, CU_DEVICE_ATTRIBUTE_ECC_ENABLED, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &ECCEnabled, CU_DEVICE_ATTRIBUTE_ECC_ENABLED, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(ECCEnabled == 0 || ECCEnabled == 1); - detail::ur::assertion((ECCEnabled == 0) | (ECCEnabled == 1)); auto Result = static_cast(ECCEnabled); return ReturnValue(Result); } case UR_DEVICE_INFO_HOST_UNIFIED_MEMORY: { int IsIntegrated = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &IsIntegrated, CU_DEVICE_ATTRIBUTE_INTEGRATED, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &IsIntegrated, CU_DEVICE_ATTRIBUTE_INTEGRATED, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(IsIntegrated == 0 || IsIntegrated == 1); - detail::ur::assertion((IsIntegrated == 0) | (IsIntegrated == 1)); auto result = static_cast(IsIntegrated); return ReturnValue(result); } @@ -589,7 +725,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_NAME: { static constexpr size_t MaxDeviceNameLength = 256u; char Name[MaxDeviceNameLength]; - UR_CHECK_ERROR(cuDeviceGetName(Name, MaxDeviceNameLength, hDevice->get())); + try { + UR_CHECK_ERROR( + cuDeviceGetName(Name, MaxDeviceNameLength, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(Name, strlen(Name) + 1); } case UR_DEVICE_INFO_VENDOR: { @@ -608,12 +749,22 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_VERSION: { std::stringstream SS; int Major; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } SS << Major; int Minor; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } SS << "." << Minor; return ReturnValue(SS.str().c_str()); } @@ -632,10 +783,21 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, int Major = 0; int Minor = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hDevice->get())); - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } if ((Major >= 6) || ((Major == 5) && (Minor >= 3))) { SupportedExtensions += "cl_khr_fp16 "; @@ -816,24 +978,34 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_GLOBAL_MEM_FREE: { size_t FreeMemory = 0; size_t TotalMemory = 0; - detail::ur::assertion(cuMemGetInfo(&FreeMemory, &TotalMemory) == - CUDA_SUCCESS, - "failed cuMemGetInfo() API."); + try { + UR_CHECK_ERROR(cuMemGetInfo(&FreeMemory, &TotalMemory)); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(FreeMemory); } case UR_DEVICE_INFO_MEMORY_CLOCK_RATE: { int Value = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Value, CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE, hDevice->get())); - detail::ur::assertion(Value >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Value, CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(Value > 0); // Convert kilohertz to megahertz when returning. return ReturnValue(Value / 1000); } case UR_DEVICE_INFO_MEMORY_BUS_WIDTH: { int Value = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Value, CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH, hDevice->get())); - detail::ur::assertion(Value >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Value, CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(Value > 0); return ReturnValue(Value); } case UR_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES: { @@ -858,30 +1030,46 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_IMAGE_PITCH_ALIGN_EXP: { int32_t tex_pitch_align = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &tex_pitch_align, CU_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT, - hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &tex_pitch_align, CU_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(tex_pitch_align); } case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_WIDTH_EXP: { int32_t tex_max_linear_width = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &tex_max_linear_width, - CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_WIDTH, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &tex_max_linear_width, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_WIDTH, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(tex_max_linear_width); } case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_HEIGHT_EXP: { int32_t tex_max_linear_height = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &tex_max_linear_height, - CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_HEIGHT, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &tex_max_linear_height, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_HEIGHT, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(tex_max_linear_height); } case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_PITCH_EXP: { int32_t tex_max_linear_pitch = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &tex_max_linear_pitch, - CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_PITCH, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &tex_max_linear_pitch, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_PITCH, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(tex_max_linear_pitch); } case UR_DEVICE_INFO_MIPMAP_SUPPORT_EXP: { @@ -927,32 +1115,48 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_DEVICE_ID: { int Value = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Value, CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID, hDevice->get())); - detail::ur::assertion(Value >= 0); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Value, CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(Value > 0); return ReturnValue(Value); } case UR_DEVICE_INFO_UUID: { CUuuid UUID; + try { #if (CUDA_VERSION >= 11040) - detail::ur::assertion(cuDeviceGetUuid_v2(&UUID, hDevice->get()) == - CUDA_SUCCESS); + UR_CHECK_ERROR(cuDeviceGetUuid_v2(&UUID, hDevice->get())); #else - detail::ur::assertion(cuDeviceGetUuid(&UUID, hDevice->get()) == - CUDA_SUCCESS); + UR_CHECK_ERROR(cuDeviceGetUuid(&UUID, hDevice->get())); #endif + } catch (ur_result_t Error) { + return Error; + } std::array Name; std::copy(UUID.bytes, UUID.bytes + 16, Name.begin()); return ReturnValue(Name.data(), 16); } case UR_DEVICE_INFO_MAX_MEMORY_BANDWIDTH: { int Major = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } int Minor = 0; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } // Some specific devices seem to need special handling. See reference // https://github.com/jeffhammond/HPCInfo/blob/master/cuda/gpu-detect.cu @@ -965,18 +1169,26 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } else if (IsOrinAGX) { MemoryClockKHz = 3200000; } else { - UR_CHECK_ERROR(cuDeviceGetAttribute(&MemoryClockKHz, - CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE, - hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MemoryClockKHz, CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } } int MemoryBusWidth = 0; if (IsOrinAGX) { MemoryBusWidth = 256; } else { - UR_CHECK_ERROR(cuDeviceGetAttribute( - &MemoryBusWidth, CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH, - hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MemoryBusWidth, CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } } uint32_t MemoryBandwidth = MemoryClockKHz * MemoryBusWidth * 250; @@ -1014,11 +1226,15 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, // Note: This number is shared by all thread blocks simultaneously resident // on a multiprocessor. int MaxRegisters{-1}; - UR_CHECK_ERROR(cuDeviceGetAttribute( - &MaxRegisters, CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK, - hDevice->get())); + try { + UR_CHECK_ERROR(cuDeviceGetAttribute( + &MaxRegisters, CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } - detail::ur::assertion(MaxRegisters >= 0); + assert(MaxRegisters > 0); return ReturnValue(static_cast(MaxRegisters)); } @@ -1029,10 +1245,16 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_PCI_ADDRESS: { constexpr size_t AddressBufferSize = 13; char AddressBuffer[AddressBufferSize]; - UR_CHECK_ERROR( - cuDeviceGetPCIBusId(AddressBuffer, AddressBufferSize, hDevice->get())); + + try { + UR_CHECK_ERROR(cuDeviceGetPCIBusId(AddressBuffer, AddressBufferSize, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + // CUDA API (8.x - 12.1) guarantees 12 bytes + \0 are written - detail::ur::assertion(strnlen(AddressBuffer, AddressBufferSize) == 12); + assert(strnlen(AddressBuffer, AddressBufferSize) != 12); return ReturnValue(AddressBuffer, strnlen(AddressBuffer, AddressBufferSize - 1) + 1); } @@ -1208,8 +1430,17 @@ ur_result_t UR_APICALL urDeviceGetGlobalTimestamps(ur_device_handle_t hDevice, ScopedContext Active(hDevice->getContext()); if (pDeviceTimestamp) { - UR_CHECK_ERROR(cuEventCreate(&Event, CU_EVENT_DEFAULT)); - UR_CHECK_ERROR(cuEventRecord(Event, 0)); + try { + UR_CHECK_ERROR(cuEventCreate(&Event, CU_EVENT_DEFAULT)); + } catch (ur_result_t Error) { + return Error; + } + + try { + UR_CHECK_ERROR(cuEventRecord(Event, 0)); + } catch (ur_result_t Error) { + return Error; + } } if (pHostTimestamp) { @@ -1220,7 +1451,11 @@ ur_result_t UR_APICALL urDeviceGetGlobalTimestamps(ur_device_handle_t hDevice, } if (pDeviceTimestamp) { - UR_CHECK_ERROR(cuEventSynchronize(Event)); + try { + UR_CHECK_ERROR(cuEventSynchronize(Event)); + } catch (ur_result_t Error) { + return Error; + } *pDeviceTimestamp = hDevice->getElapsedTime(Event); } diff --git a/source/adapters/cuda/enqueue.cpp b/source/adapters/cuda/enqueue.cpp index 7c1de98837..1038eb0036 100644 --- a/source/adapters/cuda/enqueue.cpp +++ b/source/adapters/cuda/enqueue.cpp @@ -837,23 +837,29 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferFill( } } -static size_t imageElementByteSize(CUDA_ARRAY_DESCRIPTOR ArrayDesc) { +static ur_result_t imageElementByteSize(CUDA_ARRAY_DESCRIPTOR ArrayDesc, + unsigned int *Size) { switch (ArrayDesc.Format) { case CU_AD_FORMAT_UNSIGNED_INT8: case CU_AD_FORMAT_SIGNED_INT8: - return 1; + *Size = 1; + break; case CU_AD_FORMAT_UNSIGNED_INT16: case CU_AD_FORMAT_SIGNED_INT16: case CU_AD_FORMAT_HALF: - return 2; + *Size = 2; + break; case CU_AD_FORMAT_UNSIGNED_INT32: case CU_AD_FORMAT_SIGNED_INT32: case CU_AD_FORMAT_FLOAT: - return 4; + *Size = 4; + break; default: - detail::ur::die("Invalid image format."); - return 0; + *Size = 0; + return UR_RESULT_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR; } + + return UR_RESULT_SUCCESS; } /// General ND memory copy operation for images. @@ -949,7 +955,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageRead( CUDA_ARRAY_DESCRIPTOR ArrayDesc; UR_CHECK_ERROR(cuArrayGetDescriptor(&ArrayDesc, Array)); - int ElementByteSize = imageElementByteSize(ArrayDesc); + unsigned int ElementByteSize = 0; + UR_RETURN_ON_FAILURE(imageElementByteSize(ArrayDesc, &ElementByteSize)); size_t ByteOffsetX = origin.x * ElementByteSize * ArrayDesc.NumChannels; size_t BytesToCopy = ElementByteSize * ArrayDesc.NumChannels * region.width; @@ -1021,7 +1028,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageWrite( CUDA_ARRAY_DESCRIPTOR ArrayDesc; UR_CHECK_ERROR(cuArrayGetDescriptor(&ArrayDesc, Array)); - int ElementByteSize = imageElementByteSize(ArrayDesc); + unsigned int ElementByteSize = 0; + UR_RETURN_ON_FAILURE(imageElementByteSize(ArrayDesc, &ElementByteSize)); size_t ByteOffsetX = origin.x * ElementByteSize * ArrayDesc.NumChannels; size_t BytesToCopy = ElementByteSize * ArrayDesc.NumChannels * region.width; @@ -1100,7 +1108,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageCopy( UR_ASSERT(SrcArrayDesc.NumChannels == DstArrayDesc.NumChannels, UR_RESULT_ERROR_INVALID_MEM_OBJECT); - int ElementByteSize = imageElementByteSize(SrcArrayDesc); + unsigned int ElementByteSize = 0; + UR_RETURN_ON_FAILURE(imageElementByteSize(SrcArrayDesc, &ElementByteSize)); size_t DstByteOffsetX = dstOrigin.x * ElementByteSize * SrcArrayDesc.NumChannels; diff --git a/source/adapters/cuda/event.cpp b/source/adapters/cuda/event.cpp index ac66bf479e..c1cd27ac52 100644 --- a/source/adapters/cuda/event.cpp +++ b/source/adapters/cuda/event.cpp @@ -113,10 +113,8 @@ ur_result_t ur_event_handle_t_::record() { try { EventID = Queue->getNextEventID(); - if (EventID == 0) { - detail::ur::die( - "Unrecoverable program state reached in event identifier overflow"); - } + assert(EventID == 0 && + "Unrecoverable program state reached in event identifier overflow."); UR_CHECK_ERROR(cuEventRecord(EvEnd, Stream)); } catch (ur_result_t error) { Result = error; @@ -176,10 +174,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventGetInfo(ur_event_handle_t hEvent, case UR_EVENT_INFO_CONTEXT: return ReturnValue(hEvent->getContext()); default: - detail::ur::die("Event info request not implemented"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } - - return UR_RESULT_ERROR_INVALID_ENUMERATION; } /// Obtain profiling information from PI CUDA events @@ -206,8 +202,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventGetProfilingInfo( default: break; } - detail::ur::die("Event Profiling info request not implemented"); - return {}; + + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } UR_APIEXPORT ur_result_t UR_APICALL urEventSetCallback(ur_event_handle_t, @@ -238,8 +234,8 @@ urEventWait(uint32_t numEvents, const ur_event_handle_t *phEventWaitList) { UR_APIEXPORT ur_result_t UR_APICALL urEventRetain(ur_event_handle_t hEvent) { const auto RefCount = hEvent->incrementReferenceCount(); - detail::ur::assertion(RefCount != 0, - "Reference count overflow detected in urEventRetain."); + assert(RefCount == 0 && + "Reference count overflow detected in urEventRetain."); return UR_RESULT_SUCCESS; } @@ -247,8 +243,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventRetain(ur_event_handle_t hEvent) { UR_APIEXPORT ur_result_t UR_APICALL urEventRelease(ur_event_handle_t hEvent) { // double delete or someone is messing with the ref count. // either way, cannot safely proceed. - detail::ur::assertion(hEvent->getReferenceCount() != 0, - "Reference count overflow detected in urEventRelease."); + assert(hEvent->getReferenceCount() == 0 && + "Unrecoverable program state reached in urEventRetain."); // decrement ref count. If it is 0, delete the event. if (hEvent->decrementReferenceCount() == 0) { diff --git a/source/adapters/cuda/image.cpp b/source/adapters/cuda/image.cpp index b9cc832a02..4b59de5cd4 100644 --- a/source/adapters/cuda/image.cpp +++ b/source/adapters/cuda/image.cpp @@ -898,7 +898,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( ChannelOrder = UR_IMAGE_CHANNEL_ORDER_RGBA; break; default: - die("Unexpected NumChannels returned by CUDA"); + return UR_RESULT_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR; } if (pPropValue) { ((ur_image_format_t *)pPropValue)->channelType = ChannelType; diff --git a/source/adapters/cuda/memory.cpp b/source/adapters/cuda/memory.cpp index f479522fb3..c68ab46e7e 100644 --- a/source/adapters/cuda/memory.cpp +++ b/source/adapters/cuda/memory.cpp @@ -144,13 +144,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemRelease(ur_mem_handle_t hMem) { Result = UR_RESULT_ERROR_OUT_OF_RESOURCES; } - if (Result != UR_RESULT_SUCCESS) { - // A reported CUDA error is either an implementation or an asynchronous CUDA - // error for which it is unclear if the function that reported it succeeded - // or not. Either way, the state of the program is compromised and likely - // unrecoverable. - detail::ur::die("Unrecoverable program state reached in urMemRelease"); - } + // A reported CUDA error is either an implementation or an asynchronous CUDA + // error for which it is unclear if the function that reported it succeeded + // or not. Either way, the state of the program is compromised and likely + // unrecoverable. + assert(Result != UR_RESULT_SUCCESS && + "Unrecoverable program state reached in urMemRelease."); return UR_RESULT_SUCCESS; } @@ -306,8 +305,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemImageCreate( PixelTypeSizeBytes = 4; break; default: - detail::ur::die( - "urMemImageCreate given unsupported image_channel_data_type"); + return UR_RESULT_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR; } // When a dimension isn't used pImageDesc has the size set to 1 diff --git a/source/adapters/cuda/queue.cpp b/source/adapters/cuda/queue.cpp index 120d665524..f99a6c35b3 100644 --- a/source/adapters/cuda/queue.cpp +++ b/source/adapters/cuda/queue.cpp @@ -260,12 +260,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueCreateWithNativeHandle( UR_CHECK_ERROR(cuStreamGetFlags(CuStream, &CuFlags)); ur_queue_flags_t Flags = 0; - if (CuFlags == CU_STREAM_DEFAULT) + if (CuFlags == CU_STREAM_DEFAULT) { Flags = UR_QUEUE_FLAG_USE_DEFAULT_STREAM; - else if (CuFlags == CU_STREAM_NON_BLOCKING) + } else if (CuFlags == CU_STREAM_NON_BLOCKING) { Flags = UR_QUEUE_FLAG_SYNC_WITH_DEFAULT_STREAM; - else - detail::ur::die("Unknown cuda stream"); + } else { + assert(!"Unknown cuda stream."); + } std::vector ComputeCuStreams(1, CuStream); std::vector TransferCuStreams(0); diff --git a/source/adapters/cuda/sampler.cpp b/source/adapters/cuda/sampler.cpp index 904c529fd2..245905bedb 100644 --- a/source/adapters/cuda/sampler.cpp +++ b/source/adapters/cuda/sampler.cpp @@ -94,9 +94,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urSamplerRelease(ur_sampler_handle_t hSampler) { // double delete or someone is messing with the ref count. // either way, cannot safely proceed. - detail::ur::assertion( - hSampler->getReferenceCount() != 0, - "Reference count overflow detected in urSamplerRelease."); + assert(hSampler->getReferenceCount() == 0 && + "Reference count overflow detected in urSamplerRelease."); // decrement ref count. If it is 0, delete the sampler. if (hSampler->decrementReferenceCount() == 0) { diff --git a/source/adapters/hip/common.cpp b/source/adapters/hip/common.cpp index f1f8ec4fbb..285a346f78 100644 --- a/source/adapters/hip/common.cpp +++ b/source/adapters/hip/common.cpp @@ -156,16 +156,6 @@ hipError_t getHipVersionString(std::string &Version) { return Result; } -void detail::ur::die(const char *pMessage) { - std::cerr << "ur_die: " << pMessage << '\n'; - std::terminate(); -} - -void detail::ur::assertion(bool Condition, const char *pMessage) { - if (!Condition) - die(pMessage); -} - void detail::ur::hipPrint(const char *pMessage) { std::cerr << "ur_print: " << pMessage << '\n'; } @@ -188,3 +178,24 @@ ur_result_t urGetLastResult(ur_platform_handle_t, const char **ppMessage) { *ppMessage = &ErrorMessage[0]; return ErrorMessageCode; } + +ur_result_t imageElementByteSize(hipArray_Format Format, unsigned int *Size) { + switch (Format) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + case HIP_AD_FORMAT_SIGNED_INT8: + *Size = 1; + break; + case HIP_AD_FORMAT_UNSIGNED_INT16: + case HIP_AD_FORMAT_SIGNED_INT16: + case HIP_AD_FORMAT_HALF: + *Size = 2; + break; + case HIP_AD_FORMAT_UNSIGNED_INT32: + case HIP_AD_FORMAT_SIGNED_INT32: + case HIP_AD_FORMAT_FLOAT: + *Size = 4; + break; + default: + return UR_RESULT_ERROR_IMAGE_FORMAT_NOT_SUPPORTED; + } +} diff --git a/source/adapters/hip/common.hpp b/source/adapters/hip/common.hpp index be332c280b..3435729bac 100644 --- a/source/adapters/hip/common.hpp +++ b/source/adapters/hip/common.hpp @@ -19,6 +19,16 @@ #include #include +/** + * Call a UR API and, if the result is not UR_RESULT_SUCCESS, automatically + * return from the current function. + */ +#define UR_RETURN_ON_FAILURE(urCall) \ + if (const ur_result_t ur_result_macro = urCall; \ + ur_result_macro != UR_RESULT_SUCCESS) { \ + return ur_result_macro; \ + } + // Before ROCm 6, hipify doesn't support cuArrayGetDescriptor, on AMD the // hipArray can just be indexed, but on NVidia it is an opaque type and needs to // go through cuArrayGetDescriptor so implement a utility function to get the @@ -117,17 +127,9 @@ extern thread_local char ErrorMessage[MaxMessageSize]; namespace detail { namespace ur { -// Report error and no return (keeps compiler from printing warnings). -// TODO: Probably change that to throw a catchable exception, -// but for now it is useful to see every failure. -// -[[noreturn]] void die(const char *pMessage); - // Reports error messages void hipPrint(const char *pMessage); -void assertion(bool Condition, const char *pMessage = nullptr); - } // namespace ur } // namespace detail @@ -187,7 +189,7 @@ template class ReleaseGuard { // HIP error for which it is unclear if the function that reported it // succeeded or not. Either way, the state of the program is compromised // and likely unrecoverable. - detail::ur::die("Unrecoverable program state reached in piMemRelease"); + assert(!"Unrecoverable program state reached in piMemRelease"); } } } @@ -204,3 +206,5 @@ template class ReleaseGuard { /// UR object. void dismiss() { Captive = nullptr; } }; + +ur_result_t imageElementByteSize(hipArray_Format Format, int *Size); diff --git a/source/adapters/hip/device.cpp b/source/adapters/hip/device.cpp index ae4dbe159e..e13d7be15d 100644 --- a/source/adapters/hip/device.cpp +++ b/source/adapters/hip/device.cpp @@ -45,9 +45,14 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_MAX_COMPUTE_UNITS: { int ComputeUnits = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &ComputeUnits, hipDeviceAttributeMultiprocessorCount, hDevice->get())); - detail::ur::assertion(ComputeUnits >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &ComputeUnits, hipDeviceAttributeMultiprocessorCount, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(ComputeUnits > 0); return ReturnValue(static_cast(ComputeUnits)); } case UR_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS: { @@ -59,17 +64,29 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } return_sizes; int MaxX = 0, MaxY = 0, MaxZ = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxX, hipDeviceAttributeMaxBlockDimX, - hDevice->get())); - detail::ur::assertion(MaxX >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &MaxX, hipDeviceAttributeMaxBlockDimX, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxX > 0); - UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxY, hipDeviceAttributeMaxBlockDimY, - hDevice->get())); - detail::ur::assertion(MaxY >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &MaxY, hipDeviceAttributeMaxBlockDimY, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxY > 0); - UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxZ, hipDeviceAttributeMaxBlockDimZ, - hDevice->get())); - detail::ur::assertion(MaxZ >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &MaxZ, hipDeviceAttributeMaxBlockDimZ, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxZ > 0); return_sizes.sizes[0] = size_t(MaxX); return_sizes.sizes[1] = size_t(MaxY); @@ -83,17 +100,29 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } return_sizes; int MaxX = 0, MaxY = 0, MaxZ = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxX, hipDeviceAttributeMaxGridDimX, - hDevice->get())); - detail::ur::assertion(MaxX >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxX, hipDeviceAttributeMaxGridDimX, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxX > 0); - UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxY, hipDeviceAttributeMaxGridDimY, - hDevice->get())); - detail::ur::assertion(MaxY >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxY, hipDeviceAttributeMaxGridDimY, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxY > 0); - UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxZ, hipDeviceAttributeMaxGridDimZ, - hDevice->get())); - detail::ur::assertion(MaxZ >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxZ, hipDeviceAttributeMaxGridDimZ, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxZ > 0); return_sizes.sizes[0] = size_t(MaxX); return_sizes.sizes[1] = size_t(MaxY); @@ -103,11 +132,15 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_MAX_WORK_GROUP_SIZE: { int MaxWorkGroupSize = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxWorkGroupSize, - hipDeviceAttributeMaxThreadsPerBlock, - hDevice->get())); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute(&MaxWorkGroupSize, + hipDeviceAttributeMaxThreadsPerBlock, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } - detail::ur::assertion(MaxWorkGroupSize >= 0); + assert(MaxWorkGroupSize > 0); return ReturnValue(size_t(MaxWorkGroupSize)); } @@ -156,11 +189,21 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_MAX_NUM_SUB_GROUPS: { // Number of sub-groups = max block size / warp size + possible remainder int MaxThreads = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &MaxThreads, hipDeviceAttributeMaxThreadsPerBlock, hDevice->get())); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &MaxThreads, hipDeviceAttributeMaxThreadsPerBlock, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + int WarpSize = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute(&WarpSize, hipDeviceAttributeWarpSize, - hDevice->get())); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &WarpSize, hipDeviceAttributeWarpSize, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + int MaxWarps = (MaxThreads + WarpSize - 1) / WarpSize; return ReturnValue(MaxWarps); } @@ -168,23 +211,38 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, // Volta provides independent thread scheduling // TODO: Revisit for previous generation GPUs int Major = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &Major, hipDeviceAttributeComputeCapabilityMajor, hDevice->get())); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &Major, hipDeviceAttributeComputeCapabilityMajor, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + bool IFP = (Major >= 7); return ReturnValue(IFP); } case UR_DEVICE_INFO_SUB_GROUP_SIZES_INTEL: { int WarpSize = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute(&WarpSize, hipDeviceAttributeWarpSize, - hDevice->get())); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &WarpSize, hipDeviceAttributeWarpSize, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + size_t Sizes[1] = {static_cast(WarpSize)}; return ReturnValue(Sizes, 1); } case UR_DEVICE_INFO_MAX_CLOCK_FREQUENCY: { int ClockFreq = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &ClockFreq, hipDeviceAttributeClockRate, hDevice->get())); - detail::ur::assertion(ClockFreq >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &ClockFreq, hipDeviceAttributeClockRate, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + + assert(ClockFreq > 0); return ReturnValue(static_cast(ClockFreq) / 1000u); } case UR_DEVICE_INFO_ADDRESS_BITS: { @@ -199,8 +257,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, // CL_DEVICE_TYPE_CUSTOM. size_t Global = 0; - detail::ur::assertion(hipDeviceTotalMem(&Global, hDevice->get()) == - hipSuccess); + + try { + UR_CHECK_ERROR(hipDeviceTotalMem(&Global, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } auto QuarterGlobal = static_cast(Global / 4u); @@ -233,13 +295,22 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE2D_MAX_HEIGHT: { // Take the smaller of maximum surface and maximum texture height. int TexHeight = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &TexHeight, hipDeviceAttributeMaxTexture2DHeight, hDevice->get())); - detail::ur::assertion(TexHeight >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &TexHeight, hipDeviceAttributeMaxTexture2DHeight, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexHeight > 0); + int SurfHeight = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfHeight, hipDeviceAttributeMaxTexture2DHeight, hDevice->get())); - detail::ur::assertion(SurfHeight >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &SurfHeight, hipDeviceAttributeMaxTexture2DHeight, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfHeight > 0); int Min = std::min(TexHeight, SurfHeight); @@ -248,14 +319,23 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE2D_MAX_WIDTH: { // Take the smaller of maximum surface and maximum texture width. int TexWidth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &TexWidth, hipDeviceAttributeMaxTexture2DWidth, hDevice->get())); - detail::ur::assertion(TexWidth >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &TexWidth, hipDeviceAttributeMaxTexture2DWidth, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + + assert(TexWidth > 0); int SurfWidth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfWidth, hipDeviceAttributeMaxTexture2DWidth, hDevice->get())); - detail::ur::assertion(SurfWidth >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &SurfWidth, hipDeviceAttributeMaxTexture2DWidth, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfWidth > 0); int Min = std::min(TexWidth, SurfWidth); return ReturnValue(static_cast(Min)); @@ -263,13 +343,23 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE3D_MAX_HEIGHT: { // Take the smaller of maximum surface and maximum texture height. int TexHeight = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &TexHeight, hipDeviceAttributeMaxTexture3DHeight, hDevice->get())); - detail::ur::assertion(TexHeight >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &TexHeight, hipDeviceAttributeMaxTexture3DHeight, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexHeight > 0); + int SurfHeight = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfHeight, hipDeviceAttributeMaxTexture3DHeight, hDevice->get())); - detail::ur::assertion(SurfHeight >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &SurfHeight, hipDeviceAttributeMaxTexture3DHeight, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + + assert(SurfHeight > 0); int Min = std::min(TexHeight, SurfHeight); @@ -278,13 +368,22 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE3D_MAX_WIDTH: { // Take the smaller of maximum surface and maximum texture width. int TexWidth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &TexWidth, hipDeviceAttributeMaxTexture3DWidth, hDevice->get())); - detail::ur::assertion(TexWidth >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &TexWidth, hipDeviceAttributeMaxTexture3DWidth, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexWidth > 0); + int SurfWidth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfWidth, hipDeviceAttributeMaxTexture3DWidth, hDevice->get())); - detail::ur::assertion(SurfWidth >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &SurfWidth, hipDeviceAttributeMaxTexture3DWidth, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfWidth > 0); int Min = std::min(TexWidth, SurfWidth); @@ -293,13 +392,22 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE3D_MAX_DEPTH: { // Take the smaller of maximum surface and maximum texture depth. int TexDepth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &TexDepth, hipDeviceAttributeMaxTexture3DDepth, hDevice->get())); - detail::ur::assertion(TexDepth >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &TexDepth, hipDeviceAttributeMaxTexture3DDepth, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexDepth > 0); + int SurfDepth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfDepth, hipDeviceAttributeMaxTexture3DDepth, hDevice->get())); - detail::ur::assertion(SurfDepth >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &SurfDepth, hipDeviceAttributeMaxTexture3DDepth, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfDepth > 0); int Min = std::min(TexDepth, SurfDepth); @@ -308,13 +416,22 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE: { // Take the smaller of maximum surface and maximum texture width. int TexWidth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &TexWidth, hipDeviceAttributeMaxTexture1DWidth, hDevice->get())); - detail::ur::assertion(TexWidth >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &TexWidth, hipDeviceAttributeMaxTexture1DWidth, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(TexWidth > 0); + int SurfWidth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfWidth, hipDeviceAttributeMaxTexture1DWidth, hDevice->get())); - detail::ur::assertion(SurfWidth >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &SurfWidth, hipDeviceAttributeMaxTexture1DWidth, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(SurfWidth > 0); int Min = std::min(TexWidth, SurfWidth); @@ -335,8 +452,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_MEM_BASE_ADDR_ALIGN: { int MemBaseAddrAlign = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &MemBaseAddrAlign, hipDeviceAttributeTextureAlignment, hDevice->get())); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute(&MemBaseAddrAlign, + hipDeviceAttributeTextureAlignment, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } // Multiply by 8 as clGetDeviceInfo returns this value in bits MemBaseAddrAlign *= 8; return ReturnValue(MemBaseAddrAlign); @@ -375,16 +497,24 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_GLOBAL_MEM_CACHE_SIZE: { int CacheSize = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &CacheSize, hipDeviceAttributeL2CacheSize, hDevice->get())); - detail::ur::assertion(CacheSize >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &CacheSize, hipDeviceAttributeL2CacheSize, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(CacheSize > 0); // The L2 cache is global to the GPU. return ReturnValue(static_cast(CacheSize)); } case UR_DEVICE_INFO_GLOBAL_MEM_SIZE: { size_t Bytes = 0; // Runtime API has easy access to this value, driver API info is scarse. - UR_CHECK_ERROR(hipDeviceTotalMem(&Bytes, hDevice->get())); + try { + UR_CHECK_ERROR(hipDeviceTotalMem(&Bytes, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(uint64_t{Bytes}); } case UR_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE: { @@ -394,10 +524,14 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, // memory on AMD GPU may be larger than what can fit in the positive part // of a signed integer, so use an unsigned integer and cast the pointer to // int*. - UR_CHECK_ERROR(hipDeviceGetAttribute(&ConstantMemory, - hipDeviceAttributeTotalConstantMemory, - hDevice->get())); - detail::ur::assertion(ConstantMemory >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &ConstantMemory, hipDeviceAttributeTotalConstantMemory, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(ConstantMemory > 0); return ReturnValue(static_cast(ConstantMemory)); } @@ -415,27 +549,41 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, // HIP has its own definition of "local memory", which maps to OpenCL's // "private memory". int LocalMemSize = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &LocalMemSize, hipDeviceAttributeMaxSharedMemoryPerBlock, - hDevice->get())); - detail::ur::assertion(LocalMemSize >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &LocalMemSize, hipDeviceAttributeMaxSharedMemoryPerBlock, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(LocalMemSize > 0); return ReturnValue(static_cast(LocalMemSize)); } case UR_DEVICE_INFO_ERROR_CORRECTION_SUPPORT: { int EccEnabled = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &EccEnabled, hipDeviceAttributeEccEnabled, hDevice->get())); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &EccEnabled, hipDeviceAttributeEccEnabled, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + + assert(EccEnabled == 0 || EccEnabled == 1); - detail::ur::assertion((EccEnabled == 0) | (EccEnabled == 1)); auto Result = static_cast(EccEnabled); return ReturnValue(Result); } case UR_DEVICE_INFO_HOST_UNIFIED_MEMORY: { int IsIntegrated = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &IsIntegrated, hipDeviceAttributeIntegrated, hDevice->get())); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &IsIntegrated, hipDeviceAttributeIntegrated, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + + assert(IsIntegrated == 0 || IsIntegrated == 1); - detail::ur::assertion((IsIntegrated == 0) | (IsIntegrated == 1)); auto Result = static_cast(IsIntegrated); return ReturnValue(Result); } @@ -487,14 +635,21 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_NAME: { static constexpr size_t MAX_DEVICE_NAME_LENGTH = 256u; char Name[MAX_DEVICE_NAME_LENGTH]; - UR_CHECK_ERROR( - hipDeviceGetName(Name, MAX_DEVICE_NAME_LENGTH, hDevice->get())); + try { + UR_CHECK_ERROR( + hipDeviceGetName(Name, MAX_DEVICE_NAME_LENGTH, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } // On AMD GPUs hipDeviceGetName returns an empty string, so return the arch // name instead, this is also what AMD OpenCL devices return. if (strlen(Name) == 0) { hipDeviceProp_t Props; - detail::ur::assertion(hipGetDeviceProperties(&Props, hDevice->get()) == - hipSuccess); + try { + UR_CHECK_ERROR(hipGetDeviceProperties(&Props, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(Props.gcnArchName, strlen(Props.gcnArchName) + 1); } @@ -505,7 +660,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_DRIVER_VERSION: { std::string Version; - UR_CHECK_ERROR(getHipVersionString(Version)); + try { + UR_CHECK_ERROR(getHipVersionString(Version)); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(Version.c_str()); } case UR_DEVICE_INFO_PROFILE: { @@ -518,8 +677,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, std::stringstream S; hipDeviceProp_t Props; - detail::ur::assertion(hipGetDeviceProperties(&Props, hDevice->get()) == - hipSuccess); + try { + UR_CHECK_ERROR(hipGetDeviceProperties(&Props, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } #if defined(__HIP_PLATFORM_NVIDIA__) S << Props.major << "." << Props.minor; #elif defined(__HIP_PLATFORM_AMD__) @@ -556,8 +718,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, SupportedExtensions += " "; hipDeviceProp_t Props; - detail::ur::assertion(hipGetDeviceProperties(&Props, hDevice->get()) == - hipSuccess); + try { + UR_CHECK_ERROR(hipGetDeviceProperties(&Props, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } if (Props.arch.hasDoubles) { SupportedExtensions += "cl_khr_fp64 "; @@ -710,15 +875,23 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_BACKEND_RUNTIME_VERSION: { int Major = 0, Minor = 0; - UR_CHECK_ERROR(hipDeviceComputeCapability(&Major, &Minor, hDevice->get())); + try { + UR_CHECK_ERROR( + hipDeviceComputeCapability(&Major, &Minor, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } std::string Result = std::to_string(Major) + "." + std::to_string(Minor); return ReturnValue(Result.c_str()); } case UR_DEVICE_INFO_ATOMIC_64: { hipDeviceProp_t Props; - detail::ur::assertion(hipGetDeviceProperties(&Props, hDevice->get()) == - hipSuccess); + try { + UR_CHECK_ERROR(hipGetDeviceProperties(&Props, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(Props.arch.hasGlobalInt64Atomics && Props.arch.hasSharedInt64Atomics); } @@ -726,26 +899,36 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_GLOBAL_MEM_FREE: { size_t FreeMemory = 0; size_t TotalMemory = 0; - detail::ur::assertion(hipMemGetInfo(&FreeMemory, &TotalMemory) == - hipSuccess, - "failed hipMemGetInfo() API."); + try { + UR_CHECK_ERROR(hipMemGetInfo(&FreeMemory, &TotalMemory)); + } catch (ur_result_t Error) { + return Error; + } return ReturnValue(FreeMemory); } case UR_DEVICE_INFO_MEMORY_CLOCK_RATE: { int Value = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &Value, hipDeviceAttributeMemoryClockRate, hDevice->get())); - detail::ur::assertion(Value >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &Value, hipDeviceAttributeMemoryClockRate, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(Value > 0); // Convert kilohertz to megahertz when returning. return ReturnValue(Value / 1000); } case UR_DEVICE_INFO_MEMORY_BUS_WIDTH: { int Value = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &Value, hipDeviceAttributeMemoryBusWidth, hDevice->get())); - detail::ur::assertion(Value >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &Value, hipDeviceAttributeMemoryBusWidth, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(Value > 0); return ReturnValue(Value); } case UR_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES: { @@ -783,9 +966,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_DEVICE_ID: { int Value = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute(&Value, hipDeviceAttributePciDeviceId, - hDevice->get())); - detail::ur::assertion(Value >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &Value, hipDeviceAttributePciDeviceId, hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(Value > 0); return ReturnValue(Value); } case UR_DEVICE_INFO_UUID: { @@ -793,8 +980,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, HIP_VERSION_MAJOR > 5) hipUUID UUID = {}; // Supported since 5.2+ - detail::ur::assertion(hipDeviceGetUuid(&UUID, hDevice->get()) == - hipSuccess); + UR_CHECK_ERROR(hipDeviceGetUuid(&UUID, hDevice->get())); std::array Name; std::copy(UUID.bytes, UUID.bytes + 16, Name.begin()); return ReturnValue(Name.data(), 16); @@ -806,10 +992,14 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, // Note: This number is shared by all thread blocks simultaneously resident // on a multiprocessor. int MaxRegisters{-1}; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &MaxRegisters, hipDeviceAttributeMaxRegistersPerBlock, hDevice->get())); - - detail::ur::assertion(MaxRegisters >= 0); + try { + UR_CHECK_ERROR(hipDeviceGetAttribute( + &MaxRegisters, hipDeviceAttributeMaxRegistersPerBlock, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } + assert(MaxRegisters > 0); return ReturnValue(static_cast(MaxRegisters)); } @@ -820,14 +1010,18 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_PCI_ADDRESS: { constexpr size_t AddressBufferSize = 13; char AddressBuffer[AddressBufferSize]; - UR_CHECK_ERROR( - hipDeviceGetPCIBusId(AddressBuffer, AddressBufferSize, hDevice->get())); + try { + UR_CHECK_ERROR(hipDeviceGetPCIBusId(AddressBuffer, AddressBufferSize, + hDevice->get())); + } catch (ur_result_t Error) { + return Error; + } // A typical PCI address is 12 bytes + \0: "1234:67:90.2", but the HIP API // is not guaranteed to use this format. In practice, it uses this format, // at least in 5.3-5.5. To be on the safe side, we make sure the terminating // \0 is set. AddressBuffer[AddressBufferSize - 1] = '\0'; - detail::ur::assertion(strnlen(AddressBuffer, AddressBufferSize) > 0); + assert(strnlen(AddressBuffer, AddressBufferSize) > 0); return ReturnValue(AddressBuffer, strnlen(AddressBuffer, AddressBufferSize - 1) + 1); } @@ -982,12 +1176,31 @@ ur_result_t UR_APICALL urDeviceGetGlobalTimestamps(ur_device_handle_t hDevice, ScopedContext Active(hDevice); if (pDeviceTimestamp) { - UR_CHECK_ERROR(hipEventCreateWithFlags(&Event, hipEventDefault)); - UR_CHECK_ERROR(hipEventRecord(Event)); - UR_CHECK_ERROR(hipEventSynchronize(Event)); + try { + UR_CHECK_ERROR(hipEventCreateWithFlags(&Event, hipEventDefault)); + } catch (ur_result_t Error) { + return Error; + } + + try { + UR_CHECK_ERROR(hipEventRecord(Event)); + } catch (ur_result_t Error) { + return Error; + } + + try { + UR_CHECK_ERROR(hipEventSynchronize(Event)); + } catch (ur_result_t Error) { + return Error; + } + float ElapsedTime = 0.0f; - UR_CHECK_ERROR(hipEventElapsedTime(&ElapsedTime, - ur_platform_handle_t_::EvBase, Event)); + try { + UR_CHECK_ERROR(hipEventElapsedTime(&ElapsedTime, + ur_platform_handle_t_::EvBase, Event)); + } catch (ur_result_t Error) { + return Error; + } *pDeviceTimestamp = (uint64_t)(ElapsedTime * (double)1e6); } diff --git a/source/adapters/hip/enqueue.cpp b/source/adapters/hip/enqueue.cpp index 33691ec112..78375ac12c 100644 --- a/source/adapters/hip/enqueue.cpp +++ b/source/adapters/hip/enqueue.cpp @@ -1001,7 +1001,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageRead( size_t NumChannels{}; UR_CHECK_ERROR(getArrayDesc(Array, Format, NumChannels)); - int ElementByteSize = imageElementByteSize(Format); + int ElementByteSize = 0; + UR_RETURN_ON_FAILURE(imageElementByteSize(Format, &ElementByteSize)); size_t ByteOffsetX = origin.x * ElementByteSize * NumChannels; size_t BytesToCopy = ElementByteSize * NumChannels * region.depth; @@ -1062,7 +1063,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageWrite( size_t NumChannels{}; UR_CHECK_ERROR(getArrayDesc(Array, Format, NumChannels)); - int ElementByteSize = imageElementByteSize(Format); + int ElementByteSize = 0; + UR_RETURN_ON_FAILURE(imageElementByteSize(Format, &ElementByteSize)); size_t ByteOffsetX = origin.x * ElementByteSize * NumChannels; size_t BytesToCopy = ElementByteSize * NumChannels * region.depth; @@ -1136,7 +1138,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageCopy( UR_ASSERT(SrcNumChannels == DstNumChannels, UR_RESULT_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR); - int ElementByteSize = imageElementByteSize(SrcFormat); + int ElementByteSize = 0; + UR_RETURN_ON_FAILURE(imageElementByteSize(SrcFormat, &ElementByteSize)); size_t DstByteOffsetX = dstOrigin.x * ElementByteSize * SrcNumChannels; size_t SrcByteOffsetX = srcOrigin.x * ElementByteSize * DstNumChannels; diff --git a/source/adapters/hip/event.cpp b/source/adapters/hip/event.cpp index 0af2ee02ab..5c0b0fcb4f 100644 --- a/source/adapters/hip/event.cpp +++ b/source/adapters/hip/event.cpp @@ -141,10 +141,8 @@ ur_result_t ur_event_handle_t_::record() { try { EventId = Queue->getNextEventId(); - if (EventId == 0) { - detail::ur::die( - "Unrecoverable program state reached in event identifier overflow"); - } + assert(EventId == 0 && + "Unrecoverable program state reached in event identifier overflow."); UR_CHECK_ERROR(hipEventRecord(EvEnd, Stream)); Result = UR_RESULT_SUCCESS; } catch (ur_result_t Error) { @@ -274,8 +272,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventSetCallback(ur_event_handle_t, UR_APIEXPORT ur_result_t UR_APICALL urEventRetain(ur_event_handle_t hEvent) { const auto RefCount = hEvent->incrementReferenceCount(); - detail::ur::assertion(RefCount != 0, - "Reference count overflow detected in urEventRetain."); + assert(RefCount == 0 && + "Reference count overflow detected in urEventRetain."); return UR_RESULT_SUCCESS; } @@ -283,8 +281,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventRetain(ur_event_handle_t hEvent) { UR_APIEXPORT ur_result_t UR_APICALL urEventRelease(ur_event_handle_t hEvent) { // double delete or someone is messing with the ref count. // either way, cannot safely proceed. - detail::ur::assertion(hEvent->getReferenceCount() != 0, - "Reference count overflow detected in urEventRelease."); + assert(hEvent->getReferenceCount() == 0 && + "Reference count overflow detected in urEventRelease."); // decrement ref count. If it is 0, delete the event. if (hEvent->decrementReferenceCount() == 0) { diff --git a/source/adapters/hip/kernel.cpp b/source/adapters/hip/kernel.cpp index df8c087f2b..d83651389c 100644 --- a/source/adapters/hip/kernel.cpp +++ b/source/adapters/hip/kernel.cpp @@ -296,9 +296,7 @@ urKernelSetArgMemObj(ur_kernel_handle_t hKernel, uint32_t argIndex, if (Format != HIP_AD_FORMAT_UNSIGNED_INT32 && Format != HIP_AD_FORMAT_SIGNED_INT32 && Format != HIP_AD_FORMAT_HALF && Format != HIP_AD_FORMAT_FLOAT) { - detail::ur::die( - "UR HIP kernels only support images with channel types int32, " - "uint32, float, and half."); + return UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT; } hipSurfaceObject_t hipSurf = std::get(hArgValue->Mem).getSurface(Device); diff --git a/source/adapters/hip/memory.cpp b/source/adapters/hip/memory.cpp index dcc3e34fad..d72804217d 100644 --- a/source/adapters/hip/memory.cpp +++ b/source/adapters/hip/memory.cpp @@ -16,18 +16,24 @@ size_t imageElementByteSize(hipArray_Format ArrayFormat) { switch (ArrayFormat) { case HIP_AD_FORMAT_UNSIGNED_INT8: - case HIP_AD_FORMAT_SIGNED_INT8: - return 1; + *ChannelType = UR_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; case HIP_AD_FORMAT_UNSIGNED_INT16: - case HIP_AD_FORMAT_SIGNED_INT16: - case HIP_AD_FORMAT_HALF: - return 2; + *ChannelType = UR_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16; case HIP_AD_FORMAT_UNSIGNED_INT32: + *ChannelType = UR_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; + case HIP_AD_FORMAT_SIGNED_INT8: + *ChannelType = UR_IMAGE_CHANNEL_TYPE_SIGNED_INT8; + case HIP_AD_FORMAT_SIGNED_INT16: + *ChannelType = UR_IMAGE_CHANNEL_TYPE_SIGNED_INT16; case HIP_AD_FORMAT_SIGNED_INT32: + *ChannelType = UR_IMAGE_CHANNEL_TYPE_SIGNED_INT32; + case HIP_AD_FORMAT_HALF: + *ChannelType = UR_IMAGE_CHANNEL_TYPE_HALF_FLOAT; case HIP_AD_FORMAT_FLOAT: - return 4; + *ChannelType = UR_IMAGE_CHANNEL_TYPE_FLOAT; + default: - detail::ur::die("Invalid HIP format specifier"); + return UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT; } return 0; } @@ -65,7 +71,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemRelease(ur_mem_handle_t hMem) { // error for which it is unclear if the function that reported it succeeded // or not. Either way, the state of the program is compromised and likely // unrecoverable. - detail::ur::die("Unrecoverable program state reached in urMemRelease"); + assert(!"Unrecoverable program state reached in urMemRelease"); } return UR_RESULT_SUCCESS; @@ -241,9 +247,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemGetInfo(ur_mem_handle_t hMemory, HIP_ARRAY3D_DESCRIPTOR ArrayDescriptor; UR_CHECK_ERROR( hipArray3DGetDescriptor(&ArrayDescriptor, Mem.getArray(Device))); - const auto PixelSizeBytes = - imageElementByteSize(ArrayDescriptor.Format) * - ArrayDescriptor.NumChannels; + int PixelSize = 0; + UR_RETURN_ON_FAILURE( + imageElementByteSize(ArrayDescriptor.Format, &PixelSize)); + const auto PixelSizeBytes = PixelSize * ArrayDescriptor.NumChannels; const auto ImageSizeBytes = PixelSizeBytes * (ArrayDescriptor.Width ? ArrayDescriptor.Width : 1) * @@ -393,35 +400,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemImageGetInfo(ur_mem_handle_t hMemory, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; #endif - const auto hip2urFormat = - [](hipArray_Format HipFormat) -> ur_image_channel_type_t { - switch (HipFormat) { - case HIP_AD_FORMAT_UNSIGNED_INT8: - return UR_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; - case HIP_AD_FORMAT_UNSIGNED_INT16: - return UR_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16; - case HIP_AD_FORMAT_UNSIGNED_INT32: - return UR_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; - case HIP_AD_FORMAT_SIGNED_INT8: - return UR_IMAGE_CHANNEL_TYPE_SIGNED_INT8; - case HIP_AD_FORMAT_SIGNED_INT16: - return UR_IMAGE_CHANNEL_TYPE_SIGNED_INT16; - case HIP_AD_FORMAT_SIGNED_INT32: - return UR_IMAGE_CHANNEL_TYPE_SIGNED_INT32; - case HIP_AD_FORMAT_HALF: - return UR_IMAGE_CHANNEL_TYPE_HALF_FLOAT; - case HIP_AD_FORMAT_FLOAT: - return UR_IMAGE_CHANNEL_TYPE_FLOAT; - - default: - detail::ur::die("Invalid Hip format specified."); - } - }; - switch (propName) { case UR_IMAGE_INFO_FORMAT: - return ReturnValue(ur_image_format_t{UR_IMAGE_CHANNEL_ORDER_RGBA, - hip2urFormat(ArrayInfo.Format)}); + ur_image_channel_type_t ChannelType; + UR_RETURN_ON_FAILURE(hip2urFormat(ArrayInfo.Format, &ChannelType)); + return ReturnValue( + ur_image_format_t{UR_IMAGE_CHANNEL_ORDER_RGBA, ChannelType}); case UR_IMAGE_INFO_WIDTH: return ReturnValue(ArrayInfo.Width); case UR_IMAGE_INFO_HEIGHT: diff --git a/source/adapters/hip/memory.hpp b/source/adapters/hip/memory.hpp index 7707794b3c..c0b1afa22b 100644 --- a/source/adapters/hip/memory.hpp +++ b/source/adapters/hip/memory.hpp @@ -242,7 +242,7 @@ struct SurfaceMem { break; default: // urMemImageCreate given unsupported image_channel_data_type - detail::ur::die("Bad image format given to ur_image_ constructor"); + assert(!"Bad image format given to ur_image_ constructor"); } } diff --git a/source/adapters/hip/program.cpp b/source/adapters/hip/program.cpp index 1f27cfd396..95e9ccb074 100644 --- a/source/adapters/hip/program.cpp +++ b/source/adapters/hip/program.cpp @@ -149,8 +149,10 @@ ur_result_t ur_program_handle_t_::finalizeRelocatable() { std::string ISA = "amdgcn-amd-amdhsa--"; hipDeviceProp_t Props; - detail::ur::assertion(hipGetDeviceProperties(&Props, getDevice()->get()) == - hipSuccess); + if (hipGetDeviceProperties(&Props, Context->getDevice()->get()) != + hipSuccess) { + return UR_RESULT_ERROR_INVALID_OPERATION; + } ISA += Props.gcnArchName; UR_CHECK_ERROR(amd_comgr_action_info_set_isa_name(Action, ISA.data())); @@ -259,8 +261,8 @@ ur_result_t getKernelNames(ur_program_handle_t) { UR_APIEXPORT ur_result_t UR_APICALL urProgramCreateWithIL(ur_context_handle_t, const void *, size_t, const ur_program_properties_t *, ur_program_handle_t *) { - detail::ur::die("urProgramCreateWithIL not implemented for HIP adapter" - " please use urProgramCreateWithBinary instead"); + assert(!"urProgramCreateWithIL not implemented for HIP adapter" + " please use urProgramCreateWithBinary instead"); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/hip/queue.cpp b/source/adapters/hip/queue.cpp index 6e6496fec1..dc8dc7bd09 100644 --- a/source/adapters/hip/queue.cpp +++ b/source/adapters/hip/queue.cpp @@ -287,12 +287,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueCreateWithNativeHandle( UR_CHECK_ERROR(hipStreamGetFlags(HIPStream, &HIPFlags)); ur_queue_flags_t Flags = 0; - if (HIPFlags == hipStreamDefault) + if (HIPFlags == hipStreamDefault) { Flags = UR_QUEUE_FLAG_USE_DEFAULT_STREAM; - else if (HIPFlags == hipStreamNonBlocking) + } else if (HIPFlags == hipStreamNonBlocking) { Flags = UR_QUEUE_FLAG_SYNC_WITH_DEFAULT_STREAM; - else - detail::ur::die("Unknown hip stream"); + } else { + assert(!"Unknown hip stream."); + } std::vector ComputeHIPStreams(1, HIPStream); std::vector TransferHIPStreams(0); diff --git a/source/adapters/hip/sampler.cpp b/source/adapters/hip/sampler.cpp index 1ee1996164..a948f08541 100644 --- a/source/adapters/hip/sampler.cpp +++ b/source/adapters/hip/sampler.cpp @@ -68,9 +68,8 @@ ur_result_t urSamplerRetain(ur_sampler_handle_t hSampler) { ur_result_t urSamplerRelease(ur_sampler_handle_t hSampler) { // double delete or someone is messing with the ref count. // either way, cannot safely proceed. - detail::ur::assertion( - hSampler->getReferenceCount() != 0, - "Reference count overflow detected in urSamplerRelease."); + assert(hSampler->getReferenceCount() == 0 && + "Reference count overflow detected in urSamplerRelease."); // decrement ref count. If it is 0, delete the sampler. if (hSampler->decrementReferenceCount() == 0) { diff --git a/source/adapters/level_zero/context.cpp b/source/adapters/level_zero/context.cpp index c4b5423adb..1bc57aecd2 100644 --- a/source/adapters/level_zero/context.cpp +++ b/source/adapters/level_zero/context.cpp @@ -125,7 +125,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urContextGetInfo( default: // TODO: implement other parameters - die("urGetContextInfo: unsuppported ParamName."); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } return UR_RESULT_SUCCESS; @@ -584,8 +584,8 @@ ur_context_handle_t_::decrementUnreleasedEventsInPool(ur_event_handle_t Event) { Event->isHostVisible(), Event->isProfilingEnabled(), ZeDevice); // Put the empty pool to the cache of the pools. - if (NumEventsUnreleasedInEventPool[Event->ZeEventPool] == 0) - die("Invalid event release: event pool doesn't have unreleased events"); + assert(NumEventsUnreleasedInEventPool[Event->ZeEventPool] == 0 && + "Invalid event release: event pool doesn't have unreleased events"); if (--NumEventsUnreleasedInEventPool[Event->ZeEventPool] == 0) { if (ZePoolCache->front() != Event->ZeEventPool) { ZePoolCache->push_back(Event->ZeEventPool); diff --git a/source/adapters/level_zero/device.cpp b/source/adapters/level_zero/device.cpp index b1eb1a7b1b..d57ad40404 100644 --- a/source/adapters/level_zero/device.cpp +++ b/source/adapters/level_zero/device.cpp @@ -662,7 +662,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo( case UR_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT: return ReturnValue(MapCaps(Props->sharedSystemAllocCapabilities)); default: - die("urDeviceGetInfo: unexpected ParamName."); + return UR_RESULT_ERROR_INVALID_ENUMERATION; } } diff --git a/source/adapters/level_zero/event.cpp b/source/adapters/level_zero/event.cpp index 7f611208ff..82af762b6e 100644 --- a/source/adapters/level_zero/event.cpp +++ b/source/adapters/level_zero/event.cpp @@ -594,8 +594,8 @@ ur_result_t ur_event_handle_t_::getOrCreateHostVisibleEvent( this->Mutex); if (!HostVisibleEvent) { - if (UrQueue->ZeEventsScope != OnDemandHostVisibleProxy) - die("getOrCreateHostVisibleEvent: missing host-visible event"); + assert(UrQueue->ZeEventsScope != OnDemandHostVisibleProxy && + "getOrCreateHostVisibleEvent: missing host-visible event"); // Submit the command(s) signalling the proxy event to the queue. // We have to first submit a wait for the device-only event for which this @@ -641,9 +641,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventWait( // This ensures that all signalling commands are submitted below and // thus proxy events can be waited without a deadlock. // - ur_event_handle_t_ *Event = ur_cast(e); - if (!Event->hasExternalRefs()) - die("urEventsWait must not be called for an internal event"); + ur_event_handle_t_ *Event = + ur_cast(EventWaitList[I]); + assert(!Event->hasExternalRefs() && + "urEventsWait must not be called for an internal event"); ze_event_handle_t ZeHostVisibleEvent; if (auto Res = Event->getOrCreateHostVisibleEvent(ZeHostVisibleEvent)) @@ -668,13 +669,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventWait( ur_cast(EventWaitList[I]); { std::shared_lock EventLock(Event->Mutex); - if (!Event->hasExternalRefs()) - die("urEventWait must not be called for an internal event"); + assert(!Event->hasExternalRefs() && + "urEventWait must not be called for an internal event"); if (!Event->Completed) { auto HostVisibleEvent = Event->HostVisibleEvent; - if (!HostVisibleEvent) - die("The host-visible proxy event missing"); + assert(!HostVisibleEvent && "The host-visible proxy event missing"); ze_event_handle_t ZeEvent = HostVisibleEvent->ZeEvent; urPrint("ZeEvent = %#llx\n", ur_cast(ZeEvent)); diff --git a/source/adapters/level_zero/image.cpp b/source/adapters/level_zero/image.cpp index 582121d91a..9976765746 100644 --- a/source/adapters/level_zero/image.cpp +++ b/source/adapters/level_zero/image.cpp @@ -316,8 +316,7 @@ ur_result_t ur2zeImageDesc(const ur_image_format_t *ImageFormat, } default: urPrint("format channel order = %d\n", ImageFormat->channelOrder); - die("ur2zeImageDesc: unsupported image channel order\n"); - break; + return UR_RESULT_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR; } ze_image_format_t ZeFormatDesc = { @@ -345,7 +344,7 @@ ur_result_t ur2zeImageDesc(const ur_image_format_t *ImageFormat, break; default: urPrint("ur2zeImageDesc: unsupported image type\n"); - return UR_RESULT_ERROR_INVALID_VALUE; + return UR_RESULT_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR; } ZeImageDesc.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; diff --git a/source/adapters/level_zero/kernel.cpp b/source/adapters/level_zero/kernel.cpp index c40e4ef0e3..cf5ba92288 100644 --- a/source/adapters/level_zero/kernel.cpp +++ b/source/adapters/level_zero/kernel.cpp @@ -629,8 +629,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelGetSubGroupInfo( } else if (PropName == UR_KERNEL_SUB_GROUP_INFO_SUB_GROUP_SIZE_INTEL) { ReturnValue(uint32_t{Kernel->ZeKernelProperties->requiredSubgroupSize}); } else { - die("urKernelGetSubGroupInfo: parameter not implemented"); - return {}; + return UR_RESULT_ERROR_INVALID_ENUMERATION; } return UR_RESULT_SUCCESS; } diff --git a/source/adapters/level_zero/memory.cpp b/source/adapters/level_zero/memory.cpp index 76d53c1d81..bb0636a10c 100644 --- a/source/adapters/level_zero/memory.cpp +++ b/source/adapters/level_zero/memory.cpp @@ -1443,7 +1443,7 @@ static ur_result_t ur2zeImageDesc(const ur_image_format_t *ImageFormat, } default: urPrint("format layout = %d\n", ImageFormat->channelOrder); - die("urMemImageCreate: unsupported image format layout\n"); + return UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT; break; } @@ -1498,8 +1498,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemImageCreate( ) { // TODO: implement read-only, write-only if ((Flags & UR_MEM_FLAG_READ_WRITE) == 0) { - die("urMemImageCreate: Level-Zero implements only read-write buffer," - "no read-only or write-only yet."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } std::shared_lock Lock(Context->Mutex); @@ -1655,7 +1654,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( } else if (Flags == 0 || (Flags == UR_MEM_FLAG_READ_WRITE)) { // Nothing more to do. } else - die("urMemBufferCreate: not implemented"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } *RetBuffer = reinterpret_cast(Buffer); @@ -1713,8 +1712,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferPartition( std::shared_lock Guard(Buffer->Mutex); if (Flags != UR_MEM_FLAG_READ_WRITE) { - die("urMemBufferPartition: Level-Zero implements only read-write buffer," - "no read-only or write-only yet."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } try { @@ -1780,7 +1778,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreateWithNativeHandle( // Memory allocation is unrelated to the context return UR_RESULT_ERROR_INVALID_CONTEXT; default: - die("Unexpected memory type"); + return UR_RESULT_ERROR_MEM_OBJECT_ALLOCATION_FAILURE; } ur_device_handle_t Device{}; @@ -1877,7 +1875,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemGetInfo( return ReturnValue(size_t{Buffer->Size}); } default: { - die("urMemGetInfo: Parameter is not implemented"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } } @@ -2060,8 +2058,8 @@ ur_result_t _ur_buffer::getZeHandle(char *&ZeHandle, access_mode_t AccessMode, // If some prior access invalidated this allocation then make it valid again. if (!Allocation.Valid) { // LastDeviceWithValidAllocation should always have valid allocation. - if (Device == LastDeviceWithValidAllocation) - die("getZeHandle: last used allocation is not valid"); + assert(Device == LastDeviceWithValidAllocation && + "getZeHandle: last used allocation is not valid."); // For write-only access the allocation contents is not going to be used. // So don't do anything to make it "valid". @@ -2195,7 +2193,7 @@ ur_result_t _ur_buffer::free() { ZeUSMImport.doZeUSMRelease(UrContext->getPlatform()->ZeDriver, ZeHandle); break; default: - die("_ur_buffer::free(): Unhandled release action"); + assert(!"_ur_buffer::free(): Unhandled release action"); } ZeHandle = nullptr; // don't leave hanging pointers } diff --git a/source/adapters/level_zero/program.cpp b/source/adapters/level_zero/program.cpp index d60b2a3322..12caf935ed 100644 --- a/source/adapters/level_zero/program.cpp +++ b/source/adapters/level_zero/program.cpp @@ -718,7 +718,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramGetInfo( return UR_RESULT_ERROR_UNKNOWN; } default: - die("urProgramGetInfo: not implemented"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } return UR_RESULT_SUCCESS; diff --git a/source/adapters/level_zero/queue.cpp b/source/adapters/level_zero/queue.cpp index 187f4f75f9..d0090359ed 100644 --- a/source/adapters/level_zero/queue.cpp +++ b/source/adapters/level_zero/queue.cpp @@ -162,13 +162,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueGetInfo( case UR_QUEUE_INFO_REFERENCE_COUNT: return ReturnValue(uint32_t{Queue->RefCount.load()}); case UR_QUEUE_INFO_FLAGS: - die("UR_QUEUE_INFO_FLAGS in urQueueGetInfo not implemented\n"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; break; case UR_QUEUE_INFO_SIZE: - die("UR_QUEUE_INFO_SIZE in urQueueGetInfo not implemented\n"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; break; case UR_QUEUE_INFO_DEVICE_DEFAULT: - die("UR_QUEUE_INFO_DEVICE_DEFAULT in urQueueGetInfo not implemented\n"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; break; case UR_QUEUE_INFO_EMPTY: { // We can exit early if we have in-order queue. @@ -792,7 +792,7 @@ static const zeCommandListBatchConfig ZeCommandListBatchConfig(bool IsCopy) { Config.NumTimesClosedFullThreshold = Val; break; default: - die("Unexpected batch config"); + assert(!"Unexpected batch config"); } if (IsCopy) urPrint("UR_L0_COPY_BATCH_SIZE: dynamic batch param " @@ -910,7 +910,7 @@ ur_queue_handle_t_::ur_queue_handle_t_( ComputeQueueGroup.UpperIndex = FilterUpperIndex; ComputeQueueGroup.NextIndex = ComputeQueueGroup.LowerIndex; } else { - die("No compute queue available/allowed."); + assert(!"No compute queue available/allowed."); } } if (UsingImmCmdLists) { @@ -1065,10 +1065,10 @@ ur_queue_handle_t_::executeCommandList(ur_command_list_ptr_t CommandList, if (OKToBatchCommand && this->isBatchingAllowed(UseCopyEngine) && (!ZeCommandListBatchConfig.dynamic() || !CurrentlyEmpty)) { - if (hasOpenCommandList(UseCopyEngine) && - CommandBatch.OpenCommandList != CommandList) - die("executeCommandList: OpenCommandList should be equal to" - "null or CommandList"); + assert((hasOpenCommandList(UseCopyEngine) && + CommandBatch.OpenCommandList != CommandList) && + "executeCommandList: OpenCommandList should be equal to" + "null or CommandList"); if (CommandList->second.size() < CommandBatch.QueueBatchSize) { CommandBatch.OpenCommandList = CommandList; @@ -1830,7 +1830,7 @@ ur_queue_handle_t_::ur_queue_group_t::getZeQueue(uint32_t *QueueGroupOrdinal) { zeCommandQueueCreate, (Queue->Context->ZeContext, Queue->Device->ZeDevice, &ZeCommandQueueDesc, &ZeQueue)); if (ZeResult) { - die("[L0] getZeQueue: failed to create queue"); + assert(!"[L0] getZeQueue: failed to create queue"); } return ZeQueue; diff --git a/source/adapters/level_zero/usm.cpp b/source/adapters/level_zero/usm.cpp index 4772986c1c..909c9d19d9 100644 --- a/source/adapters/level_zero/usm.cpp +++ b/source/adapters/level_zero/usm.cpp @@ -948,9 +948,8 @@ ur_result_t ZeMemFreeHelper(ur_context_handle_t Context, void *Ptr) { if (IndirectAccessTrackingEnabled) { ContextsLock.lock(); auto It = Context->MemAllocs.find(Ptr); - if (It == std::end(Context->MemAllocs)) { - die("All memory allocations must be tracked!"); - } + assert(It == std::end(Context->MemAllocs) && + "All memory allocations must be tracked!"); if (!It->second.RefCount.decrementAndTest()) { // Memory can't be deallocated yet. return UR_RESULT_SUCCESS; @@ -995,9 +994,8 @@ ur_result_t USMFreeHelper(ur_context_handle_t Context, void *Ptr, if (IndirectAccessTrackingEnabled) { auto It = Context->MemAllocs.find(Ptr); - if (It == std::end(Context->MemAllocs)) { - die("All memory allocations must be tracked!"); - } + assert(It == std::end(Context->MemAllocs) && + "All memory allocations must be tracked!"); if (!It->second.RefCount.decrementAndTest()) { // Memory can't be deallocated yet. return UR_RESULT_SUCCESS; diff --git a/source/adapters/native_cpu/command_buffer.cpp b/source/adapters/native_cpu/command_buffer.cpp index fde6c03b86..a61a88d87b 100644 --- a/source/adapters/native_cpu/command_buffer.cpp +++ b/source/adapters/native_cpu/command_buffer.cpp @@ -20,29 +20,21 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferCreateExp( ur_context_handle_t, ur_device_handle_t, const ur_exp_command_buffer_desc_t *, ur_exp_command_buffer_handle_t *) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferRetainExp(ur_exp_command_buffer_handle_t) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferReleaseExp(ur_exp_command_buffer_handle_t) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferFinalizeExp(ur_exp_command_buffer_handle_t) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -50,10 +42,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendKernelLaunchExp( ur_exp_command_buffer_handle_t, ur_kernel_handle_t, uint32_t, const size_t *, const size_t *, const size_t *, uint32_t, const ur_exp_command_buffer_sync_point_t *, - ur_exp_command_buffer_sync_point_t *, - ur_exp_command_buffer_command_handle_t *) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); + ur_exp_command_buffer_sync_point_t *) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -61,8 +50,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendUSMMemcpyExp( ur_exp_command_buffer_handle_t, void *, const void *, size_t, uint32_t, const ur_exp_command_buffer_sync_point_t *, ur_exp_command_buffer_sync_point_t *) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -70,8 +57,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendMemBufferCopyExp( ur_exp_command_buffer_handle_t, ur_mem_handle_t, ur_mem_handle_t, size_t, size_t, size_t, uint32_t, const ur_exp_command_buffer_sync_point_t *, ur_exp_command_buffer_sync_point_t *) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -80,8 +65,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendMemBufferCopyRectExp( ur_rect_offset_t, ur_rect_offset_t, ur_rect_region_t, size_t, size_t, size_t, size_t, uint32_t, const ur_exp_command_buffer_sync_point_t *, ur_exp_command_buffer_sync_point_t *) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -90,8 +73,6 @@ ur_result_t UR_APICALL urCommandBufferAppendMemBufferWriteExp( ur_exp_command_buffer_handle_t, ur_mem_handle_t, size_t, size_t, const void *, uint32_t, const ur_exp_command_buffer_sync_point_t *, ur_exp_command_buffer_sync_point_t *) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -100,8 +81,6 @@ ur_result_t UR_APICALL urCommandBufferAppendMemBufferReadExp( ur_exp_command_buffer_handle_t, ur_mem_handle_t, size_t, size_t, void *, uint32_t, const ur_exp_command_buffer_sync_point_t *, ur_exp_command_buffer_sync_point_t *) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -111,8 +90,6 @@ ur_result_t UR_APICALL urCommandBufferAppendMemBufferWriteRectExp( ur_rect_offset_t, ur_rect_region_t, size_t, size_t, size_t, size_t, void *, uint32_t, const ur_exp_command_buffer_sync_point_t *, ur_exp_command_buffer_sync_point_t *) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -122,16 +99,12 @@ ur_result_t UR_APICALL urCommandBufferAppendMemBufferReadRectExp( ur_rect_offset_t, ur_rect_region_t, size_t, size_t, size_t, size_t, void *, uint32_t, const ur_exp_command_buffer_sync_point_t *, ur_exp_command_buffer_sync_point_t *) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferEnqueueExp( ur_exp_command_buffer_handle_t, ur_queue_handle_t, uint32_t, const ur_event_handle_t *, ur_event_handle_t *) { - detail::ur::die("Experimental Command-buffer feature is not " - "implemented for the NativeCPU adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/common.cpp b/source/adapters/native_cpu/common.cpp index 90fbc1183a..1fef898b38 100644 --- a/source/adapters/native_cpu/common.cpp +++ b/source/adapters/native_cpu/common.cpp @@ -27,8 +27,3 @@ ur_result_t urGetLastResult(ur_platform_handle_t, const char **ppMessage) { *ppMessage = &ErrorMessage[0]; return ErrorMessageCode; } - -void detail::ur::die(const char *pMessage) { - std::cerr << "ur_die: " << pMessage << '\n'; - std::terminate(); -} diff --git a/source/adapters/native_cpu/common.hpp b/source/adapters/native_cpu/common.hpp index d792cbbbcf..45a756ceba 100644 --- a/source/adapters/native_cpu/common.hpp +++ b/source/adapters/native_cpu/common.hpp @@ -17,45 +17,6 @@ constexpr size_t MaxMessageSize = 256; extern thread_local ur_result_t ErrorMessageCode; extern thread_local char ErrorMessage[MaxMessageSize]; -#define DIE_NO_IMPLEMENTATION \ - if (PrintTrace) { \ - std::cerr << "Not Implemented : " << __FUNCTION__ \ - << " - File : " << __FILE__; \ - std::cerr << " / Line : " << __LINE__ << std::endl; \ - } \ - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; - -#define CONTINUE_NO_IMPLEMENTATION \ - if (PrintTrace) { \ - std::cerr << "Warning : Not Implemented : " << __FUNCTION__ \ - << " - File : " << __FILE__; \ - std::cerr << " / Line : " << __LINE__ << std::endl; \ - } \ - return UR_RESULT_SUCCESS; - -#define CASE_UR_UNSUPPORTED(not_supported) \ - case not_supported: \ - if (PrintTrace) { \ - std::cerr << std::endl \ - << "Unsupported UR case : " << #not_supported << " in " \ - << __FUNCTION__ << ":" << __LINE__ << "(" << __FILE__ << ")" \ - << std::endl; \ - } \ - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; - -/// ------ Error handling, matching OpenCL plugin semantics. -/// Taken from other adapter -namespace detail { -namespace ur { - -// Report error and no return (keeps compiler from printing warnings). -// TODO: Probably change that to throw a catchable exception, -// but for now it is useful to see every failure. -// -[[noreturn]] void die(const char *pMessage); -} // namespace ur -} // namespace detail - // Base class to store common data struct _ur_object { ur_shared_mutex Mutex; diff --git a/source/adapters/native_cpu/context.cpp b/source/adapters/native_cpu/context.cpp index c485725828..9a5d6c6f94 100644 --- a/source/adapters/native_cpu/context.cpp +++ b/source/adapters/native_cpu/context.cpp @@ -75,7 +75,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urContextGetNativeHandle( ur_context_handle_t hContext, ur_native_handle_t *phNativeContext) { std::ignore = hContext; std::ignore = phNativeContext; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urContextCreateWithNativeHandle( @@ -89,7 +89,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urContextCreateWithNativeHandle( std::ignore = pProperties; std::ignore = phContext; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urContextSetExtendedDeleter( @@ -99,5 +99,5 @@ UR_APIEXPORT ur_result_t UR_APICALL urContextSetExtendedDeleter( std::ignore = pfnDeleter; std::ignore = pUserData; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/device.cpp b/source/adapters/native_cpu/device.cpp index dfabfb81e5..664d0259dc 100644 --- a/source/adapters/native_cpu/device.cpp +++ b/source/adapters/native_cpu/device.cpp @@ -300,12 +300,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_ESIMD_SUPPORT: return ReturnValue(false); - case UR_DEVICE_INFO_COMPONENT_DEVICES: - case UR_DEVICE_INFO_COMPOSITE_DEVICE: - // These two are exclusive of L0. - return ReturnValue(0); - - CASE_UR_UNSUPPORTED(UR_DEVICE_INFO_MAX_MEMORY_BANDWIDTH); case UR_DEVICE_INFO_VIRTUAL_MEMORY_SUPPORT: return ReturnValue(false); @@ -314,9 +308,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(false); default: - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_INVALID_ENUMERATION; } - return UR_RESULT_SUCCESS; } UR_APIEXPORT ur_result_t UR_APICALL urDeviceRetain(ur_device_handle_t hDevice) { @@ -342,7 +335,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDevicePartition( std::ignore = phSubDevices; std::ignore = pNumDevicesRet; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetNativeHandle( @@ -350,7 +343,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetNativeHandle( std::ignore = hDevice; std::ignore = phNativeDevice; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urDeviceCreateWithNativeHandle( @@ -362,7 +355,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceCreateWithNativeHandle( std::ignore = pProperties; std::ignore = phDevice; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetGlobalTimestamps( diff --git a/source/adapters/native_cpu/enqueue.cpp b/source/adapters/native_cpu/enqueue.cpp index 75c2caeac0..961e86d160 100644 --- a/source/adapters/native_cpu/enqueue.cpp +++ b/source/adapters/native_cpu/enqueue.cpp @@ -56,7 +56,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch( UR_ASSERT(workDim < 4, UR_RESULT_ERROR_INVALID_WORK_DIMENSION); if (*pGlobalWorkSize == 0) { - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } // TODO: add proper error checking @@ -107,7 +107,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWait( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWaitWithBarrier( @@ -118,7 +118,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWaitWithBarrier( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } template @@ -286,7 +286,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageRead( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageWrite( @@ -306,7 +306,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageWrite( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageCopy( @@ -325,7 +325,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageCopy( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferMap( @@ -410,7 +410,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL @@ -422,7 +422,7 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size, std::ignore = advice; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFill2D( @@ -441,7 +441,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFill2D( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy2D( @@ -461,7 +461,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy2D( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueDeviceGlobalVariableWrite( @@ -480,7 +480,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueDeviceGlobalVariableWrite( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueDeviceGlobalVariableRead( @@ -499,7 +499,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueDeviceGlobalVariableRead( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueReadHostPipe( @@ -517,7 +517,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueReadHostPipe( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueWriteHostPipe( @@ -535,5 +535,5 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueWriteHostPipe( std::ignore = phEventWaitList; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/event.cpp b/source/adapters/native_cpu/event.cpp index 112bb553c0..93bd36ec13 100644 --- a/source/adapters/native_cpu/event.cpp +++ b/source/adapters/native_cpu/event.cpp @@ -23,7 +23,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventGetInfo(ur_event_handle_t hEvent, std::ignore = pPropValue; std::ignore = pPropSizeRet; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEventGetProfilingInfo( @@ -35,7 +35,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventGetProfilingInfo( std::ignore = pPropValue; std::ignore = pPropSizeRet; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL @@ -49,12 +49,12 @@ urEventWait(uint32_t numEvents, const ur_event_handle_t *phEventWaitList) { UR_APIEXPORT ur_result_t UR_APICALL urEventRetain(ur_event_handle_t hEvent) { std::ignore = hEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEventRelease(ur_event_handle_t hEvent) { std::ignore = hEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEventGetNativeHandle( @@ -62,7 +62,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventGetNativeHandle( std::ignore = hEvent; std::ignore = phNativeEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urEventCreateWithNativeHandle( @@ -74,7 +74,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventCreateWithNativeHandle( std::ignore = pProperties; std::ignore = phEvent; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL @@ -85,5 +85,5 @@ urEventSetCallback(ur_event_handle_t hEvent, ur_execution_info_t execStatus, std::ignore = pfnNotify; std::ignore = pUserData; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/kernel.cpp b/source/adapters/native_cpu/kernel.cpp index 5a7a286adc..2e4216c7f4 100644 --- a/source/adapters/native_cpu/kernel.cpp +++ b/source/adapters/native_cpu/kernel.cpp @@ -98,6 +98,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelGetInfo(ur_kernel_handle_t hKernel, default: return UR_RESULT_ERROR_INVALID_VALUE; } + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL @@ -176,7 +177,7 @@ urKernelGetSubGroupInfo(ur_kernel_handle_t hKernel, ur_device_handle_t hDevice, ur::unreachable(); } } - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urKernelRetain(ur_kernel_handle_t hKernel) { @@ -231,7 +232,7 @@ urKernelSetArgSampler(ur_kernel_handle_t hKernel, uint32_t argIndex, std::ignore = pProperties; std::ignore = hArgValue; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL @@ -262,7 +263,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelSetSpecializationConstants( std::ignore = count; std::ignore = pSpecConstants; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urKernelGetNativeHandle( @@ -270,7 +271,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelGetNativeHandle( std::ignore = hKernel; std::ignore = phNativeKernel; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urKernelCreateWithNativeHandle( @@ -284,5 +285,5 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelCreateWithNativeHandle( std::ignore = pProperties; std::ignore = phKernel; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/memory.cpp b/source/adapters/native_cpu/memory.cpp index 1f8a927c67..85f85f4588 100644 --- a/source/adapters/native_cpu/memory.cpp +++ b/source/adapters/native_cpu/memory.cpp @@ -23,7 +23,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemImageCreate( std::ignore = pHost; std::ignore = phMem; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( @@ -65,7 +65,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( UR_APIEXPORT ur_result_t UR_APICALL urMemRetain(ur_mem_handle_t hMem) { std::ignore = hMem; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urMemRelease(ur_mem_handle_t hMem) { @@ -93,8 +93,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferPartition( std::shared_lock Guard(hBuffer->Mutex); if (flags != UR_MEM_FLAG_READ_WRITE) { - die("urMemBufferPartition: NativeCPU implements only read-write buffer," - "no read-only or write-only yet."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } try { @@ -117,7 +116,7 @@ urMemGetNativeHandle(ur_mem_handle_t hMem, ur_device_handle_t hDevice, std::ignore = hDevice; std::ignore = phNativeMem; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreateWithNativeHandle( @@ -128,7 +127,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreateWithNativeHandle( std::ignore = pProperties; std::ignore = phMem; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urMemImageCreateWithNativeHandle( @@ -142,7 +141,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemImageCreateWithNativeHandle( std::ignore = pProperties; std::ignore = phMem; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urMemGetInfo(ur_mem_handle_t hMemory, @@ -156,7 +155,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemGetInfo(ur_mem_handle_t hMemory, std::ignore = pPropValue; std::ignore = pPropSizeRet; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urMemImageGetInfo(ur_mem_handle_t hMemory, @@ -170,5 +169,5 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemImageGetInfo(ur_mem_handle_t hMemory, std::ignore = pPropValue; std::ignore = pPropSizeRet; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/platform.cpp b/source/adapters/native_cpu/platform.cpp index 61093f3eed..9d58a992d7 100644 --- a/source/adapters/native_cpu/platform.cpp +++ b/source/adapters/native_cpu/platform.cpp @@ -81,7 +81,7 @@ urPlatformGetInfo(ur_platform_handle_t hPlatform, ur_platform_info_t propName, // https://github.com/oneapi-src/unified-runtime return ReturnValue(UR_PLATFORM_BACKEND_NATIVE_CPU); default: - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } return UR_RESULT_SUCCESS; @@ -94,7 +94,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urPlatformGetBackendOption( std::ignore = pFrontendOption; std::ignore = ppPlatformOption; - CONTINUE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urPlatformCreateWithNativeHandle( @@ -105,7 +105,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urPlatformCreateWithNativeHandle( std::ignore = pProperties; std::ignore = phPlatform; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urPlatformGetNativeHandle( @@ -113,5 +113,5 @@ UR_APIEXPORT ur_result_t UR_APICALL urPlatformGetNativeHandle( std::ignore = hPlatform; std::ignore = phNativePlatform; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/program.cpp b/source/adapters/native_cpu/program.cpp index f210f210c9..e2182b9b88 100644 --- a/source/adapters/native_cpu/program.cpp +++ b/source/adapters/native_cpu/program.cpp @@ -23,7 +23,7 @@ urProgramCreateWithIL(ur_context_handle_t hContext, const void *pIL, std::ignore = pProperties; std::ignore = phProgram; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urProgramCreateWithBinary( @@ -71,7 +71,7 @@ urProgramCompile(ur_context_handle_t hContext, ur_program_handle_t hProgram, std::ignore = hProgram; std::ignore = pOptions; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL @@ -84,7 +84,7 @@ urProgramLink(ur_context_handle_t hContext, uint32_t count, std::ignore = pOptions; std::ignore = phProgram; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urProgramCompileExp(ur_program_handle_t, @@ -127,7 +127,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramGetFunctionPointer( std::ignore = pFunctionName; std::ignore = ppFunctionPointer; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( @@ -139,7 +139,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramGetGlobalVariablePointer( std::ignore = pGlobalVariableSizeRet; std::ignore = ppGlobalVariablePointerRet; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL @@ -185,7 +185,7 @@ urProgramGetBuildInfo(ur_program_handle_t hProgram, ur_device_handle_t hDevice, std::ignore = pPropValue; std::ignore = pPropSizeRet; - CONTINUE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urProgramSetSpecializationConstants( @@ -195,7 +195,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramSetSpecializationConstants( std::ignore = count; std::ignore = pSpecConstants; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urProgramGetNativeHandle( @@ -203,7 +203,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramGetNativeHandle( std::ignore = hProgram; std::ignore = phNativeProgram; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urProgramCreateWithNativeHandle( @@ -215,5 +215,5 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramCreateWithNativeHandle( std::ignore = pProperties; std::ignore = phProgram; - DIE_NO_IMPLEMENTATION + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/queue.cpp b/source/adapters/native_cpu/queue.cpp index 516e66db64..2fc3a2db27 100644 --- a/source/adapters/native_cpu/queue.cpp +++ b/source/adapters/native_cpu/queue.cpp @@ -25,7 +25,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueGetInfo(ur_queue_handle_t hQueue, std::ignore = pPropValue; std::ignore = pPropSizeRet; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urQueueCreate( @@ -38,7 +38,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueCreate( auto Queue = new ur_queue_handle_t_(); *phQueue = Queue; - CONTINUE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urQueueRetain(ur_queue_handle_t hQueue) { @@ -61,7 +61,7 @@ urQueueGetNativeHandle(ur_queue_handle_t hQueue, ur_queue_native_desc_t *pDesc, std::ignore = pDesc; std::ignore = phNativeQueue; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urQueueCreateWithNativeHandle( @@ -74,7 +74,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueCreateWithNativeHandle( std::ignore = pProperties; std::ignore = phQueue; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urQueueFinish(ur_queue_handle_t hQueue) { @@ -86,5 +86,5 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueFinish(ur_queue_handle_t hQueue) { UR_APIEXPORT ur_result_t UR_APICALL urQueueFlush(ur_queue_handle_t hQueue) { std::ignore = hQueue; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/sampler.cpp b/source/adapters/native_cpu/sampler.cpp index bf9a72ce3b..2534ba4397 100644 --- a/source/adapters/native_cpu/sampler.cpp +++ b/source/adapters/native_cpu/sampler.cpp @@ -19,21 +19,21 @@ urSamplerCreate(ur_context_handle_t hContext, const ur_sampler_desc_t *pDesc, std::ignore = pDesc; std::ignore = phSampler; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urSamplerRetain(ur_sampler_handle_t hSampler) { std::ignore = hSampler; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urSamplerRelease(ur_sampler_handle_t hSampler) { std::ignore = hSampler; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL @@ -45,7 +45,7 @@ urSamplerGetInfo(ur_sampler_handle_t hSampler, ur_sampler_info_t propName, std::ignore = pPropValue; std::ignore = pPropSizeRet; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urSamplerGetNativeHandle( @@ -53,7 +53,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urSamplerGetNativeHandle( std::ignore = hSampler; std::ignore = phNativeSampler; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urSamplerCreateWithNativeHandle( @@ -65,5 +65,5 @@ UR_APIEXPORT ur_result_t UR_APICALL urSamplerCreateWithNativeHandle( std::ignore = pProperties; std::ignore = phSampler; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/usm.cpp b/source/adapters/native_cpu/usm.cpp index 7cdac0cd8f..bda0e901f2 100644 --- a/source/adapters/native_cpu/usm.cpp +++ b/source/adapters/native_cpu/usm.cpp @@ -86,7 +86,7 @@ urUSMGetMemAllocInfo(ur_context_handle_t hContext, const void *pMem, std::ignore = pPropValue; std::ignore = pPropSizeRet; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL @@ -96,21 +96,21 @@ urUSMPoolCreate(ur_context_handle_t hContext, ur_usm_pool_desc_t *pPoolDesc, std::ignore = pPoolDesc; std::ignore = ppPool; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolRetain(ur_usm_pool_handle_t pPool) { std::ignore = pPool; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urUSMPoolRelease(ur_usm_pool_handle_t pPool) { std::ignore = pPool; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL @@ -122,7 +122,7 @@ urUSMPoolGetInfo(ur_usm_pool_handle_t hPool, ur_usm_pool_info_t propName, std::ignore = pPropValue; std::ignore = pPropSizeRet; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urUSMImportExp(ur_context_handle_t Context, @@ -130,12 +130,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urUSMImportExp(ur_context_handle_t Context, std::ignore = Context; std::ignore = HostPtr; std::ignore = Size; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urUSMReleaseExp(ur_context_handle_t Context, void *HostPtr) { std::ignore = Context; std::ignore = HostPtr; - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/native_cpu/usm_p2p.cpp b/source/adapters/native_cpu/usm_p2p.cpp index dbb47d910e..7de083401a 100644 --- a/source/adapters/native_cpu/usm_p2p.cpp +++ b/source/adapters/native_cpu/usm_p2p.cpp @@ -12,16 +12,16 @@ UR_APIEXPORT ur_result_t UR_APICALL urUsmP2PEnablePeerAccessExp(ur_device_handle_t, ur_device_handle_t) { - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urUsmP2PDisablePeerAccessExp(ur_device_handle_t, ur_device_handle_t) { - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urUsmP2PPeerAccessGetInfoExp(ur_device_handle_t, ur_device_handle_t, ur_exp_peer_info_t, size_t, void *, size_t *) { - DIE_NO_IMPLEMENTATION; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/opencl/command_buffer.cpp b/source/adapters/opencl/command_buffer.cpp index ac5650b1a1..08793cbc44 100644 --- a/source/adapters/opencl/command_buffer.cpp +++ b/source/adapters/opencl/command_buffer.cpp @@ -202,9 +202,6 @@ ur_result_t UR_APICALL urCommandBufferAppendMemBufferWriteExp( [[maybe_unused]] const ur_exp_command_buffer_sync_point_t *pSyncPointWaitList, [[maybe_unused]] ur_exp_command_buffer_sync_point_t *pSyncPoint) { - - cl_adapter::die("Experimental Command-buffer feature is not " - "implemented for OpenCL adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/opencl/common.cpp b/source/adapters/opencl/common.cpp index c1668494da..d9234d5fee 100644 --- a/source/adapters/opencl/common.cpp +++ b/source/adapters/opencl/common.cpp @@ -90,11 +90,6 @@ ur_result_t mapCLErrorToUR(cl_int Result) { } } -void cl_adapter::die(const char *Message) { - logger::always("ur_die: {}", Message); - std::terminate(); -} - /// Common API for getting the native handle of a UR object /// /// \param URObj is the UR object to get the native handle of diff --git a/source/adapters/opencl/common.hpp b/source/adapters/opencl/common.hpp index 256fce0c22..cdd4ab2582 100644 --- a/source/adapters/opencl/common.hpp +++ b/source/adapters/opencl/common.hpp @@ -24,7 +24,7 @@ } /** - * Call an UR API and, if the result is not UR_RESULT_SUCCESS, automatically + * Call a UR API and, if the result is not UR_RESULT_SUCCESS, automatically * return from the current function. */ #define UR_RETURN_ON_FAILURE(urCall) \ @@ -157,8 +157,6 @@ extern thread_local char ErrorMessage[MaxMessageSize]; [[maybe_unused]] void setErrorMessage(const char *Message, ur_result_t ErrorCode); -[[noreturn]] void die(const char *Message); - template To cast(From Value) { if constexpr (std::is_pointer_v) { diff --git a/source/adapters/opencl/sampler.cpp b/source/adapters/opencl/sampler.cpp index 49f31b37fd..b49eb74788 100644 --- a/source/adapters/opencl/sampler.cpp +++ b/source/adapters/opencl/sampler.cpp @@ -12,11 +12,13 @@ namespace { -cl_sampler_info ur2CLSamplerInfo(ur_sampler_info_t URInfo) { +ur_result_t ur2CLSamplerInfo(ur_sampler_info_t URInfo, + cl_sampler_info *CLInfo) { switch (URInfo) { #define CASE(UR_INFO, CL_INFO) \ case UR_INFO: \ - return CL_INFO; + *CLInfo = CL_INFO; \ + return UR_RESULT_SUCCESS; CASE(UR_SAMPLER_INFO_REFERENCE_COUNT, CL_SAMPLER_REFERENCE_COUNT) CASE(UR_SAMPLER_INFO_CONTEXT, CL_SAMPLER_CONTEXT) @@ -27,16 +29,17 @@ cl_sampler_info ur2CLSamplerInfo(ur_sampler_info_t URInfo) { #undef CASE default: - cl_adapter::die("Unhandled: ur_sampler_info_t"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } } -cl_addressing_mode ur2CLAddressingMode(ur_sampler_addressing_mode_t Mode) { +ur_result_t ur2CLAddressingMode(ur_sampler_addressing_mode_t Mode, + cl_addressing_mode *CLMode) { switch (Mode) { - #define CASE(UR_MODE, CL_MODE) \ case UR_MODE: \ - return CL_MODE; + *CLMode = CL_MODE; \ + return UR_RESULT_SUCCESS; CASE(UR_SAMPLER_ADDRESSING_MODE_NONE, CL_ADDRESS_NONE); CASE(UR_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE, CL_ADDRESS_CLAMP_TO_EDGE); @@ -48,16 +51,17 @@ cl_addressing_mode ur2CLAddressingMode(ur_sampler_addressing_mode_t Mode) { #undef CASE default: - cl_adapter::die("Unhandled: ur_sampler_addressing_mode_t"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } } -cl_filter_mode ur2CLFilterMode(ur_sampler_filter_mode_t Mode) { +ur_result_t ur2CLFilterMode(ur_sampler_filter_mode_t Mode, + cl_filter_mode *CLMode) { switch (Mode) { - #define CASE(UR_MODE, CL_MODE) \ case UR_MODE: \ - return CL_MODE; + *CLMode = CL_MODE; \ + return UR_RESULT_SUCCESS; CASE(UR_SAMPLER_FILTER_MODE_NEAREST, CL_FILTER_NEAREST) CASE(UR_SAMPLER_FILTER_MODE_LINEAR, CL_FILTER_LINEAR) @@ -65,16 +69,17 @@ cl_filter_mode ur2CLFilterMode(ur_sampler_filter_mode_t Mode) { #undef CASE default: - cl_adapter::die("Unhandled: ur_sampler_filter_mode_t"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } } -ur_sampler_addressing_mode_t cl2URAddressingMode(cl_addressing_mode Mode) { +ur_result_t cl2URAddressingMode(cl_addressing_mode Mode, + ur_sampler_addressing_mode_t *URMode) { switch (Mode) { - #define CASE(CL_MODE, UR_MODE) \ case CL_MODE: \ - return UR_MODE; + *URMode = UR_MODE; \ + return UR_RESULT_SUCCESS; CASE(CL_ADDRESS_NONE, UR_SAMPLER_ADDRESSING_MODE_NONE); CASE(CL_ADDRESS_CLAMP_TO_EDGE, UR_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE); @@ -86,15 +91,17 @@ ur_sampler_addressing_mode_t cl2URAddressingMode(cl_addressing_mode Mode) { #undef CASE default: - cl_adapter::die("Unhandled: cl_addressing_mode"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } } -ur_sampler_filter_mode_t cl2URFilterMode(cl_filter_mode Mode) { +ur_result_t cl2URFilterMode(cl_filter_mode Mode, + ur_sampler_filter_mode_t *URMode) { switch (Mode) { #define CASE(CL_MODE, UR_MODE) \ case CL_MODE: \ - return UR_MODE; + *URMode = UR_MODE; \ + return UR_RESULT_SUCCESS; CASE(CL_FILTER_NEAREST, UR_SAMPLER_FILTER_MODE_NEAREST) CASE(CL_FILTER_LINEAR, UR_SAMPLER_FILTER_MODE_LINEAR); @@ -102,7 +109,7 @@ ur_sampler_filter_mode_t cl2URFilterMode(cl_filter_mode Mode) { #undef CASE default: - cl_adapter::die("Unhandled: cl_filter_mode"); + return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } } @@ -114,14 +121,17 @@ void cl2URSamplerInfoValue(cl_sampler_info Info, void *InfoValue) { case CL_SAMPLER_ADDRESSING_MODE: { cl_addressing_mode CLValue = *reinterpret_cast(InfoValue); + ur_sampler_addressing_mode_t AddressingMode; + cl2URAddressingMode(CLValue, &AddressingMode); *reinterpret_cast(InfoValue) = - cl2URAddressingMode(CLValue); + AddressingMode; break; } case CL_SAMPLER_FILTER_MODE: { cl_filter_mode CLMode = *reinterpret_cast(InfoValue); - *reinterpret_cast(InfoValue) = - cl2URFilterMode(CLMode); + ur_sampler_filter_mode_t FilterMode; + cl2URFilterMode(CLMode, &FilterMode); + *reinterpret_cast(InfoValue) = FilterMode; break; } @@ -138,9 +148,11 @@ ur_result_t urSamplerCreate(ur_context_handle_t hContext, // Initialize properties according to OpenCL 2.1 spec. ur_result_t ErrorCode; - cl_addressing_mode AddressingMode = - ur2CLAddressingMode(pDesc->addressingMode); - cl_filter_mode FilterMode = ur2CLFilterMode(pDesc->filterMode); + cl_addressing_mode AddressingMode; + UR_RETURN_ON_FAILURE( + ur2CLAddressingMode(pDesc->addressingMode, &AddressingMode)); + cl_filter_mode FilterMode; + UR_RETURN_ON_FAILURE(ur2CLFilterMode(pDesc->filterMode, &FilterMode)); // Always call OpenCL 1.0 API *phSampler = cl_adapter::cast(clCreateSampler( @@ -154,7 +166,8 @@ ur_result_t urSamplerCreate(ur_context_handle_t hContext, UR_APIEXPORT ur_result_t UR_APICALL urSamplerGetInfo(ur_sampler_handle_t hSampler, ur_sampler_info_t propName, size_t propSize, void *pPropValue, size_t *pPropSizeRet) { - cl_sampler_info SamplerInfo = ur2CLSamplerInfo(propName); + cl_sampler_info SamplerInfo; + UR_RETURN_ON_FAILURE(ur2CLSamplerInfo(propName, &SamplerInfo)); static_assert(sizeof(cl_addressing_mode) == sizeof(ur_sampler_addressing_mode_t)); diff --git a/source/adapters/opencl/usm_p2p.cpp b/source/adapters/opencl/usm_p2p.cpp index b0f51eac2b..9ee9920ad8 100644 --- a/source/adapters/opencl/usm_p2p.cpp +++ b/source/adapters/opencl/usm_p2p.cpp @@ -13,18 +13,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urUsmP2PEnablePeerAccessExp([[maybe_unused]] ur_device_handle_t commandDevice, [[maybe_unused]] ur_device_handle_t peerDevice) { - - cl_adapter::die( - "Experimental P2P feature is not implemented for OpenCL adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } UR_APIEXPORT ur_result_t UR_APICALL urUsmP2PDisablePeerAccessExp([[maybe_unused]] ur_device_handle_t commandDevice, [[maybe_unused]] ur_device_handle_t peerDevice) { - - cl_adapter::die( - "Experimental P2P feature is not implemented for OpenCL adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -34,8 +28,5 @@ UR_APIEXPORT ur_result_t UR_APICALL urUsmP2PPeerAccessGetInfoExp( [[maybe_unused]] ur_exp_peer_info_t propName, [[maybe_unused]] size_t propSize, [[maybe_unused]] void *pPropValue, [[maybe_unused]] size_t *pPropSizeRet) { - - cl_adapter::die( - "Experimental P2P feature is not implemented for OpenCL adapter."); return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/loader/layers/sanitizer/ur_sanddi.cpp b/source/loader/layers/sanitizer/ur_sanddi.cpp index f4e4f3ef01..b061ecda1a 100644 --- a/source/loader/layers/sanitizer/ur_sanddi.cpp +++ b/source/loader/layers/sanitizer/ur_sanddi.cpp @@ -12,6 +12,7 @@ #include "asan_interceptor.hpp" #include "ur_sanitizer_layer.hpp" +#include namespace ur_sanitizer_layer { @@ -541,11 +542,11 @@ ur_result_t context_t::init(ur_dditable_t *dditable, if (context.enabledType == SanitizerType::AddressSanitizer) { if (!(dditable->VirtualMem.pfnReserve && dditable->VirtualMem.pfnMap && dditable->VirtualMem.pfnGranularityGetInfo)) { - die("Some VirtualMem APIs are needed to enable UR_LAYER_ASAN"); + assert(!"Some VirtualMem APIs are needed to enable UR_LAYER_ASAN"); } if (!dditable->PhysicalMem.pfnCreate) { - die("Some PhysicalMem APIs are needed to enable UR_LAYER_ASAN"); + assert(!"Some PhysicalMem APIs are needed to enable UR_LAYER_ASAN"); } } diff --git a/source/ur/ur.hpp b/source/ur/ur.hpp index 48e611dda8..8a2db9ee19 100644 --- a/source/ur/ur.hpp +++ b/source/ur/ur.hpp @@ -54,12 +54,6 @@ const ur_command_t UR_EXT_COMMAND_TYPE_USER = #define __SYCL_UR_PROGRAM_METADATA_GLOBAL_ID_MAPPING "@global_id_mapping" #define __SYCL_UR_PROGRAM_METADATA_TAG_NEED_FINALIZATION "Requires finalization" -// Terminates the process with a catastrophic error message. -[[noreturn]] inline void die(const char *Message) { - std::cerr << "die: " << Message << std::endl; - std::terminate(); -} - // A single-threaded app has an opportunity to enable this mode to avoid // overhead from mutex locking. Default value is 0 which means that single // thread mode is disabled.