[llvm] [Offload] Use new error code handling mechanism and lower-case messages (PR #139275)

Ross Brunton via llvm-commits llvm-commits at lists.llvm.org
Tue May 20 02:49:18 PDT 2025


https://github.com/RossBrunton updated https://github.com/llvm/llvm-project/pull/139275

>From 149bee88daf3a97105a1d471aca6f74b8088dd3d Mon Sep 17 00:00:00 2001
From: Ross Brunton <ross at codeplay.com>
Date: Fri, 2 May 2025 12:45:47 +0100
Subject: [PATCH 1/3] [Offload] Use new error code handling mechanism

This removes the old ErrorCode-less error method and requires
every user to provide a concrete error code. All calls have been
updated.
---
 offload/include/Shared/OffloadErrcodes.inc    | 44 +++++----
 offload/liboffload/API/Common.td              | 15 +++-
 .../liboffload/include/generated/OffloadAPI.h | 47 ++++++----
 .../include/generated/OffloadPrint.hpp        | 37 ++++++--
 offload/liboffload/src/OffloadImpl.cpp        |  2 +-
 .../plugins-nextgen/amdgpu/dynamic_hsa/hsa.h  |  1 +
 offload/plugins-nextgen/amdgpu/src/rtl.cpp    | 90 +++++++++++++------
 .../common/include/PluginInterface.h          | 27 ++++--
 .../common/src/GlobalHandler.cpp              | 33 ++++---
 offload/plugins-nextgen/common/src/JIT.cpp    |  3 +-
 .../common/src/PluginInterface.cpp            | 67 +++++++++-----
 offload/plugins-nextgen/common/src/RPC.cpp    |  1 +
 offload/plugins-nextgen/cuda/src/rtl.cpp      | 72 ++++++++++-----
 offload/plugins-nextgen/host/src/rtl.cpp      | 37 +++++---
 .../OffloadAPI/kernel/olGetKernel.cpp         |  4 +-
 15 files changed, 332 insertions(+), 148 deletions(-)

diff --git a/offload/include/Shared/OffloadErrcodes.inc b/offload/include/Shared/OffloadErrcodes.inc
index 217565647c4e1..e0797e2bac194 100644
--- a/offload/include/Shared/OffloadErrcodes.inc
+++ b/offload/include/Shared/OffloadErrcodes.inc
@@ -16,22 +16,36 @@
 
 OFFLOAD_ERRC(SUCCESS, "Success", 0)
 OFFLOAD_ERRC(UNKNOWN, "Unknown or internal error", 1)
+OFFLOAD_ERRC(HOST_IO, "I/O error on host", 2)
+OFFLOAD_ERRC(INVALID_BINARY, "A provided binary image is malformed", 3)
 OFFLOAD_ERRC(INVALID_NULL_POINTER,
-             "A pointer argument is null when it should not be", 2)
-OFFLOAD_ERRC(INVALID_ARGUMENT, "An argument is invalid", 3)
-OFFLOAD_ERRC(OUT_OF_RESOURCES, "Out of resources", 4)
-OFFLOAD_ERRC(UNSUPPORTED,
-             "generic error code for unsupported features and enums", 5)
+             "A pointer argument is null when it should not be", 4)
+OFFLOAD_ERRC(INVALID_ARGUMENT, "An argument is invalid", 5)
+OFFLOAD_ERRC(NOT_FOUND, "Requested object was not found in the binary image", 6)
+OFFLOAD_ERRC(OUT_OF_RESOURCES, "Out of resources", 7)
 OFFLOAD_ERRC(
     INVALID_SIZE,
     "invalid size or dimensions (e.g., must not be zero, or is out of bounds)",
-    6)
-OFFLOAD_ERRC(INVALID_ENUMERATION, "enumerator argument is not valid", 7)
-OFFLOAD_ERRC(INVALID_KERNEL_NAME,
-             "Named kernel not found in the program binary", 8)
-OFFLOAD_ERRC(INVALID_VALUE, "Invalid Value", 9)
-OFFLOAD_ERRC(INVALID_PLATFORM, "Invalid platform", 10)
-OFFLOAD_ERRC(INVALID_DEVICE, "Invalid device", 11)
-OFFLOAD_ERRC(INVALID_QUEUE, "Invalid queue", 12)
-OFFLOAD_ERRC(INVALID_EVENT, "Invalid event", 13)
-OFFLOAD_ERRC(INVALID_NULL_HANDLE, "handle argument is not valid", 14)
+    8)
+OFFLOAD_ERRC(INVALID_ENUMERATION, "enumerator argument is not valid", 9)
+OFFLOAD_ERRC(HOST_TOOL_NOT_FOUND,
+             "A required binary (linker, etc.) was not found on the host", 10)
+OFFLOAD_ERRC(INVALID_VALUE, "Invalid Value", 11)
+OFFLOAD_ERRC(UNIMPLEMENTED,
+             "Generic error code for features currently unimplemented by the "
+             "device/backend",
+             12)
+OFFLOAD_ERRC(
+    UNSUPPORTED,
+    "Generic error code for features unsupported by the device/backend", 13)
+OFFLOAD_ERRC(ASSEMBLE_FAILURE,
+             "Assembler failure while processing binary image", 14)
+OFFLOAD_ERRC(LINK_FAILURE, "Linker failure while processing binary image", 15)
+OFFLOAD_ERRC(BACKEND_FAILURE,
+             "The plugin backend is in an invalid or unsupported state", 16)
+OFFLOAD_ERRC(INVALID_NULL_HANDLE,
+             "A handle argument is null when it should not be", 17)
+OFFLOAD_ERRC(INVALID_PLATFORM, "Invalid platform", 18)
+OFFLOAD_ERRC(INVALID_DEVICE, "Invalid device", 19)
+OFFLOAD_ERRC(INVALID_QUEUE, "Invalid queue", 20)
+OFFLOAD_ERRC(INVALID_EVENT, "Invalid event", 21)
diff --git a/offload/liboffload/API/Common.td b/offload/liboffload/API/Common.td
index d19c9c662faec..71079420a210f 100644
--- a/offload/liboffload/API/Common.td
+++ b/offload/liboffload/API/Common.td
@@ -91,21 +91,28 @@ def ErrorCode : Enum {
 
     // Universal errors
     Etor<"UNKNOWN", "Unknown or internal error">,
+    Etor<"HOST_IO", "I/O error on host">,
+    Etor<"INVALID_BINARY", "A provided binary image is malformed">,
     Etor<"INVALID_NULL_POINTER", "A pointer argument is null when it should not be">,
     Etor<"INVALID_ARGUMENT", "An argument is invalid">,
+    Etor<"NOT_FOUND", "Requested object was not found in the binary image">,
     Etor<"OUT_OF_RESOURCES", "Out of resources">,
-    Etor<"UNSUPPORTED", "generic error code for unsupported features and enums">,
     Etor<"INVALID_SIZE", "invalid size or dimensions (e.g., must not be zero, or is out of bounds)">,
     Etor<"INVALID_ENUMERATION", "enumerator argument is not valid">,
-    Etor<"INVALID_KERNEL_NAME", "Named kernel not found in the program binary">,
+    Etor<"HOST_TOOL_NOT_FOUND", "A required binary (linker, etc.) was not found on the host">,
+    Etor<"INVALID_VALUE", "Invalid Value">,
+    Etor<"UNIMPLEMENTED", "Generic error code for features currently unimplemented by the device/backend">,
+    Etor<"UNSUPPORTED", "Generic error code for features unsupported by the device/backend">,
+    Etor<"ASSEMBLE_FAILURE", "Assembler failure while processing binary image">,
+    Etor<"LINK_FAILURE", "Linker failure while processing binary image">,
+    Etor<"BACKEND_FAILURE", "The plugin backend is in an invalid or unsupported state">,
 
     // Handle related errors - only makes sense for liboffload
-    Etor<"INVALID_VALUE", "Invalid Value">,
+    Etor<"INVALID_NULL_HANDLE", "A handle argument is null when it should not be">,
     Etor<"INVALID_PLATFORM", "Invalid platform">,
     Etor<"INVALID_DEVICE", "Invalid device">,
     Etor<"INVALID_QUEUE", "Invalid queue">,
     Etor<"INVALID_EVENT", "Invalid event">,
-    Etor<"INVALID_NULL_HANDLE", "handle argument is not valid">
   ];
 }
 
