//===------------------- enqueue_kernel.cpp ---------------------*- C++ -*-===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// // // SYCL error handling of enqueue kernel operations // //===----------------------------------------------------------------------===// #include "error_handling.hpp" #include #include #include namespace sycl { inline namespace _V1 { namespace detail::enqueue_kernel_launch { void handleOutOfResources(const device_impl &DeviceImpl, ur_kernel_handle_t Kernel, const NDRDescT &NDRDesc) { sycl::platform Platform = DeviceImpl.get_platform(); sycl::backend Backend = Platform.get_backend(); if (Backend == sycl::backend::ext_oneapi_cuda) { // CUDA: // UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE is returned when the kernel // registers required for the launch config exceeds the maximum number of // registers per block (UR_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP). // This is if local_work_size[0] * ... * local_work_size[work_dim - 1] // multiplied by UR_KERNEL_INFO_NUM_REGS is greater than the value of // UR_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP. See Table 15: Technical // Specifications per Compute Capability, for limitations. const size_t TotalNumberOfWIs = NDRDesc.LocalSize[0] * NDRDesc.LocalSize[1] * NDRDesc.LocalSize[2]; const AdapterPtr &Adapter = DeviceImpl.getAdapter(); uint32_t NumRegisters = 0; Adapter->call(Kernel, UR_KERNEL_INFO_NUM_REGS, sizeof(NumRegisters), &NumRegisters, nullptr); uint32_t MaxRegistersPerBlock = DeviceImpl.get_info(); const bool HasExceededAvailableRegisters = TotalNumberOfWIs * NumRegisters > MaxRegistersPerBlock; if (HasExceededAvailableRegisters) { std::string message( "Exceeded the number of registers available on the hardware.\n"); throw sycl::exception( sycl::make_error_code(sycl::errc::nd_range), // Additional information which can be helpful to the user. message.append( "\tThe number registers per work-group cannot exceed " + std::to_string(MaxRegistersPerBlock) + " for this kernel on this device.\n" "\tThe kernel uses " + std::to_string(NumRegisters) + " registers per work-item for a total of " + std::to_string(TotalNumberOfWIs) + " work-items per work-group.\n")); } } // Fallback constexpr ur_result_t Error = UR_RESULT_ERROR_OUT_OF_RESOURCES; throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "UR backend failed. UR backend returns:" + codeToString(Error)); } void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, ur_kernel_handle_t Kernel, const NDRDescT &NDRDesc) { sycl::platform Platform = DeviceImpl.get_platform(); // Some of the error handling below is special for particular OpenCL // versions. If this is an OpenCL backend, get the version. bool IsOpenCL = false; // Backend is any OpenCL version bool IsOpenCLV1x = false; // Backend is OpenCL 1.x bool IsOpenCLVGE20 = false; // Backend is Greater or Equal to OpenCL 2.0 bool IsLevelZero = false; // Backend is any OneAPI Level 0 version auto Backend = Platform.get_backend(); if (Backend == sycl::backend::opencl) { std::string VersionString = DeviceImpl.get_info().substr(7, 3); IsOpenCL = true; IsOpenCLV1x = (VersionString.find("1.") == 0); IsOpenCLVGE20 = (VersionString.find("2.") == 0) || (VersionString.find("3.") == 0); } else if (Backend == sycl::backend::ext_oneapi_level_zero) { IsLevelZero = true; } const AdapterPtr &Adapter = DeviceImpl.getAdapter(); ur_device_handle_t Device = DeviceImpl.getHandleRef(); size_t CompileWGSize[3] = {0}; Adapter->call( Kernel, Device, UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE, sizeof(size_t) * 3, CompileWGSize, nullptr); size_t CompileMaxWGSize[3] = {0}; ur_result_t URRes = Adapter->call_nocheck( Kernel, Device, UR_KERNEL_GROUP_INFO_COMPILE_MAX_WORK_GROUP_SIZE, sizeof(size_t) * 3, CompileMaxWGSize, nullptr); if (URRes != UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION) { Adapter->checkUrResult(URRes); } size_t CompileMaxLinearWGSize = 0; URRes = Adapter->call_nocheck( Kernel, Device, UR_KERNEL_GROUP_INFO_COMPILE_MAX_LINEAR_WORK_GROUP_SIZE, sizeof(size_t), &CompileMaxLinearWGSize, nullptr); if (URRes != UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION) { Adapter->checkUrResult(URRes); } size_t MaxWGSize = 0; Adapter->call( Device, UR_DEVICE_INFO_MAX_WORK_GROUP_SIZE, sizeof(size_t), &MaxWGSize, nullptr); const bool HasLocalSize = (NDRDesc.LocalSize[0] != 0); if (CompileWGSize[0] != 0) { if (CompileWGSize[0] > MaxWGSize || CompileWGSize[1] > MaxWGSize || CompileWGSize[2] > MaxWGSize) throw sycl::exception( make_error_code(errc::kernel_not_supported), "Submitting a kernel decorated with reqd_work_group_size attribute " "to a device that does not support this work group size is invalid."); // OpenCL 1.x && 2.0: // UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is NULL and // the reqd_work_group_size attribute is used to declare the work-group size // for kernel in the program source. if (!HasLocalSize && (IsOpenCLV1x || IsOpenCLVGE20)) { throw sycl::exception( make_error_code(errc::nd_range), "OpenCL 1.x and 2.0 requires to pass local size argument even if " "required work-group size was specified in the program source"); } // UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified // and does not match the required work-group size for kernel in the program // source. if (NDRDesc.LocalSize[0] != CompileWGSize[0] || NDRDesc.LocalSize[1] != CompileWGSize[1] || NDRDesc.LocalSize[2] != CompileWGSize[2]) throw sycl::exception( make_error_code(errc::nd_range), "The specified local size {" + std::to_string(NDRDesc.LocalSize[2]) + ", " + std::to_string(NDRDesc.LocalSize[1]) + ", " + std::to_string(NDRDesc.LocalSize[0]) + "} doesn't match the required work-group size specified " "in the program source {" + std::to_string(CompileWGSize[2]) + ", " + std::to_string(CompileWGSize[1]) + ", " + std::to_string(CompileWGSize[0]) + "}"); } const size_t TotalNumberOfWIs = NDRDesc.LocalSize[0] * NDRDesc.LocalSize[1] * NDRDesc.LocalSize[2]; if (HasLocalSize) { if (CompileMaxWGSize[0] != 0) { if (NDRDesc.LocalSize[0] > CompileMaxWGSize[0] || NDRDesc.LocalSize[1] > CompileMaxWGSize[1] || NDRDesc.LocalSize[2] > CompileMaxWGSize[2]) { throw sycl::exception( make_error_code(errc::nd_range), "The specified local size {" + std::to_string(NDRDesc.LocalSize[2]) + ", " + std::to_string(NDRDesc.LocalSize[1]) + ", " + std::to_string(NDRDesc.LocalSize[0]) + "} exceeds the maximum work-group size specified " "in the program source {" + std::to_string(CompileMaxWGSize[2]) + ", " + std::to_string(CompileMaxWGSize[1]) + ", " + std::to_string(CompileMaxWGSize[0]) + "}"); } } size_t MaxThreadsPerBlock[3] = {}; Adapter->call( Device, UR_DEVICE_INFO_MAX_WORK_ITEM_SIZES, sizeof(MaxThreadsPerBlock), MaxThreadsPerBlock, nullptr); for (size_t I = 0; I < 3; ++I) { if (MaxThreadsPerBlock[I] < NDRDesc.LocalSize[I]) { throw sycl::exception(make_error_code(errc::nd_range), "The number of work-items in each dimension of a " "work-group cannot exceed {" + std::to_string(MaxThreadsPerBlock[0]) + ", " + std::to_string(MaxThreadsPerBlock[1]) + ", " + std::to_string(MaxThreadsPerBlock[2]) + "} for this device"); } } if (CompileMaxLinearWGSize && TotalNumberOfWIs > CompileMaxLinearWGSize) { throw sycl::exception( make_error_code(errc::nd_range), "The total number of work-items in the work-group (" + std::to_string(TotalNumberOfWIs) + ") exceeds the maximum specified in the program source (" + std::to_string(CompileMaxLinearWGSize) + ")"); } } if (IsOpenCLV1x) { // OpenCL 1.x: // UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified // and the total number of work-items in the work-group computed as // local_work_size[0] * ... * local_work_size[work_dim - 1] is greater // than the value specified by UR_DEVICE_INFO_MAX_WORK_GROUP_SIZE in // table 4.3 if (TotalNumberOfWIs > MaxWGSize) throw sycl::exception( make_error_code(errc::nd_range), "Total number of work-items in a work-group cannot exceed " + std::to_string(MaxWGSize)); } else if (IsOpenCLVGE20 || IsLevelZero) { // OpenCL 2.x or OneAPI Level Zero: // UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified // and the total number of work-items in the work-group computed as // local_work_size[0] * ... * local_work_size[work_dim - 1] is greater // than the value specified by UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE in // table 5.21. size_t KernelWGSize = 0; Adapter->call( Kernel, Device, UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE, sizeof(size_t), &KernelWGSize, nullptr); if (TotalNumberOfWIs > KernelWGSize) throw sycl::exception( make_error_code(errc::nd_range), "Total number of work-items in a work-group cannot exceed " + std::to_string(KernelWGSize) + " for this kernel"); } else { // TODO: Should probably have something similar for the other backends } if (HasLocalSize) { // Is the global range size evenly divisible by the local workgroup size? const bool NonUniformWGs = (NDRDesc.LocalSize[0] != 0 && NDRDesc.GlobalSize[0] % NDRDesc.LocalSize[0] != 0) || (NDRDesc.LocalSize[1] != 0 && NDRDesc.GlobalSize[1] % NDRDesc.LocalSize[1] != 0) || (NDRDesc.LocalSize[2] != 0 && NDRDesc.GlobalSize[2] % NDRDesc.LocalSize[2] != 0); // Is the local size of the workgroup greater than the global range size in // any dimension? if (IsOpenCL) { const bool LocalExceedsGlobal = NonUniformWGs && (NDRDesc.LocalSize[0] > NDRDesc.GlobalSize[0] || NDRDesc.LocalSize[1] > NDRDesc.GlobalSize[1] || NDRDesc.LocalSize[2] > NDRDesc.GlobalSize[2]); if (NonUniformWGs) { if (IsOpenCLV1x) { // OpenCL 1.x: // UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is // specified and number of workitems specified by global_work_size is // not evenly divisible by size of work-group given by local_work_size if (LocalExceedsGlobal) throw sycl::exception(make_error_code(errc::nd_range), "Local workgroup size cannot be greater than " "global range in any dimension"); else throw sycl::exception(make_error_code(errc::nd_range), "Global_work_size must be evenly divisible " "by local_work_size. Non-uniform work-groups " "are not supported by the target device"); } else { // OpenCL 2.x: // UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE if the program was compiled // with –cl-uniform-work-group-size and the number of work-items // specified by global_work_size is not evenly divisible by size of // work-group given by local_work_size ur_program_handle_t Program = nullptr; Adapter->call( Kernel, UR_KERNEL_INFO_PROGRAM, sizeof(ur_program_handle_t), &Program, nullptr); size_t OptsSize = 0; Adapter->call( Program, Device, UR_PROGRAM_BUILD_INFO_OPTIONS, 0, nullptr, &OptsSize); std::string Opts(OptsSize, '\0'); Adapter->call( Program, Device, UR_PROGRAM_BUILD_INFO_OPTIONS, OptsSize, &Opts.front(), nullptr); const bool HasStd20 = Opts.find("-cl-std=CL2.0") != std::string::npos; const bool RequiresUniformWGSize = Opts.find("-cl-uniform-work-group-size") != std::string::npos; std::string LocalWGSize = std::to_string(NDRDesc.LocalSize[0]) + ", " + std::to_string(NDRDesc.LocalSize[1]) + ", " + std::to_string(NDRDesc.LocalSize[2]); std::string GlobalWGSize = std::to_string(NDRDesc.GlobalSize[0]) + ", " + std::to_string(NDRDesc.GlobalSize[1]) + ", " + std::to_string(NDRDesc.GlobalSize[2]); std::string message = LocalExceedsGlobal ? "Local work-group size {" + LocalWGSize + "} is greater than global range size {" + GlobalWGSize + "}. " : "Global work size {" + GlobalWGSize + "} is not evenly divisible by local work-group size {" + LocalWGSize + "}. "; if (!HasStd20) throw sycl::exception( make_error_code(errc::nd_range), message.append("Non-uniform work-groups are not allowed by " "default. Underlying OpenCL 2.x implementation " "supports this feature and to enable it, build " "device program with -cl-std=CL2.0")); else if (RequiresUniformWGSize) throw sycl::exception( make_error_code(errc::nd_range), message.append("Non-uniform work-groups are not allowed when " "-cl-uniform-work-group-size flag is used. " "Underlying OpenCL 2.x implementation supports " "this feature, but it is being disabled by " "-cl-uniform-work-group-size build flag")); // else unknown. fallback (below) } } } else { // TODO: Decide what checks (if any) we need for the other backends } throw sycl::exception( make_error_code(errc::nd_range), "Non-uniform work-groups are not supported by the target device"); } // TODO: required number of sub-groups, OpenCL 2.1: // UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified and // is not consistent with the required number of sub-groups for kernel in the // program source. throw exception(make_error_code(errc::nd_range), "internal error: expected HasLocalSize"); } void handleInvalidWorkItemSize(const device_impl &DeviceImpl, const NDRDescT &NDRDesc) { const AdapterPtr &Adapter = DeviceImpl.getAdapter(); ur_device_handle_t Device = DeviceImpl.getHandleRef(); size_t MaxWISize[] = {0, 0, 0}; Adapter->call( Device, UR_DEVICE_INFO_MAX_WORK_ITEM_SIZES, sizeof(MaxWISize), &MaxWISize, nullptr); for (unsigned I = 0; I < NDRDesc.Dims; I++) { if (NDRDesc.LocalSize[I] > MaxWISize[I]) throw sycl::exception( make_error_code(errc::nd_range), "Number of work-items in a work-group exceed limit for dimension " + std::to_string(I) + " : " + std::to_string(NDRDesc.LocalSize[I]) + " > " + std::to_string(MaxWISize[I])); } } void handleInvalidValue(const device_impl &DeviceImpl, const NDRDescT &NDRDesc) { const AdapterPtr &Adapter = DeviceImpl.getAdapter(); ur_device_handle_t Device = DeviceImpl.getHandleRef(); size_t MaxNWGs[] = {0, 0, 0}; Adapter->call(Device, UR_DEVICE_INFO_MAX_WORK_GROUPS_3D, sizeof(MaxNWGs), &MaxNWGs, nullptr); for (unsigned int I = 0; I < NDRDesc.Dims; I++) { size_t NWgs = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I]; if (NWgs > MaxNWGs[I]) throw sycl::exception( make_error_code(errc::nd_range), "Number of work-groups exceed limit for dimension " + std::to_string(I) + " : " + std::to_string(NWgs) + " > " + std::to_string(MaxNWGs[I])); } // fallback throw exception(make_error_code(errc::nd_range), "unknown internal error"); } void handleErrorOrWarning(ur_result_t Error, const device_impl &DeviceImpl, ur_kernel_handle_t Kernel, const NDRDescT &NDRDesc) { assert(Error != UR_RESULT_SUCCESS && "Success is expected to be handled on caller side"); switch (Error) { case UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY: case UR_RESULT_ERROR_OUT_OF_RESOURCES: return handleOutOfResources(DeviceImpl, Kernel, NDRDesc); case UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE: return handleInvalidWorkGroupSize(DeviceImpl, Kernel, NDRDesc); case UR_RESULT_ERROR_INVALID_KERNEL_ARGS: throw detail::set_ur_error( sycl::exception( make_error_code(errc::kernel_argument), "The kernel argument values have not been specified OR a kernel " "argument declared to be a pointer to a type."), UR_RESULT_ERROR_INVALID_KERNEL_ARGS); case UR_RESULT_ERROR_INVALID_WORK_ITEM_SIZE: return handleInvalidWorkItemSize(DeviceImpl, NDRDesc); case UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT: throw detail::set_ur_error( sycl::exception( make_error_code(errc::feature_not_supported), "image object is specified as an argument value and the image " "format is not supported by device associated with queue"), UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT); case UR_RESULT_ERROR_MISALIGNED_SUB_BUFFER_OFFSET: throw detail::set_ur_error( sycl::exception(make_error_code(errc::invalid), "a sub-buffer object is specified as the value for an " "argument that is a buffer object and the offset " "specified when the sub-buffer object is created is " "not aligned to CL_DEVICE_MEM_BASE_ADDR_ALIGN value " "for device associated with queue"), UR_RESULT_ERROR_MISALIGNED_SUB_BUFFER_OFFSET); case UR_RESULT_ERROR_MEM_OBJECT_ALLOCATION_FAILURE: throw detail::set_ur_error( sycl::exception( make_error_code(errc::memory_allocation), "failure to allocate memory for data store associated with image " "or buffer objects specified as arguments to kernel"), UR_RESULT_ERROR_MEM_OBJECT_ALLOCATION_FAILURE); case UR_RESULT_ERROR_INVALID_IMAGE_SIZE: throw detail::set_ur_error( sycl::exception( make_error_code(errc::invalid), "image object is specified as an argument value and the image " "dimensions (image width, height, specified or compute row and/or " "slice pitch) are not supported by device associated with queue"), UR_RESULT_ERROR_INVALID_IMAGE_SIZE); case UR_RESULT_ERROR_INVALID_VALUE: return handleInvalidValue(DeviceImpl, NDRDesc); case UR_RESULT_ERROR_ADAPTER_SPECIFIC: // checkUrResult does all the necessary handling for // UR_RESULT_ERROR_ADAPTER_SPECIFIC_ERROR, making sure an error is thrown or // not, depending on whether UR_RESULT_ERROR_ADAPTER_SPECIFIC_ERROR contains // an error or a warning. It also ensures that the contents of the error // message buffer (used only by UR_RESULT_ERROR_ADAPTER_SPECIFIC_ERROR) get // handled correctly. return DeviceImpl.getAdapter()->checkUrResult(Error); // TODO: Handle other error codes default: throw detail::set_ur_error( exception(make_error_code(errc::runtime), "UR error"), Error); } } } // namespace detail::enqueue_kernel_launch namespace detail::kernel_get_group_info { void handleErrorOrWarning(ur_result_t Error, ur_kernel_group_info_t Descriptor, const AdapterPtr &Adapter) { assert(Error != UR_RESULT_SUCCESS && "Success is expected to be handled on caller side"); switch (Error) { case UR_RESULT_ERROR_INVALID_VALUE: if (Descriptor == UR_KERNEL_GROUP_INFO_GLOBAL_WORK_SIZE) throw sycl::exception( sycl::make_error_code(errc::invalid), "info::kernel_device_specific::global_work_size descriptor may only " "be used if the device type is device_type::custom or if the kernel " "is a built-in kernel."); break; // TODO: Handle other error codes default: Adapter->checkUrResult(Error); break; } } } // namespace detail::kernel_get_group_info } // namespace _V1 } // namespace sycl