[clang] [llvm] [LLVMOffload] Get LastError for bad Kernels (PR #213390)
Sophia Herrmann via llvm-commits
llvm-commits at lists.llvm.org
Thu Aug 13 10:57:38 PDT 2026
https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/213390
>From da9c8f52afe3fda79361381a16c3b6ac10700c8d Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Mon, 27 Jul 2026 15:48:11 -0700
Subject: [PATCH 1/4] decouple front/back end
Co-authored-by: Johannes Doerfert <jdoerfert.llvm at gmail.com>
Co-authored-by: Jonas Greifenhain <cadivus at daverkomp.de>
remove irrelevant artifacts
add unittests
---
clang/include/clang/Driver/CommonArgs.h | 6 +
clang/lib/CodeGen/CGCUDANV.cpp | 59 +++++----
clang/lib/Driver/Driver.cpp | 73 +++++++-----
clang/lib/Driver/ToolChains/AMDGPU.cpp | 13 +-
clang/lib/Driver/ToolChains/Clang.cpp | 112 +++++++++++-------
clang/lib/Driver/ToolChains/CommonArgs.cpp | 19 ++-
clang/lib/Driver/ToolChains/Cuda.cpp | 45 +++++--
clang/lib/Driver/ToolChains/Gnu.cpp | 1 +
clang/lib/Driver/ToolChains/Linux.cpp | 4 +-
clang/test/CodeGenCUDA/Inputs/cuda.h | 2 +-
clang/test/CodeGenCUDA/offload_via_llvm.cu | 60 +++++-----
clang/test/Driver/cuda-via-liboffload.cu | 15 ++-
.../ClangLinkerWrapper.cpp | 13 +-
.../include/kernel/LanguageRuntime.h | 2 +
offload/languages/kernel/CMakeLists.txt | 2 +
offload/languages/kernel/exports | 10 +-
.../kernel/include/LanguageAliases.inc | 6 -
.../languages/kernel/include/LanguageLaunch.h | 15 +--
offload/languages/kernel/include/Types.h | 13 +-
.../languages/kernel/src/LanguageLaunch.cpp | 56 ++++-----
offload/test/lit.cfg | 12 +-
offload/test/offloading/CUDA/basic_launch.cu | 25 ++--
.../CUDA/basic_launch_blocks_and_threads.cu | 22 ++--
.../offloading/CUDA/basic_launch_multi_arg.cu | 36 +++---
offload/test/offloading/CUDA/device_api.cu | 45 +++++++
.../test/offloading/CUDA/device_properties.cu | 40 +++++++
offload/test/offloading/CUDA/host_alloc.cu | 48 ++++++++
offload/test/offloading/CUDA/launch_tu.cu | 25 ++--
offload/test/offloading/CUDA/memcpy_kinds.cu | 51 ++++++++
offload/test/offloading/CUDA/stream_api.cu | 46 +++++++
offload/test/offloading/CUDA/syncthreads.cu | 40 +++++++
.../offloading/CUDA/thread_and_block_id.cu | 44 +++++++
offload/test/offloading/HIP/basic_launch.hip | 30 +++++
.../HIP/basic_launch_blocks_and_threads.hip | 31 +++++
.../offloading/HIP/basic_launch_multi_arg.hip | 39 ++++++
offload/test/offloading/HIP/device_api.hip | 45 +++++++
.../test/offloading/HIP/device_properties.hip | 40 +++++++
offload/test/offloading/HIP/host_alloc.hip | 48 ++++++++
offload/test/offloading/HIP/kernel_tu.hip.inc | 1 +
offload/test/offloading/HIP/launch_tu.hip | 30 +++++
offload/test/offloading/HIP/memcpy_kinds.hip | 51 ++++++++
offload/test/offloading/HIP/stream_api.hip | 46 +++++++
offload/test/offloading/HIP/syncthreads.hip | 40 +++++++
.../offloading/HIP/thread_and_block_id.hip | 44 +++++++
44 files changed, 1138 insertions(+), 267 deletions(-)
create mode 100644 offload/test/offloading/CUDA/device_api.cu
create mode 100644 offload/test/offloading/CUDA/device_properties.cu
create mode 100644 offload/test/offloading/CUDA/host_alloc.cu
create mode 100644 offload/test/offloading/CUDA/memcpy_kinds.cu
create mode 100644 offload/test/offloading/CUDA/stream_api.cu
create mode 100644 offload/test/offloading/CUDA/syncthreads.cu
create mode 100644 offload/test/offloading/CUDA/thread_and_block_id.cu
create mode 100644 offload/test/offloading/HIP/basic_launch.hip
create mode 100644 offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
create mode 100644 offload/test/offloading/HIP/basic_launch_multi_arg.hip
create mode 100644 offload/test/offloading/HIP/device_api.hip
create mode 100644 offload/test/offloading/HIP/device_properties.hip
create mode 100644 offload/test/offloading/HIP/host_alloc.hip
create mode 100644 offload/test/offloading/HIP/kernel_tu.hip.inc
create mode 100644 offload/test/offloading/HIP/launch_tu.hip
create mode 100644 offload/test/offloading/HIP/memcpy_kinds.hip
create mode 100644 offload/test/offloading/HIP/stream_api.hip
create mode 100644 offload/test/offloading/HIP/syncthreads.hip
create mode 100644 offload/test/offloading/HIP/thread_and_block_id.hip
diff --git a/clang/include/clang/Driver/CommonArgs.h b/clang/include/clang/Driver/CommonArgs.h
index 01e358a1d0717..3b53df94fc79f 100644
--- a/clang/include/clang/Driver/CommonArgs.h
+++ b/clang/include/clang/Driver/CommonArgs.h
@@ -148,6 +148,12 @@ void addArchSpecificRPath(const ToolChain &TC, const llvm::opt::ArgList &Args,
void addOpenMPRuntimeLibraryPath(const ToolChain &TC,
const llvm::opt::ArgList &Args,
llvm::opt::ArgStringList &CmdArgs);
+
+bool addLLVMOffloadingRuntime(const Compilation &C,
+ llvm::opt::ArgStringList &CmdArgs,
+ const ToolChain &TC,
+ const llvm::opt::ArgList &Args);
+
/// Returns true, if an OpenMP runtime has been added.
bool addOpenMPRuntime(const Compilation &C, llvm::opt::ArgStringList &CmdArgs,
const ToolChain &TC, const llvm::opt::ArgList &Args,
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 1e688d29d15a5..88a8a9df73581 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -345,41 +345,52 @@ void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
emitDeviceStubBodyLegacy(CGF, Args);
}
-/// Build the input as a sized array of pointers so that it can be launched by
-/// the offloading runtime.
+/// CUDA passes the arguments with a level of indirection. For example, a
+/// (void*, short, void*) is passed as {void **, short *, void **} to the launch
+/// function. For the LLVM/Offload launch we include the number of arguments and
+/// their size. Thus, we pass {{void **, short*, void **}, 3, {sizeof(void*),
+/// sizeof(short), sizeof(void*)}}.
Address CGNVCUDARuntime::prepareKernelArgsLLVMOffload(CodeGenFunction &CGF,
FunctionArgList &Args) {
- SmallVector<llvm::Type *> ArgTypes, KernelLaunchParamsTypes;
- for (auto &Arg : Args)
- ArgTypes.push_back(CGF.ConvertTypeForMem(Arg->getType()));
- llvm::StructType *KernelArgsTy = llvm::StructType::create(ArgTypes);
- llvm::Type *KernelArgsPtrsTy = llvm::ArrayType::get(PtrTy, Args.size());
-
- auto *Int32Ty = CGF.Builder.getInt32Ty();
- KernelLaunchParamsTypes.push_back(Int32Ty);
+ SmallVector<llvm::Type *> KernelLaunchParamsTypes;
+
+ auto *Int64Ty = CGF.Builder.getInt64Ty();
+ KernelLaunchParamsTypes.push_back(PtrTy);
+ KernelLaunchParamsTypes.push_back(Int64Ty);
KernelLaunchParamsTypes.push_back(PtrTy);
llvm::StructType *KernelLaunchParamsTy =
llvm::StructType::create(KernelLaunchParamsTypes);
- Address KernelArgs = CGF.CreateTempAllocaWithoutCast(
- KernelArgsTy, CharUnits::fromQuantity(16), "kernel_args");
- Address KernelArgsPtrs = CGF.CreateTempAllocaWithoutCast(
- KernelArgsPtrsTy, CharUnits::fromQuantity(16), "kernel_args_ptrs");
Address KernelLaunchParams = CGF.CreateTempAllocaWithoutCast(
KernelLaunchParamsTy, CharUnits::fromQuantity(16),
"kernel_launch_params");
+ Address KernelArgs = CGF.CreateTempAlloca(
+ PtrTy, LangAS::Default, CharUnits::fromQuantity(16), "kernel_args",
+ llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
+ Address KernelArgSizes = CGF.CreateTempAlloca(
+ SizeTy, LangAS::Default, CharUnits::fromQuantity(16), "kernel_arg_sizes",
+ llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
- CGF.Builder.CreateStore(llvm::ConstantInt::get(Int32Ty, Args.size()),
+ CGF.Builder.CreateStore(KernelArgs.emitRawPointer(CGF),
CGF.Builder.CreateStructGEP(KernelLaunchParams, 0));
- CGF.Builder.CreateStore(KernelArgsPtrs.emitRawPointer(CGF),
+ CGF.Builder.CreateStore(llvm::ConstantInt::get(Int64Ty, Args.size()),
CGF.Builder.CreateStructGEP(KernelLaunchParams, 1));
+ CGF.Builder.CreateStore(KernelArgSizes.emitRawPointer(CGF),
+ CGF.Builder.CreateStructGEP(KernelLaunchParams, 2));
for (unsigned i = 0; i < Args.size(); ++i) {
- auto *ArgVal = CGF.Builder.CreateLoad(CGF.GetAddrOfLocalVar(Args[i]));
- Address ArgAddr = CGF.Builder.CreateStructGEP(KernelArgs, i);
- CGF.Builder.CreateStore(ArgVal, ArgAddr);
- CGF.Builder.CreateStore(ArgAddr.emitRawPointer(CGF),
- CGF.Builder.CreateConstArrayGEP(KernelArgsPtrs, i));
+ llvm::Value *VarPtr = CGF.GetAddrOfLocalVar(Args[i]).emitRawPointer(CGF);
+ llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, PtrTy);
+ CGF.Builder.CreateDefaultAlignedStore(
+ VoidVarPtr, CGF.Builder.CreateConstGEP1_32(
+ PtrTy, KernelArgs.emitRawPointer(CGF), i));
+
+ auto ArgSize = CGM.getDataLayout().getTypeAllocSize(
+ CGM.getTypes().ConvertType(Args[i]->getType()));
+ CGF.Builder.CreateDefaultAlignedStore(
+ llvm::ConstantInt::get(SizeTy, ArgSize),
+ CGF.Builder.CreateConstGEP1_32(SizeTy,
+ KernelArgSizes.emitRawPointer(CGF), i));
}
return KernelLaunchParams;
@@ -408,8 +419,9 @@ Address CGNVCUDARuntime::prepareKernelArgs(CodeGenFunction &CGF,
// array and kernels are launched using cudaLaunchKernel().
void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
FunctionArgList &Args) {
+ bool UsesLLVMOffloading = CGF.getLangOpts().OffloadViaLLVM;
// Build the shadow stack entry at the very start of the function.
- Address KernelArgs = CGF.getLangOpts().OffloadViaLLVM
+ Address KernelArgs = UsesLLVMOffloading
? prepareKernelArgsLLVMOffload(CGF, Args)
: prepareKernelArgs(CGF, Args);
@@ -1282,9 +1294,6 @@ void CGNVCUDARuntime::createOffloadingEntries() {
llvm::object::OffloadKind Kind = CGM.getLangOpts().HIP
? llvm::object::OffloadKind::OFK_HIP
: llvm::object::OffloadKind::OFK_Cuda;
- // For now, just spoof this as OpenMP because that's the runtime it uses.
- if (CGM.getLangOpts().OffloadViaLLVM)
- Kind = llvm::object::OffloadKind::OFK_OpenMP;
llvm::Module &M = CGM.getModule();
for (KernelInfo &I : EmittedKernels)
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index 39691538c355a..c2b670d8edd11 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -928,10 +928,15 @@ getSystemOffloadArchs(Compilation &C, Action::OffloadKind Kind) {
if (llvm::ErrorOr<std::string> Executable =
llvm::sys::findProgramByName(Program, {C.getDriver().Dir})) {
llvm::SmallVector<StringRef> Args{*Executable};
- if (Kind == Action::OFK_HIP)
- Args.push_back("--only=amdgpu");
- else if (Kind == Action::OFK_Cuda)
- Args.push_back("--only=nvptx");
+ bool UsesLLVMOffloading =
+ C.getArgs().hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false);
+ if (!UsesLLVMOffloading) {
+ if (Kind == Action::OFK_HIP)
+ Args.push_back("--only=amdgpu");
+ else if (Kind == Action::OFK_Cuda)
+ Args.push_back("--only=nvptx");
+ }
auto StdoutOrErr = C.getDriver().executeProgram(Args);
if (!StdoutOrErr) {
@@ -988,15 +993,20 @@ static TripleSet inferOffloadToolchains(Compilation &C,
ID = StringToOffloadArch(
getProcessorFromTargetID(llvm::Triple("amdgcn-amd-amdhsa"), Arch));
- if (Kind == Action::OFK_HIP && !ID.isAMDGPU() && !ID.isSPIRV()) {
- C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
- << "HIP" << Arch;
- return {};
- }
- if (Kind == Action::OFK_Cuda && !ID.isNVPTX()) {
- C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
- << "CUDA" << Arch;
- return {};
+ bool UsesLLVMOffloading =
+ C.getArgs().hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false);
+ if (!UsesLLVMOffloading) {
+ if (Kind == Action::OFK_HIP && !ID.isAMDGPU() && !ID.isSPIRV()) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "HIP" << Arch;
+ return {};
+ }
+ if (Kind == Action::OFK_Cuda && !ID.isNVPTX()) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "CUDA" << Arch;
+ return {};
+ }
}
if (Kind == Action::OFK_OpenMP && (ID.isUnknown() || ID.isUnused())) {
C.getDriver().Diag(clang::diag::err_drv_failed_to_deduce_target_from_arch)
@@ -1011,6 +1021,8 @@ static TripleSet inferOffloadToolchains(Compilation &C,
llvm::Triple Triple =
OffloadArchToTriple(C.getDefaultToolChain().getTriple(), ID);
+ if (UsesLLVMOffloading)
+ Triple.setEnvironment(llvm::Triple::LLVM);
// Make a new argument that dispatches this argument to the appropriate
// toolchain. This is required when we infer it and create potentially
@@ -1054,26 +1066,20 @@ static TripleSet inferOffloadToolchains(Compilation &C,
void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
InputList &Inputs) {
- bool UseLLVMOffload = C.getInputArgs().hasArg(
- options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
bool IsCuda =
- llvm::any_of(Inputs,
- [](std::pair<types::ID, const llvm::opt::Arg *> &I) {
- return types::isCuda(I.first);
- }) &&
- !UseLLVMOffload;
+ llvm::any_of(Inputs, [](std::pair<types::ID, const llvm::opt::Arg *> &I) {
+ return types::isCuda(I.first);
+ });
bool IsHIP =
(llvm::any_of(Inputs,
[](std::pair<types::ID, const llvm::opt::Arg *> &I) {
return types::isHIP(I.first);
}) ||
C.getInputArgs().hasArg(options::OPT_hip_link) ||
- C.getInputArgs().hasArg(options::OPT_hipstdpar)) &&
- !UseLLVMOffload;
+ C.getInputArgs().hasArg(options::OPT_hipstdpar));
bool IsSYCL = C.getInputArgs().hasFlag(options::OPT_fsycl,
options::OPT_fno_sycl, false);
bool IsOpenMPOffloading =
- UseLLVMOffload ||
(C.getInputArgs().hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ,
options::OPT_fno_openmp, false) &&
(C.getInputArgs().hasArg(options::OPT_offload_targets_EQ) ||
@@ -1165,7 +1171,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
C.getDefaultToolChain().getTriple());
// Emit a warning if the detected CUDA version is too new.
- if (Kind == Action::OFK_Cuda) {
+ if (Kind == Action::OFK_Cuda && Target.getOS() == llvm::Triple::CUDA) {
auto &CudaInstallation =
static_cast<const toolchains::CudaToolChain &>(TC).CudaInstallation;
if (CudaInstallation.isValid())
@@ -5090,6 +5096,9 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
getFinalPhase(Args) == phases::Preprocess))
return HostAction;
+ bool UsesLLVMOffloading = Args.hasArg(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+
ActionList OffloadActions;
OffloadAction::DeviceDependences DDeps;
@@ -5205,9 +5214,12 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
OffloadAction::DeviceDependences DDep;
DDep.add(*A, *TCAndArch->first, TCAndArch->second, Kind);
- // Compiling CUDA in non-RDC mode uses the PTX output if available.
+ // The legacy CUDA fatbinary path can include PTX alongside the cubin.
+ // The LLVM offload wrapper path feeds these images through a device
+ // linker first, and clang-nvlink-wrapper does not accept PTX as input.
for (Action *Input : A->getInputs())
- if (Kind == Action::OFK_Cuda && A->getType() == types::TY_Object &&
+ if (!UsesLLVMOffloading && Kind == Action::OFK_Cuda &&
+ A->getType() == types::TY_Object &&
!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
false))
DDep.add(*Input, *TCAndArch->first, TCAndArch->second, Kind);
@@ -5235,7 +5247,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
return HostAction;
OffloadAction::DeviceDependences DDep;
- if (C.isOffloadingHostKind(Action::OFK_Cuda) &&
+ if (!UsesLLVMOffloading && C.isOffloadingHostKind(Action::OFK_Cuda) &&
(!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false) ||
Args.hasArg(options::OPT_cuda_emit_nvcc_abi))) {
// If we are not in RDC-mode or are targeting the NVCC ABI we just emit the
@@ -5244,7 +5256,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
C.MakeAction<LinkJobAction>(OffloadActions, types::TY_CUDA_FATBIN);
DDep.add(*FatbinAction, *C.getSingleOffloadToolChain<Action::OFK_Cuda>(),
/*BA=*/{}, Action::OFK_Cuda);
- } else if (HIPNoRDC && offloadDeviceOnly()) {
+ } else if (!UsesLLVMOffloading && HIPNoRDC && offloadDeviceOnly()) {
// If we are in device-only non-RDC-mode we just emit the final HIP
// fatbinary for each translation unit, linking each input individually.
Action *FatbinAction =
@@ -5252,7 +5264,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
DDep.add(*FatbinAction,
*C.getOffloadToolChains<Action::OFK_HIP>().first->second,
/*BA=*/{}, Action::OFK_HIP);
- } else if (HIPNoRDC) {
+ } else if (!UsesLLVMOffloading && HIPNoRDC) {
// Host + device assembly: defer to clang-offload-bundler (see
// BuildActions).
if (HIPAsmBundleDeviceOut &&
@@ -7110,7 +7122,8 @@ const ToolChain &Driver::getOffloadToolChain(
// For AMDHSA offloading (HIP, OpenMP), use the unified AMDGPUToolChain
// This handles both amdgpu-amd-amdhsa and spirv64-amd-amdhsa
// FIXME: This should not key off language or OS.
- if (Kind == Action::OFK_HIP || Kind == Action::OFK_OpenMP)
+ if (Kind == Action::OFK_HIP || Kind == Action::OFK_OpenMP ||
+ Kind == Action::OFK_Cuda)
TC = std::make_unique<toolchains::AMDGPUToolChain>(*this, Target, Args,
HostTC.get(), Kind);
break;
diff --git a/clang/lib/Driver/ToolChains/AMDGPU.cpp b/clang/lib/Driver/ToolChains/AMDGPU.cpp
index 7a7d5be48443e..82f9158a69f04 100644
--- a/clang/lib/Driver/ToolChains/AMDGPU.cpp
+++ b/clang/lib/Driver/ToolChains/AMDGPU.cpp
@@ -533,6 +533,10 @@ void RocmInstallationDetector::AddHIPIncludeArgs(const ArgList &DriverArgs,
!DriverArgs.hasArg(options::OPT_nohipwrapperinc);
bool HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
+ return;
+
if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
// HIP header includes standard library wrapper headers under clang
// cuda_wrappers directory. Since these wrapper headers include_next
@@ -733,8 +737,10 @@ AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
// each tool invocation.
checkAMDGPUCodeObjectVersion(D, Args);
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
if (Triple.getOS() == llvm::Triple::AMDHSA &&
- Triple.getEnvironment() != llvm::Triple::LLVM)
+ Triple.getEnvironment() != llvm::Triple::LLVM && !UsesLLVMOffloading)
RocmInstallation->detectDeviceLibrary();
if (HostTC)
@@ -913,7 +919,10 @@ bool AMDGPUToolChain::isWave64(const llvm::opt::ArgList &DriverArgs,
void AMDGPUToolChain::addClangTargetOptions(
const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
BoundArch BA, Action::OffloadKind DeviceOffloadingKind) const {
- if (DeviceOffloadingKind == Action::OFK_HIP) {
+ bool UsesLLVMOffloading = DriverArgs.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ if (DeviceOffloadingKind == Action::OFK_HIP ||
+ (DeviceOffloadingKind == Action::OFK_Cuda && UsesLLVMOffloading)) {
CC1Args.append({"-fcuda-is-device", "-fno-threadsafe-statics"});
if (!DriverArgs.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 79127fa5e2db1..19d37b2ed7017 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -915,48 +915,71 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
Args.AddLastArg(CmdArgs, options::OPT_MP);
Args.AddLastArg(CmdArgs, options::OPT_MV);
- // Add offload include arguments specific for CUDA/HIP/SYCL. This must happen
- // before we -I or -include anything else, because we must pick up the
- // CUDA/HIP/SYCL headers from the particular CUDA/ROCm/SYCL installation,
- // rather than from e.g. /usr/local/include.
- if (JA.isOffloading(Action::OFK_Cuda))
- getToolChain().AddCudaIncludeArgs(Args, CmdArgs);
- if (JA.isOffloading(Action::OFK_HIP))
- getToolChain().AddHIPIncludeArgs(Args, CmdArgs);
- if (JA.isOffloading(Action::OFK_SYCL))
- getToolChain().addSYCLIncludeArgs(Args, CmdArgs);
-
- // If we are offloading to a target via OpenMP we need to include the
- // openmp_wrappers folder which contains alternative system headers.
- if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&
- !Args.hasArg(options::OPT_nostdinc) &&
- Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
- true) &&
- getToolChain().getTriple().isGPU()) {
- if (!Args.hasArg(options::OPT_nobuiltininc)) {
- // Add openmp_wrappers/* to our system include path. This lets us wrap
- // standard library headers.
- SmallString<128> P(D.ResourceDir);
- llvm::sys::path::append(P, "include");
- llvm::sys::path::append(P, "openmp_wrappers");
- CmdArgs.push_back("-internal-isystem");
- CmdArgs.push_back(Args.MakeArgString(P));
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ bool UsesOffloadInclude =
+ Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc, true);
+ bool NoBuiltinInc = Args.hasArg(options::OPT_nobuiltininc);
+
+ // Add offload include arguments for CUDA/HIP when using LLVM offloading. We
+ // want to pull in our wrappers instead of the vendor headers.
+ if (UsesLLVMOffloading) {
+ if (UsesOffloadInclude && !NoBuiltinInc) {
+ auto AddOffloadHeadersInclude = [&](StringRef IncludeSubdir,
+ StringRef RuntimeHeader) {
+ SmallString<128> OffloadInclude(D.Dir);
+ llvm::sys::path::append(OffloadInclude, "..", "include", "offload");
+ if (!IncludeSubdir.empty())
+ llvm::sys::path::append(OffloadInclude, IncludeSubdir);
+ CmdArgs.append({"-internal-isystem", Args.MakeArgString(OffloadInclude),
+ "-include", Args.MakeArgString(RuntimeHeader)});
+ };
+ auto AddForcedInclude = [&](StringRef Header) {
+ CmdArgs.push_back("-include");
+ CmdArgs.push_back(Args.MakeArgString(Header));
+ };
+ AddForcedInclude("__clang_gpu_runtime_wrapper.h");
+ AddForcedInclude("__clang_gpu_builtin_vars.h");
+ AddForcedInclude("__clang_gpu_device_functions.h");
+ AddForcedInclude("__clang_gpu_intrinsics.h");
+ if (JA.isOffloading(Action::OFK_Cuda))
+ AddOffloadHeadersInclude("cuda", "cuda_runtime.h");
+ if (JA.isOffloading(Action::OFK_HIP) &&
+ !Args.hasArg(options::OPT_nohipwrapperinc)) {
+ // HIP code commonly includes this as "hip/hip_runtime.h".
+ AddOffloadHeadersInclude("", "hip/hip_runtime.h");
+ }
}
+ } else {
+ // Add offload include arguments specific for CUDA/HIP/SYCL. This must
+ // happen before we -I or -include anything else, because we must pick up
+ // the CUDA/HIP/SYCL headers from the particular CUDA/ROCm/SYCL
+ // installation, rather than from e.g. /usr/local/include.
+ if (JA.isOffloading(Action::OFK_Cuda))
+ getToolChain().AddCudaIncludeArgs(Args, CmdArgs);
+ if (JA.isOffloading(Action::OFK_HIP))
+ getToolChain().AddHIPIncludeArgs(Args, CmdArgs);
+ if (JA.isOffloading(Action::OFK_SYCL))
+ getToolChain().addSYCLIncludeArgs(Args, CmdArgs);
+
+ // If we are offloading to a target via OpenMP we need to include the
+ // openmp_wrappers folder which contains alternative system headers.
+ if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&
+ !Args.hasArg(options::OPT_nostdinc) && UsesOffloadInclude &&
+ getToolChain().getTriple().isGPU()) {
+ if (!NoBuiltinInc) {
+ // Add openmp_wrappers/* to our system include path. This lets us
+ // wrap standard library headers.
+ SmallString<128> P(D.ResourceDir);
+ llvm::sys::path::append(P, "include");
+ llvm::sys::path::append(P, "openmp_wrappers");
+ CmdArgs.push_back("-internal-isystem");
+ CmdArgs.push_back(Args.MakeArgString(P));
+ }
- CmdArgs.push_back("-include");
- CmdArgs.push_back("__clang_openmp_device_functions.h");
- }
-
- if (Args.hasArg(options::OPT_foffload_via_llvm)) {
- // Add llvm_wrappers/* to our system include path. This lets us wrap
- // standard library headers and other headers.
- SmallString<128> P(D.ResourceDir);
- llvm::sys::path::append(P, "include", "llvm_offload_wrappers");
- CmdArgs.append({"-internal-isystem", Args.MakeArgString(P), "-include"});
- if (JA.isDeviceOffloading(Action::OFK_OpenMP))
- CmdArgs.push_back("__llvm_offload_device.h");
- else
- CmdArgs.push_back("__llvm_offload_host.h");
+ CmdArgs.push_back("-include");
+ CmdArgs.push_back("__clang_openmp_device_functions.h");
+ }
}
// Add -i* options, and automatically translate to
@@ -5166,6 +5189,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
bool IsSYCLDevice = JA.isDeviceOffloading(Action::OFK_SYCL);
bool IsOpenMPDevice = JA.isDeviceOffloading(Action::OFK_OpenMP);
bool IsExtractAPI = isa<ExtractAPIJobAction>(JA);
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
bool IsDeviceOffloadAction = !(JA.isDeviceOffloading(Action::OFK_None) ||
JA.isDeviceOffloading(Action::OFK_Host));
bool IsHostOffloadingAction =
@@ -5285,7 +5310,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
}
}
- if (IsCuda && !IsCudaDevice) {
+ if (IsCuda && !IsCudaDevice && !UsesLLVMOffloading) {
// We need to figure out which CUDA version we're compiling for, as that
// determines how we load and launch GPU kernels.
auto *CTC = static_cast<const toolchains::CudaToolChain *>(
@@ -8293,12 +8318,13 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
// Host-side offloading compilation receives all device-side outputs. Include
// them in the host compilation depending on the target. If the host inputs
// are not empty we use the new-driver scheme, otherwise use the old scheme.
- if ((IsCuda || IsHIP) && CudaDeviceInput) {
+ if ((IsCuda || IsHIP) && !UsesLLVMOffloading && CudaDeviceInput) {
CmdArgs.push_back("-fcuda-include-gpubinary");
CmdArgs.push_back(CudaDeviceInput->getFilename());
} else if (!HostOffloadingInputs.empty()) {
if ((IsCuda || IsHIP) &&
- (!IsRDCMode || Args.hasArg(options::OPT_cuda_emit_nvcc_abi))) {
+ (!IsRDCMode || Args.hasArg(options::OPT_cuda_emit_nvcc_abi)) &&
+ !UsesLLVMOffloading) {
assert(HostOffloadingInputs.size() == 1 && "Only one input expected");
CmdArgs.push_back("-fcuda-include-gpubinary");
CmdArgs.push_back(HostOffloadingInputs.front().getFilename());
diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp b/clang/lib/Driver/ToolChains/CommonArgs.cpp
index 883296e43111b..a76f4aa6ae853 100644
--- a/clang/lib/Driver/ToolChains/CommonArgs.cpp
+++ b/clang/lib/Driver/ToolChains/CommonArgs.cpp
@@ -1491,18 +1491,25 @@ void tools::addArchSpecificRPath(const ToolChain &TC, const ArgList &Args,
}
}
+bool tools::addLLVMOffloadingRuntime(const Compilation &C,
+ ArgStringList &CmdArgs,
+ const ToolChain &TC, const ArgList &Args) {
+
+ if (!Args.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
+ return false;
+
+ CmdArgs.push_back("-lLLVMOffloadKernel");
+ return true;
+}
+
bool tools::addOpenMPRuntime(const Compilation &C, ArgStringList &CmdArgs,
const ToolChain &TC, const ArgList &Args,
bool ForceStaticHostRuntime, bool IsOffloadingHost,
bool GompNeedsRT) {
if (!Args.hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ,
- options::OPT_fno_openmp, false)) {
- // We need libomptarget (liboffload) if it's the choosen offloading runtime.
- if (Args.hasFlag(options::OPT_foffload_via_llvm,
- options::OPT_fno_offload_via_llvm, false))
- CmdArgs.push_back("-lomptarget");
+ options::OPT_fno_openmp, false))
return false;
- }
Driver::OpenMPRuntimeKind RTKind = TC.getDriver().getOpenMPRuntime(Args);
diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp
index 54585105373da..77d1f6bb556d5 100644
--- a/clang/lib/Driver/ToolChains/Cuda.cpp
+++ b/clang/lib/Driver/ToolChains/Cuda.cpp
@@ -303,6 +303,10 @@ CudaInstallationDetector::CudaInstallationDetector(
void CudaInstallationDetector::AddCudaIncludeArgs(
const ArgList &DriverArgs, ArgStringList &CC1Args) const {
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
+ return;
+
if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
// Add cuda_wrappers/* to our system include path. This lets us wrap
// standard library headers.
@@ -398,7 +402,10 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
const char *LinkingOutput) const {
const auto &TC =
static_cast<const toolchains::NVPTXToolChain &>(getToolChain());
- assert(TC.getTriple().isNVPTX() && "Wrong platform");
+
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ assert((TC.getTriple().isNVPTX() || UsesLLVMOffloading) && "Wrong platform");
BoundArch GPUArch;
// If this is a CUDA action we need to extract the device architecture
@@ -421,7 +428,7 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
"Device action expected to have an architecture.");
// Check that our installation's ptxas supports gpu_arch.
- if (!Args.hasArg(options::OPT_no_cuda_version_check)) {
+ if (!UsesLLVMOffloading && !Args.hasArg(options::OPT_no_cuda_version_check)) {
TC.CudaInstallation.CheckCudaVersionSupportsArch(GPUArch.Arch);
}
@@ -494,7 +501,8 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
/*Default=*/true);
else if (JA.isOffloading(Action::OFK_Cuda))
// In CUDA we generate relocatable code by default.
- Relocatable = Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
+ Relocatable = UsesLLVMOffloading ||
+ Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
/*Default=*/false);
else
// Otherwise, we are compiling directly and should create linkable output.
@@ -543,7 +551,9 @@ void NVPTX::FatBinary::ConstructJob(Compilation &C, const JobAction &JA,
const char *LinkingOutput) const {
const auto &TC =
static_cast<const toolchains::CudaToolChain &>(getToolChain());
- assert(TC.getTriple().isNVPTX() && "Wrong platform");
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ assert((UsesLLVMOffloading || TC.getTriple().isNVPTX()) && "Wrong platform");
ArgStringList CmdArgs;
if (TC.CudaInstallation.version() <= CudaVersion::CUDA_100)
@@ -591,7 +601,9 @@ void NVPTX::Linker::ConstructJob(Compilation &C, const JobAction &JA,
static_cast<const toolchains::NVPTXToolChain &>(getToolChain());
ArgStringList CmdArgs;
- assert(TC.getTriple().isNVPTX() && "Wrong platform");
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ assert((UsesLLVMOffloading || TC.getTriple().isNVPTX()) && "Wrong platform");
assert((Output.isFilename() || Output.isNothing()) && "Invalid output.");
if (Output.isFilename()) {
@@ -897,9 +909,12 @@ void CudaToolChain::addClangTargetOptions(
BoundArch BA, Action::OffloadKind DeviceOffloadingKind) const {
HostTC.addClangTargetOptions(DriverArgs, CC1Args, BA, DeviceOffloadingKind);
+ bool UsesLLVMOffloading = DriverArgs.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+
StringRef GpuArch = DriverArgs.getLastArgValue(options::OPT_march_EQ);
assert((DeviceOffloadingKind == Action::OFK_OpenMP ||
- DeviceOffloadingKind == Action::OFK_Cuda) &&
+ DeviceOffloadingKind == Action::OFK_Cuda || UsesLLVMOffloading) &&
"Only OpenMP or CUDA offloading kinds are supported for NVIDIA GPUs.");
CC1Args.append({"-fcuda-is-device", "-mllvm",
@@ -918,6 +933,9 @@ void CudaToolChain::addClangTargetOptions(
DriverArgs.hasArg(options::OPT_S))
return;
+ if (UsesLLVMOffloading)
+ return;
+
std::string LibDeviceFile = CudaInstallation.getLibDeviceFile(GpuArch);
if (LibDeviceFile.empty()) {
getDriver().Diag(diag::err_drv_no_cuda_libdevice) << GpuArch;
@@ -927,13 +945,6 @@ void CudaToolChain::addClangTargetOptions(
CC1Args.push_back("-mlink-builtin-bitcode");
CC1Args.push_back(DriverArgs.MakeArgString(LibDeviceFile));
- // For now, we don't use any Offload/OpenMP device runtime when we offload
- // CUDA via LLVM/Offload. We should split the Offload/OpenMP device runtime
- // and include the "generic" (or CUDA-specific) parts.
- if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
- options::OPT_fno_offload_via_llvm, false))
- return;
-
clang::CudaVersion CudaInstallationVersion = CudaInstallation.version();
if (CudaInstallationVersion >= CudaVersion::UNKNOWN)
@@ -974,6 +985,10 @@ llvm::DenormalMode CudaToolChain::getDefaultDenormalModeForType(
void CudaToolChain::AddCudaIncludeArgs(const ArgList &DriverArgs,
ArgStringList &CC1Args) const {
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
+ return;
+
// Check our CUDA version if we're going to include the CUDA headers.
if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
true) &&
@@ -1046,6 +1061,10 @@ CudaToolChain::GetCXXStdlibType(const ArgList &Args) const {
void CudaToolChain::AddClangSystemIncludeArgs(const ArgList &DriverArgs,
ArgStringList &CC1Args) const {
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
+ return;
+
HostTC.AddClangSystemIncludeArgs(DriverArgs, CC1Args);
if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
diff --git a/clang/lib/Driver/ToolChains/Gnu.cpp b/clang/lib/Driver/ToolChains/Gnu.cpp
index 24076d8814322..72affac131701 100644
--- a/clang/lib/Driver/ToolChains/Gnu.cpp
+++ b/clang/lib/Driver/ToolChains/Gnu.cpp
@@ -510,6 +510,7 @@ void tools::gnutools::Linker::ConstructJob(Compilation &C, const JobAction &JA,
// FIXME: Does this really make sense for all GNU toolchains?
WantPthread = true;
+ addLLVMOffloadingRuntime(C, CmdArgs, ToolChain, Args);
AddRunTimeLibs(ToolChain, D, CmdArgs, Args);
// LLVM support for atomics on 32-bit SPARC V8+ is incomplete, so
diff --git a/clang/lib/Driver/ToolChains/Linux.cpp b/clang/lib/Driver/ToolChains/Linux.cpp
index 89af9847e5ae5..e295b2516da16 100644
--- a/clang/lib/Driver/ToolChains/Linux.cpp
+++ b/clang/lib/Driver/ToolChains/Linux.cpp
@@ -886,7 +886,9 @@ void Linux::addOffloadRTLibs(unsigned ActiveKinds, const ArgList &Args,
if (!Args.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
true) ||
Args.hasArg(options::OPT_nostdlib) ||
- Args.hasArg(options::OPT_no_hip_rt) || Args.hasArg(options::OPT_r))
+ Args.hasArg(options::OPT_no_hip_rt) || Args.hasArg(options::OPT_r) ||
+ Args.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
return;
llvm::SmallVector<std::pair<StringRef, StringRef>> Libraries;
diff --git a/clang/test/CodeGenCUDA/Inputs/cuda.h b/clang/test/CodeGenCUDA/Inputs/cuda.h
index 421fa4dd7dbae..0968c8978ea9f 100644
--- a/clang/test/CodeGenCUDA/Inputs/cuda.h
+++ b/clang/test/CodeGenCUDA/Inputs/cuda.h
@@ -56,7 +56,7 @@ extern "C" hipError_t hipLaunchKernel_spt(const void *func, dim3 gridDim,
extern "C" unsigned __llvmPushCallConfiguration(dim3 gridDim, dim3 blockDim,
size_t sharedMem = 0, void *stream = 0);
extern "C" unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
- void **args, size_t sharedMem = 0, void *stream = 0);
+ void *args, size_t sharedMem = 0, void *stream = 0);
#else
typedef struct cudaStream *cudaStream_t;
typedef enum cudaError {} cudaError_t;
diff --git a/clang/test/CodeGenCUDA/offload_via_llvm.cu b/clang/test/CodeGenCUDA/offload_via_llvm.cu
index b13a64c81b775..c99b3e5e3b334 100644
--- a/clang/test/CodeGenCUDA/offload_via_llvm.cu
+++ b/clang/test/CodeGenCUDA/offload_via_llvm.cu
@@ -14,9 +14,7 @@
// HST-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 2
// HST-NEXT: [[DOTADDR2:%.*]] = alloca ptr, align 4
// HST-NEXT: [[DOTADDR3:%.*]] = alloca ptr, align 4
-// HST-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[TMP0]], align 16
-// HST-NEXT: [[KERNEL_ARGS_PTRS:%.*]] = alloca [4 x ptr], align 16
-// HST-NEXT: [[KERNEL_LAUNCH_PARAMS:%.*]] = alloca [[TMP1]], align 16
+// HST-NEXT: [[KERNEL_LAUNCH_PARAMS:%.*]] = alloca [[TMP0]], align 16
// HST-NEXT: [[GRID_DIM:%.*]] = alloca [[STRUCT_DIM3:%.*]], align 8
// HST-NEXT: [[BLOCK_DIM:%.*]] = alloca [[STRUCT_DIM3]], align 8
// HST-NEXT: [[SHMEM_SIZE:%.*]] = alloca i32, align 4
@@ -25,34 +23,34 @@
// HST-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 2
// HST-NEXT: store ptr [[TMP2]], ptr [[DOTADDR2]], align 4
// HST-NEXT: store ptr [[TMP3]], ptr [[DOTADDR3]], align 4
-// HST-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 0
-// HST-NEXT: store i32 4, ptr [[TMP4]], align 16
-// HST-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 1
-// HST-NEXT: store ptr [[KERNEL_ARGS_PTRS]], ptr [[TMP5]], align 4
-// HST-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTADDR]], align 4
-// HST-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 0
-// HST-NEXT: store i32 [[TMP6]], ptr [[TMP7]], align 16
-// HST-NEXT: [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 0
-// HST-NEXT: store ptr [[TMP7]], ptr [[TMP8]], align 16
-// HST-NEXT: [[TMP9:%.*]] = load i16, ptr [[DOTADDR1]], align 2
-// HST-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 1
-// HST-NEXT: store i16 [[TMP9]], ptr [[TMP10]], align 4
-// HST-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 1
-// HST-NEXT: store ptr [[TMP10]], ptr [[TMP11]], align 4
-// HST-NEXT: [[TMP12:%.*]] = load ptr, ptr [[DOTADDR2]], align 4
-// HST-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 2
-// HST-NEXT: store ptr [[TMP12]], ptr [[TMP13]], align 8
-// HST-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 2
-// HST-NEXT: store ptr [[TMP13]], ptr [[TMP14]], align 8
-// HST-NEXT: [[TMP15:%.*]] = load ptr, ptr [[DOTADDR3]], align 4
-// HST-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 3
-// HST-NEXT: store ptr [[TMP15]], ptr [[TMP16]], align 4
-// HST-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 3
-// HST-NEXT: store ptr [[TMP16]], ptr [[TMP17]], align 4
-// HST-NEXT: [[TMP18:%.*]] = call i32 @__llvmPopCallConfiguration(ptr [[GRID_DIM]], ptr [[BLOCK_DIM]], ptr [[SHMEM_SIZE]], ptr [[STREAM]])
-// HST-NEXT: [[TMP19:%.*]] = load i32, ptr [[SHMEM_SIZE]], align 4
-// HST-NEXT: [[TMP20:%.*]] = load ptr, ptr [[STREAM]], align 4
-// HST-NEXT: [[CALL:%.*]] = call noundef i32 @llvmLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP19]], ptr noundef [[TMP20]]) #[[ATTR3:[0-9]+]]
+// HST-NEXT: [[KERNEL_ARGS:%.*]] = alloca ptr, i32 4, align 16
+// HST-NEXT: [[KERNEL_ARG_SIZES:%.*]] = alloca i32, i32 4, align 16
+// HST-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 0
+// HST-NEXT: store ptr [[KERNEL_ARGS]], ptr [[TMP4]], align 16
+// HST-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 1
+// HST-NEXT: store i64 4, ptr [[TMP5]], align 8
+// HST-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 2
+// HST-NEXT: store ptr [[KERNEL_ARG_SIZES]], ptr [[TMP6]], align 16
+// HST-NEXT: [[TMP7:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 0
+// HST-NEXT: store ptr [[DOTADDR]], ptr [[TMP7]], align 4
+// HST-NEXT: [[TMP8:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 0
+// HST-NEXT: store i32 4, ptr [[TMP8]], align 4
+// HST-NEXT: [[TMP9:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 1
+// HST-NEXT: store ptr [[DOTADDR1]], ptr [[TMP9]], align 4
+// HST-NEXT: [[TMP10:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 1
+// HST-NEXT: store i32 2, ptr [[TMP10]], align 4
+// HST-NEXT: [[TMP11:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 2
+// HST-NEXT: store ptr [[DOTADDR2]], ptr [[TMP11]], align 4
+// HST-NEXT: [[TMP12:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 2
+// HST-NEXT: store i32 4, ptr [[TMP12]], align 4
+// HST-NEXT: [[TMP13:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 3
+// HST-NEXT: store ptr [[DOTADDR3]], ptr [[TMP13]], align 4
+// HST-NEXT: [[TMP14:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 3
+// HST-NEXT: store i32 4, ptr [[TMP14]], align 4
+// HST-NEXT: [[TMP15:%.*]] = call i32 @__llvmPopCallConfiguration(ptr [[GRID_DIM]], ptr [[BLOCK_DIM]], ptr [[SHMEM_SIZE]], ptr [[STREAM]])
+// HST-NEXT: [[TMP16:%.*]] = load i32, ptr [[SHMEM_SIZE]], align 4
+// HST-NEXT: [[TMP17:%.*]] = load ptr, ptr [[STREAM]], align 4
+// HST-NEXT: [[CALL:%.*]] = call noundef i32 @llvmLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP16]], ptr noundef [[TMP17]]) #[[ATTR3:[0-9]+]]
// HST-NEXT: br label %[[SETUP_END:.*]]
// HST: [[SETUP_END]]:
// HST-NEXT: ret void
diff --git a/clang/test/Driver/cuda-via-liboffload.cu b/clang/test/Driver/cuda-via-liboffload.cu
index 68dc963e906b2..d30e529f0ce12 100644
--- a/clang/test/Driver/cuda-via-liboffload.cu
+++ b/clang/test/Driver/cuda-via-liboffload.cu
@@ -2,21 +2,20 @@
// RUN: --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
// RUN: | FileCheck -check-prefix BINDINGS %s
-// BINDINGS: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[HOST_BC:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT]]", "[[HOST_BC]]"], output: "[[PTX_SM_35:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX_SM_35]]"], output: "[[CUBIN_SM_35:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT]]", "[[HOST_BC]]"], output: "[[PTX_SM_70:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX_SM_70:.+]]"], output: "[[CUBIN_SM_70:.+]]"
+// BINDINGS: "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX_SM_35:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX_SM_35]]"], output: "[[CUBIN_SM_35:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT]]"], output: "[[PTX_SM_70:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX_SM_70:.+]]"], output: "[[CUBIN_SM_70:.+]]"
// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Packager", inputs: ["[[CUBIN_SM_35]]", "[[CUBIN_SM_70]]"], output: "[[BINARY:.+]]"
-// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[HOST_BC]]", "[[BINARY]]"], output: "[[HOST_OBJ:.+]]"
+// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT]]", "[[BINARY]]"], output: "[[HOST_OBJ:.+]]"
// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Linker", inputs: ["[[HOST_OBJ]]"], output: "a.out"
// RUN: %clang -### -target x86_64-linux-gnu -foffload-via-llvm -ccc-print-bindings \
// RUN: --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
// RUN: | FileCheck -check-prefix BINDINGS-DEVICE %s
-// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX:.+]]"
-// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX]]"], output: "[[CUBIN:.+]]"
+// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX:.+]]"
+// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX]]"], output: "[[CUBIN:.+]]"
// RUN: %clang -### -target x86_64-linux-gnu -ccc-print-bindings --offload-link -foffload-via-llvm %s 2>&1 | FileCheck -check-prefix DEVICE-LINK %s
diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
index f2a58774e99af..21aee2121f255 100644
--- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
+++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
@@ -139,6 +139,13 @@ static bool CanonicalPrefixes = true;
using OffloadingImage = OffloadBinary::OffloadingImage;
+static bool usesLLVMOffloadWrapper(ArrayRef<OffloadingImage> Images) {
+ return llvm::any_of(Images, [](const OffloadingImage &Image) {
+ return Triple(Image.StringData.lookup("triple")).getEnvironment() ==
+ Triple::LLVM;
+ });
+}
+
namespace llvm {
// Provide DenseMapInfo so that OffloadKind can be used in a DenseMap.
template <> struct DenseMapInfo<OffloadKind> {
@@ -977,6 +984,9 @@ Expected<SmallVector<std::unique_ptr<MemoryBuffer>>>
bundleLinkedOutput(ArrayRef<OffloadingImage> Images, const ArgList &Args,
OffloadKind Kind) {
llvm::TimeTraceScope TimeScope("Bundle linked output");
+ if (usesLLVMOffloadWrapper(Images))
+ return bundleOpenMP(Images);
+
switch (Kind) {
case OFK_OpenMP:
return (Verbose && SaveTemps) ? bundleOpenMPVerbose(Images)
@@ -1220,7 +1230,8 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
continue;
}
- auto OutputOrErr = wrapDeviceImages(*BundledImagesOrErr, Args, Kind);
+ OffloadKind WrapperKind = usesLLVMOffloadWrapper(Input) ? OFK_OpenMP : Kind;
+ auto OutputOrErr = wrapDeviceImages(*BundledImagesOrErr, Args, WrapperKind);
if (!OutputOrErr)
return OutputOrErr.takeError();
WrappedOutput.push_back(*OutputOrErr);
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index b072087ba0e7d..9b0d9c98e8e5f 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -9,6 +9,8 @@
#ifndef LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_RUNTIME_H
#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_RUNTIME_H
+#include "LanguageLaunch.h"
+#include "Types.h"
#include <cstddef>
#include <cstdint>
#include <cstdio>
diff --git a/offload/languages/kernel/CMakeLists.txt b/offload/languages/kernel/CMakeLists.txt
index e64852b940702..293df7becdc93 100644
--- a/offload/languages/kernel/CMakeLists.txt
+++ b/offload/languages/kernel/CMakeLists.txt
@@ -78,4 +78,6 @@ install(FILES
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/DefineLanguageNames.inc
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageRuntime.h
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/UndefineLanguageNames.inc
+ ${CMAKE_CURRENT_SOURCE_DIR}/include/LanguageLaunch.h
+ ${CMAKE_CURRENT_SOURCE_DIR}/include/Types.h
DESTINATION ${CMAKE_INSTALL_PREFIX}/include/offload/kernel/)
diff --git a/offload/languages/kernel/exports b/offload/languages/kernel/exports
index 3372890ff42c6..d58962564c008 100644
--- a/offload/languages/kernel/exports
+++ b/offload/languages/kernel/exports
@@ -2,8 +2,14 @@ VERS1.0 {
global:
*cuda*;
*hip*;
- llvmLaunchKernel*;
- __llvm*;
+ llvmLaunchKernel;
+ __llvmPushCallConfiguration;
+ __llvmPopCallConfiguration;
+ __llvmRegisterFunction;
+ __llvmRegisterVar;
+ __llvmRegisterManagedVar;
+ __llvmRegisterSurface;
+ __llvmRegisterTexture;
__tgt_register_lib;
__tgt_unregister_lib;
local:
diff --git a/offload/languages/kernel/include/LanguageAliases.inc b/offload/languages/kernel/include/LanguageAliases.inc
index 551f64a5b8fc8..bebd5c9d78f46 100644
--- a/offload/languages/kernel/include/LanguageAliases.inc
+++ b/offload/languages/kernel/include/LanguageAliases.inc
@@ -63,12 +63,6 @@ extern "C" void LANGUAGE_NAME(__, RegisterTexture)(
__llvmRegisterTexture(Data, TexRef, DevPtr, Name, Dim, Norm, Ext);
}
-extern "C" unsigned
-LANGUAGE_NAME(__, PopCallConfiguration)(dim3 *GridSize, dim3 *BlockSize,
- size_t *SharedMemory, void **Stream) {
- return __llvmPopCallConfiguration(GridSize, BlockSize, SharedMemory, Stream);
-}
-
#undef LANGUAGE_NAME
#undef LA_IMPL1
#undef LA_IMPL2
diff --git a/offload/languages/kernel/include/LanguageLaunch.h b/offload/languages/kernel/include/LanguageLaunch.h
index c72f3ec215ac1..de6d3e25d77ca 100644
--- a/offload/languages/kernel/include/LanguageLaunch.h
+++ b/offload/languages/kernel/include/LanguageLaunch.h
@@ -9,7 +9,6 @@
#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_LAUNCH_H
#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_LAUNCH_H
-#include "OffloadAPI.h"
#include "Types.h"
#include <cstddef>
@@ -19,21 +18,17 @@ extern "C" {
/// Push call configuration for kernel launch
unsigned __llvmPushCallConfiguration(dim3 __grid_size, dim3 __block_size,
- size_t __shared_memory, void *__stream);
+ size_t __shared_memory = 0,
+ void *__stream = 0);
/// Pop call configuration for kernel launch
unsigned __llvmPopCallConfiguration(dim3 *__grid_size, dim3 *__block_size,
size_t *__shared_memory, void **__stream);
-/// Internal kernel launch implementation
-ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
- dim3 BlockDim, void *KernelArgsPtr,
- size_t DynamicSharedMem, void *Stream);
-
/// LLVM-style kernel launch entry point
-unsigned __llvmLaunchKernel(const char *KernelID, dim3 GridDim, dim3 BlockDim,
- void *KernelArgsPtr, size_t DynamicSharedMem,
- void *Stream);
+unsigned llvmLaunchKernel(const char *KernelID, dim3 GridDim, dim3 BlockDim,
+ void *KernelArgsPtr, size_t DynamicSharedMem = 0,
+ void *Stream = 0);
}
#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_LAUNCH_H
diff --git a/offload/languages/kernel/include/Types.h b/offload/languages/kernel/include/Types.h
index 56435daf602fc..582a1a0e08c01 100644
--- a/offload/languages/kernel/include/Types.h
+++ b/offload/languages/kernel/include/Types.h
@@ -13,10 +13,19 @@
#include <cstdint>
struct uint3 {
- unsigned x = 0, y = 0, z = 0;
+ // CUDA/HIP uint3 is plain vector storage; default init does not zero it.
+ unsigned int x, y, z;
};
-using dim3 = uint3;
+struct dim3 {
+ // CUDA/HIP dim3 model launch dimensions; omitted axes default to one.
+ dim3(unsigned int X = 1, unsigned int Y = 1, unsigned int Z = 1)
+ : x(X), y(Y), z(Z) {}
+ dim3(uint3 V) : x(V.x), y(V.y), z(V.z) {}
+ operator uint3() const { return {x, y, z}; }
+
+ unsigned int x, y, z;
+};
struct CallConfigurationTy {
dim3 GridSize;
diff --git a/offload/languages/kernel/src/LanguageLaunch.cpp b/offload/languages/kernel/src/LanguageLaunch.cpp
index eff3d11668c17..61003b44330a7 100644
--- a/offload/languages/kernel/src/LanguageLaunch.cpp
+++ b/offload/languages/kernel/src/LanguageLaunch.cpp
@@ -15,31 +15,6 @@
using RuntimeState = llvm::offload::StateTy;
using ThreadState = llvm::offload::ThreadStateTy;
-extern "C" {
-
-/// Push call configuration for kernel launch
-unsigned __llvmPushCallConfiguration(dim3 GridSize, dim3 BlockSize,
- size_t SharedMemory, void *Stream) {
- CallConfigurationTy &CC = ThreadState::getCallConfiguration();
-
- CC.GridSize = GridSize;
- CC.BlockSize = BlockSize;
- CC.SharedMemory = SharedMemory;
- CC.Stream = Stream;
- return 0;
-}
-
-/// Pop call configuration for kernel launch
-unsigned __llvmPopCallConfiguration(dim3 *GridSize, dim3 *BlockSize,
- size_t *SharedMemory, void **Stream) {
- CallConfigurationTy &CC = ThreadState::getCallConfiguration();
- *GridSize = CC.GridSize;
- *BlockSize = CC.BlockSize;
- *SharedMemory = CC.SharedMemory;
- *Stream = CC.Stream;
- return 0;
-}
-
/// Internal kernel launch implementation
ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
dim3 BlockDim, void *KernelArgsPtr,
@@ -73,9 +48,34 @@ ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
OKA->ArgSizes);
}
-unsigned __llvmLaunchKernel(const char *KernelID, dim3 GridDim, dim3 BlockDim,
- void *KernelArgsPtr, size_t DynamicSharedMem,
- void *Stream) {
+extern "C" {
+
+/// Push call configuration for kernel launch
+unsigned __llvmPushCallConfiguration(dim3 GridSize, dim3 BlockSize,
+ size_t SharedMemory, void *Stream) {
+ CallConfigurationTy &CC = ThreadState::getCallConfiguration();
+
+ CC.GridSize = GridSize;
+ CC.BlockSize = BlockSize;
+ CC.SharedMemory = SharedMemory;
+ CC.Stream = Stream;
+ return 0;
+}
+
+/// Pop call configuration for kernel launch
+unsigned __llvmPopCallConfiguration(dim3 *GridSize, dim3 *BlockSize,
+ size_t *SharedMemory, void **Stream) {
+ CallConfigurationTy &CC = ThreadState::getCallConfiguration();
+ *GridSize = CC.GridSize;
+ *BlockSize = CC.BlockSize;
+ *SharedMemory = CC.SharedMemory;
+ *Stream = CC.Stream;
+ return 0;
+}
+
+unsigned llvmLaunchKernel(const char *KernelID, dim3 GridDim, dim3 BlockDim,
+ void *KernelArgsPtr, size_t DynamicSharedMem,
+ void *Stream) {
ol_result_t Result = __llvmLaunchKernelImpl(
KernelID, GridDim, BlockDim, KernelArgsPtr, DynamicSharedMem, Stream);
return Result ? Result->Code : 0;
diff --git a/offload/test/lit.cfg b/offload/test/lit.cfg
index ace2b1ea8a749..37255f14ec124 100644
--- a/offload/test/lit.cfg
+++ b/offload/test/lit.cfg
@@ -83,7 +83,7 @@ def remove_suffix_if_present(name):
config.name = 'libomptarget :: ' + config.libomptarget_current_target
# suffixes: A list of file extensions to treat as test files.
-config.suffixes = ['.c', '.cpp', '.cc', '.f90', '.cu', '.td']
+config.suffixes = ['.c', '.cpp', '.cc', '.f90', '.cu', '.hip', '.td']
# excludes: A list of directories to exclude from the testuites.
config.excludes = ['Inputs', 'unit']
@@ -91,6 +91,12 @@ config.excludes = ['Inputs', 'unit']
# test_source_root: The root path where tests are located.
config.test_source_root = os.path.dirname(__file__)
+# language includes
+config.test_language_includes = os.path.join(config.test_source_root, "../languages/include")
+config.test_language_cuda_includes = os.path.join(config.test_language_includes, "cuda")
+config.test_language_hip_includes = os.path.join(config.test_language_includes, "hip")
+config.test_language_kernel_includes = os.path.join(config.test_source_root, "../languages/kernel/include")
+
# test_exec_root: The root object directory where output is placed
config.test_exec_root = config.libomptarget_obj_root
@@ -100,6 +106,10 @@ config.test_format = lit.formats.ShTest()
# compiler flags
config.test_flags = " -I " + config.test_source_root + \
" -I " + config.omp_header_directory + \
+ " -I " + config.test_language_includes + \
+ " -I " + config.test_language_cuda_includes + \
+ " -I " + config.test_language_hip_includes + \
+ " -I " + config.test_language_kernel_includes + \
" -L " + config.library_dir + \
" -L " + config.llvm_library_intdir + \
" -L " + config.llvm_lib_directory
diff --git a/offload/test/offloading/CUDA/basic_launch.cu b/offload/test/offloading/CUDA/basic_launch.cu
index e017241bb9a74..5ecfc3e9d5601 100644
--- a/offload/test/offloading/CUDA/basic_launch.cu
+++ b/offload/test/offloading/CUDA/basic_launch.cu
@@ -7,25 +7,24 @@
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
__global__ void square(int *A) { *A = 42; }
int main(int argc, char **argv) {
- int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- *Ptr = 7;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
+ int *Ptr;
+ cudaMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
square<<<1, 1>>>(Ptr);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- llvm_omp_target_free_shared(Ptr, DevNo);
+ int I = 0;
+ cudaDeviceSynchronize();
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
}
diff --git a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
index a428e25d82359..0bfc9e231ffde 100644
--- a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
+++ b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
@@ -7,27 +7,25 @@
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
__global__ void square(int *A) {
__scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
}
int main(int argc, char **argv) {
int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- *Ptr = 0;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 0
+ int *Ptr, I;
+ cudaMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
square<<<7, 6>>>(Ptr);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- llvm_omp_target_free_shared(Ptr, DevNo);
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
}
diff --git a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
index db2a1e48371b0..505a9f9379c08 100644
--- a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
+++ b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
@@ -6,15 +6,13 @@
// clang-format on
// REQUIRES: gpu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
__global__ void square(int *Dst, short Q, int *Src, short P) {
*Dst = (Src[0] + Src[1]) * (Q + P);
Src[0] = Q;
@@ -23,19 +21,19 @@ __global__ void square(int *Dst, short Q, int *Src, short P) {
int main(int argc, char **argv) {
int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- int *Src = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(8, DevNo));
- *Ptr = 7;
- Src[0] = -2;
- Src[1] = 8;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
- printf("Src: %i : %i\n", Src[0], Src[1]);
- // CHECK: Src: -2 : 8
+ int *Src, *Ptr;
+ cudaMalloc(&Ptr, 4);
+ cudaMalloc(&Src, 8);
+
+ int I = 7;
+ int HostSrc[2] = {-2, 8};
+ cudaMemcpy(Ptr, &I, sizeof(int), cudaMemcpyHostToDevice);
+ cudaMemcpy(Src, &HostSrc[0], 2 * sizeof(int), cudaMemcpyHostToDevice);
square<<<1, 1>>>(Ptr, 3, Src, 4);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- printf("Src: %i : %i\n", Src[0], Src[1]);
- // CHECK: Src: 3 : 4
- llvm_omp_target_free_shared(Ptr, DevNo);
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ cudaMemcpy(&HostSrc[0], Src, 2 * sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+ printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
+ // CHECK: Src: 3, 4
}
diff --git a/offload/test/offloading/CUDA/device_api.cu b/offload/test/offloading/CUDA/device_api.cu
new file mode 100644
index 0000000000000..184167f2d17e4
--- /dev/null
+++ b/offload/test/offloading/CUDA/device_api.cu
@@ -0,0 +1,45 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int Count = 0;
+ if (cudaGetDeviceCount(&Count) != cudaSuccess)
+ return 1;
+
+ printf("device count: %d\n", Count);
+ // CHECK: device count: {{[1-9][0-9]*}}
+
+ int Device = -1;
+ if (cudaGetDevice(&Device) != cudaSuccess)
+ return 1;
+
+ printf("device: %d\n", Device);
+ // CHECK: device: {{[0-9]+}}
+
+ if (cudaSetDevice(Device) != cudaSuccess)
+ return 1;
+
+ int After = -1;
+ if (cudaGetDevice(&After) != cudaSuccess)
+ return 1;
+
+ printf("device after set: %d\n", After);
+ // CHECK: device after set: {{[0-9]+}}
+
+ cudaError_t Err = cudaSetDevice(-1);
+ printf("set invalid device: %u\n", Err);
+ // CHECK: set invalid device: 1
+}
diff --git a/offload/test/offloading/CUDA/device_properties.cu b/offload/test/offloading/CUDA/device_properties.cu
new file mode 100644
index 0000000000000..8f625f6ccabe1
--- /dev/null
+++ b/offload/test/offloading/CUDA/device_properties.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ cudaDeviceProp Prop = {};
+ cudaError_t Err = cudaGetDeviceProperties(&Prop, 0);
+ if (Err != cudaSuccess) {
+ printf("cudaGetDeviceProperties failed: %u\n", Err);
+ return 1;
+ }
+
+ printf("Device name: %s\n", Prop.name);
+ // CHECK: Device name:
+ printf("Total global memory: %zu\n", Prop.totalGlobalMem);
+ // CHECK: Total global memory:
+ printf("Multiprocessors: %i\n", Prop.multiProcessorCount);
+ // CHECK: Multiprocessors:
+ printf("Warp size: %i\n", Prop.warpSize);
+ // CHECK: Warp size:
+
+ if (!Prop.name[0] || !Prop.totalGlobalMem || !Prop.multiProcessorCount ||
+ !Prop.warpSize)
+ return 1;
+
+ printf("Device properties are populated.\n");
+ // CHECK: Device properties are populated.
+}
diff --git a/offload/test/offloading/CUDA/host_alloc.cu b/offload/test/offloading/CUDA/host_alloc.cu
new file mode 100644
index 0000000000000..1440b276c6602
--- /dev/null
+++ b/offload/test/offloading/CUDA/host_alloc.cu
@@ -0,0 +1,48 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void add(int *Ptr, int Value) { *Ptr += Value; }
+
+int main(int argc, char **argv) {
+ int *HostAllocPtr = nullptr;
+ if (cudaHostAlloc(&HostAllocPtr, sizeof(int), cudaHostAllocDefault) !=
+ cudaSuccess)
+ return 1;
+
+ *HostAllocPtr = 17;
+ add<<<1, 1>>>(HostAllocPtr, 5);
+ if (cudaDeviceSynchronize() != cudaSuccess)
+ return 1;
+ printf("cudaHostAlloc value: %d\n", *HostAllocPtr);
+ // CHECK: cudaHostAlloc value: 22
+
+ if (cudaFreeHost(HostAllocPtr) != cudaSuccess)
+ return 1;
+
+ int *MallocHostPtr = nullptr;
+ if (cudaMallocHost(&MallocHostPtr, sizeof(int)) != cudaSuccess)
+ return 1;
+
+ *MallocHostPtr = 23;
+ add<<<1, 1>>>(MallocHostPtr, 7);
+ if (cudaDeviceSynchronize() != cudaSuccess)
+ return 1;
+ printf("cudaMallocHost value: %d\n", *MallocHostPtr);
+ // CHECK: cudaMallocHost value: 30
+
+ if (cudaFreeHost(MallocHostPtr) != cudaSuccess)
+ return 1;
+}
diff --git a/offload/test/offloading/CUDA/launch_tu.cu b/offload/test/offloading/CUDA/launch_tu.cu
index a46472b514a6c..8b92194ba435e 100644
--- a/offload/test/offloading/CUDA/launch_tu.cu
+++ b/offload/test/offloading/CUDA/launch_tu.cu
@@ -1,31 +1,30 @@
// clang-format off
// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t.launch_tu.o -c
// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda %S/kernel_tu.cu.inc -o %t.kernel_tu.o -c
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %t.launch_tu.o %t.kernel_tu.o -o %t
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.launch_tu.o %t.kernel_tu.o -o %t
// RUN: %t | %fcheck-generic
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
extern __global__ void square(int *A);
int main(int argc, char **argv) {
int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- *Ptr = 7;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
+ int *Ptr;
+ cudaMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
square<<<1, 1>>>(Ptr);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- llvm_omp_target_free_shared(Ptr, DevNo);
+ int I;
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
}
diff --git a/offload/test/offloading/CUDA/memcpy_kinds.cu b/offload/test/offloading/CUDA/memcpy_kinds.cu
new file mode 100644
index 0000000000000..a4288ee51ee3b
--- /dev/null
+++ b/offload/test/offloading/CUDA/memcpy_kinds.cu
@@ -0,0 +1,51 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int HostSrc = 11;
+ int HostDst = 0;
+ if (cudaMemcpy(&HostDst, &HostSrc, sizeof(int), cudaMemcpyHostToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("host to host: %d\n", HostDst);
+ // CHECK: host to host: 11
+
+ int *DevSrc = nullptr;
+ int *DevDst = nullptr;
+ int Result = 0;
+ if (cudaMalloc(&DevSrc, sizeof(int)) != cudaSuccess)
+ return 1;
+ if (cudaMalloc(&DevDst, sizeof(int)) != cudaSuccess)
+ return 1;
+
+ HostSrc = 42;
+ if (cudaMemcpy(DevSrc, &HostSrc, sizeof(int), cudaMemcpyHostToDevice) !=
+ cudaSuccess)
+ return 1;
+ if (cudaMemcpy(DevDst, DevSrc, sizeof(int), cudaMemcpyDeviceToDevice) !=
+ cudaSuccess)
+ return 1;
+ if (cudaMemcpy(&Result, DevDst, sizeof(int), cudaMemcpyDeviceToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("device to device: %d\n", Result);
+ // CHECK: device to device: 42
+
+ cudaFree(DevSrc);
+ cudaFree(DevDst);
+}
diff --git a/offload/test/offloading/CUDA/stream_api.cu b/offload/test/offloading/CUDA/stream_api.cu
new file mode 100644
index 0000000000000..7202751f8207e
--- /dev/null
+++ b/offload/test/offloading/CUDA/stream_api.cu
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void setValue(int *Out) { *Out = 42; }
+
+int main(int argc, char **argv) {
+ cudaStream_t Stream = nullptr;
+ if (cudaStreamCreate(&Stream) != cudaSuccess)
+ return 1;
+
+ printf("stream created: %d\n", Stream != nullptr);
+ // CHECK: stream created: 1
+
+ int *DevPtr = nullptr;
+ int Result = 0;
+ if (cudaMalloc(&DevPtr, sizeof(int)) != cudaSuccess)
+ return 1;
+
+ setValue<<<1, 1, 0, Stream>>>(DevPtr);
+
+ if (cudaStreamSynchronize(Stream) != cudaSuccess)
+ return 1;
+ if (cudaMemcpy(&Result, DevPtr, sizeof(int), cudaMemcpyDeviceToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("stream result: %d\n", Result);
+ // CHECK: stream result: 42
+
+ if (cudaStreamDestroy(Stream) != cudaSuccess)
+ return 1;
+ cudaFree(DevPtr);
+}
diff --git a/offload/test/offloading/CUDA/syncthreads.cu b/offload/test/offloading/CUDA/syncthreads.cu
new file mode 100644
index 0000000000000..0c6048c32f824
--- /dev/null
+++ b/offload/test/offloading/CUDA/syncthreads.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void reduceBlock(int *Out) {
+ __shared__ int Scratch[64];
+ int Tid = threadIdx.x;
+ Scratch[Tid] = Tid;
+ __syncthreads();
+
+ if (Tid == 0) {
+ int Sum = 0;
+ for (int I = 0; I < 64; ++I)
+ Sum += Scratch[I];
+ Out[0] = Sum;
+ }
+}
+
+int main(int argc, char **argv) {
+ int *DevPtr;
+ int Result = 0;
+ cudaMalloc(&DevPtr, sizeof(int));
+ reduceBlock<<<1, 64>>>(DevPtr);
+ cudaMemcpy(&Result, DevPtr, sizeof(int), cudaMemcpyDeviceToHost);
+
+ printf("sum: %i\n", Result);
+ // CHECK: sum: 2016
+}
diff --git a/offload/test/offloading/CUDA/thread_and_block_id.cu b/offload/test/offloading/CUDA/thread_and_block_id.cu
new file mode 100644
index 0000000000000..30b87659d2eed
--- /dev/null
+++ b/offload/test/offloading/CUDA/thread_and_block_id.cu
@@ -0,0 +1,44 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+
+#include <stdio.h>
+#include <stdlib.h>
+
+__global__ void fill(int *A) {
+ int tid = threadIdx.x + blockDim.x * blockIdx.x;
+ A[tid] = 42;
+}
+
+int main(int argc, char **argv) {
+ int NThreads = 128;
+ int NBlocks = 512;
+ int Size = sizeof(int) * NThreads * NBlocks;
+ int *Ptr = (int *)calloc(1, Size);
+ int *DevPtr;
+ cudaMalloc(&DevPtr, Size);
+ cudaMemcpy(DevPtr, Ptr, Size, cudaMemcpyHostToDevice);
+ printf("DevPtr %p\n", DevPtr);
+ // CHECK: DevPtr [[DevPtr:0x.*]]
+ fill<<<NBlocks, NThreads>>>(DevPtr);
+ cudaMemcpy(Ptr, DevPtr, Size, cudaMemcpyDeviceToHost);
+
+ for (int I = 0; I < NBlocks * NThreads; ++I) {
+ if (Ptr[I] == 42)
+ continue;
+ printf("Error at %i: %i vs %i\n", I, Ptr[I], 42);
+ return 1;
+ }
+ return 0;
+}
diff --git a/offload/test/offloading/HIP/basic_launch.hip b/offload/test/offloading/HIP/basic_launch.hip
new file mode 100644
index 0000000000000..bd2f2a6078671
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch.hip
@@ -0,0 +1,30 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *A) { *A = 42; }
+
+int main(int argc, char **argv) {
+ int *Ptr;
+ hipMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
+ square<<<1, 1>>>(Ptr);
+ int I = 0;
+ hipDeviceSynchronize();
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
new file mode 100644
index 0000000000000..344b98b1636f1
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
@@ -0,0 +1,31 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *A) {
+ __scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
+}
+
+int main(int argc, char **argv) {
+ int DevNo = 0;
+ int *Ptr, I;
+ hipMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
+ square<<<7, 6>>>(Ptr);
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/basic_launch_multi_arg.hip b/offload/test/offloading/HIP/basic_launch_multi_arg.hip
new file mode 100644
index 0000000000000..6e599d6704598
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch_multi_arg.hip
@@ -0,0 +1,39 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// REQUIRES: gpu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *Dst, short Q, int *Src, short P) {
+ *Dst = (Src[0] + Src[1]) * (Q + P);
+ Src[0] = Q;
+ Src[1] = P;
+}
+
+int main(int argc, char **argv) {
+ int DevNo = 0;
+ int *Src, *Ptr;
+ hipMalloc(&Ptr, 4);
+ hipMalloc(&Src, 8);
+
+ int I = 7;
+ int HostSrc[2] = {-2,8};
+ hipMemcpy(Ptr, &I, sizeof(int), hipMemcpyHostToDevice);
+ hipMemcpy(Src, &HostSrc[0], 2*sizeof(int), hipMemcpyHostToDevice);
+ square<<<1, 1>>>(Ptr, 3, Src, 4);
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ hipMemcpy(&HostSrc[0], Src, 2 * sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+ printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
+ // CHECK: Src: 3, 4
+}
diff --git a/offload/test/offloading/HIP/device_api.hip b/offload/test/offloading/HIP/device_api.hip
new file mode 100644
index 0000000000000..031e3703e66c1
--- /dev/null
+++ b/offload/test/offloading/HIP/device_api.hip
@@ -0,0 +1,45 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int Count = 0;
+ if (hipGetDeviceCount(&Count) != hipSuccess)
+ return 1;
+
+ printf("device count: %d\n", Count);
+ // CHECK: device count: {{[1-9][0-9]*}}
+
+ int Device = -1;
+ if (hipGetDevice(&Device) != hipSuccess)
+ return 1;
+
+ printf("device: %d\n", Device);
+ // CHECK: device: {{[0-9]+}}
+
+ if (hipSetDevice(Device) != hipSuccess)
+ return 1;
+
+ int After = -1;
+ if (hipGetDevice(&After) != hipSuccess)
+ return 1;
+
+ printf("device after set: %d\n", After);
+ // CHECK: device after set: {{[0-9]+}}
+
+ hipError_t Err = hipSetDevice(-1);
+ printf("set invalid device: %u\n", Err);
+ // CHECK: set invalid device: 1
+}
diff --git a/offload/test/offloading/HIP/device_properties.hip b/offload/test/offloading/HIP/device_properties.hip
new file mode 100644
index 0000000000000..1a9b9a70f8ea9
--- /dev/null
+++ b/offload/test/offloading/HIP/device_properties.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ hipDeviceProp_t Prop = {};
+ hipError_t Err = hipGetDeviceProperties(&Prop, 0);
+ if (Err != hipSuccess) {
+ printf("hipGetDeviceProperties failed: %u\n", Err);
+ return 1;
+ }
+
+ printf("Device name: %s\n", Prop.name);
+ // CHECK: Device name:
+ printf("Total global memory: %zu\n", Prop.totalGlobalMem);
+ // CHECK: Total global memory:
+ printf("Multiprocessors: %i\n", Prop.multiProcessorCount);
+ // CHECK: Multiprocessors:
+ printf("Warp size: %i\n", Prop.warpSize);
+ // CHECK: Warp size:
+
+ if (!Prop.name[0] || !Prop.totalGlobalMem || !Prop.multiProcessorCount ||
+ !Prop.warpSize)
+ return 1;
+
+ printf("Device properties are populated.\n");
+ // CHECK: Device properties are populated.
+}
diff --git a/offload/test/offloading/HIP/host_alloc.hip b/offload/test/offloading/HIP/host_alloc.hip
new file mode 100644
index 0000000000000..38a83ee64f333
--- /dev/null
+++ b/offload/test/offloading/HIP/host_alloc.hip
@@ -0,0 +1,48 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void add(int *Ptr, int Value) { *Ptr += Value; }
+
+int main(int argc, char **argv) {
+ int *HostAllocPtr = nullptr;
+ if (hipHostAlloc(&HostAllocPtr, sizeof(int), hipHostAllocDefault) !=
+ hipSuccess)
+ return 1;
+
+ *HostAllocPtr = 17;
+ add<<<1, 1>>>(HostAllocPtr, 5);
+ if (hipDeviceSynchronize() != hipSuccess)
+ return 1;
+ printf("hipHostAlloc value: %d\n", *HostAllocPtr);
+ // CHECK: hipHostAlloc value: 22
+
+ if (hipFreeHost(HostAllocPtr) != hipSuccess)
+ return 1;
+
+ int *MallocHostPtr = nullptr;
+ if (hipMallocHost(&MallocHostPtr, sizeof(int)) != hipSuccess)
+ return 1;
+
+ *MallocHostPtr = 23;
+ add<<<1, 1>>>(MallocHostPtr, 7);
+ if (hipDeviceSynchronize() != hipSuccess)
+ return 1;
+ printf("hipMallocHost value: %d\n", *MallocHostPtr);
+ // CHECK: hipMallocHost value: 30
+
+ if (hipFreeHost(MallocHostPtr) != hipSuccess)
+ return 1;
+}
diff --git a/offload/test/offloading/HIP/kernel_tu.hip.inc b/offload/test/offloading/HIP/kernel_tu.hip.inc
new file mode 100644
index 0000000000000..d7d28a109dfc5
--- /dev/null
+++ b/offload/test/offloading/HIP/kernel_tu.hip.inc
@@ -0,0 +1 @@
+__global__ void square(int *A) { *A = 42; }
diff --git a/offload/test/offloading/HIP/launch_tu.hip b/offload/test/offloading/HIP/launch_tu.hip
new file mode 100644
index 0000000000000..03073029ca211
--- /dev/null
+++ b/offload/test/offloading/HIP/launch_tu.hip
@@ -0,0 +1,30 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t.launch_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip %S/kernel_tu.hip.inc -o %t.kernel_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.launch_tu.o %t.kernel_tu.o -o %t
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+extern __global__ void square(int *A);
+
+int main(int argc, char **argv) {
+ int DevNo = 0;
+ int *Ptr;
+ hipMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
+ square<<<1, 1>>>(Ptr);
+ int I;
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/memcpy_kinds.hip b/offload/test/offloading/HIP/memcpy_kinds.hip
new file mode 100644
index 0000000000000..6755a55aa0794
--- /dev/null
+++ b/offload/test/offloading/HIP/memcpy_kinds.hip
@@ -0,0 +1,51 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int HostSrc = 11;
+ int HostDst = 0;
+ if (hipMemcpy(&HostDst, &HostSrc, sizeof(int), hipMemcpyHostToHost) !=
+ hipSuccess)
+ return 1;
+
+ printf("host to host: %d\n", HostDst);
+ // CHECK: host to host: 11
+
+ int *DevSrc = nullptr;
+ int *DevDst = nullptr;
+ int Result = 0;
+ if (hipMalloc(&DevSrc, sizeof(int)) != hipSuccess)
+ return 1;
+ if (hipMalloc(&DevDst, sizeof(int)) != hipSuccess)
+ return 1;
+
+ HostSrc = 42;
+ if (hipMemcpy(DevSrc, &HostSrc, sizeof(int), hipMemcpyHostToDevice) !=
+ hipSuccess)
+ return 1;
+ if (hipMemcpy(DevDst, DevSrc, sizeof(int), hipMemcpyDeviceToDevice) !=
+ hipSuccess)
+ return 1;
+ if (hipMemcpy(&Result, DevDst, sizeof(int), hipMemcpyDeviceToHost) !=
+ hipSuccess)
+ return 1;
+
+ printf("device to device: %d\n", Result);
+ // CHECK: device to device: 42
+
+ hipFree(DevSrc);
+ hipFree(DevDst);
+}
diff --git a/offload/test/offloading/HIP/stream_api.hip b/offload/test/offloading/HIP/stream_api.hip
new file mode 100644
index 0000000000000..c0e2699822814
--- /dev/null
+++ b/offload/test/offloading/HIP/stream_api.hip
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void setValue(int *Out) { *Out = 42; }
+
+int main(int argc, char **argv) {
+ hipStream_t Stream = nullptr;
+ if (hipStreamCreate(&Stream) != hipSuccess)
+ return 1;
+
+ printf("stream created: %d\n", Stream != nullptr);
+ // CHECK: stream created: 1
+
+ int *DevPtr = nullptr;
+ int Result = 0;
+ if (hipMalloc(&DevPtr, sizeof(int)) != hipSuccess)
+ return 1;
+
+ setValue<<<1, 1, 0, Stream>>>(DevPtr);
+
+ if (hipStreamSynchronize(Stream) != hipSuccess)
+ return 1;
+ if (hipMemcpy(&Result, DevPtr, sizeof(int), hipMemcpyDeviceToHost) !=
+ hipSuccess)
+ return 1;
+
+ printf("stream result: %d\n", Result);
+ // CHECK: stream result: 42
+
+ if (hipStreamDestroy(Stream) != hipSuccess)
+ return 1;
+ hipFree(DevPtr);
+}
diff --git a/offload/test/offloading/HIP/syncthreads.hip b/offload/test/offloading/HIP/syncthreads.hip
new file mode 100644
index 0000000000000..5962ab5468b86
--- /dev/null
+++ b/offload/test/offloading/HIP/syncthreads.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void reduceBlock(int *Out) {
+ __shared__ int Scratch[64];
+ int Tid = threadIdx.x;
+ Scratch[Tid] = Tid;
+ __syncthreads();
+
+ if (Tid == 0) {
+ int Sum = 0;
+ for (int I = 0; I < 64; ++I)
+ Sum += Scratch[I];
+ Out[0] = Sum;
+ }
+}
+
+int main(int argc, char **argv) {
+ int *DevPtr;
+ int Result = 0;
+ hipMalloc(&DevPtr, sizeof(int));
+ reduceBlock<<<1, 64>>>(DevPtr);
+ hipMemcpy(&Result, DevPtr, sizeof(int), hipMemcpyDeviceToHost);
+
+ printf("sum: %i\n", Result);
+ // CHECK: sum: 2016
+}
diff --git a/offload/test/offloading/HIP/thread_and_block_id.hip b/offload/test/offloading/HIP/thread_and_block_id.hip
new file mode 100644
index 0000000000000..c9c33c55d72fb
--- /dev/null
+++ b/offload/test/offloading/HIP/thread_and_block_id.hip
@@ -0,0 +1,44 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+
+#include <stdio.h>
+#include <stdlib.h>
+
+__global__ void fill(int *A) {
+ int tid = threadIdx.x + blockDim.x * blockIdx.x;
+ A[tid] = 42;
+}
+
+int main(int argc, char **argv) {
+ int NThreads = 128;
+ int NBlocks = 512;
+ int Size = sizeof(int) * NThreads * NBlocks;
+ int *Ptr = (int*)calloc(1, Size);
+ int *DevPtr;
+ hipMalloc(&DevPtr, Size);
+ hipMemcpy(DevPtr, Ptr, Size, hipMemcpyHostToDevice);
+ printf("DevPtr %p\n", DevPtr);
+ // CHECK: DevPtr [[DevPtr:0x.*]]
+ fill<<<NBlocks, NThreads>>>(DevPtr);
+ hipMemcpy(Ptr, DevPtr, Size, hipMemcpyDeviceToHost);
+
+ for (int I = 0; I < NBlocks * NThreads; ++I) {
+ if (Ptr[I] == 42)
+ continue;
+ printf("Error at %i: %i vs %i\n", I, Ptr[I], 42);
+ return 1;
+ }
+ return 0;
+}
>From 5156495e392d6a242c57395e320df56b89e56c4e Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Wed, 29 Jul 2026 15:12:50 -0700
Subject: [PATCH 2/4] add getErrorName + getErrorString + Test
---
.../include/kernel/DefineLanguageNames.inc | 4 ++
.../languages/include/kernel/LanguageErrors.h | 25 ++++++++
.../include/kernel/LanguageRuntime.h | 5 +-
.../include/kernel/UndefineLanguageNames.inc | 4 ++
offload/languages/kernel/CMakeLists.txt | 2 +
.../languages/kernel/include/LanguageUtils.h | 43 +++++++++++++
.../languages/kernel/src/LanguageErrors.cpp | 51 ++++++++++++++++
.../languages/kernel/src/LanguageRuntime.cpp | 30 +++------
offload/test/offloading/CUDA/device_api.cu | 2 +-
offload/test/offloading/CUDA/error_kinds.cu | 61 +++++++++++++++++++
offload/test/offloading/HIP/device_api.hip | 2 +-
offload/test/offloading/HIP/error_kinds.hip | 61 +++++++++++++++++++
12 files changed, 261 insertions(+), 29 deletions(-)
create mode 100644 offload/languages/include/kernel/LanguageErrors.h
create mode 100644 offload/languages/kernel/include/LanguageUtils.h
create mode 100644 offload/languages/kernel/src/LanguageErrors.cpp
create mode 100644 offload/test/offloading/CUDA/error_kinds.cu
create mode 100644 offload/test/offloading/HIP/error_kinds.hip
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index cfab6998c28e6..989d119ae6ed7 100644
--- a/offload/languages/include/kernel/DefineLanguageNames.inc
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -17,6 +17,10 @@
#define DeviceSynchronize COMBINE(LANGUAGE, DeviceSynchronize)
#define Success COMBINE(LANGUAGE, Success)
#define ErrorInvalidValue COMBINE(LANGUAGE, ErrorInvalidValue)
+#define ErrorInvalidDevice COMBINE(LANGUAGE, ErrorInvalidDevice)
+#define ErrorUnknown COMBINE(LANGUAGE, ErrorUnknown)
+#define GetErrorName COMBINE(LANGUAGE, GetErrorName)
+#define GetErrorString COMBINE(LANGUAGE, GetErrorString)
#define MemcpyKind COMBINE(LANGUAGE, MemcpyKind)
#define MemcpyHostToHost COMBINE(LANGUAGE, MemcpyHostToHost)
#define MemcpyHostToDevice COMBINE(LANGUAGE, MemcpyHostToDevice)
diff --git a/offload/languages/include/kernel/LanguageErrors.h b/offload/languages/include/kernel/LanguageErrors.h
new file mode 100644
index 0000000000000..0c2a83cf5ef5b
--- /dev/null
+++ b/offload/languages/include/kernel/LanguageErrors.h
@@ -0,0 +1,25 @@
+//===-- LanguageErrors.h - Kernel language error API declarations ---------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_ERRORS_H
+#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_ERRORS_H
+
+#include <cstdint>
+
+enum Error_t : uint32_t {
+ Success = 0,
+ ErrorInvalidValue = 1,
+ ErrorInvalidDevice = 2,
+ ErrorUnknown = 3,
+};
+
+const char *GetErrorName(Error_t Error);
+
+const char *GetErrorString(Error_t Error);
+
+#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_ERRORS_H
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index 9b0d9c98e8e5f..3fbc4c0af3e2c 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -16,10 +16,7 @@
#include <cstdio>
#include <cstdlib>
-enum Error_t : uint32_t {
- Success = 0,
- ErrorInvalidValue = 1,
-};
+#include "LanguageErrors.h"
struct DeviceProp_t {
char name[256];
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
index c1cc75719a18c..b2a9f2fd846ae 100644
--- a/offload/languages/include/kernel/UndefineLanguageNames.inc
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -14,6 +14,10 @@
#undef DeviceSynchronize
#undef Success
#undef ErrorInvalidValue
+#undef ErrorInvalidDevice
+#undef ErrorUnknown
+#undef GetErrorName
+#undef GetErrorString
#undef MemcpyKind
#undef MemcpyHostToHost
#undef MemcpyHostToDevice
diff --git a/offload/languages/kernel/CMakeLists.txt b/offload/languages/kernel/CMakeLists.txt
index 293df7becdc93..ff0bb81c7f017 100644
--- a/offload/languages/kernel/CMakeLists.txt
+++ b/offload/languages/kernel/CMakeLists.txt
@@ -13,6 +13,7 @@ set(LLVM_OFFLOAD_KERNEL_INCLUDE_DIRS
function(add_llvm_offload_kernel_language_runtime_objects target language)
add_library(${target} OBJECT
src/LanguageRuntime.cpp
+ src/LanguageErrors.cpp
)
add_dependencies(${target} OffloadAPI)
target_include_directories(${target} PRIVATE ${LLVM_OFFLOAD_KERNEL_INCLUDE_DIRS})
@@ -76,6 +77,7 @@ install(TARGETS LLVMOffloadKernel
install(FILES
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/DefineLanguageNames.inc
+ ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageErrors.h
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageRuntime.h
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/UndefineLanguageNames.inc
${CMAKE_CURRENT_SOURCE_DIR}/include/LanguageLaunch.h
diff --git a/offload/languages/kernel/include/LanguageUtils.h b/offload/languages/kernel/include/LanguageUtils.h
new file mode 100644
index 0000000000000..3b937dce16356
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageUtils.h
@@ -0,0 +1,43 @@
+//===-- LanguageUtils.h - Kernel Language utility functions ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_UTILS_H
+#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_UTILS_H
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+#include "OffloadAPI.h"
+
+/// Convert an ol_result_t to the active language's Error_t.
+static inline Error_t convertResult(ol_result_t Result) {
+ if (Result == OL_SUCCESS)
+ return Success;
+ switch (Result->Code) {
+ case OL_ERRC_INVALID_VALUE:
+ case OL_ERRC_INVALID_ARGUMENT:
+ case OL_ERRC_INVALID_NULL_POINTER:
+ return ErrorInvalidValue;
+ case OL_ERRC_INVALID_DEVICE:
+ return ErrorInvalidDevice;
+ default:
+ return ErrorUnknown;
+ }
+}
+
+/// Convert a Stream_t to an ol_queue_handle_t.
+static inline Error_t getQueueFromStream(Stream_t Stream,
+ ol_queue_handle_t *Queue) {
+ if (!Stream)
+ return ErrorInvalidValue;
+ *Queue = reinterpret_cast<ol_queue_handle_t>(Stream);
+ return Success;
+}
+
+#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_UTILS_H
diff --git a/offload/languages/kernel/src/LanguageErrors.cpp b/offload/languages/kernel/src/LanguageErrors.cpp
new file mode 100644
index 0000000000000..f725e5f029d28
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageErrors.cpp
@@ -0,0 +1,51 @@
+//===-- LanguageErrors.cpp - Kernel language error API implementation -----===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+// Rename the generic error API before declaring or defining language symbols.
+// clang-format off
+#include "DefineLanguageNames.inc"
+#include "LanguageErrors.h"
+// clang-format on
+
+const char *GetErrorName(Error_t Error) {
+ switch (Error) {
+#define LLVM_OFFLOAD_STRINGIFY_IMPL(NAME) #NAME
+#define LLVM_OFFLOAD_STRINGIFY(NAME) LLVM_OFFLOAD_STRINGIFY_IMPL(NAME)
+#define LLVM_OFFLOAD_ERR_STR(NAME) \
+ case NAME: \
+ return LLVM_OFFLOAD_STRINGIFY(NAME);
+ LLVM_OFFLOAD_ERR_STR(Success)
+ LLVM_OFFLOAD_ERR_STR(ErrorInvalidValue)
+ LLVM_OFFLOAD_ERR_STR(ErrorInvalidDevice)
+#undef LLVM_OFFLOAD_ERR_STR
+#undef LLVM_OFFLOAD_STRINGIFY
+#undef LLVM_OFFLOAD_STRINGIFY_IMPL
+ default:
+ return "Unrecognized error";
+ };
+}
+
+const char *GetErrorString(Error_t Error) {
+ switch (Error) {
+ case Success:
+ return "No error";
+ case ErrorInvalidValue:
+ return "Invalid argument value";
+ case ErrorInvalidDevice:
+ return "Invalid device number";
+ case ErrorUnknown:
+ return "Unknown error";
+ }
+ return "Unrecognized error";
+}
+
+#include "UndefineLanguageNames.inc"
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 9c599c68436b0..0bd82bf0bb243 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -16,6 +16,7 @@
#include "LanguageRuntime.h"
// clang-format on
+#include "LanguageUtils.h"
#include "State.h"
#include "Types.h"
@@ -32,17 +33,6 @@
using RuntimeState = llvm::offload::StateTy;
using ThreadState = llvm::offload::ThreadStateTy;
-static Error_t convertResult(ol_result_t Result) {
- if (Result == OL_SUCCESS)
- return Success;
- switch (Result->Code) {
- case OL_ERRC_INVALID_VALUE:
- return ErrorInvalidValue;
- default:
- return ErrorInvalidValue;
- }
-}
-
Error_t Malloc(void **DevPtr, size_t Size) {
ol_device_handle_t Device = ThreadState::getDefaultDevice();
ol_result_t Result = olMemAlloc(Device, OL_ALLOC_TYPE_DEVICE, Size, DevPtr);
@@ -105,7 +95,7 @@ Error_t DeviceSynchronize() {
Error_t GetDevice(int *DeviceNo) {
ol_device_handle_t Device = ThreadState::getDevice(DeviceNo);
if (!Device)
- return ErrorInvalidValue;
+ return ErrorInvalidDevice;
return Success;
}
@@ -117,7 +107,7 @@ Error_t GetDeviceCount(int *Count) {
Error_t SetDevice(int DeviceNo) {
ol_device_handle_t Device = ThreadState::setDefaultDevice(DeviceNo);
if (!Device)
- return ErrorInvalidValue;
+ return ErrorInvalidDevice;
assert(Device == ThreadState::getDefaultDevice() &&
"Set Device is not Default Device");
return Success;
@@ -154,19 +144,13 @@ Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
return Success;
}
-static Error_t getQueueFromStream(Stream_t Stream, ol_queue_handle_t *Queue) {
- if (!Stream)
- return ErrorInvalidValue;
- *Queue = reinterpret_cast<ol_queue_handle_t>(Stream);
- return Success;
-}
-
Error_t StreamCreate(Stream_t *Stream) {
ol_queue_handle_t Queue;
- olCreateQueue(RuntimeState::getContext(), ThreadState::getDefaultDevice(),
+ ol_result_t Result = olCreateQueue(RuntimeState::getContext(), ThreadState::getDefaultDevice(),
&Queue);
- *Stream = reinterpret_cast<Stream_t>(Queue);
- return Success;
+ if (Result == OL_SUCCESS)
+ *Stream = reinterpret_cast<Stream_t>(Queue);
+ return convertResult(Result);
}
Error_t StreamDestroy(Stream_t Stream) {
diff --git a/offload/test/offloading/CUDA/device_api.cu b/offload/test/offloading/CUDA/device_api.cu
index 184167f2d17e4..af2b046eee397 100644
--- a/offload/test/offloading/CUDA/device_api.cu
+++ b/offload/test/offloading/CUDA/device_api.cu
@@ -41,5 +41,5 @@ int main(int argc, char **argv) {
cudaError_t Err = cudaSetDevice(-1);
printf("set invalid device: %u\n", Err);
- // CHECK: set invalid device: 1
+ // CHECK: set invalid device: 2
}
diff --git a/offload/test/offloading/CUDA/error_kinds.cu b/offload/test/offloading/CUDA/error_kinds.cu
new file mode 100644
index 0000000000000..5fa3035fe52f0
--- /dev/null
+++ b/offload/test/offloading/CUDA/error_kinds.cu
@@ -0,0 +1,61 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+static void print_error(const char *Label, cudaError_t Error) {
+ printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
+ printf("%s name: %s\n", Label, cudaGetErrorName(Error));
+ printf("%s string: %s\n", Label, cudaGetErrorString(Error));
+}
+
+int main() {
+ print_error("success", cudaSuccess);
+ // CHECK: success value: 0
+ // CHECK: success name: cudaSuccess
+ // CHECK: success string: No error
+
+ print_error("invalid value", cudaErrorInvalidValue);
+ // CHECK: invalid value value: 1
+ // CHECK: invalid value name: cudaErrorInvalidValue
+ // CHECK: invalid value string: Invalid argument value
+
+ print_error("invalid device", cudaErrorInvalidDevice);
+ // CHECK: invalid device value: 2
+ // CHECK: invalid device name: cudaErrorInvalidDevice
+ // CHECK: invalid device string: Invalid device number
+
+ print_error("unknown", cudaErrorUnknown);
+ // CHECK: unknown value: 3
+ // CHECK: unknown name: Unrecognized error
+ // CHECK: unknown string: Unknown error
+
+ cudaError_t Unrecognized = static_cast<cudaError_t>(999);
+ print_error("unrecognized", Unrecognized);
+ // CHECK: unrecognized value: 999
+ // CHECK: unrecognized name: Unrecognized error
+ // CHECK: unrecognized string: Unrecognized error
+
+ print_error("set invalid device", cudaSetDevice(-1));
+ // CHECK: set invalid device value: 2
+ // CHECK: set invalid device name: cudaErrorInvalidDevice
+ // CHECK: set invalid device string: Invalid device number
+
+ print_error("null stream destroy", cudaStreamDestroy(nullptr));
+ // CHECK: null stream destroy value: 1
+ // CHECK: null stream destroy name: cudaErrorInvalidValue
+ // CHECK: null stream destroy string: Invalid argument value
+
+ return 0;
+}
diff --git a/offload/test/offloading/HIP/device_api.hip b/offload/test/offloading/HIP/device_api.hip
index 031e3703e66c1..5fb66e6e45eeb 100644
--- a/offload/test/offloading/HIP/device_api.hip
+++ b/offload/test/offloading/HIP/device_api.hip
@@ -41,5 +41,5 @@ int main(int argc, char **argv) {
hipError_t Err = hipSetDevice(-1);
printf("set invalid device: %u\n", Err);
- // CHECK: set invalid device: 1
+ // CHECK: set invalid device: 2
}
diff --git a/offload/test/offloading/HIP/error_kinds.hip b/offload/test/offloading/HIP/error_kinds.hip
new file mode 100644
index 0000000000000..af760a4cd4399
--- /dev/null
+++ b/offload/test/offloading/HIP/error_kinds.hip
@@ -0,0 +1,61 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+static void print_error(const char *Label, hipError_t Error) {
+ printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
+ printf("%s name: %s\n", Label, hipGetErrorName(Error));
+ printf("%s string: %s\n", Label, hipGetErrorString(Error));
+}
+
+int main() {
+ print_error("success", hipSuccess);
+ // CHECK: success value: 0
+ // CHECK: success name: hipSuccess
+ // CHECK: success string: No error
+
+ print_error("invalid value", hipErrorInvalidValue);
+ // CHECK: invalid value value: 1
+ // CHECK: invalid value name: hipErrorInvalidValue
+ // CHECK: invalid value string: Invalid argument value
+
+ print_error("invalid device", hipErrorInvalidDevice);
+ // CHECK: invalid device value: 2
+ // CHECK: invalid device name: hipErrorInvalidDevice
+ // CHECK: invalid device string: Invalid device number
+
+ print_error("unknown", hipErrorUnknown);
+ // CHECK: unknown value: 3
+ // CHECK: unknown name: Unrecognized error
+ // CHECK: unknown string: Unknown error
+
+ hipError_t Unrecognized = static_cast<hipError_t>(999);
+ print_error("unrecognized", Unrecognized);
+ // CHECK: unrecognized value: 999
+ // CHECK: unrecognized name: Unrecognized error
+ // CHECK: unrecognized string: Unrecognized error
+
+ print_error("set invalid device", hipSetDevice(-1));
+ // CHECK: set invalid device value: 2
+ // CHECK: set invalid device name: hipErrorInvalidDevice
+ // CHECK: set invalid device string: Invalid device number
+
+ print_error("null stream destroy", hipStreamDestroy(nullptr));
+ // CHECK: null stream destroy value: 1
+ // CHECK: null stream destroy name: hipErrorInvalidValue
+ // CHECK: null stream destroy string: Invalid argument value
+
+ return 0;
+}
>From 8693843f5aa26d894844336c91eb9211a791e977 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Fri, 31 Jul 2026 15:22:08 -0700
Subject: [PATCH 3/4] add new errors and threadlocal last error
remove guards, split cpp files
---
offload/languages/README.md | 3 +-
.../include/kernel/DefineLanguageNames.inc | 4 ++
.../languages/include/kernel/LanguageErrors.h | 6 +++
.../include/kernel/UndefineLanguageNames.inc | 4 ++
.../languages/kernel/include/LanguageUtils.h | 13 +++--
offload/languages/kernel/include/State.h | 16 +++++-
.../languages/kernel/src/LanguageErrors.cpp | 19 +++++++
.../languages/kernel/src/LanguageRuntime.cpp | 49 ++++++++++---------
offload/languages/kernel/src/State.cpp | 9 ++++
offload/test/offloading/CUDA/error_kinds.cu | 20 ++++++++
offload/test/offloading/HIP/error_kinds.hip | 20 ++++++++
11 files changed, 134 insertions(+), 29 deletions(-)
diff --git a/offload/languages/README.md b/offload/languages/README.md
index 2ef13f49321f0..ae79e22beb78d 100644
--- a/offload/languages/README.md
+++ b/offload/languages/README.md
@@ -13,7 +13,8 @@ code calls them, but they are for compiler-generated code rather than for users
to call directly.
The rest of the runtime is internal. This includes device lookup, queues,
-registered programs, registered kernels, and error conversion.
+registered programs, registered kernels, error conversion, and per-thread
+last-error storage.
Some source files are shared by CUDA and HIP. CMake compiles those files once
with `LANGUAGE=cuda` and once with `LANGUAGE=hip`. The language-name includes
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index 989d119ae6ed7..790fd3ad1a359 100644
--- a/offload/languages/include/kernel/DefineLanguageNames.inc
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -19,8 +19,12 @@
#define ErrorInvalidValue COMBINE(LANGUAGE, ErrorInvalidValue)
#define ErrorInvalidDevice COMBINE(LANGUAGE, ErrorInvalidDevice)
#define ErrorUnknown COMBINE(LANGUAGE, ErrorUnknown)
+#define ErrorInvalidResourceHandle COMBINE(LANGUAGE, ErrorInvalidResourceHandle)
+#define ErrorInvalidConfiguration COMBINE(LANGUAGE, ErrorInvalidConfiguration)
#define GetErrorName COMBINE(LANGUAGE, GetErrorName)
#define GetErrorString COMBINE(LANGUAGE, GetErrorString)
+#define GetLastError COMBINE(LANGUAGE, GetLastError)
+#define PeekAtLastError COMBINE(LANGUAGE, PeekAtLastError)
#define MemcpyKind COMBINE(LANGUAGE, MemcpyKind)
#define MemcpyHostToHost COMBINE(LANGUAGE, MemcpyHostToHost)
#define MemcpyHostToDevice COMBINE(LANGUAGE, MemcpyHostToDevice)
diff --git a/offload/languages/include/kernel/LanguageErrors.h b/offload/languages/include/kernel/LanguageErrors.h
index 0c2a83cf5ef5b..4e97363ab6582 100644
--- a/offload/languages/include/kernel/LanguageErrors.h
+++ b/offload/languages/include/kernel/LanguageErrors.h
@@ -16,10 +16,16 @@ enum Error_t : uint32_t {
ErrorInvalidValue = 1,
ErrorInvalidDevice = 2,
ErrorUnknown = 3,
+ ErrorInvalidResourceHandle = 4,
+ ErrorInvalidConfiguration = 5,
};
const char *GetErrorName(Error_t Error);
const char *GetErrorString(Error_t Error);
+Error_t GetLastError();
+
+Error_t PeekAtLastError();
+
#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_ERRORS_H
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
index b2a9f2fd846ae..cd0e6b33aafc6 100644
--- a/offload/languages/include/kernel/UndefineLanguageNames.inc
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -16,8 +16,12 @@
#undef ErrorInvalidValue
#undef ErrorInvalidDevice
#undef ErrorUnknown
+#undef ErrorInvalidResourceHandle
+#undef ErrorInvalidConfiguration
#undef GetErrorName
#undef GetErrorString
+#undef GetLastError
+#undef PeekAtLastError
#undef MemcpyKind
#undef MemcpyHostToHost
#undef MemcpyHostToDevice
diff --git a/offload/languages/kernel/include/LanguageUtils.h b/offload/languages/kernel/include/LanguageUtils.h
index 3b937dce16356..874643a81ccd6 100644
--- a/offload/languages/kernel/include/LanguageUtils.h
+++ b/offload/languages/kernel/include/LanguageUtils.h
@@ -9,11 +9,9 @@
#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_UTILS_H
#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_UTILS_H
-#ifndef LANGUAGE
-#error This file should be included, or used, with a LANGUAGE macro set.
-#endif
-
+#include "LanguageRuntime.h"
#include "OffloadAPI.h"
+#include "State.h"
/// Convert an ol_result_t to the active language's Error_t.
static inline Error_t convertResult(ol_result_t Result) {
@@ -31,6 +29,13 @@ static inline Error_t convertResult(ol_result_t Result) {
}
}
+/// Set the last error for the current thread and return it.
+static inline Error_t setLastError(Error_t Error) {
+ // TODO: find a more efficient way to set last error
+ return static_cast<Error_t>(
+ llvm::offload::ThreadStateTy::setLastError(Error));
+}
+
/// Convert a Stream_t to an ol_queue_handle_t.
static inline Error_t getQueueFromStream(Stream_t Stream,
ol_queue_handle_t *Queue) {
diff --git a/offload/languages/kernel/include/State.h b/offload/languages/kernel/include/State.h
index 8a981ddff8ed9..7315b2e710df5 100644
--- a/offload/languages/kernel/include/State.h
+++ b/offload/languages/kernel/include/State.h
@@ -16,6 +16,7 @@
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/Support/raw_ostream.h"
+#include <cstdint>
#define CHECK_FATAL(ResultExpr, ...) \
do { \
@@ -29,6 +30,12 @@
} \
} while (false)
+#define FATAL_UNIMPLEMENTED(...) \
+ do { \
+ llvm::errs() << __VA_ARGS__ << '\n'; \
+ abort(); \
+ } while (false)
+
namespace llvm {
namespace offload {
@@ -40,7 +47,7 @@ using KernelIDTy = const void *;
/// Per-thread state used by the language runtime entry points.
///
/// Tracks the current thread's default device, optional per-thread queue,
-/// and pending kernel launch configuration.
+/// last-error code, and pending kernel launch configuration.
struct ThreadStateTy {
~ThreadStateTy();
@@ -59,6 +66,12 @@ struct ThreadStateTy {
/// \returns the selected device, or nullptr if \p DeviceNo is invalid.
static ol_device_handle_t setDefaultDevice(int DeviceNo);
+ /// Return the last language-runtime error code for this thread.
+ static uint32_t getLastError();
+
+ /// Set the last language-runtime error code for this thread.
+ static uint32_t setLastError(uint32_t Error);
+
/// Return the pending kernel launch configuration for this thread.
static CallConfigurationTy &getCallConfiguration();
@@ -71,6 +84,7 @@ struct ThreadStateTy {
void createDefaultQueue(ol_device_handle_t Device);
int DefaultDevice = 0;
+ uint32_t LastError = 0;
ol_queue_handle_t DefaultQueue = nullptr;
CallConfigurationTy CC = {};
diff --git a/offload/languages/kernel/src/LanguageErrors.cpp b/offload/languages/kernel/src/LanguageErrors.cpp
index f725e5f029d28..823b2e34e2124 100644
--- a/offload/languages/kernel/src/LanguageErrors.cpp
+++ b/offload/languages/kernel/src/LanguageErrors.cpp
@@ -10,12 +10,15 @@
#error This file should be included, or used, with a LANGUAGE macro set.
#endif
+#include "State.h"
// Rename the generic error API before declaring or defining language symbols.
// clang-format off
#include "DefineLanguageNames.inc"
#include "LanguageErrors.h"
// clang-format on
+using ThreadState = llvm::offload::ThreadStateTy;
+
const char *GetErrorName(Error_t Error) {
switch (Error) {
#define LLVM_OFFLOAD_STRINGIFY_IMPL(NAME) #NAME
@@ -26,6 +29,8 @@ const char *GetErrorName(Error_t Error) {
LLVM_OFFLOAD_ERR_STR(Success)
LLVM_OFFLOAD_ERR_STR(ErrorInvalidValue)
LLVM_OFFLOAD_ERR_STR(ErrorInvalidDevice)
+ LLVM_OFFLOAD_ERR_STR(ErrorInvalidResourceHandle)
+ LLVM_OFFLOAD_ERR_STR(ErrorInvalidConfiguration)
#undef LLVM_OFFLOAD_ERR_STR
#undef LLVM_OFFLOAD_STRINGIFY
#undef LLVM_OFFLOAD_STRINGIFY_IMPL
@@ -44,8 +49,22 @@ const char *GetErrorString(Error_t Error) {
return "Invalid device number";
case ErrorUnknown:
return "Unknown error";
+ case ErrorInvalidResourceHandle:
+ return "Invalid resource handle";
+ case ErrorInvalidConfiguration:
+ return "Invalid configuration argument";
}
return "Unrecognized error";
}
+Error_t GetLastError() {
+ Error_t Error = static_cast<Error_t>(ThreadState::getLastError());
+ ThreadState::setLastError(Success);
+ return Error;
+}
+
+Error_t PeekAtLastError() {
+ return static_cast<Error_t>(ThreadState::getLastError());
+}
+
#include "UndefineLanguageNames.inc"
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 0bd82bf0bb243..549333251aa9c 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -13,6 +13,7 @@
// Rename the generic runtime API before declaring or defining language symbols.
// clang-format off
#include "DefineLanguageNames.inc"
+#include "LanguageErrors.h"
#include "LanguageRuntime.h"
// clang-format on
@@ -27,21 +28,22 @@
#include <cstdlib>
#include <cstring>
-#define STR(X) #X
-#define LANGUAGE_STR STR(LANGUAGE)
-
using RuntimeState = llvm::offload::StateTy;
using ThreadState = llvm::offload::ThreadStateTy;
+static inline Error_t convertAndSetLastError(ol_result_t Result) {
+ return setLastError(convertResult(Result));
+}
+
Error_t Malloc(void **DevPtr, size_t Size) {
ol_device_handle_t Device = ThreadState::getDefaultDevice();
ol_result_t Result = olMemAlloc(Device, OL_ALLOC_TYPE_DEVICE, Size, DevPtr);
- return convertResult(Result);
+ return convertAndSetLastError(Result);
}
Error_t Free(void *DevPtr) {
ol_result_t Result = olMemFree(DevPtr);
- return convertResult(Result);
+ return convertAndSetLastError(Result);
}
Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
@@ -75,13 +77,14 @@ Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
break;
}
case MemcpyDefault:
- fprintf(stderr, LANGUAGE_STR "MemcpyDefault is not implemented yet");
- abort();
+ FATAL_UNIMPLEMENTED("MemcpyDefault is not implemented yet");
};
- Result = olSyncQueue(Queue);
+ if (Result != OL_SUCCESS)
+ return convertAndSetLastError(Result);
- return convertResult(Result);
+ Result = olSyncQueue(Queue);
+ return convertAndSetLastError(Result);
}
Error_t DeviceSynchronize() {
@@ -89,34 +92,34 @@ Error_t DeviceSynchronize() {
// plugins.
ol_queue_handle_t Queue = ThreadState::getDefaultQueue();
ol_result_t Result = olSyncQueue(Queue);
- return convertResult(Result);
+ return convertAndSetLastError(Result);
}
Error_t GetDevice(int *DeviceNo) {
ol_device_handle_t Device = ThreadState::getDevice(DeviceNo);
if (!Device)
- return ErrorInvalidDevice;
- return Success;
+ return setLastError(ErrorInvalidDevice);
+ return setLastError(Success);
}
Error_t GetDeviceCount(int *Count) {
*Count = RuntimeState::getDeviceCount();
- return Success;
+ return setLastError(Success);
}
Error_t SetDevice(int DeviceNo) {
ol_device_handle_t Device = ThreadState::setDefaultDevice(DeviceNo);
if (!Device)
- return ErrorInvalidDevice;
+ return setLastError(ErrorInvalidDevice);
assert(Device == ThreadState::getDefaultDevice() &&
"Set Device is not Default Device");
- return Success;
+ return setLastError(Success);
}
Error_t HostAlloc(void **Ptr, size_t Size, unsigned int Flags) {
ol_device_handle_t Device = ThreadState::getDefaultDevice();
ol_result_t Result = olMemAllocHost(Device, Size, Ptr);
- return convertResult(Result);
+ return convertAndSetLastError(Result);
}
Error_t MallocHost(void **Ptr, size_t Size) {
@@ -125,7 +128,7 @@ Error_t MallocHost(void **Ptr, size_t Size) {
Error_t FreeHost(void *Ptr) {
ol_result_t Result = olMemFree(Ptr);
- return convertResult(Result);
+ return convertAndSetLastError(Result);
}
Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
@@ -141,7 +144,7 @@ Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
&DeviceProp->multiProcessorCount);
olGetDeviceInfo(Device, OL_DEVICE_INFO_NUM_LANES, sizeof(uint32_t),
&DeviceProp->warpSize);
- return Success;
+ return setLastError(Success);
}
Error_t StreamCreate(Stream_t *Stream) {
@@ -150,25 +153,25 @@ Error_t StreamCreate(Stream_t *Stream) {
&Queue);
if (Result == OL_SUCCESS)
*Stream = reinterpret_cast<Stream_t>(Queue);
- return convertResult(Result);
+ return convertAndSetLastError(Result);
}
Error_t StreamDestroy(Stream_t Stream) {
ol_queue_handle_t Queue;
Error_t Err = getQueueFromStream(Stream, &Queue);
if (Err != Success)
- return Err;
+ return setLastError(Err);
ol_result_t Result = olDestroyQueue(Queue);
- return convertResult(Result);
+ return convertAndSetLastError(Result);
}
Error_t StreamSynchronize(Stream_t Stream) {
ol_queue_handle_t Queue;
Error_t Err = getQueueFromStream(Stream, &Queue);
if (Err != Success)
- return Err;
+ return setLastError(Err);
ol_result_t Result = olSyncQueue(Queue);
- return convertResult(Result);
+ return convertAndSetLastError(Result);
}
#include "UndefineLanguageNames.inc"
diff --git a/offload/languages/kernel/src/State.cpp b/offload/languages/kernel/src/State.cpp
index da4de0d5da8d4..28af7c6e385d0 100644
--- a/offload/languages/kernel/src/State.cpp
+++ b/offload/languages/kernel/src/State.cpp
@@ -17,6 +17,7 @@
#include <atomic>
#include <cassert>
+#include <cstdint>
#include <cstdio>
#include <mutex>
@@ -122,6 +123,14 @@ ol_device_handle_t ThreadStateTy::getDevice(int *DeviceNo) {
return ThreadStateTy::getDefaultDevice();
}
+uint32_t ThreadStateTy::getLastError() {
+ return ThreadStateTy::get().LastError;
+}
+
+uint32_t ThreadStateTy::setLastError(uint32_t Error) {
+ return ThreadStateTy::get().LastError = Error;
+}
+
void ThreadStateTy::createDefaultQueue(ol_device_handle_t Device) {
if (DefaultQueue)
olDestroyQueue(DefaultQueue);
diff --git a/offload/test/offloading/CUDA/error_kinds.cu b/offload/test/offloading/CUDA/error_kinds.cu
index 5fa3035fe52f0..c5c2a1d83680d 100644
--- a/offload/test/offloading/CUDA/error_kinds.cu
+++ b/offload/test/offloading/CUDA/error_kinds.cu
@@ -41,6 +41,16 @@ int main() {
// CHECK: unknown name: Unrecognized error
// CHECK: unknown string: Unknown error
+ print_error("invalid resource handle", cudaErrorInvalidResourceHandle);
+ // CHECK: invalid resource handle value: 4
+ // CHECK: invalid resource handle name: cudaErrorInvalidResourceHandle
+ // CHECK: invalid resource handle string: Invalid resource handle
+
+ print_error("invalid configuration", cudaErrorInvalidConfiguration);
+ // CHECK: invalid configuration value: 5
+ // CHECK: invalid configuration name: cudaErrorInvalidConfiguration
+ // CHECK: invalid configuration string: Invalid configuration argument
+
cudaError_t Unrecognized = static_cast<cudaError_t>(999);
print_error("unrecognized", Unrecognized);
// CHECK: unrecognized value: 999
@@ -52,6 +62,16 @@ int main() {
// CHECK: set invalid device name: cudaErrorInvalidDevice
// CHECK: set invalid device string: Invalid device number
+ print_error("get last error", cudaGetLastError());
+ // CHECK: get last error value: 2
+ // CHECK: get last error name: cudaErrorInvalidDevice
+ // CHECK: get last error string: Invalid device number
+
+ print_error("cleared last error", cudaGetLastError());
+ // CHECK: cleared last error value: 0
+ // CHECK: cleared last error name: cudaSuccess
+ // CHECK: cleared last error string: No error
+
print_error("null stream destroy", cudaStreamDestroy(nullptr));
// CHECK: null stream destroy value: 1
// CHECK: null stream destroy name: cudaErrorInvalidValue
diff --git a/offload/test/offloading/HIP/error_kinds.hip b/offload/test/offloading/HIP/error_kinds.hip
index af760a4cd4399..2830d9de9b9dc 100644
--- a/offload/test/offloading/HIP/error_kinds.hip
+++ b/offload/test/offloading/HIP/error_kinds.hip
@@ -41,6 +41,16 @@ int main() {
// CHECK: unknown name: Unrecognized error
// CHECK: unknown string: Unknown error
+ print_error("invalid resource handle", hipErrorInvalidResourceHandle);
+ // CHECK: invalid resource handle value: 4
+ // CHECK: invalid resource handle name: hipErrorInvalidResourceHandle
+ // CHECK: invalid resource handle string: Invalid resource handle
+
+ print_error("invalid configuration", hipErrorInvalidConfiguration);
+ // CHECK: invalid configuration value: 5
+ // CHECK: invalid configuration name: hipErrorInvalidConfiguration
+ // CHECK: invalid configuration string: Invalid configuration argument
+
hipError_t Unrecognized = static_cast<hipError_t>(999);
print_error("unrecognized", Unrecognized);
// CHECK: unrecognized value: 999
@@ -52,6 +62,16 @@ int main() {
// CHECK: set invalid device name: hipErrorInvalidDevice
// CHECK: set invalid device string: Invalid device number
+ print_error("get last error", hipGetLastError());
+ // CHECK: get last error value: 2
+ // CHECK: get last error name: hipErrorInvalidDevice
+ // CHECK: get last error string: Invalid device number
+
+ print_error("cleared last error", hipGetLastError());
+ // CHECK: cleared last error value: 0
+ // CHECK: cleared last error name: hipSuccess
+ // CHECK: cleared last error string: No error
+
print_error("null stream destroy", hipStreamDestroy(nullptr));
// CHECK: null stream destroy value: 1
// CHECK: null stream destroy name: hipErrorInvalidValue
>From 7c00ee0ac9e82a76fe516385ebd67b0b3e0d980a Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Fri, 31 Jul 2026 15:23:38 -0700
Subject: [PATCH 4/4] add language-based KernelLaunches and preserve error
---
offload/languages/include/hip/hip_runtime.h | 15 ----
.../languages/kernel/include/LanguageUtils.h | 13 +++
.../languages/kernel/src/LanguageLaunch.cpp | 44 ++++++++--
.../languages/kernel/src/LanguageRuntime.cpp | 4 -
offload/test/offloading/CUDA/get_errs.cu | 82 +++++++++++++++++++
offload/test/offloading/HIP/get_errs.hip | 81 ++++++++++++++++++
6 files changed, 214 insertions(+), 25 deletions(-)
create mode 100644 offload/test/offloading/CUDA/get_errs.cu
create mode 100644 offload/test/offloading/HIP/get_errs.hip
diff --git a/offload/languages/include/hip/hip_runtime.h b/offload/languages/include/hip/hip_runtime.h
index b56e295c904e3..5c58f829daa84 100644
--- a/offload/languages/include/hip/hip_runtime.h
+++ b/offload/languages/include/hip/hip_runtime.h
@@ -43,19 +43,4 @@ template <class T> static inline hipError_t hipHostFree(T *Ptr) {
}
#endif
-#if defined(__AMDGPU__) || defined(__NVPTX__)
-#define HIP_KERNEL_NAME(...) __VA_ARGS__
-
-extern "C" hipError_t hipLaunchKernel(const char *Kernel, dim3 GridDim,
- dim3 BlockDim, void **KernelArgs,
- size_t DynamicSharedMem, void *Stream);
-
-template <typename... AT, typename FT = void (*)(AT...)>
-static inline void hipLaunchKernelGGL(FT Kernel, dim3 GridDim, dim3 BlockDim,
- size_t DynamicSharedMem, void *Stream,
- AT... KernelArgs) {
- Kernel<<<GridDim, BlockDim, DynamicSharedMem, Stream>>>(KernelArgs...);
-}
-#endif
-
#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_HIP_HIP_RUNTIME_H
diff --git a/offload/languages/kernel/include/LanguageUtils.h b/offload/languages/kernel/include/LanguageUtils.h
index 874643a81ccd6..9726d0599a9f4 100644
--- a/offload/languages/kernel/include/LanguageUtils.h
+++ b/offload/languages/kernel/include/LanguageUtils.h
@@ -22,6 +22,13 @@ static inline Error_t convertResult(ol_result_t Result) {
case OL_ERRC_INVALID_ARGUMENT:
case OL_ERRC_INVALID_NULL_POINTER:
return ErrorInvalidValue;
+ case OL_ERRC_INVALID_SIZE:
+ return ErrorInvalidConfiguration;
+ case OL_ERRC_INVALID_NULL_HANDLE:
+ case OL_ERRC_INVALID_QUEUE:
+ case OL_ERRC_INVALID_EVENT:
+ case OL_ERRC_INVALID_CONTEXT:
+ return ErrorInvalidResourceHandle;
case OL_ERRC_INVALID_DEVICE:
return ErrorInvalidDevice;
default:
@@ -36,6 +43,12 @@ static inline Error_t setLastError(Error_t Error) {
llvm::offload::ThreadStateTy::setLastError(Error));
}
+/// Convert an ol_result_t to the active language's Error_t and set it as the
+/// last error for the current thread.
+static inline Error_t convertAndSetLastError(ol_result_t Result) {
+ return setLastError(convertResult(Result));
+}
+
/// Convert a Stream_t to an ol_queue_handle_t.
static inline Error_t getQueueFromStream(Stream_t Stream,
ol_queue_handle_t *Queue) {
diff --git a/offload/languages/kernel/src/LanguageLaunch.cpp b/offload/languages/kernel/src/LanguageLaunch.cpp
index 61003b44330a7..abec0c8e1507b 100644
--- a/offload/languages/kernel/src/LanguageLaunch.cpp
+++ b/offload/languages/kernel/src/LanguageLaunch.cpp
@@ -7,30 +7,54 @@
//===----------------------------------------------------------------------===//
#include "LanguageLaunch.h"
+#include "LanguageUtils.h"
#include "State.h"
-
-#include <algorithm>
#include <cstdio>
using RuntimeState = llvm::offload::StateTy;
using ThreadState = llvm::offload::ThreadStateTy;
+static constexpr ol_error_struct_t InvalidKernelError = {
+ OL_ERRC_INVALID_NULL_HANDLE, "kernel is not registered"};
+
+static constexpr ol_error_struct_t InvalidDeviceError = {OL_ERRC_INVALID_DEVICE,
+ "invalid device"};
+
+static constexpr ol_error_struct_t InvalidArgumentError = {
+ OL_ERRC_INVALID_ARGUMENT, "invalid argument to kernel launch"};
+
+static constexpr ol_error_struct_t InvalidConfigurationError = {
+ OL_ERRC_INVALID_SIZE, "invalid kernel launch configuration"};
+
/// Internal kernel launch implementation
ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
dim3 BlockDim, void *KernelArgsPtr,
size_t DynamicSharedMem, void *Stream) {
ol_device_handle_t Device = ThreadState::getDefaultDevice();
ol_symbol_handle_t Kernel = RuntimeState::getKernel(KernelID);
+ if (!Device)
+ return &InvalidDeviceError;
+ if (!KernelID)
+ return &InvalidArgumentError;
+ if (!Kernel)
+ return &InvalidKernelError;
+
+ if (GridDim.x == 0 || GridDim.y == 0 || GridDim.z == 0 || BlockDim.x == 0 ||
+ BlockDim.y == 0 || BlockDim.z == 0)
+ return &InvalidConfigurationError;
+
+ if (!KernelArgsPtr)
+ return &InvalidArgumentError;
ol_kernel_launch_size_args_t LaunchSizeArgs;
LaunchSizeArgs.Dimensions =
1 + (GridDim.y > 1 || BlockDim.y > 1) + (GridDim.z > 1 || BlockDim.z > 1);
LaunchSizeArgs.NumGroups.x = GridDim.x;
- LaunchSizeArgs.NumGroups.y = std::max(GridDim.y, 1u);
- LaunchSizeArgs.NumGroups.z = std::max(GridDim.z, 1u);
+ LaunchSizeArgs.NumGroups.y = GridDim.y;
+ LaunchSizeArgs.NumGroups.z = GridDim.z;
LaunchSizeArgs.GroupSize.x = BlockDim.x;
- LaunchSizeArgs.GroupSize.y = std::max(BlockDim.y, 1u);
- LaunchSizeArgs.GroupSize.z = std::max(BlockDim.z, 1u);
+ LaunchSizeArgs.GroupSize.y = BlockDim.y;
+ LaunchSizeArgs.GroupSize.z = BlockDim.z;
LaunchSizeArgs.DynSharedMemory = DynamicSharedMem;
ol_queue_handle_t Queue = Stream ? reinterpret_cast<ol_queue_handle_t>(Stream)
@@ -42,6 +66,13 @@ ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
size_t *ArgSizes;
};
OffloadKernelArgs *OKA = static_cast<OffloadKernelArgs *>(KernelArgsPtr);
+ if ((!OKA->Args) != (!OKA->ArgSizes))
+ return &InvalidArgumentError;
+ if (OKA->NumArgs > 0 && !OKA->Args)
+ return &InvalidArgumentError;
+ for (size_t I = 0; I < OKA->NumArgs; ++I)
+ if (!OKA->Args[I] || OKA->ArgSizes[I] == 0)
+ return &InvalidArgumentError;
return olLaunchKernel(Queue, Device, Kernel, &LaunchSizeArgs,
/*Properties=*/nullptr, OKA->NumArgs, OKA->Args,
@@ -78,6 +109,7 @@ unsigned llvmLaunchKernel(const char *KernelID, dim3 GridDim, dim3 BlockDim,
void *Stream) {
ol_result_t Result = __llvmLaunchKernelImpl(
KernelID, GridDim, BlockDim, KernelArgsPtr, DynamicSharedMem, Stream);
+ convertAndSetLastError(Result);
return Result ? Result->Code : 0;
}
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 549333251aa9c..78bdf5276a46b 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -31,10 +31,6 @@
using RuntimeState = llvm::offload::StateTy;
using ThreadState = llvm::offload::ThreadStateTy;
-static inline Error_t convertAndSetLastError(ol_result_t Result) {
- return setLastError(convertResult(Result));
-}
-
Error_t Malloc(void **DevPtr, size_t Size) {
ol_device_handle_t Device = ThreadState::getDefaultDevice();
ol_result_t Result = olMemAlloc(Device, OL_ALLOC_TYPE_DEVICE, Size, DevPtr);
diff --git a/offload/test/offloading/CUDA/get_errs.cu b/offload/test/offloading/CUDA/get_errs.cu
new file mode 100644
index 0000000000000..3e000e7782eb5
--- /dev/null
+++ b/offload/test/offloading/CUDA/get_errs.cu
@@ -0,0 +1,82 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -pthread -std=c++17
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp -pthread -std=c++17
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <cstdio>
+#include <cuda_runtime.h>
+#include <mutex>
+#include <thread>
+
+static std::mutex PrintMutex;
+
+static void printError(int ThreadId, const char *Label, cudaError_t Error) {
+ std::lock_guard<std::mutex> Lock(PrintMutex);
+ printf("thread %d %s: %s\n", ThreadId, Label, cudaGetErrorName(Error));
+ std::fflush(stdout);
+}
+
+__global__ void errorKernel(float *d_out) {
+ int idx = blockIdx.x * blockDim.x + threadIdx.x;
+ d_out[idx] = idx * 0.5f;
+}
+
+void runTask(int thread_id) {
+ const int N = 1 << 20;
+ size_t bytes = N * sizeof(float);
+
+ float *d_data;
+ cudaMalloc(&d_data, bytes);
+ printError(thread_id, "cudaMalloc", cudaGetLastError());
+
+ thread_id == 1 ? errorKernel<<<4096, 256>>>(d_data)
+ : errorKernel<<<4096, 0>>>(d_data);
+ printError(thread_id, "kernel launch", cudaPeekAtLastError());
+
+ printError(thread_id, "kernel launch get", cudaGetLastError());
+
+ printError(thread_id, "kernel launch get again", cudaGetLastError());
+
+ cudaDeviceSynchronize();
+ cudaFree(d_data);
+}
+
+int main() {
+ printError(0, "initial", cudaPeekAtLastError());
+ // CHECK: thread 0 initial: cudaSuccess
+
+ std::thread t1(runTask, 1);
+ std::thread t2(runTask, 2);
+
+ t1.join();
+ t2.join();
+ // CHECK-DAG: thread 1 cudaMalloc: cudaSuccess
+ // CHECK-DAG: thread 2 cudaMalloc: cudaSuccess
+ // CHECK-DAG: thread 1 kernel launch: cudaSuccess
+ // CHECK-DAG: thread 2 kernel launch: cudaErrorInvalidConfiguration
+ // CHECK-DAG: thread 1 kernel launch get: cudaSuccess
+ // CHECK-DAG: thread 2 kernel launch get: cudaErrorInvalidConfiguration
+ // CHECK-DAG: thread 1 kernel launch get again: cudaSuccess
+ // CHECK-DAG: thread 2 kernel launch get again: cudaSuccess
+
+ std::thread t3(runTask, 3);
+ t3.join();
+ // CHECK: thread 3 cudaMalloc: cudaSuccess
+ // CHECK: thread 3 kernel launch: cudaErrorInvalidConfiguration
+ // CHECK: thread 3 kernel launch get: cudaErrorInvalidConfiguration
+ // CHECK: thread 3 kernel launch get again: cudaSuccess
+
+ printError(0, "joined", cudaGetLastError());
+ // CHECK: thread 0 joined: cudaSuccess
+
+ return 0;
+}
diff --git a/offload/test/offloading/HIP/get_errs.hip b/offload/test/offloading/HIP/get_errs.hip
new file mode 100644
index 0000000000000..c400c5e371cda
--- /dev/null
+++ b/offload/test/offloading/HIP/get_errs.hip
@@ -0,0 +1,81 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -pthread -std=c++17
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp -pthread -std=c++17
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <cstdio>
+#include <mutex>
+#include <thread>
+
+static std::mutex PrintMutex;
+
+static void printError(int ThreadId, const char *Label, hipError_t Error) {
+ std::lock_guard<std::mutex> Lock(PrintMutex);
+ printf("thread %d %s: %s\n", ThreadId, Label, hipGetErrorName(Error));
+ std::fflush(stdout);
+}
+
+__global__ void errorKernel(float *d_out) {
+ int idx = blockIdx.x * blockDim.x + threadIdx.x;
+ d_out[idx] = idx * 0.5f;
+}
+
+void runTask(int ThreadId) {
+ const int N = 1 << 20;
+ size_t Bytes = N * sizeof(float);
+
+ float *d_data;
+ hipMalloc(&d_data, Bytes);
+ printError(ThreadId, "hipMalloc", hipGetLastError());
+
+ ThreadId == 1 ? errorKernel<<<4096, 256>>>(d_data)
+ : errorKernel<<<4096, 0>>>(d_data);
+ printError(ThreadId, "kernel launch", hipPeekAtLastError());
+
+ printError(ThreadId, "kernel launch get", hipGetLastError());
+
+ printError(ThreadId, "kernel launch get again", hipGetLastError());
+
+ hipDeviceSynchronize();
+ hipFree(d_data);
+}
+
+int main() {
+ printError(0, "initial", hipPeekAtLastError());
+ // CHECK: thread 0 initial: hipSuccess
+
+ std::thread t1(runTask, 1);
+ std::thread t2(runTask, 2);
+
+ t1.join();
+ t2.join();
+ // CHECK-DAG: thread 1 hipMalloc: hipSuccess
+ // CHECK-DAG: thread 2 hipMalloc: hipSuccess
+ // CHECK-DAG: thread 1 kernel launch: hipSuccess
+ // CHECK-DAG: thread 2 kernel launch: hipErrorInvalidConfiguration
+ // CHECK-DAG: thread 1 kernel launch get: hipSuccess
+ // CHECK-DAG: thread 2 kernel launch get: hipErrorInvalidConfiguration
+ // CHECK-DAG: thread 1 kernel launch get again: hipSuccess
+ // CHECK-DAG: thread 2 kernel launch get again: hipSuccess
+
+ std::thread t3(runTask, 3);
+ t3.join();
+ // CHECK: thread 3 hipMalloc: hipSuccess
+ // CHECK: thread 3 kernel launch: hipErrorInvalidConfiguration
+ // CHECK: thread 3 kernel launch get: hipErrorInvalidConfiguration
+ // CHECK: thread 3 kernel launch get again: hipSuccess
+
+ printError(0, "joined", hipGetLastError());
+ // CHECK: thread 0 joined: hipSuccess
+
+ return 0;
+}
More information about the llvm-commits
mailing list