diff --git a/offload/liboffload/include/generated/OffloadAPI.h b/offload/liboffload/include/generated/OffloadAPI.h
index 66ea48c8ca14f..d38ed06b91559 100644
--- a/offload/liboffload/include/generated/OffloadAPI.h
+++ b/offload/liboffload/include/generated/OffloadAPI.h
@@ -24,32 +24,47 @@ typedef enum ol_errc_t {
   OL_ERRC_SUCCESS = 0,
   /// Unknown or internal error
   OL_ERRC_UNKNOWN = 1,
+  /// I/O error on host
+  OL_ERRC_HOST_IO = 2,
+  /// A provided binary image is malformed
+  OL_ERRC_INVALID_BINARY = 3,
   /// A pointer argument is null when it should not be
-  OL_ERRC_INVALID_NULL_POINTER = 2,
+  OL_ERRC_INVALID_NULL_POINTER = 4,
   /// An argument is invalid
-  OL_ERRC_INVALID_ARGUMENT = 3,
+  OL_ERRC_INVALID_ARGUMENT = 5,
+  /// Requested object was not found in the binary image
+  OL_ERRC_NOT_FOUND = 6,
   /// Out of resources
-  OL_ERRC_OUT_OF_RESOURCES = 4,
-  /// generic error code for unsupported features and enums
-  OL_ERRC_UNSUPPORTED = 5,
+  OL_ERRC_OUT_OF_RESOURCES = 7,
   /// invalid size or dimensions (e.g., must not be zero, or is out of bounds)
-  OL_ERRC_INVALID_SIZE = 6,
+  OL_ERRC_INVALID_SIZE = 8,
   /// enumerator argument is not valid
-  OL_ERRC_INVALID_ENUMERATION = 7,
-  /// Named kernel not found in the program binary
-  OL_ERRC_INVALID_KERNEL_NAME = 8,
+  OL_ERRC_INVALID_ENUMERATION = 9,
+  /// A required binary (linker, etc.) was not found on the host
+  OL_ERRC_HOST_TOOL_NOT_FOUND = 10,
   /// Invalid Value
-  OL_ERRC_INVALID_VALUE = 9,
+  OL_ERRC_INVALID_VALUE = 11,
+  /// Generic error code for features currently unimplemented by the
+  /// device/backend
+  OL_ERRC_UNIMPLEMENTED = 12,
+  /// Generic error code for features unsupported by the device/backend
+  OL_ERRC_UNSUPPORTED = 13,
+  /// Assembler failure while processing binary image
+  OL_ERRC_ASSEMBLE_FAILURE = 14,
+  /// Linker failure while processing binary image
+  OL_ERRC_LINK_FAILURE = 15,
+  /// The plugin backend is in an invalid or unsupported state
+  OL_ERRC_BACKEND_FAILURE = 16,
+  /// A handle argument is null when it should not be
+  OL_ERRC_INVALID_NULL_HANDLE = 17,
   /// Invalid platform
-  OL_ERRC_INVALID_PLATFORM = 10,
+  OL_ERRC_INVALID_PLATFORM = 18,
   /// Invalid device
-  OL_ERRC_INVALID_DEVICE = 11,
+  OL_ERRC_INVALID_DEVICE = 19,
   /// Invalid queue
-  OL_ERRC_INVALID_QUEUE = 12,
+  OL_ERRC_INVALID_QUEUE = 20,
   /// Invalid event
-  OL_ERRC_INVALID_EVENT = 13,
-  /// handle argument is not valid
-  OL_ERRC_INVALID_NULL_HANDLE = 14,
+  OL_ERRC_INVALID_EVENT = 21,
   /// @cond
   OL_ERRC_FORCE_UINT32 = 0x7fffffff
   /// @endcond
diff --git a/offload/liboffload/include/generated/OffloadPrint.hpp b/offload/liboffload/include/generated/OffloadPrint.hpp
index af5ca0a96b509..9b916543eec0d 100644
--- a/offload/liboffload/include/generated/OffloadPrint.hpp
+++ b/offload/liboffload/include/generated/OffloadPrint.hpp
@@ -52,30 +52,54 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &os,
   case OL_ERRC_UNKNOWN:
     os << "OL_ERRC_UNKNOWN";
     break;
+  case OL_ERRC_HOST_IO:
+    os << "OL_ERRC_HOST_IO";
+    break;
+  case OL_ERRC_INVALID_BINARY:
+    os << "OL_ERRC_INVALID_BINARY";
+    break;
   case OL_ERRC_INVALID_NULL_POINTER:
     os << "OL_ERRC_INVALID_NULL_POINTER";
     break;
   case OL_ERRC_INVALID_ARGUMENT:
     os << "OL_ERRC_INVALID_ARGUMENT";
     break;
+  case OL_ERRC_NOT_FOUND:
+    os << "OL_ERRC_NOT_FOUND";
+    break;
   case OL_ERRC_OUT_OF_RESOURCES:
     os << "OL_ERRC_OUT_OF_RESOURCES";
     break;
-  case OL_ERRC_UNSUPPORTED:
-    os << "OL_ERRC_UNSUPPORTED";
-    break;
   case OL_ERRC_INVALID_SIZE:
     os << "OL_ERRC_INVALID_SIZE";
     break;
   case OL_ERRC_INVALID_ENUMERATION:
     os << "OL_ERRC_INVALID_ENUMERATION";
     break;
-  case OL_ERRC_INVALID_KERNEL_NAME:
-    os << "OL_ERRC_INVALID_KERNEL_NAME";
+  case OL_ERRC_HOST_TOOL_NOT_FOUND:
+    os << "OL_ERRC_HOST_TOOL_NOT_FOUND";
     break;
   case OL_ERRC_INVALID_VALUE:
     os << "OL_ERRC_INVALID_VALUE";
     break;
+  case OL_ERRC_UNIMPLEMENTED:
+    os << "OL_ERRC_UNIMPLEMENTED";
+    break;
+  case OL_ERRC_UNSUPPORTED:
+    os << "OL_ERRC_UNSUPPORTED";
+    break;
+  case OL_ERRC_ASSEMBLE_FAILURE:
+    os << "OL_ERRC_ASSEMBLE_FAILURE";
+    break;
+  case OL_ERRC_LINK_FAILURE:
+    os << "OL_ERRC_LINK_FAILURE";
+    break;
+  case OL_ERRC_BACKEND_FAILURE:
+    os << "OL_ERRC_BACKEND_FAILURE";
+    break;
+  case OL_ERRC_INVALID_NULL_HANDLE:
+    os << "OL_ERRC_INVALID_NULL_HANDLE";
+    break;
   case OL_ERRC_INVALID_PLATFORM:
     os << "OL_ERRC_INVALID_PLATFORM";
     break;
@@ -88,9 +112,6 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &os,
   case OL_ERRC_INVALID_EVENT:
     os << "OL_ERRC_INVALID_EVENT";
     break;
-  case OL_ERRC_INVALID_NULL_HANDLE:
-    os << "OL_ERRC_INVALID_NULL_HANDLE";
-    break;
   default:
     os << "unknown enumerator";
     break;
diff --git a/offload/liboffload/src/OffloadImpl.cpp b/offload/liboffload/src/OffloadImpl.cpp
index b50c7e0f87b7c..ea65282e3ba52 100644
--- a/offload/liboffload/src/OffloadImpl.cpp
+++ b/offload/liboffload/src/OffloadImpl.cpp
@@ -482,7 +482,7 @@ ol_impl_result_t olGetKernel_impl(ol_program_handle_t Program,
   auto &Device = Program->Image->getDevice();
   auto KernelImpl = Device.constructKernel(KernelName);
   if (!KernelImpl)
-    return OL_ERRC_INVALID_KERNEL_NAME;
+    return ol_impl_result_t::fromError(KernelImpl.takeError());
 
   auto Err = KernelImpl->init(Device, *Program->Image);
   if (Err)
diff --git a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h
index 27e4541481301..61f680bab3a07 100644
--- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h
+++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h
@@ -30,6 +30,7 @@ typedef enum {
   HSA_STATUS_INFO_BREAK = 0x1,
   HSA_STATUS_ERROR = 0x1000,
   HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010,
+  HSA_STATUS_ERROR_INVALID_SYMBOL_NAME = 0x1013,
   HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B,
   HSA_STATUS_ERROR_EXCEPTION = 0x1016,
 } hsa_status_t;
diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
index ed575f2213f28..ad2ba9b23081d 100644
--- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -75,6 +75,8 @@
 #include "hsa/hsa_ext_amd.h"
 #endif
 
+using namespace error;
+
 namespace llvm {
 namespace omp {
 namespace target {
@@ -167,7 +169,8 @@ Error asyncMemCopy(bool UseMultipleSdmaEngines, void *Dst, hsa_agent_t DstAgent,
 // This solution is probably not the best
 #if !(HSA_AMD_INTERFACE_VERSION_MAJOR >= 1 &&                                  \
       HSA_AMD_INTERFACE_VERSION_MINOR >= 2)
-  return Plugin::error("Async copy on selected SDMA requires ROCm 5.7");
+  return Plugin::error(ErrorCode::UNSUPPORTED,
+                       "Async copy on selected SDMA requires ROCm 5.7");
 #else
   static std::atomic<int> SdmaEngine{1};
 
@@ -237,7 +240,8 @@ struct AMDGPUResourceRef : public GenericDeviceResourceRef {
   /// reference must be to a valid resource before calling to this function.
   Error destroy(GenericDeviceTy &Device) override {
     if (!Resource)
-      return Plugin::error("Destroying an invalid resource");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Destroying an invalid resource");
 
     if (auto Err = Resource->deinit())
       return Err;
@@ -335,7 +339,8 @@ struct AMDGPUMemoryPoolTy {
       // The agent is not allowed to access the memory pool in any case. Do not
       // continue because otherwise it result in undefined behavior.
       if (Access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED)
-        return Plugin::error("An agent is not allowed to access a memory pool");
+        return Plugin::error(ErrorCode::INVALID_VALUE,
+                             "An agent is not allowed to access a memory pool");
     }
 #endif
 
@@ -416,7 +421,8 @@ struct AMDGPUMemoryManagerTy : public DeviceAllocatorTy {
 
     *PtrStorage = MemoryManager->allocate(Size, nullptr);
     if (*PtrStorage == nullptr)
-      return Plugin::error("Failure to allocate from AMDGPU memory manager");
+      return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
+                           "Failure to allocate from AMDGPU memory manager");
 
     return Plugin::success();
   }
@@ -426,7 +432,8 @@ struct AMDGPUMemoryManagerTy : public DeviceAllocatorTy {
     assert(Ptr && "Invalid pointer");
 
     if (MemoryManager->free(Ptr))
-      return Plugin::error("Failure to deallocate from AMDGPU memory manager");
+      return Plugin::error(ErrorCode::UNKNOWN,
+                           "Failure to deallocate from AMDGPU memory manager");
 
     return Plugin::success();
   }
@@ -540,7 +547,8 @@ struct AMDGPUKernelTy : public GenericKernelTy {
 
     // Make sure it is a kernel symbol.
     if (SymbolType != HSA_SYMBOL_KIND_KERNEL)
-      return Plugin::error("Symbol %s is not a kernel function");
+      return Plugin::error(ErrorCode::INVALID_BINARY,
+                           "Symbol %s is not a kernel function");
 
     // TODO: Read the kernel descriptor for the max threads per block. May be
     // read from the image.
@@ -1115,7 +1123,8 @@ struct AMDGPUStreamTy {
   /// Use a barrier packet with two input signals.
   Error waitOnStreamOperation(AMDGPUStreamTy &OtherStream, uint32_t Slot) {
     if (Queue == nullptr)
-      return Plugin::error("Target queue was nullptr");
+      return Plugin::error(ErrorCode::INVALID_NULL_POINTER,
+                           "Target queue was nullptr");
 
     /// The signal that we must wait from the other stream.
     AMDGPUSignalTy *OtherSignal = OtherStream.Slots[Slot].Signal;
@@ -1236,7 +1245,8 @@ struct AMDGPUStreamTy {
                          uint32_t GroupSize, uint64_t StackSize,
                          AMDGPUMemoryManagerTy &MemoryManager) {
     if (Queue == nullptr)
-      return Plugin::error("Target queue was nullptr");
+      return Plugin::error(ErrorCode::INVALID_NULL_POINTER,
+                           "Target queue was nullptr");
 
     // Retrieve an available signal for the operation's output.
     AMDGPUSignalTy *OutputSignal = nullptr;
@@ -1553,7 +1563,8 @@ struct AMDGPUEventTy {
     std::lock_guard<std::mutex> Lock(Mutex);
 
     if (!RecordedStream)
-      return Plugin::error("Event does not have any recorded stream");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Event does not have any recorded stream");
 
     // Synchronizing the same stream. Do nothing.
     if (RecordedStream == &Stream)
@@ -1942,7 +1953,8 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     else if (WavefrontSize == 64)
       GridValues = getAMDGPUGridValues<64>();
     else
-      return Plugin::error("Unexpected AMDGPU wavefront %d", WavefrontSize);
+      return Plugin::error(ErrorCode::UNSUPPORTED,
+                           "Unexpected AMDGPU wavefront %d", WavefrontSize);
 
     // Get maximum number of workitems per workgroup.
     uint16_t WorkgroupMaxDim[3];
@@ -1958,7 +1970,8 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
 
     GridValues.GV_Max_Teams = GridMaxDim.x / GridValues.GV_Max_WG_Size;
     if (GridValues.GV_Max_Teams == 0)
-      return Plugin::error("Maximum number of teams cannot be zero");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Maximum number of teams cannot be zero");
 
     // Compute the default number of teams.
     uint32_t ComputeUnits = 0;
@@ -2071,7 +2084,8 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     std::error_code EC = sys::fs::createTemporaryFile("amdgpu-pre-link-jit",
                                                       "o", LinkerInputFilePath);
     if (EC)
-      return Plugin::error("Failed to create temporary file for linker");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to create temporary file for linker");
 
     // Write the file's contents to the output file.
     Expected<std::unique_ptr<FileOutputBuffer>> OutputOrErr =
@@ -2087,11 +2101,12 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     EC = sys::fs::createTemporaryFile("amdgpu-pre-link-jit", "so",
                                       LinkerOutputFilePath);
     if (EC)
-      return Plugin::error("Failed to create temporary file for linker");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to create temporary file for linker");
 
     const auto &ErrorOrPath = sys::findProgramByName("lld");
     if (!ErrorOrPath)
-      return createStringError(inconvertibleErrorCode(),
+      return createStringError(ErrorCode::HOST_TOOL_NOT_FOUND,
                                "Failed to find `lld` on the PATH.");
 
     std::string LLDPath = ErrorOrPath.get();
@@ -2112,18 +2127,22 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     std::string Error;
     int RC = sys::ExecuteAndWait(LLDPath, Args, std::nullopt, {}, 0, 0, &Error);
     if (RC)
-      return Plugin::error("Linking optimized bitcode failed: %s",
+      return Plugin::error(ErrorCode::LINK_FAILURE,
+                           "Linking optimized bitcode failed: %s",
                            Error.c_str());
 
     auto BufferOrErr = MemoryBuffer::getFileOrSTDIN(LinkerOutputFilePath);
     if (!BufferOrErr)
-      return Plugin::error("Failed to open temporary file for lld");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to open temporary file for lld");
 
     // Clean up the temporary files afterwards.
     if (sys::fs::remove(LinkerOutputFilePath))
-      return Plugin::error("Failed to remove temporary output file for lld");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to remove temporary output file for lld");
     if (sys::fs::remove(LinkerInputFilePath))
-      return Plugin::error("Failed to remove temporary input file for lld");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to remove temporary input file for lld");
 
     return std::move(*BufferOrErr);
   }
@@ -2139,7 +2158,8 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     // Allocate and construct the AMDGPU kernel.
     AMDGPUKernelTy *AMDGPUKernel = Plugin.allocate<AMDGPUKernelTy>();
     if (!AMDGPUKernel)
-      return Plugin::error("Failed to allocate memory for AMDGPU kernel");
+      return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
+                           "Failed to allocate memory for AMDGPU kernel");
 
     new (AMDGPUKernel) AMDGPUKernelTy(Name);
 
@@ -2529,7 +2549,8 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
 
   /// Synchronize the current thread with the event.
   Error syncEventImpl(void *EventPtr) override {
-    return Plugin::error("Synchronize event not implemented");
+    return Plugin::error(ErrorCode::UNIMPLEMENTED,
+                         "Synchronize event not implemented");
   }
 
   /// Print information about the device.
@@ -2777,7 +2798,8 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
         return Plugin::check(Status, "Error in getting device memory size: %s");
       }
     }
-    return Plugin::error("getDeviceMemorySize:: no global pool");
+    return Plugin::error(ErrorCode::UNSUPPORTED,
+                         "getDeviceMemorySize:: no global pool");
   }
 
   /// AMDGPU-specific function to get device attributes.
@@ -2981,7 +3003,8 @@ Error AMDGPUDeviceImageTy::loadExecutable(const AMDGPUDeviceTy &Device) {
     return Err;
 
   if (Result)
-    return Plugin::error("Loaded HSA executable does not validate");
+    return Plugin::error(ErrorCode::INVALID_BINARY,
+                         "Loaded HSA executable does not validate");
 
   Status = hsa_code_object_reader_destroy(Reader);
   if (auto Err =
@@ -3016,7 +3039,8 @@ AMDGPUDeviceImageTy::findDeviceSymbol(GenericDeviceTy &Device,
 template <typename ResourceTy>
 Error AMDGPUResourceRef<ResourceTy>::create(GenericDeviceTy &Device) {
   if (Resource)
-    return Plugin::error("Creating an existing resource");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Creating an existing resource");
 
   AMDGPUDeviceTy &AMDGPUDevice = static_cast<AMDGPUDeviceTy &>(Device);
 
@@ -3072,6 +3096,7 @@ struct AMDGPUGlobalHandlerTy final : public GenericGlobalHandlerTy {
     // Check the size of the symbol.
     if (SymbolSize != DeviceGlobal.getSize())
       return Plugin::error(
+          ErrorCode::INVALID_BINARY,
           "Failed to load global '%s' due to size mismatch (%zu != %zu)",
           DeviceGlobal.getName().data(), SymbolSize,
           (size_t)DeviceGlobal.getSize());
@@ -3151,7 +3176,7 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
     // There are kernel agents but there is no host agent. That should be
     // treated as an error.
     if (HostAgents.empty())
-      return Plugin::error("No AMDGPU host agents");
+      return Plugin::error(ErrorCode::BACKEND_FAILURE, "No AMDGPU host agents");
 
     // Initialize the host device using host agents.
     HostDevice = allocate<AMDHostDeviceTy>();
@@ -3341,7 +3366,8 @@ Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
                                  AsyncInfoWrapperTy &AsyncInfoWrapper) const {
   if (ArgsSize != LaunchParams.Size &&
       ArgsSize != LaunchParams.Size + getImplicitArgsSize())
-    return Plugin::error("Mismatch of kernel arguments size");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Mismatch of kernel arguments size");
 
   AMDGPUPluginTy &AMDGPUPlugin =
       static_cast<AMDGPUPluginTy &>(GenericDevice.Plugin);
@@ -3461,8 +3487,18 @@ static Error Plugin::check(int32_t Code, const char *ErrFmt, ArgsTy... Args) {
   if (Ret != HSA_STATUS_SUCCESS)
     REPORT("Unrecognized " GETNAME(TARGET_NAME) " error code %d\n", Code);
 
-  return createStringError<ArgsTy..., const char *>(inconvertibleErrorCode(),
-                                                    ErrFmt, Args..., Desc);
+  // TODO: Add more entries to this switch
+  ErrorCode OffloadErrCode;
+  switch (ResultCode) {
+  case HSA_STATUS_ERROR_INVALID_SYMBOL_NAME:
+    OffloadErrCode = ErrorCode::NOT_FOUND;
+    break;
+  default:
+    OffloadErrCode = ErrorCode::UNKNOWN;
+  }
+
+  return createStringError<ArgsTy..., const char *>(OffloadErrCode, ErrFmt,
+                                                    Args..., Desc);
 }
 
 void *AMDGPUMemoryManagerTy::allocate(size_t Size, void *HstPtr,
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index fe603c6269b3b..09b5c8ed3e8ab 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -1388,17 +1388,30 @@ static Error error(error::ErrorCode Code, const char *ErrFmt, ArgsTy... Args) {
   return make_error<error::OffloadError>(Code, Buffer);
 }
 
-template <typename... ArgsTy>
-static Error error(const char *ErrFmt, ArgsTy... Args) {
-  return error(error::ErrorCode::UNKNOWN, ErrFmt, Args...);
-}
-
 inline Error error(error::ErrorCode Code, const char *S) {
   return make_error<error::OffloadError>(Code, S);
 }
 
-inline Error error(const char *S) {
-  return make_error<error::OffloadError>(error::ErrorCode::UNKNOWN, S);
+// The OffloadError will have a message of either:
+// * "{Context}: {Message}" if the other error is a StringError
+// * "{Context}" otherwise
+inline Error error(error::ErrorCode Code, Error &&OtherError,
+                   const char *Context) {
+  std::string Buffer{Context};
+  raw_string_ostream buffer(Buffer);
+
+  handleAllErrors(
+      std::move(OtherError),
+      [&](llvm::StringError &Err) {
+        buffer << ": ";
+        buffer << Err.getMessage();
+      },
+      [&](llvm::ErrorInfoBase &Err) {
+        // Non-string error message don't add anything to the offload error's
+        // error message
+      });
+
+  return make_error<error::OffloadError>(Code, Buffer);
 }
 
 /// Check the plugin-specific error code and return an error or success
diff --git a/offload/plugins-nextgen/common/src/GlobalHandler.cpp b/offload/plugins-nextgen/common/src/GlobalHandler.cpp
index 35a70d8eff901..4166980de1606 100644
--- a/offload/plugins-nextgen/common/src/GlobalHandler.cpp
+++ b/offload/plugins-nextgen/common/src/GlobalHandler.cpp
@@ -26,13 +26,20 @@ using namespace llvm;
 using namespace omp;
 using namespace target;
 using namespace plugin;
+using namespace error;
 
 Expected<std::unique_ptr<ObjectFile>>
 GenericGlobalHandlerTy::getELFObjectFile(DeviceImageTy &Image) {
   assert(utils::elf::isELF(Image.getMemoryBuffer().getBuffer()) &&
          "Input is not an ELF file");
 
-  return ELFObjectFileBase::createELFObjectFile(Image.getMemoryBuffer());
+  auto Expected =
+      ELFObjectFileBase::createELFObjectFile(Image.getMemoryBuffer());
+  if (!Expected) {
+    return Plugin::error(ErrorCode::INVALID_BINARY, Expected.takeError(),
+                         "Error parsing binary");
+  }
+  return Expected;
 }
 
 Error GenericGlobalHandlerTy::moveGlobalBetweenDeviceAndHost(
@@ -112,20 +119,21 @@ Error GenericGlobalHandlerTy::getGlobalMetadataFromImage(
   // Search the ELF symbol using the symbol name.
   auto SymOrErr = utils::elf::getSymbol(**ELFObj, ImageGlobal.getName());
   if (!SymOrErr)
-    return Plugin::error("Failed ELF lookup of global '%s': %s",
-                         ImageGlobal.getName().data(),
-                         toString(SymOrErr.takeError()).data());
+    return Plugin::error(
+        ErrorCode::NOT_FOUND, "Failed ELF lookup of global '%s': %s",
+        ImageGlobal.getName().data(), toString(SymOrErr.takeError()).data());
 
   if (!SymOrErr->has_value())
-    return Plugin::error("Failed to find global symbol '%s' in the ELF image",
+    return Plugin::error(ErrorCode::NOT_FOUND,
+                         "Failed to find global symbol '%s' in the ELF image",
                          ImageGlobal.getName().data());
 
   auto AddrOrErr = utils::elf::getSymbolAddress(**SymOrErr);
   // Get the section to which the symbol belongs.
   if (!AddrOrErr)
-    return Plugin::error("Failed to get ELF symbol from global '%s': %s",
-                         ImageGlobal.getName().data(),
-                         toString(AddrOrErr.takeError()).data());
+    return Plugin::error(
+        ErrorCode::NOT_FOUND, "Failed to get ELF symbol from global '%s': %s",
+        ImageGlobal.getName().data(), toString(AddrOrErr.takeError()).data());
 
   // Setup the global symbol's address and size.
   ImageGlobal.setPtr(const_cast<void *>(*AddrOrErr));
@@ -143,7 +151,8 @@ Error GenericGlobalHandlerTy::readGlobalFromImage(GenericDeviceTy &Device,
     return Err;
 
   if (ImageGlobal.getSize() != HostGlobal.getSize())
-    return Plugin::error("Transfer failed because global symbol '%s' has "
+    return Plugin::error(ErrorCode::INVALID_BINARY,
+                         "Transfer failed because global symbol '%s' has "
                          "%u bytes in the ELF image but %u bytes on the host",
                          HostGlobal.getName().data(), ImageGlobal.getSize(),
                          HostGlobal.getSize());
@@ -274,7 +283,8 @@ void GPUProfGlobals::dump() const {
 
 Error GPUProfGlobals::write() const {
   if (!__llvm_write_custom_profile)
-    return Plugin::error("Could not find symbol __llvm_write_custom_profile. "
+    return Plugin::error(ErrorCode::INVALID_BINARY,
+                         "Could not find symbol __llvm_write_custom_profile. "
                          "The compiler-rt profiling library must be linked for "
                          "GPU PGO to work.");
 
@@ -307,7 +317,8 @@ Error GPUProfGlobals::write() const {
       TargetTriple.str().c_str(), DataBegin, DataEnd, CountersBegin,
       CountersEnd, NamesBegin, NamesEnd, &Version);
   if (result != 0)
-    return Plugin::error("Error writing GPU PGO data to file");
+    return Plugin::error(ErrorCode::HOST_IO,
+                         "Error writing GPU PGO data to file");
 
   return Plugin::success();
 }
diff --git a/offload/plugins-nextgen/common/src/JIT.cpp b/offload/plugins-nextgen/common/src/JIT.cpp
index affedb1a33687..093ebe16cc851 100644
--- a/offload/plugins-nextgen/common/src/JIT.cpp
+++ b/offload/plugins-nextgen/common/src/JIT.cpp
@@ -221,7 +221,8 @@ JITEngine::backend(Module &M, const std::string &ComputeUnitKind,
     raw_fd_stream FD(PostOptIRModuleFileName.get(), EC);
     if (EC)
       return createStringError(
-          EC, "Could not open %s to write the post-opt IR module\n",
+          error::ErrorCode::HOST_IO,
+          "Could not open %s to write the post-opt IR module\n",
           PostOptIRModuleFileName.get().c_str());
     M.print(FD, nullptr);
   }
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index 059f14f59c38b..f35961627b974 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -42,6 +42,7 @@ using namespace llvm;
 using namespace omp;
 using namespace target;
 using namespace plugin;
+using namespace error;
 
 // TODO: Fix any thread safety issues for multi-threaded kernel recording.
 namespace llvm::omp::target::plugin {
@@ -94,7 +95,8 @@ struct RecordReplayTy {
       return Err;
 
     if (isReplaying() && VAddr != MemoryStart) {
-      return Plugin::error("Record-Replay cannot assign the"
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Record-Replay cannot assign the"
                            "requested recorded address (%p, %p)",
                            VAddr, MemoryStart);
     }
@@ -121,7 +123,8 @@ struct RecordReplayTy {
         break;
     }
     if (!MemoryStart)
-      return Plugin::error("Allocating record/replay memory");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Allocating record/replay memory");
 
     if (VAddr && VAddr != MemoryStart)
       MemoryOffset = uintptr_t(VAddr) - uintptr_t(MemoryStart);
@@ -166,7 +169,8 @@ struct RecordReplayTy {
 
     uint64_t DevMemSize;
     if (Device->getDeviceMemorySize(DevMemSize))
-      return Plugin::error("Cannot determine Device Memory Size");
+      return Plugin::error(ErrorCode::UNKNOWN,
+                           "Cannot determine Device Memory Size");
 
     return preAllocateHeuristic(DevMemSize, DeviceMemorySize, ReqVAddr);
   }
@@ -1078,7 +1082,8 @@ Error PinnedAllocationMapTy::insertEntry(void *HstPtr, void *DevAccessiblePtr,
   // Insert the new entry into the map.
   auto Res = Allocs.insert({HstPtr, DevAccessiblePtr, Size, ExternallyLocked});
   if (!Res.second)
-    return Plugin::error("Cannot insert locked buffer entry");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Cannot insert locked buffer entry");
 
   // Check whether the next entry overlaps with the inserted entry.
   auto It = std::next(Res.first);
@@ -1087,7 +1092,8 @@ Error PinnedAllocationMapTy::insertEntry(void *HstPtr, void *DevAccessiblePtr,
 
   const EntryTy *NextEntry = &(*It);
   if (intersects(NextEntry->HstPtr, NextEntry->Size, HstPtr, Size))
-    return Plugin::error("Partial overlapping not allowed in locked buffers");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Partial overlapping not allowed in locked buffers");
 
   return Plugin::success();
 }
@@ -1098,14 +1104,16 @@ Error PinnedAllocationMapTy::eraseEntry(const EntryTy &Entry) {
   // the code more difficult to read.
   size_t Erased = Allocs.erase({Entry.HstPtr});
   if (!Erased)
-    return Plugin::error("Cannot erase locked buffer entry");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Cannot erase locked buffer entry");
   return Plugin::success();
 }
 
 Error PinnedAllocationMapTy::registerEntryUse(const EntryTy &Entry,
                                               void *HstPtr, size_t Size) {
   if (!contains(Entry.HstPtr, Entry.Size, HstPtr, Size))
-    return Plugin::error("Partial overlapping not allowed in locked buffers");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Partial overlapping not allowed in locked buffers");
 
   ++Entry.References;
   return Plugin::success();
@@ -1113,7 +1121,8 @@ Error PinnedAllocationMapTy::registerEntryUse(const EntryTy &Entry,
 
 Expected<bool> PinnedAllocationMapTy::unregisterEntryUse(const EntryTy &Entry) {
   if (Entry.References == 0)
-    return Plugin::error("Invalid number of references");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Invalid number of references");
 
   // Return whether this was the last user.
   return (--Entry.References == 0);
@@ -1131,7 +1140,8 @@ Error PinnedAllocationMapTy::registerHostBuffer(void *HstPtr,
   // No pinned allocation should intersect.
   const EntryTy *Entry = findIntersecting(HstPtr);
   if (Entry)
-    return Plugin::error("Cannot insert entry due to an existing one");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Cannot insert entry due to an existing one");
 
   // Now insert the new entry.
   return insertEntry(HstPtr, DevAccessiblePtr, Size);
@@ -1144,11 +1154,13 @@ Error PinnedAllocationMapTy::unregisterHostBuffer(void *HstPtr) {
 
   const EntryTy *Entry = findIntersecting(HstPtr);
   if (!Entry)
-    return Plugin::error("Cannot find locked buffer");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Cannot find locked buffer");
 
   // The address in the entry should be the same we are unregistering.
   if (Entry->HstPtr != HstPtr)
-    return Plugin::error("Unexpected host pointer in locked buffer entry");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Unexpected host pointer in locked buffer entry");
 
   // Unregister from the entry.
   auto LastUseOrErr = unregisterEntryUse(*Entry);
@@ -1157,7 +1169,8 @@ Error PinnedAllocationMapTy::unregisterHostBuffer(void *HstPtr) {
 
   // There should be no other references to the pinned allocation.
   if (!(*LastUseOrErr))
-    return Plugin::error("The locked buffer is still being used");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "The locked buffer is still being used");
 
   // Erase the entry from the map.
   return eraseEntry(*Entry);
@@ -1203,7 +1216,8 @@ Error PinnedAllocationMapTy::unlockHostBuffer(void *HstPtr) {
 
   const EntryTy *Entry = findIntersecting(HstPtr);
   if (!Entry)
-    return Plugin::error("Cannot find locked buffer");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Cannot find locked buffer");
 
   // Unregister from the locked buffer. No need to do anything if there are
   // others using the allocation.
@@ -1289,7 +1303,8 @@ Error PinnedAllocationMapTy::unlockUnmappedHostBuffer(void *HstPtr) {
 
   // No entry, but the automatic locking is enabled, so this is an error.
   if (!Entry)
-    return Plugin::error("Locked buffer not found");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Locked buffer not found");
 
   // There is entry, so unregister a user and check whether it was the last one.
   auto LastUseOrErr = unregisterEntryUse(*Entry);
@@ -1312,7 +1327,8 @@ Error PinnedAllocationMapTy::unlockUnmappedHostBuffer(void *HstPtr) {
 
 Error GenericDeviceTy::synchronize(__tgt_async_info *AsyncInfo) {
   if (!AsyncInfo || !AsyncInfo->Queue)
-    return Plugin::error("Invalid async info queue");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Invalid async info queue");
 
   if (auto Err = synchronizeImpl(*AsyncInfo))
     return Err;
@@ -1327,21 +1343,25 @@ Error GenericDeviceTy::synchronize(__tgt_async_info *AsyncInfo) {
 
 Error GenericDeviceTy::queryAsync(__tgt_async_info *AsyncInfo) {
   if (!AsyncInfo || !AsyncInfo->Queue)
-    return Plugin::error("Invalid async info queue");
+    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                         "Invalid async info queue");
 
   return queryAsyncImpl(*AsyncInfo);
 }
 
 Error GenericDeviceTy::memoryVAMap(void **Addr, void *VAddr, size_t *RSize) {
-  return Plugin::error("Device does not support VA Management");
+  return Plugin::error(ErrorCode::UNSUPPORTED,
+                       "Device does not support VA Management");
 }
 
 Error GenericDeviceTy::memoryVAUnMap(void *VAddr, size_t Size) {
-  return Plugin::error("Device does not support VA Management");
+  return Plugin::error(ErrorCode::UNSUPPORTED,
+                       "Device does not support VA Management");
 }
 
 Error GenericDeviceTy::getDeviceMemorySize(uint64_t &DSize) {
   return Plugin::error(
+      ErrorCode::UNIMPLEMENTED,
       "Missing getDeviceMemorySize implementation (required by RR-heuristic");
 }
 
@@ -1359,7 +1379,8 @@ Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr,
     if (MemoryManager) {
       Alloc = MemoryManager->allocate(Size, HostPtr);
       if (!Alloc)
-        return Plugin::error("Failed to allocate from memory manager");
+        return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
+                             "Failed to allocate from memory manager");
       break;
     }
     [[fallthrough]];
@@ -1367,13 +1388,15 @@ Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr,
   case TARGET_ALLOC_SHARED:
     Alloc = allocate(Size, HostPtr, Kind);
     if (!Alloc)
-      return Plugin::error("Failed to allocate from device allocator");
+      return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
+                           "Failed to allocate from device allocator");
   }
 
   // Report error if the memory manager or the device allocator did not return
   // any memory buffer.
   if (!Alloc)
-    return Plugin::error("Invalid target data allocation kind or requested "
+    return Plugin::error(ErrorCode::UNIMPLEMENTED,
+                         "Invalid target data allocation kind or requested "
                          "allocator not implemented yet");
 
   // Register allocated buffer as pinned memory if the type is host memory.
@@ -1448,6 +1471,7 @@ Error GenericDeviceTy::dataDelete(void *TgtPtr, TargetAllocTy Kind) {
       Res = MemoryManager->free(TgtPtr);
       if (Res)
         return Plugin::error(
+            ErrorCode::OUT_OF_RESOURCES,
             "Failure to deallocate device pointer %p via memory manager",
             TgtPtr);
       break;
@@ -1458,6 +1482,7 @@ Error GenericDeviceTy::dataDelete(void *TgtPtr, TargetAllocTy Kind) {
     Res = free(TgtPtr, Kind);
     if (Res)
       return Plugin::error(
+          ErrorCode::UNKNOWN,
           "Failure to deallocate device pointer %p via device deallocator",
           TgtPtr);
   }
diff --git a/offload/plugins-nextgen/common/src/RPC.cpp b/offload/plugins-nextgen/common/src/RPC.cpp
index fc90bb2e032f2..c4afe10586cf5 100644
--- a/offload/plugins-nextgen/common/src/RPC.cpp
+++ b/offload/plugins-nextgen/common/src/RPC.cpp
@@ -176,6 +176,7 @@ Error RPCServerTy::initDevice(plugin::GenericDeviceTy &Device,
       TARGET_ALLOC_HOST);
   if (!RPCBuffer)
     return plugin::Plugin::error(
+        error::ErrorCode::UNKNOWN,
         "Failed to initialize RPC server for device %d", Device.getDeviceId());
 
   // Get the address of the RPC client from the device.
diff --git a/offload/plugins-nextgen/cuda/src/rtl.cpp b/offload/plugins-nextgen/cuda/src/rtl.cpp
index 0d0c4858aa7fa..2219cd88cf93d 100644
--- a/offload/plugins-nextgen/cuda/src/rtl.cpp
+++ b/offload/plugins-nextgen/cuda/src/rtl.cpp
@@ -33,6 +33,8 @@
 #include "llvm/Support/FileSystem.h"
 #include "llvm/Support/Program.h"
 
+using namespace error;
+
 namespace llvm {
 namespace omp {
 namespace target {
@@ -134,7 +136,8 @@ struct CUDAKernelTy : public GenericKernelTy {
 
     // Check that the function pointer is valid.
     if (!Func)
-      return Plugin::error("Invalid function for kernel %s", getName());
+      return Plugin::error(ErrorCode::INVALID_BINARY,
+                           "Invalid function for kernel %s", getName());
 
     int MaxThreads;
     Res = cuFuncGetAttribute(&MaxThreads,
@@ -175,7 +178,8 @@ struct CUDAStreamRef final : public GenericDeviceResourceRef {
   /// before calling to this function.
   Error create(GenericDeviceTy &Device) override {
     if (Stream)
-      return Plugin::error("Creating an existing stream");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Creating an existing stream");
 
     CUresult Res = cuStreamCreate(&Stream, CU_STREAM_NON_BLOCKING);
     if (auto Err = Plugin::check(Res, "Error in cuStreamCreate: %s"))
@@ -188,7 +192,8 @@ struct CUDAStreamRef final : public GenericDeviceResourceRef {
   /// must be to a valid stream before calling to this function.
   Error destroy(GenericDeviceTy &Device) override {
     if (!Stream)
-      return Plugin::error("Destroying an invalid stream");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Destroying an invalid stream");
 
     CUresult Res = cuStreamDestroy(Stream);
     if (auto Err = Plugin::check(Res, "Error in cuStreamDestroy: %s"))
@@ -222,7 +227,8 @@ struct CUDAEventRef final : public GenericDeviceResourceRef {
   /// before calling to this function.
   Error create(GenericDeviceTy &Device) override {
     if (Event)
-      return Plugin::error("Creating an existing event");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Creating an existing event");
 
     CUresult Res = cuEventCreate(&Event, CU_EVENT_DEFAULT);
     if (auto Err = Plugin::check(Res, "Error in cuEventCreate: %s"))
@@ -235,7 +241,8 @@ struct CUDAEventRef final : public GenericDeviceResourceRef {
   /// must be to a valid event before calling to this function.
   Error destroy(GenericDeviceTy &Device) override {
     if (!Event)
-      return Plugin::error("Destroying an invalid event");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Destroying an invalid event");
 
     CUresult Res = cuEventDestroy(Event);
     if (auto Err = Plugin::check(Res, "Error in cuEventDestroy: %s"))
@@ -419,7 +426,8 @@ struct CUDADeviceTy : public GenericDeviceTy {
     std::error_code EC = sys::fs::createTemporaryFile("nvptx-pre-link-jit", "s",
                                                       PTXInputFilePath);
     if (EC)
-      return Plugin::error("Failed to create temporary file for ptxas");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to create temporary file for ptxas");
 
     // Write the file's contents to the output file.
     Expected<std::unique_ptr<FileOutputBuffer>> OutputOrErr =
@@ -435,12 +443,14 @@ struct CUDADeviceTy : public GenericDeviceTy {
     EC = sys::fs::createTemporaryFile("nvptx-post-link-jit", "cubin",
                                       PTXOutputFilePath);
     if (EC)
-      return Plugin::error("Failed to create temporary file for ptxas");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to create temporary file for ptxas");
 
     // Try to find `ptxas` in the path to compile the PTX to a binary.
     const auto ErrorOrPath = sys::findProgramByName("ptxas");
     if (!ErrorOrPath)
-      return Plugin::error("Failed to find 'ptxas' on the PATH.");
+      return Plugin::error(ErrorCode::HOST_TOOL_NOT_FOUND,
+                           "Failed to find 'ptxas' on the PATH.");
 
     std::string Arch = getComputeUnitKind();
     StringRef Args[] = {*ErrorOrPath,
@@ -455,17 +465,21 @@ struct CUDADeviceTy : public GenericDeviceTy {
     std::string ErrMsg;
     if (sys::ExecuteAndWait(*ErrorOrPath, Args, std::nullopt, {}, 0, 0,
                             &ErrMsg))
-      return Plugin::error("Running 'ptxas' failed: %s\n", ErrMsg.c_str());
+      return Plugin::error(ErrorCode::ASSEMBLE_FAILURE,
+                           "Running 'ptxas' failed: %s\n", ErrMsg.c_str());
 
     auto BufferOrErr = MemoryBuffer::getFileOrSTDIN(PTXOutputFilePath.data());
     if (!BufferOrErr)
-      return Plugin::error("Failed to open temporary file for ptxas");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to open temporary file for ptxas");
 
     // Clean up the temporary files afterwards.
     if (sys::fs::remove(PTXOutputFilePath))
-      return Plugin::error("Failed to remove temporary file for ptxas");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to remove temporary file for ptxas");
     if (sys::fs::remove(PTXInputFilePath))
-      return Plugin::error("Failed to remove temporary file for ptxas");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to remove temporary file for ptxas");
 
     return std::move(*BufferOrErr);
   }
@@ -475,7 +489,8 @@ struct CUDADeviceTy : public GenericDeviceTy {
     // Allocate and construct the CUDA kernel.
     CUDAKernelTy *CUDAKernel = Plugin.allocate<CUDAKernelTy>();
     if (!CUDAKernel)
-      return Plugin::error("Failed to allocate memory for CUDA kernel");
+      return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
+                           "Failed to allocate memory for CUDA kernel");
 
     new (CUDAKernel) CUDAKernelTy(Name);
 
@@ -658,11 +673,13 @@ struct CUDADeviceTy : public GenericDeviceTy {
     size_t Size = *RSize;
 
     if (Size == 0)
-      return Plugin::error("Memory Map Size must be larger than 0");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Memory Map Size must be larger than 0");
 
     // Check if we have already mapped this address
     if (IHandle != DeviceMMaps.end())
-      return Plugin::error("Address already memory mapped");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Address already memory mapped");
 
     CUmemAllocationProp Prop = {};
     size_t Granularity = 0;
@@ -675,6 +692,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
     if (Size >= Free) {
       *Addr = nullptr;
       return Plugin::error(
+          ErrorCode::OUT_OF_RESOURCES,
           "Cannot map memory size larger than the available device memory");
     }
 
@@ -690,7 +708,8 @@ struct CUDADeviceTy : public GenericDeviceTy {
       return Err;
 
     if (Granularity == 0)
-      return Plugin::error("Wrong device Page size");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Wrong device Page size");
 
     // Ceil to page size.
     Size = utils::roundUp(Size, Granularity);
@@ -732,11 +751,13 @@ struct CUDADeviceTy : public GenericDeviceTy {
     auto IHandle = DeviceMMaps.find(DVAddr);
     // Mapping does not exist
     if (IHandle == DeviceMMaps.end()) {
-      return Plugin::error("Addr is not MemoryMapped");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Addr is not MemoryMapped");
     }
 
     if (IHandle == DeviceMMaps.end())
-      return Plugin::error("Addr is not MemoryMapped");
+      return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+                           "Addr is not MemoryMapped");
 
     CUmemGenericAllocationHandle &AllocHandle = IHandle->second;
 
@@ -1156,7 +1177,8 @@ struct CUDADeviceTy : public GenericDeviceTy {
 
       uint16_t Priority;
       if (NameOrErr->rsplit('_').second.getAsInteger(10, Priority))
-        return Plugin::error("Invalid priority for constructor or destructor");
+        return Plugin::error(ErrorCode::INVALID_BINARY,
+                             "Invalid priority for constructor or destructor");
 
       Funcs.emplace_back(*NameOrErr, Priority);
     }
@@ -1169,7 +1191,8 @@ struct CUDADeviceTy : public GenericDeviceTy {
     void *Buffer =
         allocate(Funcs.size() * sizeof(void *), nullptr, TARGET_ALLOC_DEVICE);
     if (!Buffer)
-      return Plugin::error("Failed to allocate memory for global buffer");
+      return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
+                           "Failed to allocate memory for global buffer");
 
     auto *GlobalPtrStart = reinterpret_cast<uintptr_t *>(Buffer);
     auto *GlobalPtrStop = reinterpret_cast<uintptr_t *>(Buffer) + Funcs.size();
@@ -1217,7 +1240,8 @@ struct CUDADeviceTy : public GenericDeviceTy {
     AsyncInfoWrapper.finalize(Err);
 
     if (free(Buffer, TARGET_ALLOC_DEVICE) != OFFLOAD_SUCCESS)
-      return Plugin::error("Failed to free memory for global buffer");
+      return Plugin::error(ErrorCode::UNKNOWN,
+                           "Failed to free memory for global buffer");
 
     return Err;
   }
@@ -1316,6 +1340,7 @@ class CUDAGlobalHandlerTy final : public GenericGlobalHandlerTy {
 
     if (CUSize != DeviceGlobal.getSize())
       return Plugin::error(
+          ErrorCode::INVALID_BINARY,
           "Failed to load global '%s' due to size mismatch (%zu != %zu)",
           GlobalName, CUSize, (size_t)DeviceGlobal.getSize());
 
@@ -1501,8 +1526,9 @@ static Error Plugin::check(int32_t Code, const char *ErrFmt, ArgsTy... Args) {
   if (Ret != CUDA_SUCCESS)
     REPORT("Unrecognized " GETNAME(TARGET_NAME) " error code %d\n", Code);
 
-  return createStringError<ArgsTy..., const char *>(inconvertibleErrorCode(),
-                                                    ErrFmt, Args..., Desc);
+  // TODO: Create a map for CUDA error codes to Offload error codes
+  return createStringError<ArgsTy..., const char *>(ErrorCode::UNKNOWN, ErrFmt,
+                                                    Args..., Desc);
 }
 
 } // namespace plugin
diff --git a/offload/plugins-nextgen/host/src/rtl.cpp b/offload/plugins-nextgen/host/src/rtl.cpp
index 171c202e3e39e..8dfb06fec976e 100644
--- a/offload/plugins-nextgen/host/src/rtl.cpp
+++ b/offload/plugins-nextgen/host/src/rtl.cpp
@@ -56,6 +56,7 @@ struct GenELF64DeviceTy;
 struct GenELF64PluginTy;
 
 using llvm::sys::DynamicLibrary;
+using namespace error;
 
 /// Class implementing kernel functionalities for GenELF64.
 struct GenELF64KernelTy : public GenericKernelTy {
@@ -74,7 +75,8 @@ struct GenELF64KernelTy : public GenericKernelTy {
 
     // Check that the function pointer is valid.
     if (!Global.getPtr())
-      return Plugin::error("Invalid function for kernel %s", getName());
+      return Plugin::error(ErrorCode::INVALID_BINARY,
+                           "Invalid function for kernel %s", getName());
 
     // Save the function pointer.
     Func = (void (*)())Global.getPtr();
@@ -102,7 +104,8 @@ struct GenELF64KernelTy : public GenericKernelTy {
     ffi_status Status = ffi_prep_cif(&Cif, FFI_DEFAULT_ABI, KernelArgs.NumArgs,
                                      &ffi_type_void, ArgTypesPtr);
     if (Status != FFI_OK)
-      return Plugin::error("Error in ffi_prep_cif: %d", Status);
+      return Plugin::error(ErrorCode::UNKNOWN, "Error in ffi_prep_cif: %d",
+                           Status);
 
     // Call the kernel function through libffi.
     long Return;
@@ -155,7 +158,8 @@ struct GenELF64DeviceTy : public GenericDeviceTy {
     // Allocate and construct the kernel.
     GenELF64KernelTy *GenELF64Kernel = Plugin.allocate<GenELF64KernelTy>();
     if (!GenELF64Kernel)
-      return Plugin::error("Failed to allocate memory for GenELF64 kernel");
+      return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
+                           "Failed to allocate memory for GenELF64 kernel");
 
     new (GenELF64Kernel) GenELF64KernelTy(Name);
 
@@ -176,24 +180,28 @@ struct GenELF64DeviceTy : public GenericDeviceTy {
     char TmpFileName[] = "/tmp/tmpfile_XXXXXX";
     int TmpFileFd = mkstemp(TmpFileName);
     if (TmpFileFd == -1)
-      return Plugin::error("Failed to create tmpfile for loading target image");
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to create tmpfile for loading target image");
 
     // Open the temporary file.
     FILE *TmpFile = fdopen(TmpFileFd, "wb");
     if (!TmpFile)
-      return Plugin::error("Failed to open tmpfile %s for loading target image",
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to open tmpfile %s for loading target image",
                            TmpFileName);
 
     // Write the image into the temporary file.
     size_t Written = fwrite(Image->getStart(), Image->getSize(), 1, TmpFile);
     if (Written != 1)
-      return Plugin::error("Failed to write target image to tmpfile %s",
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to write target image to tmpfile %s",
                            TmpFileName);
 
     // Close the temporary file.
     int Ret = fclose(TmpFile);
     if (Ret)
-      return Plugin::error("Failed to close tmpfile %s with the target image",
+      return Plugin::error(ErrorCode::HOST_IO,
+                           "Failed to close tmpfile %s with the target image",
                            TmpFileName);
 
     // Load the temporary file as a dynamic library.
@@ -203,7 +211,8 @@ struct GenELF64DeviceTy : public GenericDeviceTy {
 
     // Check if the loaded library is valid.
     if (!DynLib.isValid())
-      return Plugin::error("Failed to load target image: %s", ErrMsg.c_str());
+      return Plugin::error(ErrorCode::INVALID_BINARY,
+                           "Failed to load target image: %s", ErrMsg.c_str());
 
     // Save a reference of the image's dynamic library.
     Image->setDynamicLibrary(DynLib);
@@ -272,7 +281,8 @@ struct GenELF64DeviceTy : public GenericDeviceTy {
                          AsyncInfoWrapperTy &AsyncInfoWrapper) override {
     // This function should never be called because the function
     // GenELF64PluginTy::isDataExchangable() returns false.
-    return Plugin::error("dataExchangeImpl not supported");
+    return Plugin::error(ErrorCode::UNSUPPORTED,
+                         "dataExchangeImpl not supported");
   }
 
   /// All functions are already synchronous. No need to do anything on this
@@ -289,12 +299,14 @@ struct GenELF64DeviceTy : public GenericDeviceTy {
 
   /// This plugin does not support interoperability
   Error initAsyncInfoImpl(AsyncInfoWrapperTy &AsyncInfoWrapper) override {
-    return Plugin::error("initAsyncInfoImpl not supported");
+    return Plugin::error(ErrorCode::UNSUPPORTED,
+                         "initAsyncInfoImpl not supported");
   }
 
   /// This plugin does not support interoperability
   Error initDeviceInfoImpl(__tgt_device_info *DeviceInfo) override {
-    return Plugin::error("initDeviceInfoImpl not supported");
+    return Plugin::error(ErrorCode::UNSUPPORTED,
+                         "initDeviceInfoImpl not supported");
   }
 
   /// This plugin does not support the event API. Do nothing without failing.
@@ -365,7 +377,8 @@ class GenELF64GlobalHandlerTy final : public GenericGlobalHandlerTy {
     // Get the address of the symbol.
     void *Addr = DynLib.getAddressOfSymbol(GlobalName);
     if (Addr == nullptr) {
-      return Plugin::error("Failed to load global '%s'", GlobalName);
+      return Plugin::error(ErrorCode::NOT_FOUND, "Failed to load global '%s'",
+                           GlobalName);
     }
 
     // Save the pointer to the symbol.
diff --git a/offload/unittests/OffloadAPI/kernel/olGetKernel.cpp b/offload/unittests/OffloadAPI/kernel/olGetKernel.cpp
index bd1b562eac71e..83755e554ebe9 100644
--- a/offload/unittests/OffloadAPI/kernel/olGetKernel.cpp
+++ b/offload/unittests/OffloadAPI/kernel/olGetKernel.cpp
@@ -31,8 +31,8 @@ TEST_P(olGetKernelTest, InvalidNullKernelPointer) {
 }
 
 // Error code returning from plugin interface not yet supported
-TEST_F(olGetKernelTest, DISABLED_InvalidKernelName) {
+TEST_P(olGetKernelTest, InvalidKernelName) {
   ol_kernel_handle_t Kernel = nullptr;
-  ASSERT_ERROR(OL_ERRC_INVALID_KERNEL_NAME,
+  ASSERT_ERROR(OL_ERRC_NOT_FOUND,
                olGetKernel(Program, "invalid_kernel_name", &Kernel));
 }

>From 9cab985b56cdcca03de86d948b886a34560b209c Mon Sep 17 00:00:00 2001
From: Ross Brunton <ross at codeplay.com>
Date: Tue, 20 May 2025 10:43:33 +0100
Subject: [PATCH 2/3] Lowercase error messages and port over more calls

Error messages should be lower case, so all of them have been changed.
In addition, a few more string error creations have been ported to
use offload errors, which also meant moving some reporting things to
OffloadError.h.
---
 offload/include/Shared/OffloadErrcodes.inc    |  38 ++---
 offload/include/Shared/OffloadError.h         |  34 +++++
 offload/liboffload/API/Common.td              |  38 ++---
 .../liboffload/include/generated/OffloadAPI.h |  38 ++---
 offload/libomptarget/PluginManager.cpp        |  12 +-
 offload/libomptarget/device.cpp               |  10 +-
 offload/plugins-nextgen/amdgpu/src/rtl.cpp    | 135 +++++++++--------
 .../common/include/PluginInterface.h          |  23 +--
 .../common/src/GlobalHandler.cpp              |  14 +-
 offload/plugins-nextgen/common/src/JIT.cpp    |  11 +-
 .../common/src/PluginInterface.cpp            |  48 +++---
 offload/plugins-nextgen/common/src/RPC.cpp    |   2 +-
 offload/plugins-nextgen/cuda/src/rtl.cpp      | 141 +++++++++---------
 offload/plugins-nextgen/host/src/rtl.cpp      |  26 ++--
 14 files changed, 292 insertions(+), 278 deletions(-)

diff --git a/offload/include/Shared/OffloadErrcodes.inc b/offload/include/Shared/OffloadErrcodes.inc
index e0797e2bac194..130e553aa70a4 100644
--- a/offload/include/Shared/OffloadErrcodes.inc
+++ b/offload/include/Shared/OffloadErrcodes.inc
@@ -14,38 +14,38 @@
 // To add new error codes, add them to offload/liboffload/API/Common.td and run
 // the GenerateOffload target.
 
-OFFLOAD_ERRC(SUCCESS, "Success", 0)
-OFFLOAD_ERRC(UNKNOWN, "Unknown or internal error", 1)
+OFFLOAD_ERRC(SUCCESS, "success", 0)
+OFFLOAD_ERRC(UNKNOWN, "unknown or internal error", 1)
 OFFLOAD_ERRC(HOST_IO, "I/O error on host", 2)
-OFFLOAD_ERRC(INVALID_BINARY, "A provided binary image is malformed", 3)
+OFFLOAD_ERRC(INVALID_BINARY, "a provided binary image is malformed", 3)
 OFFLOAD_ERRC(INVALID_NULL_POINTER,
-             "A pointer argument is null when it should not be", 4)
-OFFLOAD_ERRC(INVALID_ARGUMENT, "An argument is invalid", 5)
-OFFLOAD_ERRC(NOT_FOUND, "Requested object was not found in the binary image", 6)
-OFFLOAD_ERRC(OUT_OF_RESOURCES, "Out of resources", 7)
+             "a pointer argument is null when it should not be", 4)
+OFFLOAD_ERRC(INVALID_ARGUMENT, "an argument is invalid", 5)
+OFFLOAD_ERRC(NOT_FOUND, "requested object was not found in the binary image", 6)
+OFFLOAD_ERRC(OUT_OF_RESOURCES, "out of resources", 7)
 OFFLOAD_ERRC(
     INVALID_SIZE,
     "invalid size or dimensions (e.g., must not be zero, or is out of bounds)",
     8)
 OFFLOAD_ERRC(INVALID_ENUMERATION, "enumerator argument is not valid", 9)
 OFFLOAD_ERRC(HOST_TOOL_NOT_FOUND,
-             "A required binary (linker, etc.) was not found on the host", 10)
-OFFLOAD_ERRC(INVALID_VALUE, "Invalid Value", 11)
+             "a required binary (linker, etc.) was not found on the host", 10)
+OFFLOAD_ERRC(INVALID_VALUE, "invalid value", 11)
 OFFLOAD_ERRC(UNIMPLEMENTED,
-             "Generic error code for features currently unimplemented by the "
+             "generic error code for features currently unimplemented by the "
              "device/backend",
              12)
 OFFLOAD_ERRC(
     UNSUPPORTED,
-    "Generic error code for features unsupported by the device/backend", 13)
+    "generic error code for features unsupported by the device/backend", 13)
 OFFLOAD_ERRC(ASSEMBLE_FAILURE,
-             "Assembler failure while processing binary image", 14)
-OFFLOAD_ERRC(LINK_FAILURE, "Linker failure while processing binary image", 15)
+             "assembler failure while processing binary image", 14)
+OFFLOAD_ERRC(LINK_FAILURE, "linker failure while processing binary image", 15)
 OFFLOAD_ERRC(BACKEND_FAILURE,
-             "The plugin backend is in an invalid or unsupported state", 16)
+             "the plugin backend is in an invalid or unsupported state", 16)
 OFFLOAD_ERRC(INVALID_NULL_HANDLE,
-             "A handle argument is null when it should not be", 17)
-OFFLOAD_ERRC(INVALID_PLATFORM, "Invalid platform", 18)
-OFFLOAD_ERRC(INVALID_DEVICE, "Invalid device", 19)
-OFFLOAD_ERRC(INVALID_QUEUE, "Invalid queue", 20)
-OFFLOAD_ERRC(INVALID_EVENT, "Invalid event", 21)
+             "a handle argument is null when it should not be", 17)
+OFFLOAD_ERRC(INVALID_PLATFORM, "invalid platform", 18)
+OFFLOAD_ERRC(INVALID_DEVICE, "invalid device", 19)
+OFFLOAD_ERRC(INVALID_QUEUE, "invalid queue", 20)
+OFFLOAD_ERRC(INVALID_EVENT, "invalid event", 21)
diff --git a/offload/include/Shared/OffloadError.h b/offload/include/Shared/OffloadError.h
index b999705893ad5..ab9bf136b444c 100644
--- a/offload/include/Shared/OffloadError.h
+++ b/offload/include/Shared/OffloadError.h
@@ -46,6 +46,40 @@ class OffloadError : public llvm::ErrorInfo<OffloadError, llvm::StringError> {
   // The definition for this resides in the plugin static library
   static char ID;
 };
+
+/// Create an Offload error.
+template <typename... ArgsTy>
+static llvm::Error createOffloadError(error::ErrorCode Code, const char *ErrFmt, ArgsTy... Args) {
+  std::string Buffer;
+  llvm::raw_string_ostream(Buffer) << llvm::format(ErrFmt, Args...);
+  return llvm::make_error<error::OffloadError>(Code, Buffer);
+}
+
+inline llvm::Error createOffloadError(error::ErrorCode Code, const char *S) {
+  return llvm::make_error<error::OffloadError>(Code, S);
+}
+
+// The OffloadError will have a message of either:
+// * "{Context}: {Message}" if the other error is a StringError
+// * "{Context}" otherwise
+inline llvm::Error createOffloadError(error::ErrorCode Code, llvm::Error &&OtherError,
+                   const char *Context) {
+  std::string Buffer{Context};
+  llvm::raw_string_ostream buffer(Buffer);
+
+  handleAllErrors(
+      std::move(OtherError),
+      [&](llvm::StringError &Err) {
+        buffer << ": ";
+        buffer << Err.getMessage();
+      },
+      [&](llvm::ErrorInfoBase &Err) {
+        // Non-string error message don't add anything to the offload error's
+        // error message
+      });
+
+  return llvm::make_error<error::OffloadError>(Code, Buffer);
+}
 } // namespace error
 
 #endif
diff --git a/offload/liboffload/API/Common.td b/offload/liboffload/API/Common.td
index 71079420a210f..7674da0438c29 100644
--- a/offload/liboffload/API/Common.td
+++ b/offload/liboffload/API/Common.td
@@ -87,32 +87,32 @@ def ErrorCode : Enum {
   let name = "ol_errc_t";
   let desc = "Defines Return/Error codes";
   let etors =[
-    Etor<"SUCCESS", "Success">,
+    Etor<"SUCCESS", "success">,
 
     // Universal errors
-    Etor<"UNKNOWN", "Unknown or internal error">,
+    Etor<"UNKNOWN", "unknown or internal error">,
     Etor<"HOST_IO", "I/O error on host">,
-    Etor<"INVALID_BINARY", "A provided binary image is malformed">,
-    Etor<"INVALID_NULL_POINTER", "A pointer argument is null when it should not be">,
-    Etor<"INVALID_ARGUMENT", "An argument is invalid">,
-    Etor<"NOT_FOUND", "Requested object was not found in the binary image">,
-    Etor<"OUT_OF_RESOURCES", "Out of resources">,
+    Etor<"INVALID_BINARY", "a provided binary image is malformed">,
+    Etor<"INVALID_NULL_POINTER", "a pointer argument is null when it should not be">,
+    Etor<"INVALID_ARGUMENT", "an argument is invalid">,
+    Etor<"NOT_FOUND", "requested object was not found in the binary image">,
+    Etor<"OUT_OF_RESOURCES", "out of resources">,
     Etor<"INVALID_SIZE", "invalid size or dimensions (e.g., must not be zero, or is out of bounds)">,
     Etor<"INVALID_ENUMERATION", "enumerator argument is not valid">,
-    Etor<"HOST_TOOL_NOT_FOUND", "A required binary (linker, etc.) was not found on the host">,
-    Etor<"INVALID_VALUE", "Invalid Value">,
-    Etor<"UNIMPLEMENTED", "Generic error code for features currently unimplemented by the device/backend">,
-    Etor<"UNSUPPORTED", "Generic error code for features unsupported by the device/backend">,
-    Etor<"ASSEMBLE_FAILURE", "Assembler failure while processing binary image">,
-    Etor<"LINK_FAILURE", "Linker failure while processing binary image">,
-    Etor<"BACKEND_FAILURE", "The plugin backend is in an invalid or unsupported state">,
+    Etor<"HOST_TOOL_NOT_FOUND", "a required binary (linker, etc.) was not found on the host">,
+    Etor<"INVALID_VALUE", "invalid value">,
+    Etor<"UNIMPLEMENTED", "generic error code for features currently unimplemented by the device/backend">,
+    Etor<"UNSUPPORTED", "generic error code for features unsupported by the device/backend">,
+    Etor<"ASSEMBLE_FAILURE", "assembler failure while processing binary image">,
+    Etor<"LINK_FAILURE", "linker failure while processing binary image">,
+    Etor<"BACKEND_FAILURE", "the plugin backend is in an invalid or unsupported state">,
 
     // Handle related errors - only makes sense for liboffload
-    Etor<"INVALID_NULL_HANDLE", "A handle argument is null when it should not be">,
-    Etor<"INVALID_PLATFORM", "Invalid platform">,
-    Etor<"INVALID_DEVICE", "Invalid device">,
-    Etor<"INVALID_QUEUE", "Invalid queue">,
-    Etor<"INVALID_EVENT", "Invalid event">,
+    Etor<"INVALID_NULL_HANDLE", "a handle argument is null when it should not be">,
+    Etor<"INVALID_PLATFORM", "invalid platform">,
+    Etor<"INVALID_DEVICE", "invalid device">,
+    Etor<"INVALID_QUEUE", "invalid queue">,
+    Etor<"INVALID_EVENT", "invalid event">,
   ];
 }
 
diff --git a/offload/liboffload/include/generated/OffloadAPI.h b/offload/liboffload/include/generated/OffloadAPI.h
index d38ed06b91559..f7ec749f6efa3 100644
--- a/offload/liboffload/include/generated/OffloadAPI.h
+++ b/offload/liboffload/include/generated/OffloadAPI.h
@@ -20,50 +20,50 @@ extern "C" {
 ///////////////////////////////////////////////////////////////////////////////
 /// @brief Defines Return/Error codes
 typedef enum ol_errc_t {
-  /// Success
+  /// success
   OL_ERRC_SUCCESS = 0,
-  /// Unknown or internal error
+  /// unknown or internal error
   OL_ERRC_UNKNOWN = 1,
   /// I/O error on host
   OL_ERRC_HOST_IO = 2,
-  /// A provided binary image is malformed
+  /// a provided binary image is malformed
   OL_ERRC_INVALID_BINARY = 3,
-  /// A pointer argument is null when it should not be
+  /// a pointer argument is null when it should not be
   OL_ERRC_INVALID_NULL_POINTER = 4,
-  /// An argument is invalid
+  /// an argument is invalid
   OL_ERRC_INVALID_ARGUMENT = 5,
-  /// Requested object was not found in the binary image
+  /// requested object was not found in the binary image
   OL_ERRC_NOT_FOUND = 6,
-  /// Out of resources
+  /// out of resources
   OL_ERRC_OUT_OF_RESOURCES = 7,
   /// invalid size or dimensions (e.g., must not be zero, or is out of bounds)
   OL_ERRC_INVALID_SIZE = 8,
   /// enumerator argument is not valid
   OL_ERRC_INVALID_ENUMERATION = 9,
-  /// A required binary (linker, etc.) was not found on the host
+  /// a required binary (linker, etc.) was not found on the host
   OL_ERRC_HOST_TOOL_NOT_FOUND = 10,
-  /// Invalid Value
+  /// invalid value
   OL_ERRC_INVALID_VALUE = 11,
-  /// Generic error code for features currently unimplemented by the
+  /// generic error code for features currently unimplemented by the
   /// device/backend
   OL_ERRC_UNIMPLEMENTED = 12,
-  /// Generic error code for features unsupported by the device/backend
+  /// generic error code for features unsupported by the device/backend
   OL_ERRC_UNSUPPORTED = 13,
-  /// Assembler failure while processing binary image
+  /// assembler failure while processing binary image
   OL_ERRC_ASSEMBLE_FAILURE = 14,
-  /// Linker failure while processing binary image
+  /// linker failure while processing binary image
   OL_ERRC_LINK_FAILURE = 15,
-  /// The plugin backend is in an invalid or unsupported state
+  /// the plugin backend is in an invalid or unsupported state
   OL_ERRC_BACKEND_FAILURE = 16,
-  /// A handle argument is null when it should not be
+  /// a handle argument is null when it should not be
   OL_ERRC_INVALID_NULL_HANDLE = 17,
-  /// Invalid platform
+  /// invalid platform
   OL_ERRC_INVALID_PLATFORM = 18,
-  /// Invalid device
+  /// invalid device
   OL_ERRC_INVALID_DEVICE = 19,
-  /// Invalid queue
+  /// invalid queue
   OL_ERRC_INVALID_QUEUE = 20,
-  /// Invalid event
+  /// invalid event
   OL_ERRC_INVALID_EVENT = 21,
   /// @cond
   OL_ERRC_FORCE_UINT32 = 0x7fffffff
diff --git a/offload/libomptarget/PluginManager.cpp b/offload/libomptarget/PluginManager.cpp
index d99d6ad7b017c..93589960a426d 100644
--- a/offload/libomptarget/PluginManager.cpp
+++ b/offload/libomptarget/PluginManager.cpp
@@ -538,9 +538,9 @@ Expected<DeviceTy &> PluginManager::getDevice(uint32_t DeviceNo) {
   {
     auto ExclusiveDevicesAccessor = getExclusiveDevicesAccessor();
     if (DeviceNo >= ExclusiveDevicesAccessor->size())
-      return createStringError(
-          inconvertibleErrorCode(),
-          "Device number '%i' out of range, only %i devices available",
+      return error::createOffloadError(
+          error::ErrorCode::INVALID_VALUE,
+          "device number '%i' out of range, only %i devices available",
           DeviceNo, ExclusiveDevicesAccessor->size());
 
     DevicePtr = &*(*ExclusiveDevicesAccessor)[DeviceNo];
@@ -549,8 +549,8 @@ Expected<DeviceTy &> PluginManager::getDevice(uint32_t DeviceNo) {
   // Check whether global data has been mapped for this device
   if (DevicePtr->hasPendingImages())
     if (loadImagesOntoDevice(*DevicePtr) != OFFLOAD_SUCCESS)
-      return createStringError(inconvertibleErrorCode(),
-                               "Failed to load images on device '%i'",
-                               DeviceNo);
+      return error::createOffloadError(error::ErrorCode::BACKEND_FAILURE,
+                                       "failed to load images on device '%i'",
+                                       DeviceNo);
   return *DevicePtr;
 }
diff --git a/offload/libomptarget/device.cpp b/offload/libomptarget/device.cpp
index 2beb4093572da..f88e30ae9e76b 100644
--- a/offload/libomptarget/device.cpp
+++ b/offload/libomptarget/device.cpp
@@ -79,9 +79,9 @@ DeviceTy::~DeviceTy() {
 llvm::Error DeviceTy::init() {
   int32_t Ret = RTL->init_device(RTLDeviceID);
   if (Ret != OFFLOAD_SUCCESS)
-    return llvm::createStringError(llvm::inconvertibleErrorCode(),
-                                   "Failed to initialize device %d\n",
-                                   DeviceID);
+    return error::createOffloadError(error::ErrorCode::BACKEND_FAILURE,
+                                     "failed to initialize device %d\n",
+                                     DeviceID);
 
   // Enables recording kernels if set.
   BoolEnvar OMPX_RecordKernel("LIBOMPTARGET_RECORD", false);
@@ -103,8 +103,8 @@ DeviceTy::loadBinary(__tgt_device_image *Img) {
   __tgt_device_binary Binary;
 
   if (RTL->load_binary(RTLDeviceID, Img, &Binary) != OFFLOAD_SUCCESS)
-    return llvm::createStringError(llvm::inconvertibleErrorCode(),
-                                   "Failed to load binary %p", Img);
+    return error::createOffloadError(error::ErrorCode::INVALID_BINARY,
+                                     "failed to load binary %p", Img);
   return Binary;
 }
 
diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
index ad2ba9b23081d..2733796611d9b 100644
--- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -134,14 +134,14 @@ hsa_status_t iterate(IterFuncTy Func, IterFuncArgTy FuncArg, CallbackTy Cb) {
 /// Iterate agents.
 template <typename CallbackTy> Error iterateAgents(CallbackTy Callback) {
   hsa_status_t Status = iterate<hsa_agent_t>(hsa_iterate_agents, Callback);
-  return Plugin::check(Status, "Error in hsa_iterate_agents: %s");
+  return Plugin::check(Status, "error in hsa_iterate_agents: %s");
 }
 
 /// Iterate ISAs of an agent.
 template <typename CallbackTy>
 Error iterateAgentISAs(hsa_agent_t Agent, CallbackTy Cb) {
   hsa_status_t Status = iterate<hsa_isa_t>(hsa_agent_iterate_isas, Agent, Cb);
-  return Plugin::check(Status, "Error in hsa_agent_iterate_isas: %s");
+  return Plugin::check(Status, "error in hsa_agent_iterate_isas: %s");
 }
 
 /// Iterate memory pools of an agent.
@@ -150,7 +150,7 @@ Error iterateAgentMemoryPools(hsa_agent_t Agent, CallbackTy Cb) {
   hsa_status_t Status = iterate<hsa_amd_memory_pool_t>(
       hsa_amd_agent_iterate_memory_pools, Agent, Cb);
   return Plugin::check(Status,
-                       "Error in hsa_amd_agent_iterate_memory_pools: %s");
+                       "error in hsa_amd_agent_iterate_memory_pools: %s");
 }
 
 /// Dispatches an asynchronous memory copy.
@@ -163,14 +163,14 @@ Error asyncMemCopy(bool UseMultipleSdmaEngines, void *Dst, hsa_agent_t DstAgent,
     hsa_status_t S =
         hsa_amd_memory_async_copy(Dst, DstAgent, Src, SrcAgent, Size,
                                   NumDepSignals, DepSignals, CompletionSignal);
-    return Plugin::check(S, "Error in hsa_amd_memory_async_copy: %s");
+    return Plugin::check(S, "error in hsa_amd_memory_async_copy: %s");
   }
 
 // This solution is probably not the best
 #if !(HSA_AMD_INTERFACE_VERSION_MAJOR >= 1 &&                                  \
       HSA_AMD_INTERFACE_VERSION_MINOR >= 2)
   return Plugin::error(ErrorCode::UNSUPPORTED,
-                       "Async copy on selected SDMA requires ROCm 5.7");
+                       "async copy on selected SDMA requires ROCm 5.7");
 #else
   static std::atomic<int> SdmaEngine{1};
 
@@ -189,7 +189,7 @@ Error asyncMemCopy(bool UseMultipleSdmaEngines, void *Dst, hsa_agent_t DstAgent,
   LocalSdmaEngine = (LocalSdmaEngine << 1) % 3;
   SdmaEngine.store(LocalSdmaEngine, std::memory_order_relaxed);
 
-  return Plugin::check(S, "Error in hsa_amd_memory_async_copy_on_engine: %s");
+  return Plugin::check(S, "error in hsa_amd_memory_async_copy_on_engine: %s");
 #endif
 }
 
@@ -241,7 +241,7 @@ struct AMDGPUResourceRef : public GenericDeviceResourceRef {
   Error destroy(GenericDeviceTy &Device) override {
     if (!Resource)
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Destroying an invalid resource");
+                           "destroying an invalid resource");
 
     if (auto Err = Resource->deinit())
       return Err;
@@ -308,13 +308,13 @@ struct AMDGPUMemoryPoolTy {
   Error allocate(size_t Size, void **PtrStorage) {
     hsa_status_t Status =
         hsa_amd_memory_pool_allocate(MemoryPool, Size, 0, PtrStorage);
-    return Plugin::check(Status, "Error in hsa_amd_memory_pool_allocate: %s");
+    return Plugin::check(Status, "error in hsa_amd_memory_pool_allocate: %s");
   }
 
   /// Return memory to the memory pool.
   Error deallocate(void *Ptr) {
     hsa_status_t Status = hsa_amd_memory_pool_free(Ptr);
-    return Plugin::check(Status, "Error in hsa_amd_memory_pool_free: %s");
+    return Plugin::check(Status, "error in hsa_amd_memory_pool_free: %s");
   }
 
   /// Returns if the \p Agent can access the memory pool.
@@ -340,14 +340,14 @@ struct AMDGPUMemoryPoolTy {
       // continue because otherwise it result in undefined behavior.
       if (Access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED)
         return Plugin::error(ErrorCode::INVALID_VALUE,
-                             "An agent is not allowed to access a memory pool");
+                             "an agent is not allowed to access a memory pool");
     }
 #endif
 
     // We can access but it is disabled by default. Enable the access then.
     hsa_status_t Status =
         hsa_amd_agents_allow_access(Agents.size(), Agents.data(), nullptr, Ptr);
-    return Plugin::check(Status, "Error in hsa_amd_agents_allow_access: %s");
+    return Plugin::check(Status, "error in hsa_amd_agents_allow_access: %s");
   }
 
   /// Get attribute from the memory pool.
@@ -355,7 +355,7 @@ struct AMDGPUMemoryPoolTy {
   Error getAttr(hsa_amd_memory_pool_info_t Kind, Ty &Value) const {
     hsa_status_t Status;
     Status = hsa_amd_memory_pool_get_info(MemoryPool, Kind, &Value);
-    return Plugin::check(Status, "Error in hsa_amd_memory_pool_get_info: %s");
+    return Plugin::check(Status, "error in hsa_amd_memory_pool_get_info: %s");
   }
 
   template <typename Ty>
@@ -371,7 +371,7 @@ struct AMDGPUMemoryPoolTy {
     Status =
         hsa_amd_agent_memory_pool_get_info(Agent, MemoryPool, Kind, &Value);
     return Plugin::check(Status,
-                         "Error in hsa_amd_agent_memory_pool_get_info: %s");
+                         "error in hsa_amd_agent_memory_pool_get_info: %s");
   }
 
 private:
@@ -422,7 +422,7 @@ struct AMDGPUMemoryManagerTy : public DeviceAllocatorTy {
     *PtrStorage = MemoryManager->allocate(Size, nullptr);
     if (*PtrStorage == nullptr)
       return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
-                           "Failure to allocate from AMDGPU memory manager");
+                           "failure to allocate from AMDGPU memory manager");
 
     return Plugin::success();
   }
@@ -433,7 +433,7 @@ struct AMDGPUMemoryManagerTy : public DeviceAllocatorTy {
 
     if (MemoryManager->free(Ptr))
       return Plugin::error(ErrorCode::UNKNOWN,
-                           "Failure to deallocate from AMDGPU memory manager");
+                           "failure to deallocate from AMDGPU memory manager");
 
     return Plugin::success();
   }
@@ -475,7 +475,7 @@ struct AMDGPUDeviceImageTy : public DeviceImageTy {
   /// Unload the executable.
   Error unloadExecutable() {
     hsa_status_t Status = hsa_executable_destroy(Executable);
-    return Plugin::check(Status, "Error in hsa_executable_destroy: %s");
+    return Plugin::check(Status, "error in hsa_executable_destroy: %s");
   }
 
   /// Get the executable.
@@ -541,14 +541,14 @@ struct AMDGPUKernelTy : public GenericKernelTy {
     for (auto &Info : RequiredInfos) {
       Status = hsa_executable_symbol_get_info(Symbol, Info.first, Info.second);
       if (auto Err = Plugin::check(
-              Status, "Error in hsa_executable_symbol_get_info: %s"))
+              Status, "error in hsa_executable_symbol_get_info: %s"))
         return Err;
     }
 
     // Make sure it is a kernel symbol.
     if (SymbolType != HSA_SYMBOL_KIND_KERNEL)
       return Plugin::error(ErrorCode::INVALID_BINARY,
-                           "Symbol %s is not a kernel function");
+                           "symbol %s is not a kernel function");
 
     // TODO: Read the kernel descriptor for the max threads per block. May be
     // read from the image.
@@ -618,13 +618,13 @@ struct AMDGPUSignalTy {
   Error init(uint32_t InitialValue = 1) {
     hsa_status_t Status =
         hsa_amd_signal_create(InitialValue, 0, nullptr, 0, &HSASignal);
-    return Plugin::check(Status, "Error in hsa_signal_create: %s");
+    return Plugin::check(Status, "error in hsa_signal_create: %s");
   }
 
   /// Deinitialize the signal.
   Error deinit() {
     hsa_status_t Status = hsa_signal_destroy(HSASignal);
-    return Plugin::check(Status, "Error in hsa_signal_destroy: %s");
+    return Plugin::check(Status, "error in hsa_signal_destroy: %s");
   }
 
   /// Wait until the signal gets a zero value.
@@ -696,7 +696,7 @@ struct AMDGPUQueueTy {
     hsa_status_t Status =
         hsa_queue_create(Agent, QueueSize, HSA_QUEUE_TYPE_MULTI, callbackError,
                          &Device, UINT32_MAX, UINT32_MAX, &Queue);
-    return Plugin::check(Status, "Error in hsa_queue_create: %s");
+    return Plugin::check(Status, "error in hsa_queue_create: %s");
   }
 
   /// Deinitialize the queue and destroy its resources.
@@ -705,7 +705,7 @@ struct AMDGPUQueueTy {
     if (!Queue)
       return Plugin::success();
     hsa_status_t Status = hsa_queue_destroy(Queue);
-    return Plugin::check(Status, "Error in hsa_queue_destroy: %s");
+    return Plugin::check(Status, "error in hsa_queue_destroy: %s");
   }
 
   /// Returns the number of streams, this queue is currently assigned to.
@@ -1124,7 +1124,7 @@ struct AMDGPUStreamTy {
   Error waitOnStreamOperation(AMDGPUStreamTy &OtherStream, uint32_t Slot) {
     if (Queue == nullptr)
       return Plugin::error(ErrorCode::INVALID_NULL_POINTER,
-                           "Target queue was nullptr");
+                           "target queue was nullptr");
 
     /// The signal that we must wait from the other stream.
     AMDGPUSignalTy *OtherSignal = OtherStream.Slots[Slot].Signal;
@@ -1246,7 +1246,7 @@ struct AMDGPUStreamTy {
                          AMDGPUMemoryManagerTy &MemoryManager) {
     if (Queue == nullptr)
       return Plugin::error(ErrorCode::INVALID_NULL_POINTER,
-                           "Target queue was nullptr");
+                           "target queue was nullptr");
 
     // Retrieve an available signal for the operation's output.
     AMDGPUSignalTy *OutputSignal = nullptr;
@@ -1377,7 +1377,7 @@ struct AMDGPUStreamTy {
         InputSignal->get(), HSA_SIGNAL_CONDITION_EQ, 0, asyncActionCallback,
         (void *)&Slots[Curr]);
 
-    return Plugin::check(Status, "Error in hsa_amd_signal_async_handler: %s");
+    return Plugin::check(Status, "error in hsa_amd_signal_async_handler: %s");
   }
 
   /// Push an asynchronous memory copy host-to-device involving an unpinned
@@ -1421,7 +1421,7 @@ struct AMDGPUStreamTy {
           (void *)&Slots[Curr]);
 
       if (auto Err = Plugin::check(Status,
-                                   "Error in hsa_amd_signal_async_handler: %s"))
+                                   "error in hsa_amd_signal_async_handler: %s"))
         return Err;
 
       // Let's use now the second output signal.
@@ -1564,7 +1564,7 @@ struct AMDGPUEventTy {
 
     if (!RecordedStream)
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Event does not have any recorded stream");
+                           "event does not have any recorded stream");
 
     // Synchronizing the same stream. Do nothing.
     if (RecordedStream == &Stream)
@@ -1954,7 +1954,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
       GridValues = getAMDGPUGridValues<64>();
     else
       return Plugin::error(ErrorCode::UNSUPPORTED,
-                           "Unexpected AMDGPU wavefront %d", WavefrontSize);
+                           "unexpected AMDGPU wavefront %d", WavefrontSize);
 
     // Get maximum number of workitems per workgroup.
     uint16_t WorkgroupMaxDim[3];
@@ -1971,7 +1971,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     GridValues.GV_Max_Teams = GridMaxDim.x / GridValues.GV_Max_WG_Size;
     if (GridValues.GV_Max_Teams == 0)
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Maximum number of teams cannot be zero");
+                           "maximum number of teams cannot be zero");
 
     // Compute the default number of teams.
     uint32_t ComputeUnits = 0;
@@ -2085,7 +2085,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
                                                       "o", LinkerInputFilePath);
     if (EC)
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to create temporary file for linker");
+                           "failed to create temporary file for linker");
 
     // Write the file's contents to the output file.
     Expected<std::unique_ptr<FileOutputBuffer>> OutputOrErr =
@@ -2102,12 +2102,12 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
                                       LinkerOutputFilePath);
     if (EC)
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to create temporary file for linker");
+                           "failed to create temporary file for linker");
 
     const auto &ErrorOrPath = sys::findProgramByName("lld");
     if (!ErrorOrPath)
       return createStringError(ErrorCode::HOST_TOOL_NOT_FOUND,
-                               "Failed to find `lld` on the PATH.");
+                               "failed to find `lld` on the PATH.");
 
     std::string LLDPath = ErrorOrPath.get();
     INFO(OMP_INFOTYPE_PLUGIN_KERNEL, getDeviceId(),
@@ -2128,21 +2128,21 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     int RC = sys::ExecuteAndWait(LLDPath, Args, std::nullopt, {}, 0, 0, &Error);
     if (RC)
       return Plugin::error(ErrorCode::LINK_FAILURE,
-                           "Linking optimized bitcode failed: %s",
+                           "linking optimized bitcode failed: %s",
                            Error.c_str());
 
     auto BufferOrErr = MemoryBuffer::getFileOrSTDIN(LinkerOutputFilePath);
     if (!BufferOrErr)
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to open temporary file for lld");
+                           "failed to open temporary file for lld");
 
     // Clean up the temporary files afterwards.
     if (sys::fs::remove(LinkerOutputFilePath))
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to remove temporary output file for lld");
+                           "failed to remove temporary output file for lld");
     if (sys::fs::remove(LinkerInputFilePath))
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to remove temporary input file for lld");
+                           "failed to remove temporary input file for lld");
 
     return std::move(*BufferOrErr);
   }
@@ -2159,7 +2159,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     AMDGPUKernelTy *AMDGPUKernel = Plugin.allocate<AMDGPUKernelTy>();
     if (!AMDGPUKernel)
       return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
-                           "Failed to allocate memory for AMDGPU kernel");
+                           "failed to allocate memory for AMDGPU kernel");
 
     new (AMDGPUKernel) AMDGPUKernelTy(Name);
 
@@ -2295,7 +2295,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
 
     hsa_status_t Status =
         hsa_amd_memory_lock(HstPtr, Size, nullptr, 0, &PinnedPtr);
-    if (auto Err = Plugin::check(Status, "Error in hsa_amd_memory_lock: %s\n"))
+    if (auto Err = Plugin::check(Status, "error in hsa_amd_memory_lock: %s\n"))
       return std::move(Err);
 
     return PinnedPtr;
@@ -2304,7 +2304,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
   /// Unpin the host buffer.
   Error dataUnlockImpl(void *HstPtr) override {
     hsa_status_t Status = hsa_amd_memory_unlock(HstPtr);
-    return Plugin::check(Status, "Error in hsa_amd_memory_unlock: %s\n");
+    return Plugin::check(Status, "error in hsa_amd_memory_unlock: %s\n");
   }
 
   /// Check through the HSA runtime whether the \p HstPtr buffer is pinned.
@@ -2317,7 +2317,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     hsa_status_t Status = hsa_amd_pointer_info(
         HstPtr, &Info, /*Allocator=*/nullptr, /*num_agents_accessible=*/nullptr,
         /*accessible=*/nullptr);
-    if (auto Err = Plugin::check(Status, "Error in hsa_amd_pointer_info: %s"))
+    if (auto Err = Plugin::check(Status, "error in hsa_amd_pointer_info: %s"))
       return std::move(Err);
 
     // The buffer may be locked or allocated through HSA allocators. Assume that
@@ -2362,7 +2362,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
       Status = hsa_amd_memory_lock(const_cast<void *>(HstPtr), Size, nullptr, 0,
                                    &PinnedPtr);
       if (auto Err =
-              Plugin::check(Status, "Error in hsa_amd_memory_lock: %s\n"))
+              Plugin::check(Status, "error in hsa_amd_memory_lock: %s\n"))
         return Err;
 
       AMDGPUSignalTy Signal;
@@ -2381,7 +2381,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
         return Err;
 
       Status = hsa_amd_memory_unlock(const_cast<void *>(HstPtr));
-      return Plugin::check(Status, "Error in hsa_amd_memory_unlock: %s\n");
+      return Plugin::check(Status, "error in hsa_amd_memory_unlock: %s\n");
     }
 
     // Otherwise, use two-step copy with an intermediate pinned host buffer.
@@ -2422,7 +2422,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
       Status = hsa_amd_memory_lock(const_cast<void *>(HstPtr), Size, nullptr, 0,
                                    &PinnedPtr);
       if (auto Err =
-              Plugin::check(Status, "Error in hsa_amd_memory_lock: %s\n"))
+              Plugin::check(Status, "error in hsa_amd_memory_lock: %s\n"))
         return Err;
 
       AMDGPUSignalTy Signal;
@@ -2441,7 +2441,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
         return Err;
 
       Status = hsa_amd_memory_unlock(const_cast<void *>(HstPtr));
-      return Plugin::check(Status, "Error in hsa_amd_memory_unlock: %s\n");
+      return Plugin::check(Status, "error in hsa_amd_memory_unlock: %s\n");
     }
 
     // Otherwise, use two-step copy with an intermediate pinned host buffer.
@@ -2550,7 +2550,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
   /// Synchronize the current thread with the event.
   Error syncEventImpl(void *EventPtr) override {
     return Plugin::error(ErrorCode::UNIMPLEMENTED,
-                         "Synchronize event not implemented");
+                         "synchronize event not implemented");
   }
 
   /// Print information about the device.
@@ -2795,7 +2795,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
       if (Pool->isGlobal()) {
         hsa_status_t Status =
             Pool->getAttrRaw(HSA_AMD_MEMORY_POOL_INFO_SIZE, Value);
-        return Plugin::check(Status, "Error in getting device memory size: %s");
+        return Plugin::check(Status, "error in getting device memory size: %s");
       }
     }
     return Plugin::error(ErrorCode::UNSUPPORTED,
@@ -2977,38 +2977,38 @@ Error AMDGPUDeviceImageTy::loadExecutable(const AMDGPUDeviceTy &Device) {
   hsa_status_t Status =
       hsa_code_object_reader_create_from_memory(getStart(), getSize(), &Reader);
   if (auto Err = Plugin::check(
-          Status, "Error in hsa_code_object_reader_create_from_memory: %s"))
+          Status, "error in hsa_code_object_reader_create_from_memory: %s"))
     return Err;
 
   Status = hsa_executable_create_alt(
       HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO, "", &Executable);
   if (auto Err =
-          Plugin::check(Status, "Error in hsa_executable_create_alt: %s"))
+          Plugin::check(Status, "error in hsa_executable_create_alt: %s"))
     return Err;
 
   hsa_loaded_code_object_t Object;
   Status = hsa_executable_load_agent_code_object(Executable, Device.getAgent(),
                                                  Reader, "", &Object);
   if (auto Err = Plugin::check(
-          Status, "Error in hsa_executable_load_agent_code_object: %s"))
+          Status, "error in hsa_executable_load_agent_code_object: %s"))
     return Err;
 
   Status = hsa_executable_freeze(Executable, "");
-  if (auto Err = Plugin::check(Status, "Error in hsa_executable_freeze: %s"))
+  if (auto Err = Plugin::check(Status, "error in hsa_executable_freeze: %s"))
     return Err;
 
   uint32_t Result;
   Status = hsa_executable_validate(Executable, &Result);
-  if (auto Err = Plugin::check(Status, "Error in hsa_executable_validate: %s"))
+  if (auto Err = Plugin::check(Status, "error in hsa_executable_validate: %s"))
     return Err;
 
   if (Result)
     return Plugin::error(ErrorCode::INVALID_BINARY,
-                         "Loaded HSA executable does not validate");
+                         "loaded HSA executable does not validate");
 
   Status = hsa_code_object_reader_destroy(Reader);
   if (auto Err =
-          Plugin::check(Status, "Error in hsa_code_object_reader_destroy: %s"))
+          Plugin::check(Status, "error in hsa_code_object_reader_destroy: %s"))
     return Err;
 
   if (auto Err = hsa_utils::readAMDGPUMetaDataFromImage(
@@ -3029,7 +3029,7 @@ AMDGPUDeviceImageTy::findDeviceSymbol(GenericDeviceTy &Device,
   hsa_status_t Status = hsa_executable_get_symbol_by_name(
       Executable, SymbolName.data(), &Agent, &Symbol);
   if (auto Err = Plugin::check(
-          Status, "Error in hsa_executable_get_symbol_by_name(%s): %s",
+          Status, "error in hsa_executable_get_symbol_by_name(%s): %s",
           SymbolName.data()))
     return std::move(Err);
 
@@ -3040,7 +3040,7 @@ template <typename ResourceTy>
 Error AMDGPUResourceRef<ResourceTy>::create(GenericDeviceTy &Device) {
   if (Resource)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Creating an existing resource");
+                         "creating an existing resource");
 
   AMDGPUDeviceTy &AMDGPUDevice = static_cast<AMDGPUDeviceTy &>(Device);
 
@@ -3089,7 +3089,7 @@ struct AMDGPUGlobalHandlerTy final : public GenericGlobalHandlerTy {
     for (auto &Info : RequiredInfos) {
       Status = hsa_executable_symbol_get_info(Symbol, Info.first, Info.second);
       if (auto Err = Plugin::check(
-              Status, "Error in hsa_executable_symbol_get_info: %s"))
+              Status, "error in hsa_executable_symbol_get_info: %s"))
         return Err;
     }
 
@@ -3097,7 +3097,7 @@ struct AMDGPUGlobalHandlerTy final : public GenericGlobalHandlerTy {
     if (SymbolSize != DeviceGlobal.getSize())
       return Plugin::error(
           ErrorCode::INVALID_BINARY,
-          "Failed to load global '%s' due to size mismatch (%zu != %zu)",
+          "failed to load global '%s' due to size mismatch (%zu != %zu)",
           DeviceGlobal.getName().data(), SymbolSize,
           (size_t)DeviceGlobal.getSize());
 
@@ -3135,7 +3135,7 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
     // Register event handler to detect memory errors on the devices.
     Status = hsa_amd_register_system_event_handler(eventHandler, this);
     if (auto Err = Plugin::check(
-            Status, "Error in hsa_amd_register_system_event_handler: %s"))
+            Status, "error in hsa_amd_register_system_event_handler: %s"))
       return std::move(Err);
 
     // List of host (CPU) agents.
@@ -3176,7 +3176,7 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
     // There are kernel agents but there is no host agent. That should be
     // treated as an error.
     if (HostAgents.empty())
-      return Plugin::error(ErrorCode::BACKEND_FAILURE, "No AMDGPU host agents");
+      return Plugin::error(ErrorCode::BACKEND_FAILURE, "no AMDGPU host agents");
 
     // Initialize the host device using host agents.
     HostDevice = allocate<AMDHostDeviceTy>();
@@ -3202,7 +3202,7 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
 
     // Finalize the HSA runtime.
     hsa_status_t Status = hsa_shut_down();
-    return Plugin::check(Status, "Error in hsa_shut_down: %s");
+    return Plugin::check(Status, "error in hsa_shut_down: %s");
   }
 
   /// Creates an AMDGPU device.
@@ -3322,7 +3322,7 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
       void *DevicePtr = (void *)Event->memory_fault.virtual_address;
       std::string S;
       llvm::raw_string_ostream OS(S);
-      OS << llvm::format("Memory access fault by GPU %" PRIu32
+      OS << llvm::format("memory access fault by GPU %" PRIu32
                          " (agent 0x%" PRIx64
                          ") at virtual address %p. Reasons: %s",
                          Node, Event->memory_fault.agent.handle,
@@ -3335,7 +3335,7 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
 
     // Abort the execution since we do not recover from this error.
     FATAL_MESSAGE(1,
-                  "Memory access fault by GPU %" PRIu32 " (agent 0x%" PRIx64
+                  "memory access fault by GPU %" PRIu32 " (agent 0x%" PRIx64
                   ") at virtual address %p. Reasons: %s",
                   Node, Event->memory_fault.agent.handle,
                   (void *)Event->memory_fault.virtual_address,
@@ -3367,7 +3367,7 @@ Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
   if (ArgsSize != LaunchParams.Size &&
       ArgsSize != LaunchParams.Size + getImplicitArgsSize())
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Mismatch of kernel arguments size");
+                         "mismatch of kernel arguments size");
 
   AMDGPUPluginTy &AMDGPUPlugin =
       static_cast<AMDGPUPluginTy &>(GenericDevice.Plugin);
@@ -3480,9 +3480,9 @@ template <typename... ArgsTy>
 static Error Plugin::check(int32_t Code, const char *ErrFmt, ArgsTy... Args) {
   hsa_status_t ResultCode = static_cast<hsa_status_t>(Code);
   if (ResultCode == HSA_STATUS_SUCCESS || ResultCode == HSA_STATUS_INFO_BREAK)
-    return Error::success();
+    return Plugin::success();
 
-  const char *Desc = "Unknown error";
+  const char *Desc = "unknown error";
   hsa_status_t Ret = hsa_status_string(ResultCode, &Desc);
   if (Ret != HSA_STATUS_SUCCESS)
     REPORT("Unrecognized " GETNAME(TARGET_NAME) " error code %d\n", Code);
@@ -3497,8 +3497,7 @@ static Error Plugin::check(int32_t Code, const char *ErrFmt, ArgsTy... Args) {
     OffloadErrCode = ErrorCode::UNKNOWN;
   }
 
-  return createStringError<ArgsTy..., const char *>(OffloadErrCode, ErrFmt,
-                                                    Args..., Desc);
+  return Plugin::error(OffloadErrCode, ErrFmt, Args..., Desc);
 }
 
 void *AMDGPUMemoryManagerTy::allocate(size_t Size, void *HstPtr,
@@ -3595,7 +3594,7 @@ void AMDGPUQueueTy::callbackError(hsa_status_t Status, hsa_queue_t *Source,
                                       AsyncInfoWrapperMatcher);
   }
 
-  auto Err = Plugin::check(Status, "Received error in queue %p: %s", Source);
+  auto Err = Plugin::check(Status, "received error in queue %p: %s", Source);
   FATAL_MESSAGE(1, "%s", toString(std::move(Err)).data());
 }
 
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index 09b5c8ed3e8ab..fa13e26bc3483 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -1383,35 +1383,16 @@ static inline Error success() { return Error::success(); }
 /// Create an Offload error.
 template <typename... ArgsTy>
 static Error error(error::ErrorCode Code, const char *ErrFmt, ArgsTy... Args) {
-  std::string Buffer;
-  raw_string_ostream(Buffer) << format(ErrFmt, Args...);
-  return make_error<error::OffloadError>(Code, Buffer);
+  return error::createOffloadError(Code, ErrFmt, Args...);
 }
 
 inline Error error(error::ErrorCode Code, const char *S) {
   return make_error<error::OffloadError>(Code, S);
 }
 
-// The OffloadError will have a message of either:
-// * "{Context}: {Message}" if the other error is a StringError
-// * "{Context}" otherwise
 inline Error error(error::ErrorCode Code, Error &&OtherError,
                    const char *Context) {
-  std::string Buffer{Context};
-  raw_string_ostream buffer(Buffer);
-
-  handleAllErrors(
-      std::move(OtherError),
-      [&](llvm::StringError &Err) {
-        buffer << ": ";
-        buffer << Err.getMessage();
-      },
-      [&](llvm::ErrorInfoBase &Err) {
-        // Non-string error message don't add anything to the offload error's
-        // error message
-      });
-
-  return make_error<error::OffloadError>(Code, Buffer);
+  return error::createOffloadError(Code, std::move(OtherError), Context);
 }
 
 /// Check the plugin-specific error code and return an error or success
diff --git a/offload/plugins-nextgen/common/src/GlobalHandler.cpp b/offload/plugins-nextgen/common/src/GlobalHandler.cpp
index 4166980de1606..27d7e8ee2fdf3 100644
--- a/offload/plugins-nextgen/common/src/GlobalHandler.cpp
+++ b/offload/plugins-nextgen/common/src/GlobalHandler.cpp
@@ -37,7 +37,7 @@ GenericGlobalHandlerTy::getELFObjectFile(DeviceImageTy &Image) {
       ELFObjectFileBase::createELFObjectFile(Image.getMemoryBuffer());
   if (!Expected) {
     return Plugin::error(ErrorCode::INVALID_BINARY, Expected.takeError(),
-                         "Error parsing binary");
+                         "error parsing binary");
   }
   return Expected;
 }
@@ -120,19 +120,19 @@ Error GenericGlobalHandlerTy::getGlobalMetadataFromImage(
   auto SymOrErr = utils::elf::getSymbol(**ELFObj, ImageGlobal.getName());
   if (!SymOrErr)
     return Plugin::error(
-        ErrorCode::NOT_FOUND, "Failed ELF lookup of global '%s': %s",
+        ErrorCode::NOT_FOUND, "failed ELF lookup of global '%s': %s",
         ImageGlobal.getName().data(), toString(SymOrErr.takeError()).data());
 
   if (!SymOrErr->has_value())
     return Plugin::error(ErrorCode::NOT_FOUND,
-                         "Failed to find global symbol '%s' in the ELF image",
+                         "failed to find global symbol '%s' in the ELF image",
                          ImageGlobal.getName().data());
 
   auto AddrOrErr = utils::elf::getSymbolAddress(**SymOrErr);
   // Get the section to which the symbol belongs.
   if (!AddrOrErr)
     return Plugin::error(
-        ErrorCode::NOT_FOUND, "Failed to get ELF symbol from global '%s': %s",
+        ErrorCode::NOT_FOUND, "failed to get ELF symbol from global '%s': %s",
         ImageGlobal.getName().data(), toString(AddrOrErr.takeError()).data());
 
   // Setup the global symbol's address and size.
@@ -152,7 +152,7 @@ Error GenericGlobalHandlerTy::readGlobalFromImage(GenericDeviceTy &Device,
 
   if (ImageGlobal.getSize() != HostGlobal.getSize())
     return Plugin::error(ErrorCode::INVALID_BINARY,
-                         "Transfer failed because global symbol '%s' has "
+                         "transfer failed because global symbol '%s' has "
                          "%u bytes in the ELF image but %u bytes on the host",
                          HostGlobal.getName().data(), ImageGlobal.getSize(),
                          HostGlobal.getSize());
@@ -284,7 +284,7 @@ void GPUProfGlobals::dump() const {
 Error GPUProfGlobals::write() const {
   if (!__llvm_write_custom_profile)
     return Plugin::error(ErrorCode::INVALID_BINARY,
-                         "Could not find symbol __llvm_write_custom_profile. "
+                         "could not find symbol __llvm_write_custom_profile. "
                          "The compiler-rt profiling library must be linked for "
                          "GPU PGO to work.");
 
@@ -318,7 +318,7 @@ Error GPUProfGlobals::write() const {
       CountersEnd, NamesBegin, NamesEnd, &Version);
   if (result != 0)
     return Plugin::error(ErrorCode::HOST_IO,
-                         "Error writing GPU PGO data to file");
+                         "error writing GPU PGO data to file");
 
   return Plugin::success();
 }
diff --git a/offload/plugins-nextgen/common/src/JIT.cpp b/offload/plugins-nextgen/common/src/JIT.cpp
index 093ebe16cc851..c82a06e36d8f9 100644
--- a/offload/plugins-nextgen/common/src/JIT.cpp
+++ b/offload/plugins-nextgen/common/src/JIT.cpp
@@ -62,8 +62,8 @@ createModuleFromMemoryBuffer(std::unique_ptr<MemoryBuffer> &MB,
   SMDiagnostic Err;
   auto Mod = parseIR(*MB, Err, Context);
   if (!Mod)
-    return make_error<StringError>("Failed to create module",
-                                   inconvertibleErrorCode());
+    return error::createOffloadError(error::ErrorCode::UNKNOWN,
+                                     "failed to create module");
   return std::move(Mod);
 }
 Expected<std::unique_ptr<Module>>
@@ -100,7 +100,8 @@ createTargetMachine(Module &M, std::string CPU, unsigned OptLevel) {
   std::string Msg;
   const Target *T = TargetRegistry::lookupTarget(M.getTargetTriple(), Msg);
   if (!T)
-    return make_error<StringError>(Msg, inconvertibleErrorCode());
+    return error::createOffloadError(error::ErrorCode::INVALID_BINARY,
+                                     Msg.data());
 
   SubtargetFeatures Features;
   Features.getDefaultSubtargetFeatures(TT);
@@ -118,8 +119,8 @@ createTargetMachine(Module &M, std::string CPU, unsigned OptLevel) {
       T->createTargetMachine(M.getTargetTriple(), CPU, Features.getString(),
                              Options, RelocModel, CodeModel, CGOptLevel));
   if (!TM)
-    return make_error<StringError>("Failed to create target machine",
-                                   inconvertibleErrorCode());
+    return error::createOffloadError(error::ErrorCode::INVALID_BINARY,
+                                     "failed to create target machine");
   return std::move(TM);
 }
 
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index f35961627b974..f9e316adad8f4 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -96,7 +96,7 @@ struct RecordReplayTy {
 
     if (isReplaying() && VAddr != MemoryStart) {
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Record-Replay cannot assign the"
+                           "record-Replay cannot assign the"
                            "requested recorded address (%p, %p)",
                            VAddr, MemoryStart);
     }
@@ -124,7 +124,7 @@ struct RecordReplayTy {
     }
     if (!MemoryStart)
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Allocating record/replay memory");
+                           "allocating record/replay memory");
 
     if (VAddr && VAddr != MemoryStart)
       MemoryOffset = uintptr_t(VAddr) - uintptr_t(MemoryStart);
@@ -170,7 +170,7 @@ struct RecordReplayTy {
     uint64_t DevMemSize;
     if (Device->getDeviceMemorySize(DevMemSize))
       return Plugin::error(ErrorCode::UNKNOWN,
-                           "Cannot determine Device Memory Size");
+                           "cannot determine Device Memory Size");
 
     return preAllocateHeuristic(DevMemSize, DeviceMemorySize, ReqVAddr);
   }
@@ -1083,7 +1083,7 @@ Error PinnedAllocationMapTy::insertEntry(void *HstPtr, void *DevAccessiblePtr,
   auto Res = Allocs.insert({HstPtr, DevAccessiblePtr, Size, ExternallyLocked});
   if (!Res.second)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Cannot insert locked buffer entry");
+                         "cannot insert locked buffer entry");
 
   // Check whether the next entry overlaps with the inserted entry.
   auto It = std::next(Res.first);
@@ -1093,7 +1093,7 @@ Error PinnedAllocationMapTy::insertEntry(void *HstPtr, void *DevAccessiblePtr,
   const EntryTy *NextEntry = &(*It);
   if (intersects(NextEntry->HstPtr, NextEntry->Size, HstPtr, Size))
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Partial overlapping not allowed in locked buffers");
+                         "partial overlapping not allowed in locked buffers");
 
   return Plugin::success();
 }
@@ -1105,7 +1105,7 @@ Error PinnedAllocationMapTy::eraseEntry(const EntryTy &Entry) {
   size_t Erased = Allocs.erase({Entry.HstPtr});
   if (!Erased)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Cannot erase locked buffer entry");
+                         "cannot erase locked buffer entry");
   return Plugin::success();
 }
 
@@ -1113,7 +1113,7 @@ Error PinnedAllocationMapTy::registerEntryUse(const EntryTy &Entry,
                                               void *HstPtr, size_t Size) {
   if (!contains(Entry.HstPtr, Entry.Size, HstPtr, Size))
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Partial overlapping not allowed in locked buffers");
+                         "partial overlapping not allowed in locked buffers");
 
   ++Entry.References;
   return Plugin::success();
@@ -1122,7 +1122,7 @@ Error PinnedAllocationMapTy::registerEntryUse(const EntryTy &Entry,
 Expected<bool> PinnedAllocationMapTy::unregisterEntryUse(const EntryTy &Entry) {
   if (Entry.References == 0)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Invalid number of references");
+                         "invalid number of references");
 
   // Return whether this was the last user.
   return (--Entry.References == 0);
@@ -1141,7 +1141,7 @@ Error PinnedAllocationMapTy::registerHostBuffer(void *HstPtr,
   const EntryTy *Entry = findIntersecting(HstPtr);
   if (Entry)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Cannot insert entry due to an existing one");
+                         "cannot insert entry due to an existing one");
 
   // Now insert the new entry.
   return insertEntry(HstPtr, DevAccessiblePtr, Size);
@@ -1155,12 +1155,12 @@ Error PinnedAllocationMapTy::unregisterHostBuffer(void *HstPtr) {
   const EntryTy *Entry = findIntersecting(HstPtr);
   if (!Entry)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Cannot find locked buffer");
+                         "cannot find locked buffer");
 
   // The address in the entry should be the same we are unregistering.
   if (Entry->HstPtr != HstPtr)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Unexpected host pointer in locked buffer entry");
+                         "unexpected host pointer in locked buffer entry");
 
   // Unregister from the entry.
   auto LastUseOrErr = unregisterEntryUse(*Entry);
@@ -1170,7 +1170,7 @@ Error PinnedAllocationMapTy::unregisterHostBuffer(void *HstPtr) {
   // There should be no other references to the pinned allocation.
   if (!(*LastUseOrErr))
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "The locked buffer is still being used");
+                         "the locked buffer is still being used");
 
   // Erase the entry from the map.
   return eraseEntry(*Entry);
@@ -1217,7 +1217,7 @@ Error PinnedAllocationMapTy::unlockHostBuffer(void *HstPtr) {
   const EntryTy *Entry = findIntersecting(HstPtr);
   if (!Entry)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Cannot find locked buffer");
+                         "cannot find locked buffer");
 
   // Unregister from the locked buffer. No need to do anything if there are
   // others using the allocation.
@@ -1304,7 +1304,7 @@ Error PinnedAllocationMapTy::unlockUnmappedHostBuffer(void *HstPtr) {
   // No entry, but the automatic locking is enabled, so this is an error.
   if (!Entry)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Locked buffer not found");
+                         "locked buffer not found");
 
   // There is entry, so unregister a user and check whether it was the last one.
   auto LastUseOrErr = unregisterEntryUse(*Entry);
@@ -1328,7 +1328,7 @@ Error PinnedAllocationMapTy::unlockUnmappedHostBuffer(void *HstPtr) {
 Error GenericDeviceTy::synchronize(__tgt_async_info *AsyncInfo) {
   if (!AsyncInfo || !AsyncInfo->Queue)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Invalid async info queue");
+                         "invalid async info queue");
 
   if (auto Err = synchronizeImpl(*AsyncInfo))
     return Err;
@@ -1344,25 +1344,25 @@ Error GenericDeviceTy::synchronize(__tgt_async_info *AsyncInfo) {
 Error GenericDeviceTy::queryAsync(__tgt_async_info *AsyncInfo) {
   if (!AsyncInfo || !AsyncInfo->Queue)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                         "Invalid async info queue");
+                         "invalid async info queue");
 
   return queryAsyncImpl(*AsyncInfo);
 }
 
 Error GenericDeviceTy::memoryVAMap(void **Addr, void *VAddr, size_t *RSize) {
   return Plugin::error(ErrorCode::UNSUPPORTED,
-                       "Device does not support VA Management");
+                       "device does not support VA Management");
 }
 
 Error GenericDeviceTy::memoryVAUnMap(void *VAddr, size_t Size) {
   return Plugin::error(ErrorCode::UNSUPPORTED,
-                       "Device does not support VA Management");
+                       "device does not support VA Management");
 }
 
 Error GenericDeviceTy::getDeviceMemorySize(uint64_t &DSize) {
   return Plugin::error(
       ErrorCode::UNIMPLEMENTED,
-      "Missing getDeviceMemorySize implementation (required by RR-heuristic");
+      "missing getDeviceMemorySize implementation (required by RR-heuristic");
 }
 
 Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr,
@@ -1380,7 +1380,7 @@ Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr,
       Alloc = MemoryManager->allocate(Size, HostPtr);
       if (!Alloc)
         return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
-                             "Failed to allocate from memory manager");
+                             "failed to allocate from memory manager");
       break;
     }
     [[fallthrough]];
@@ -1389,14 +1389,14 @@ Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr,
     Alloc = allocate(Size, HostPtr, Kind);
     if (!Alloc)
       return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
-                           "Failed to allocate from device allocator");
+                           "failed to allocate from device allocator");
   }
 
   // Report error if the memory manager or the device allocator did not return
   // any memory buffer.
   if (!Alloc)
     return Plugin::error(ErrorCode::UNIMPLEMENTED,
-                         "Invalid target data allocation kind or requested "
+                         "invalid target data allocation kind or requested "
                          "allocator not implemented yet");
 
   // Register allocated buffer as pinned memory if the type is host memory.
@@ -1472,7 +1472,7 @@ Error GenericDeviceTy::dataDelete(void *TgtPtr, TargetAllocTy Kind) {
       if (Res)
         return Plugin::error(
             ErrorCode::OUT_OF_RESOURCES,
-            "Failure to deallocate device pointer %p via memory manager",
+            "failure to deallocate device pointer %p via memory manager",
             TgtPtr);
       break;
     }
@@ -1483,7 +1483,7 @@ Error GenericDeviceTy::dataDelete(void *TgtPtr, TargetAllocTy Kind) {
     if (Res)
       return Plugin::error(
           ErrorCode::UNKNOWN,
-          "Failure to deallocate device pointer %p via device deallocator",
+          "failure to deallocate device pointer %p via device deallocator",
           TgtPtr);
   }
 
diff --git a/offload/plugins-nextgen/common/src/RPC.cpp b/offload/plugins-nextgen/common/src/RPC.cpp
index c4afe10586cf5..678be78b56af6 100644
--- a/offload/plugins-nextgen/common/src/RPC.cpp
+++ b/offload/plugins-nextgen/common/src/RPC.cpp
@@ -177,7 +177,7 @@ Error RPCServerTy::initDevice(plugin::GenericDeviceTy &Device,
   if (!RPCBuffer)
     return plugin::Plugin::error(
         error::ErrorCode::UNKNOWN,
-        "Failed to initialize RPC server for device %d", Device.getDeviceId());
+        "failed to initialize RPC server for device %d", Device.getDeviceId());
 
   // Get the address of the RPC client from the device.
   plugin::GlobalTy ClientGlobal("__llvm_rpc_client", sizeof(rpc::Client));
diff --git a/offload/plugins-nextgen/cuda/src/rtl.cpp b/offload/plugins-nextgen/cuda/src/rtl.cpp
index 2219cd88cf93d..e47cc42e4bde1 100644
--- a/offload/plugins-nextgen/cuda/src/rtl.cpp
+++ b/offload/plugins-nextgen/cuda/src/rtl.cpp
@@ -89,7 +89,7 @@ struct CUDADeviceImageTy : public DeviceImageTy {
     assert(!Module && "Module already loaded");
 
     CUresult Res = cuModuleLoadDataEx(&Module, getStart(), 0, nullptr, nullptr);
-    if (auto Err = Plugin::check(Res, "Error in cuModuleLoadDataEx: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuModuleLoadDataEx: %s"))
       return Err;
 
     return Plugin::success();
@@ -100,7 +100,7 @@ struct CUDADeviceImageTy : public DeviceImageTy {
     assert(Module && "Module not loaded");
 
     CUresult Res = cuModuleUnload(Module);
-    if (auto Err = Plugin::check(Res, "Error in cuModuleUnload: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuModuleUnload: %s"))
       return Err;
 
     Module = nullptr;
@@ -130,19 +130,19 @@ struct CUDAKernelTy : public GenericKernelTy {
 
     // Retrieve the function pointer of the kernel.
     Res = cuModuleGetFunction(&Func, CUDAImage.getModule(), getName());
-    if (auto Err = Plugin::check(Res, "Error in cuModuleGetFunction('%s'): %s",
+    if (auto Err = Plugin::check(Res, "error in cuModuleGetFunction('%s'): %s",
                                  getName()))
       return Err;
 
     // Check that the function pointer is valid.
     if (!Func)
       return Plugin::error(ErrorCode::INVALID_BINARY,
-                           "Invalid function for kernel %s", getName());
+                           "invalid function for kernel %s", getName());
 
     int MaxThreads;
     Res = cuFuncGetAttribute(&MaxThreads,
                              CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, Func);
-    if (auto Err = Plugin::check(Res, "Error in cuFuncGetAttribute: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuFuncGetAttribute: %s"))
       return Err;
 
     // The maximum number of threads cannot exceed the maximum of the kernel.
@@ -179,10 +179,10 @@ struct CUDAStreamRef final : public GenericDeviceResourceRef {
   Error create(GenericDeviceTy &Device) override {
     if (Stream)
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Creating an existing stream");
+                           "creating an existing stream");
 
     CUresult Res = cuStreamCreate(&Stream, CU_STREAM_NON_BLOCKING);
-    if (auto Err = Plugin::check(Res, "Error in cuStreamCreate: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuStreamCreate: %s"))
       return Err;
 
     return Plugin::success();
@@ -193,10 +193,10 @@ struct CUDAStreamRef final : public GenericDeviceResourceRef {
   Error destroy(GenericDeviceTy &Device) override {
     if (!Stream)
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Destroying an invalid stream");
+                           "destroying an invalid stream");
 
     CUresult Res = cuStreamDestroy(Stream);
-    if (auto Err = Plugin::check(Res, "Error in cuStreamDestroy: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuStreamDestroy: %s"))
       return Err;
 
     Stream = nullptr;
@@ -228,10 +228,10 @@ struct CUDAEventRef final : public GenericDeviceResourceRef {
   Error create(GenericDeviceTy &Device) override {
     if (Event)
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Creating an existing event");
+                           "creating an existing event");
 
     CUresult Res = cuEventCreate(&Event, CU_EVENT_DEFAULT);
-    if (auto Err = Plugin::check(Res, "Error in cuEventCreate: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuEventCreate: %s"))
       return Err;
 
     return Plugin::success();
@@ -242,10 +242,10 @@ struct CUDAEventRef final : public GenericDeviceResourceRef {
   Error destroy(GenericDeviceTy &Device) override {
     if (!Event)
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Destroying an invalid event");
+                           "destroying an invalid event");
 
     CUresult Res = cuEventDestroy(Event);
-    if (auto Err = Plugin::check(Res, "Error in cuEventDestroy: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuEventDestroy: %s"))
       return Err;
 
     Event = nullptr;
@@ -273,7 +273,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
   /// Initialize the device, its resources and get its properties.
   Error initImpl(GenericPluginTy &Plugin) override {
     CUresult Res = cuDeviceGet(&Device, DeviceId);
-    if (auto Err = Plugin::check(Res, "Error in cuDeviceGet: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuDeviceGet: %s"))
       return Err;
 
     // Query the current flags of the primary context and set its flags if
@@ -283,7 +283,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
     Res = cuDevicePrimaryCtxGetState(Device, &FormerPrimaryCtxFlags,
                                      &FormerPrimaryCtxIsActive);
     if (auto Err =
-            Plugin::check(Res, "Error in cuDevicePrimaryCtxGetState: %s"))
+            Plugin::check(Res, "error in cuDevicePrimaryCtxGetState: %s"))
       return Err;
 
     if (FormerPrimaryCtxIsActive) {
@@ -299,14 +299,14 @@ struct CUDADeviceTy : public GenericDeviceTy {
            "CU_CTX_SCHED_BLOCKING_SYNC\n");
       Res = cuDevicePrimaryCtxSetFlags(Device, CU_CTX_SCHED_BLOCKING_SYNC);
       if (auto Err =
-              Plugin::check(Res, "Error in cuDevicePrimaryCtxSetFlags: %s"))
+              Plugin::check(Res, "error in cuDevicePrimaryCtxSetFlags: %s"))
         return Err;
     }
 
     // Retain the per device primary context and save it to use whenever this
     // device is selected.
     Res = cuDevicePrimaryCtxRetain(&Context, Device);
-    if (auto Err = Plugin::check(Res, "Error in cuDevicePrimaryCtxRetain: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuDevicePrimaryCtxRetain: %s"))
       return Err;
 
     if (auto Err = setContext())
@@ -389,7 +389,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
     if (Context) {
       CUresult Res = cuDevicePrimaryCtxRelease(Device);
       if (auto Err =
-              Plugin::check(Res, "Error in cuDevicePrimaryCtxRelease: %s"))
+              Plugin::check(Res, "error in cuDevicePrimaryCtxRelease: %s"))
         return Err;
     }
 
@@ -427,7 +427,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
                                                       PTXInputFilePath);
     if (EC)
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to create temporary file for ptxas");
+                           "failed to create temporary file for ptxas");
 
     // Write the file's contents to the output file.
     Expected<std::unique_ptr<FileOutputBuffer>> OutputOrErr =
@@ -444,13 +444,13 @@ struct CUDADeviceTy : public GenericDeviceTy {
                                       PTXOutputFilePath);
     if (EC)
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to create temporary file for ptxas");
+                           "failed to create temporary file for ptxas");
 
     // Try to find `ptxas` in the path to compile the PTX to a binary.
     const auto ErrorOrPath = sys::findProgramByName("ptxas");
     if (!ErrorOrPath)
       return Plugin::error(ErrorCode::HOST_TOOL_NOT_FOUND,
-                           "Failed to find 'ptxas' on the PATH.");
+                           "failed to find 'ptxas' on the PATH.");
 
     std::string Arch = getComputeUnitKind();
     StringRef Args[] = {*ErrorOrPath,
@@ -466,20 +466,20 @@ struct CUDADeviceTy : public GenericDeviceTy {
     if (sys::ExecuteAndWait(*ErrorOrPath, Args, std::nullopt, {}, 0, 0,
                             &ErrMsg))
       return Plugin::error(ErrorCode::ASSEMBLE_FAILURE,
-                           "Running 'ptxas' failed: %s\n", ErrMsg.c_str());
+                           "running 'ptxas' failed: %s\n", ErrMsg.c_str());
 
     auto BufferOrErr = MemoryBuffer::getFileOrSTDIN(PTXOutputFilePath.data());
     if (!BufferOrErr)
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to open temporary file for ptxas");
+                           "failed to open temporary file for ptxas");
 
     // Clean up the temporary files afterwards.
     if (sys::fs::remove(PTXOutputFilePath))
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to remove temporary file for ptxas");
+                           "failed to remove temporary file for ptxas");
     if (sys::fs::remove(PTXInputFilePath))
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to remove temporary file for ptxas");
+                           "failed to remove temporary file for ptxas");
 
     return std::move(*BufferOrErr);
   }
@@ -490,7 +490,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
     CUDAKernelTy *CUDAKernel = Plugin.allocate<CUDAKernelTy>();
     if (!CUDAKernel)
       return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
-                           "Failed to allocate memory for CUDA kernel");
+                           "failed to allocate memory for CUDA kernel");
 
     new (CUDAKernel) CUDAKernelTy(Name);
 
@@ -500,7 +500,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
   /// Set the current context to this device's context.
   Error setContext() override {
     CUresult Res = cuCtxSetCurrent(Context);
-    return Plugin::check(Res, "Error in cuCtxSetCurrent: %s");
+    return Plugin::check(Res, "error in cuCtxSetCurrent: %s");
   }
 
   /// NVIDIA returns the product of the SM count and the number of warps that
@@ -594,7 +594,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
     }
 
     if (auto Err =
-            Plugin::check(Res, "Error in cuMemAlloc[Host|Managed]: %s")) {
+            Plugin::check(Res, "error in cuMemAlloc[Host|Managed]: %s")) {
       REPORT("Failure to alloc memory: %s\n", toString(std::move(Err)).data());
       return nullptr;
     }
@@ -632,7 +632,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
     }
     }
 
-    if (auto Err = Plugin::check(Res, "Error in cuMemFree[Host]: %s")) {
+    if (auto Err = Plugin::check(Res, "error in cuMemFree[Host]: %s")) {
       REPORT("Failure to free memory: %s\n", toString(std::move(Err)).data());
       return OFFLOAD_FAIL;
     }
@@ -652,7 +652,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
     if (auto Err = CUDAStreamManager.returnResource(Stream))
       return Err;
 
-    return Plugin::check(Res, "Error in cuStreamSynchronize: %s");
+    return Plugin::check(Res, "error in cuStreamSynchronize: %s");
   }
 
   /// CUDA support VA management
@@ -674,12 +674,12 @@ struct CUDADeviceTy : public GenericDeviceTy {
 
     if (Size == 0)
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Memory Map Size must be larger than 0");
+                           "memory Map Size must be larger than 0");
 
     // Check if we have already mapped this address
     if (IHandle != DeviceMMaps.end())
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Address already memory mapped");
+                           "address already memory mapped");
 
     CUmemAllocationProp Prop = {};
     size_t Granularity = 0;
@@ -693,7 +693,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
       *Addr = nullptr;
       return Plugin::error(
           ErrorCode::OUT_OF_RESOURCES,
-          "Cannot map memory size larger than the available device memory");
+          "cannot map memory size larger than the available device memory");
     }
 
     // currently NVidia only supports pinned device types
@@ -704,12 +704,12 @@ struct CUDADeviceTy : public GenericDeviceTy {
     cuMemGetAllocationGranularity(&Granularity, &Prop,
                                   CU_MEM_ALLOC_GRANULARITY_MINIMUM);
     if (auto Err =
-            Plugin::check(Res, "Error in cuMemGetAllocationGranularity: %s"))
+            Plugin::check(Res, "error in cuMemGetAllocationGranularity: %s"))
       return Err;
 
     if (Granularity == 0)
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Wrong device Page size");
+                           "wrong device Page size");
 
     // Ceil to page size.
     Size = utils::roundUp(Size, Granularity);
@@ -717,16 +717,16 @@ struct CUDADeviceTy : public GenericDeviceTy {
     // Create a handler of our allocation
     CUmemGenericAllocationHandle AHandle;
     Res = cuMemCreate(&AHandle, Size, &Prop, 0);
-    if (auto Err = Plugin::check(Res, "Error in cuMemCreate: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuMemCreate: %s"))
       return Err;
 
     CUdeviceptr DevPtr = 0;
     Res = cuMemAddressReserve(&DevPtr, Size, 0, DVAddr, 0);
-    if (auto Err = Plugin::check(Res, "Error in cuMemAddressReserve: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuMemAddressReserve: %s"))
       return Err;
 
     Res = cuMemMap(DevPtr, Size, 0, AHandle, 0);
-    if (auto Err = Plugin::check(Res, "Error in cuMemMap: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuMemMap: %s"))
       return Err;
 
     CUmemAccessDesc ADesc = {};
@@ -736,7 +736,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
 
     // Sets address
     Res = cuMemSetAccess(DevPtr, Size, &ADesc, 1);
-    if (auto Err = Plugin::check(Res, "Error in cuMemSetAccess: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuMemSetAccess: %s"))
       return Err;
 
     *Addr = reinterpret_cast<void *>(DevPtr);
@@ -752,25 +752,25 @@ struct CUDADeviceTy : public GenericDeviceTy {
     // Mapping does not exist
     if (IHandle == DeviceMMaps.end()) {
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Addr is not MemoryMapped");
+                           "addr is not MemoryMapped");
     }
 
     if (IHandle == DeviceMMaps.end())
       return Plugin::error(ErrorCode::INVALID_ARGUMENT,
-                           "Addr is not MemoryMapped");
+                           "addr is not MemoryMapped");
 
     CUmemGenericAllocationHandle &AllocHandle = IHandle->second;
 
     CUresult Res = cuMemUnmap(DVAddr, Size);
-    if (auto Err = Plugin::check(Res, "Error in cuMemUnmap: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuMemUnmap: %s"))
       return Err;
 
     Res = cuMemRelease(AllocHandle);
-    if (auto Err = Plugin::check(Res, "Error in cuMemRelease: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuMemRelease: %s"))
       return Err;
 
     Res = cuMemAddressFree(DVAddr, Size);
-    if (auto Err = Plugin::check(Res, "Error in cuMemAddressFree: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuMemAddressFree: %s"))
       return Err;
 
     DeviceMMaps.erase(IHandle);
@@ -793,7 +793,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
     if (auto Err = CUDAStreamManager.returnResource(Stream))
       return Err;
 
-    return Plugin::check(Res, "Error in cuStreamQuery: %s");
+    return Plugin::check(Res, "error in cuStreamQuery: %s");
   }
 
   Expected<void *> dataLockImpl(void *HstPtr, int64_t Size) override {
@@ -821,7 +821,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
       return Err;
 
     CUresult Res = cuMemcpyHtoDAsync((CUdeviceptr)TgtPtr, HstPtr, Size, Stream);
-    return Plugin::check(Res, "Error in cuMemcpyHtoDAsync: %s");
+    return Plugin::check(Res, "error in cuMemcpyHtoDAsync: %s");
   }
 
   /// Retrieve data from the device (device to host transfer).
@@ -835,7 +835,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
       return Err;
 
     CUresult Res = cuMemcpyDtoHAsync(HstPtr, (CUdeviceptr)TgtPtr, Size, Stream);
-    return Plugin::check(Res, "Error in cuMemcpyDtoHAsync: %s");
+    return Plugin::check(Res, "error in cuMemcpyDtoHAsync: %s");
   }
 
   /// Exchange data between two devices directly. We may use peer access if
@@ -895,7 +895,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
       return Err;
 
     CUresult Res = cuEventRecord(Event, Stream);
-    return Plugin::check(Res, "Error in cuEventRecord: %s");
+    return Plugin::check(Res, "error in cuEventRecord: %s");
   }
 
   /// Make the stream wait on the event.
@@ -911,14 +911,14 @@ struct CUDADeviceTy : public GenericDeviceTy {
     // specific CUDA version, and defined as 0x0. In previous version, per CUDA
     // API document, that argument has to be 0x0.
     CUresult Res = cuStreamWaitEvent(Stream, Event, 0);
-    return Plugin::check(Res, "Error in cuStreamWaitEvent: %s");
+    return Plugin::check(Res, "error in cuStreamWaitEvent: %s");
   }
 
   /// Synchronize the current thread with the event.
   Error syncEventImpl(void *EventPtr) override {
     CUevent Event = reinterpret_cast<CUevent>(EventPtr);
     CUresult Res = cuEventSynchronize(Event);
-    return Plugin::check(Res, "Error in cuEventSynchronize: %s");
+    return Plugin::check(Res, "error in cuEventSynchronize: %s");
   }
 
   /// Print information about the device.
@@ -1110,17 +1110,17 @@ struct CUDADeviceTy : public GenericDeviceTy {
   }
   Error getDeviceMemorySize(uint64_t &Value) override {
     CUresult Res = cuDeviceTotalMem(&Value, Device);
-    return Plugin::check(Res, "Error in getDeviceMemorySize %s");
+    return Plugin::check(Res, "error in getDeviceMemorySize %s");
   }
 
   /// CUDA-specific functions for getting and setting context limits.
   Error setCtxLimit(CUlimit Kind, uint64_t Value) {
     CUresult Res = cuCtxSetLimit(Kind, Value);
-    return Plugin::check(Res, "Error in cuCtxSetLimit: %s");
+    return Plugin::check(Res, "error in cuCtxSetLimit: %s");
   }
   Error getCtxLimit(CUlimit Kind, uint64_t &Value) {
     CUresult Res = cuCtxGetLimit(&Value, Kind);
-    return Plugin::check(Res, "Error in cuCtxGetLimit: %s");
+    return Plugin::check(Res, "error in cuCtxGetLimit: %s");
   }
 
   /// CUDA-specific function to get device attributes.
@@ -1128,7 +1128,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
     // TODO: Warn if the new value is larger than the old.
     CUresult Res =
         cuDeviceGetAttribute((int *)&Value, (CUdevice_attribute)Kind, Device);
-    return Plugin::check(Res, "Error in cuDeviceGetAttribute: %s");
+    return Plugin::check(Res, "error in cuDeviceGetAttribute: %s");
   }
 
   CUresult getDeviceAttrRaw(uint32_t Kind, int &Value) {
@@ -1178,7 +1178,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
       uint16_t Priority;
       if (NameOrErr->rsplit('_').second.getAsInteger(10, Priority))
         return Plugin::error(ErrorCode::INVALID_BINARY,
-                             "Invalid priority for constructor or destructor");
+                             "invalid priority for constructor or destructor");
 
       Funcs.emplace_back(*NameOrErr, Priority);
     }
@@ -1192,7 +1192,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
         allocate(Funcs.size() * sizeof(void *), nullptr, TARGET_ALLOC_DEVICE);
     if (!Buffer)
       return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
-                           "Failed to allocate memory for global buffer");
+                           "failed to allocate memory for global buffer");
 
     auto *GlobalPtrStart = reinterpret_cast<uintptr_t *>(Buffer);
     auto *GlobalPtrStop = reinterpret_cast<uintptr_t *>(Buffer) + Funcs.size();
@@ -1241,7 +1241,7 @@ struct CUDADeviceTy : public GenericDeviceTy {
 
     if (free(Buffer, TARGET_ALLOC_DEVICE) != OFFLOAD_SUCCESS)
       return Plugin::error(ErrorCode::UNKNOWN,
-                           "Failed to free memory for global buffer");
+                           "failed to free memory for global buffer");
 
     return Err;
   }
@@ -1314,7 +1314,7 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
         },
         &GenericDevice.Plugin);
 
-  return Plugin::check(Res, "Error in cuLaunchKernel for '%s': %s", getName());
+  return Plugin::check(Res, "error in cuLaunchKernel for '%s': %s", getName());
 }
 
 /// Class implementing the CUDA-specific functionalities of the global handler.
@@ -1334,14 +1334,14 @@ class CUDAGlobalHandlerTy final : public GenericGlobalHandlerTy {
     CUdeviceptr CUPtr;
     CUresult Res =
         cuModuleGetGlobal(&CUPtr, &CUSize, CUDAImage.getModule(), GlobalName);
-    if (auto Err = Plugin::check(Res, "Error in cuModuleGetGlobal for '%s': %s",
+    if (auto Err = Plugin::check(Res, "error in cuModuleGetGlobal for '%s': %s",
                                  GlobalName))
       return Err;
 
     if (CUSize != DeviceGlobal.getSize())
       return Plugin::error(
           ErrorCode::INVALID_BINARY,
-          "Failed to load global '%s' due to size mismatch (%zu != %zu)",
+          "failed to load global '%s' due to size mismatch (%zu != %zu)",
           GlobalName, CUSize, (size_t)DeviceGlobal.getSize());
 
     DeviceGlobal.setPtr(reinterpret_cast<void *>(CUPtr));
@@ -1373,13 +1373,13 @@ struct CUDAPluginTy final : public GenericPluginTy {
       return 0;
     }
 
-    if (auto Err = Plugin::check(Res, "Error in cuInit: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuInit: %s"))
       return std::move(Err);
 
     // Get the number of devices.
     int NumDevices;
     Res = cuDeviceGetCount(&NumDevices);
-    if (auto Err = Plugin::check(Res, "Error in cuDeviceGetCount: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuDeviceGetCount: %s"))
       return std::move(Err);
 
     // Do not initialize if there are no devices.
@@ -1427,18 +1427,18 @@ struct CUDAPluginTy final : public GenericPluginTy {
 
     CUdevice Device;
     CUresult Res = cuDeviceGet(&Device, DeviceId);
-    if (auto Err = Plugin::check(Res, "Error in cuDeviceGet: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuDeviceGet: %s"))
       return std::move(Err);
 
     int32_t Major, Minor;
     Res = cuDeviceGetAttribute(
         &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, Device);
-    if (auto Err = Plugin::check(Res, "Error in cuDeviceGetAttribute: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuDeviceGetAttribute: %s"))
       return std::move(Err);
 
     Res = cuDeviceGetAttribute(
         &Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, Device);
-    if (auto Err = Plugin::check(Res, "Error in cuDeviceGetAttribute: %s"))
+    if (auto Err = Plugin::check(Res, "error in cuDeviceGetAttribute: %s"))
       return std::move(Err);
 
     int32_t ImageMajor = SM / 10;
@@ -1490,7 +1490,7 @@ Error CUDADeviceTy::dataExchangeImpl(const void *SrcPtr,
           CanAccessPeer = 0;
           DP("Too many P2P so fall back to D2D memcpy");
         } else if (auto Err =
-                       Plugin::check(Res, "Error in cuCtxEnablePeerAccess: %s"))
+                       Plugin::check(Res, "error in cuCtxEnablePeerAccess: %s"))
           return Err;
       }
       PeerAccesses[DstDeviceId] = (CanAccessPeer)
@@ -1507,19 +1507,19 @@ Error CUDADeviceTy::dataExchangeImpl(const void *SrcPtr,
     // TODO: Should we fallback to D2D if peer access fails?
     Res = cuMemcpyPeerAsync(CUDstPtr, Context, CUSrcPtr, DstDevice.Context,
                             Size, Stream);
-    return Plugin::check(Res, "Error in cuMemcpyPeerAsync: %s");
+    return Plugin::check(Res, "error in cuMemcpyPeerAsync: %s");
   }
 
   // Fallback to D2D copy.
   Res = cuMemcpyDtoDAsync(CUDstPtr, CUSrcPtr, Size, Stream);
-  return Plugin::check(Res, "Error in cuMemcpyDtoDAsync: %s");
+  return Plugin::check(Res, "error in cuMemcpyDtoDAsync: %s");
 }
 
 template <typename... ArgsTy>
 static Error Plugin::check(int32_t Code, const char *ErrFmt, ArgsTy... Args) {
   CUresult ResultCode = static_cast<CUresult>(Code);
   if (ResultCode == CUDA_SUCCESS)
-    return Error::success();
+    return Plugin::success();
 
   const char *Desc = "Unknown error";
   CUresult Ret = cuGetErrorString(ResultCode, &Desc);
@@ -1527,8 +1527,7 @@ static Error Plugin::check(int32_t Code, const char *ErrFmt, ArgsTy... Args) {
     REPORT("Unrecognized " GETNAME(TARGET_NAME) " error code %d\n", Code);
 
   // TODO: Create a map for CUDA error codes to Offload error codes
-  return createStringError<ArgsTy..., const char *>(ErrorCode::UNKNOWN, ErrFmt,
-                                                    Args..., Desc);
+  return Plugin::error(ErrorCode::UNKNOWN, ErrFmt, Args..., Desc);
 }
 
 } // namespace plugin
diff --git a/offload/plugins-nextgen/host/src/rtl.cpp b/offload/plugins-nextgen/host/src/rtl.cpp
index 8dfb06fec976e..9916f4d0ab250 100644
--- a/offload/plugins-nextgen/host/src/rtl.cpp
+++ b/offload/plugins-nextgen/host/src/rtl.cpp
@@ -76,7 +76,7 @@ struct GenELF64KernelTy : public GenericKernelTy {
     // Check that the function pointer is valid.
     if (!Global.getPtr())
       return Plugin::error(ErrorCode::INVALID_BINARY,
-                           "Invalid function for kernel %s", getName());
+                           "invalid function for kernel %s", getName());
 
     // Save the function pointer.
     Func = (void (*)())Global.getPtr();
@@ -104,7 +104,7 @@ struct GenELF64KernelTy : public GenericKernelTy {
     ffi_status Status = ffi_prep_cif(&Cif, FFI_DEFAULT_ABI, KernelArgs.NumArgs,
                                      &ffi_type_void, ArgTypesPtr);
     if (Status != FFI_OK)
-      return Plugin::error(ErrorCode::UNKNOWN, "Error in ffi_prep_cif: %d",
+      return Plugin::error(ErrorCode::UNKNOWN, "error in ffi_prep_cif: %d",
                            Status);
 
     // Call the kernel function through libffi.
@@ -159,7 +159,7 @@ struct GenELF64DeviceTy : public GenericDeviceTy {
     GenELF64KernelTy *GenELF64Kernel = Plugin.allocate<GenELF64KernelTy>();
     if (!GenELF64Kernel)
       return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
-                           "Failed to allocate memory for GenELF64 kernel");
+                           "failed to allocate memory for GenELF64 kernel");
 
     new (GenELF64Kernel) GenELF64KernelTy(Name);
 
@@ -181,27 +181,27 @@ struct GenELF64DeviceTy : public GenericDeviceTy {
     int TmpFileFd = mkstemp(TmpFileName);
     if (TmpFileFd == -1)
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to create tmpfile for loading target image");
+                           "failed to create tmpfile for loading target image");
 
     // Open the temporary file.
     FILE *TmpFile = fdopen(TmpFileFd, "wb");
     if (!TmpFile)
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to open tmpfile %s for loading target image",
+                           "failed to open tmpfile %s for loading target image",
                            TmpFileName);
 
     // Write the image into the temporary file.
     size_t Written = fwrite(Image->getStart(), Image->getSize(), 1, TmpFile);
     if (Written != 1)
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to write target image to tmpfile %s",
+                           "failed to write target image to tmpfile %s",
                            TmpFileName);
 
     // Close the temporary file.
     int Ret = fclose(TmpFile);
     if (Ret)
       return Plugin::error(ErrorCode::HOST_IO,
-                           "Failed to close tmpfile %s with the target image",
+                           "failed to close tmpfile %s with the target image",
                            TmpFileName);
 
     // Load the temporary file as a dynamic library.
@@ -212,7 +212,7 @@ struct GenELF64DeviceTy : public GenericDeviceTy {
     // Check if the loaded library is valid.
     if (!DynLib.isValid())
       return Plugin::error(ErrorCode::INVALID_BINARY,
-                           "Failed to load target image: %s", ErrMsg.c_str());
+                           "failed to load target image: %s", ErrMsg.c_str());
 
     // Save a reference of the image's dynamic library.
     Image->setDynamicLibrary(DynLib);
@@ -377,7 +377,7 @@ class GenELF64GlobalHandlerTy final : public GenericGlobalHandlerTy {
     // Get the address of the symbol.
     void *Addr = DynLib.getAddressOfSymbol(GlobalName);
     if (Addr == nullptr) {
-      return Plugin::error(ErrorCode::NOT_FOUND, "Failed to load global '%s'",
+      return Plugin::error(ErrorCode::NOT_FOUND, "failed to load global '%s'",
                            GlobalName);
     }
 
@@ -400,7 +400,7 @@ struct GenELF64PluginTy final : public GenericPluginTy {
   /// Initialize the plugin and return the number of devices.
   Expected<int32_t> initImpl() override {
 #ifdef USES_DYNAMIC_FFI
-    if (auto Err = Plugin::check(ffi_init(), "Failed to initialize libffi"))
+    if (auto Err = Plugin::check(ffi_init(), "failed to initialize libffi"))
       return std::move(Err);
 #endif
 
@@ -468,10 +468,10 @@ struct GenELF64PluginTy final : public GenericPluginTy {
 template <typename... ArgsTy>
 static Error Plugin::check(int32_t Code, const char *ErrMsg, ArgsTy... Args) {
   if (Code == 0)
-    return Error::success();
+    return Plugin::success();
 
-  return createStringError<ArgsTy..., const char *>(
-      inconvertibleErrorCode(), ErrMsg, Args..., std::to_string(Code).data());
+  return Plugin::error(ErrorCode::UNKNOWN, ErrMsg, Args...,
+                       std::to_string(Code).data());
 }
 
 } // namespace plugin

>From 5efadf169beef6c6050a5e9c6736c70158a93d84 Mon Sep 17 00:00:00 2001
From: Ross Brunton <ross at codeplay.com>
Date: Tue, 20 May 2025 10:49:05 +0100
Subject: [PATCH 3/3] Formatting

---
 offload/include/Shared/OffloadError.h | 8 +++++---
 1 file changed, 5 insertions(+), 3 deletions(-)

diff --git a/offload/include/Shared/OffloadError.h b/offload/include/Shared/OffloadError.h
index ab9bf136b444c..01b2c59d7f6d6 100644
--- a/offload/include/Shared/OffloadError.h
+++ b/offload/include/Shared/OffloadError.h
@@ -49,7 +49,8 @@ class OffloadError : public llvm::ErrorInfo<OffloadError, llvm::StringError> {
 
 /// Create an Offload error.
 template <typename... ArgsTy>
-static llvm::Error createOffloadError(error::ErrorCode Code, const char *ErrFmt, ArgsTy... Args) {
+static llvm::Error createOffloadError(error::ErrorCode Code, const char *ErrFmt,
+                                      ArgsTy... Args) {
   std::string Buffer;
   llvm::raw_string_ostream(Buffer) << llvm::format(ErrFmt, Args...);
   return llvm::make_error<error::OffloadError>(Code, Buffer);
@@ -62,8 +63,9 @@ inline llvm::Error createOffloadError(error::ErrorCode Code, const char *S) {
 // The OffloadError will have a message of either:
 // * "{Context}: {Message}" if the other error is a StringError
 // * "{Context}" otherwise
-inline llvm::Error createOffloadError(error::ErrorCode Code, llvm::Error &&OtherError,
-                   const char *Context) {
+inline llvm::Error createOffloadError(error::ErrorCode Code,
+                                      llvm::Error &&OtherError,
+                                      const char *Context) {
   std::string Buffer{Context};
   llvm::raw_string_ostream buffer(Buffer);
 



More information about the llvm-commits mailing list