[clang] [llvm] [LLVMOffload] Get LastError for bad Kernels (PR #213390)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 31 17:47:38 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang-codegen
@llvm/pr-subscribers-backend-amdgpu
Author: Sophia Herrmann (jellytabby)
<details>
<summary>Changes</summary>
This PR moves Kernel launches into the language specific layer in order to retrieve errors generated by them. It also adds extra guards to the internal kernel launch. It also adds tests for #<!-- -->213389.
This PR depends on #<!-- -->213389 and its predecessors but because I do not have commit access I cannot stack the PR. For review only consider the LAST ONE commit.
---
Patch is 161.79 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/213390.diff
72 Files Affected:
- (modified) clang/include/clang/Driver/CommonArgs.h (+6)
- (modified) clang/lib/CodeGen/CGCUDANV.cpp (+35-28)
- (modified) clang/lib/Driver/Driver.cpp (+47-31)
- (modified) clang/lib/Driver/ToolChains/AMDGPU.cpp (+30-3)
- (modified) clang/lib/Driver/ToolChains/Clang.cpp (+40-16)
- (modified) clang/lib/Driver/ToolChains/CommonArgs.cpp (+13-6)
- (modified) clang/lib/Driver/ToolChains/Cuda.cpp (+57-13)
- (modified) clang/lib/Driver/ToolChains/Gnu.cpp (+1)
- (modified) clang/lib/Driver/ToolChains/Linux.cpp (+3-1)
- (modified) clang/lib/Headers/__clang_gpu_builtin_vars.h (+19)
- (modified) clang/test/CodeGenCUDA/Inputs/cuda.h (+6-2)
- (modified) clang/test/CodeGenCUDA/offload_via_llvm.cu (+29-31)
- (modified) clang/test/Driver/cuda-via-liboffload.cu (+7-8)
- (modified) clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c (+20-18)
- (modified) clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp (+12-1)
- (modified) llvm/lib/Frontend/Offloading/OffloadWrapper.cpp (+3-1)
- (modified) offload/CMakeLists.txt (+1)
- (added) offload/languages/CMakeLists.txt (+3)
- (added) offload/languages/cuda/CMakeLists.txt (+1)
- (added) offload/languages/cuda/src/cuda_runtime.cpp (+33)
- (added) offload/languages/hip/CMakeLists.txt (+1)
- (added) offload/languages/hip/src/hip_runtime.cpp (+33)
- (added) offload/languages/include/cuda/cuda_runtime.h (+22)
- (added) offload/languages/include/hip/hip_runtime.h (+55)
- (added) offload/languages/include/kernel/DefineLanguageNames.inc (+53)
- (added) offload/languages/include/kernel/LanguageRuntime.h (+207)
- (added) offload/languages/include/kernel/Types.h (+35)
- (added) offload/languages/include/kernel/UndefineLanguageNames.inc (+51)
- (added) offload/languages/kernel/CMakeLists.txt (+55)
- (added) offload/languages/kernel/exports (+11)
- (added) offload/languages/kernel/include/LanguageAliases.h (+40)
- (added) offload/languages/kernel/include/LanguageLaunch.h (+49)
- (added) offload/languages/kernel/include/LanguageRegistration.h (+41)
- (added) offload/languages/kernel/include/LanguageUtils.h (+21)
- (added) offload/languages/kernel/include/Registration.h (+17)
- (added) offload/languages/kernel/include/RuntimeAPI.h (+47)
- (added) offload/languages/kernel/include/State.h (+106)
- (added) offload/languages/kernel/src/LanguageCommon.cpp (+18)
- (added) offload/languages/kernel/src/LanguageLaunch.cpp (+129)
- (added) offload/languages/kernel/src/LanguageRegistration.cpp (+130)
- (added) offload/languages/kernel/src/LanguageRuntime.cpp (+171)
- (added) offload/languages/kernel/src/LanguageUtils.cpp (+90)
- (added) offload/languages/kernel/src/RuntimeAPI.cpp (+101)
- (added) offload/languages/kernel/src/State.cpp (+198)
- (modified) offload/test/lit.cfg (+9-1)
- (modified) offload/test/offloading/CUDA/basic_launch.cu (+12-13)
- (modified) offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu (+10-12)
- (modified) offload/test/offloading/CUDA/basic_launch_multi_arg.cu (+17-19)
- (added) offload/test/offloading/CUDA/device_api.cu (+45)
- (added) offload/test/offloading/CUDA/device_properties.cu (+40)
- (added) offload/test/offloading/CUDA/error_kinds.cu (+71)
- (added) offload/test/offloading/CUDA/get_errs.cu (+82)
- (added) offload/test/offloading/CUDA/host_alloc.cu (+40)
- (modified) offload/test/offloading/CUDA/launch_tu.cu (+12-13)
- (added) offload/test/offloading/CUDA/memcpy_kinds.cu (+51)
- (added) offload/test/offloading/CUDA/stream_api.cu (+46)
- (added) offload/test/offloading/CUDA/syncthreads.cu (+40)
- (added) offload/test/offloading/CUDA/thread_and_block_id.cu (+44)
- (added) offload/test/offloading/HIP/basic_launch.hip (+30)
- (added) offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip (+31)
- (added) offload/test/offloading/HIP/basic_launch_multi_arg.hip (+39)
- (added) offload/test/offloading/HIP/device_api.hip (+45)
- (added) offload/test/offloading/HIP/device_properties.hip (+40)
- (added) offload/test/offloading/HIP/error_kinds.hip (+71)
- (added) offload/test/offloading/HIP/get_errs.hip (+81)
- (added) offload/test/offloading/HIP/host_alloc.hip (+40)
- (added) offload/test/offloading/HIP/kernel_tu.hip.inc (+1)
- (added) offload/test/offloading/HIP/launch_tu.hip (+30)
- (added) offload/test/offloading/HIP/memcpy_kinds.hip (+51)
- (added) offload/test/offloading/HIP/stream_api.hip (+46)
- (added) offload/test/offloading/HIP/syncthreads.hip (+40)
- (added) offload/test/offloading/HIP/thread_and_block_id.hip (+44)
``````````diff
diff --git a/clang/include/clang/Driver/CommonArgs.h b/clang/include/clang/Driver/CommonArgs.h
index 8c861df793311..ad1912247e01f 100644
--- a/clang/include/clang/Driver/CommonArgs.h
+++ b/clang/include/clang/Driver/CommonArgs.h
@@ -144,6 +144,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 416ed935c1b30..f976aac60a062 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -254,9 +254,7 @@ CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
VoidTy = CGM.VoidTy;
PtrTy = CGM.DefaultPtrTy;
- if (CGM.getLangOpts().OffloadViaLLVM)
- Prefix = "llvm";
- else if (CGM.getLangOpts().HIP)
+ if (CGM.getLangOpts().HIP)
Prefix = "hip";
else
Prefix = "cuda";
@@ -345,41 +343,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 +417,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 +1292,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 38795f7c2ae7a..f27557c09fca3 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -905,10 +905,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) {
@@ -965,15 +970,20 @@ static TripleSet inferOffloadToolchains(Compilation &C,
ID = StringToOffloadArch(
getProcessorFromTargetID(llvm::Triple("amdgcn-amd-amdhsa"), Arch));
- if (Kind == Action::OFK_HIP && !IsAMDOffloadArch(ID)) {
- C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
- << "HIP" << Arch;
- return {};
- }
- if (Kind == Action::OFK_Cuda && !IsNVIDIAOffloadArch(ID)) {
- 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 && !IsAMDOffloadArch(ID)) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "HIP" << Arch;
+ return {};
+ }
+ if (Kind == Action::OFK_Cuda && !IsNVIDIAOffloadArch(ID)) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "CUDA" << Arch;
+ return {};
+ }
}
if (Kind == Action::OFK_OpenMP &&
(ID == OffloadArch::Unknown || ID == OffloadArch::Unused)) {
@@ -989,6 +999,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
@@ -1032,32 +1044,30 @@ 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) ||
(C.getInputArgs().hasArg(options::OPT_offload_arch_EQ) &&
!(IsCuda || IsHIP))));
+ // We currently don't support any kind of mixed offloading.
+ if (IsOpenMPOffloading)
+ IsCuda = IsHIP = IsSYCL = false;
+
llvm::SmallSet<Action::OffloadKind, 4> Kinds;
const std::pair<bool, Action::OffloadKind> ActiveKinds[] = {
{IsCuda, Action::OFK_Cuda},
@@ -1143,7 +1153,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())
@@ -5069,6 +5079,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;
@@ -5089,7 +5102,6 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
types::ID InputType = Input.first;
const Arg *InputArg = Input.second;
- // The toolchain can be active for unsupported file types.
if ((Kind == Action::OFK_Cuda && !types::isCuda(InputType)) ||
(Kind == Action::OFK_HIP && !types::isHIP(InputType)))
continue;
@@ -5184,9 +5196,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);
@@ -5214,7 +5229,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)) {
// If we are not in RDC-mode we just emit the final CUDA fatbinary for
// each translation unit without requiring any linking.
@@ -5222,7 +5237,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 =
@@ -5230,7 +5245,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 &&
@@ -7091,7 +7106,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 7bce060de0596..5c30417365b94 100644
--- a/clang/lib/Driver/ToolChains/AMDGPU.cpp
+++ b/clang/lib/Driver/ToolChains/AMDGPU.cpp
@@ -515,6 +515,24 @@ 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)) {
+ if (DriverArgs.hasFlag(options::OPT_offload_inc,
+ options::OPT_no_offload_inc, true) &&
+ !DriverArgs.hasArg(options::OPT_nohipwrapperinc) &&
+ !DriverArgs.hasArg(options::OPT_nobuiltininc)) {
+ CC1Args.append({"-include", "__clang_gpu_device_functions.h"});
+
+ SmallString<128> HIPIncludePath(D.ResourceDir);
+ llvm::sys::path::append(HIPIncludePath, "..", "..", "..");
+ llvm::sys::path::append(HIPIncludePath, "include", "offload");
+ CC1Args.push_back("-internal-isystem");
+ CC1Args.push_back(DriverArgs.MakeArgString(HIPIncludePath));
+ CC1Args.append({"-include", "hip/hip_runtime.h"});
+ }
+ 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
@@ -699,7 +717,11 @@ AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
: Generic_ELF(D, Triple, Args),
OptionsDefault(
{{options::OPT_O, "3"}, {options::OPT_cl_std_EQ, "CL1.2"}}),
- HostTC(HostTC_), UseHIPLinker(Kind == Action::OFK_HIP),
+ HostTC(HostTC_),
+ UseHIPLinker(Kind == Action::OFK_HIP ||
+ (Kind == Action::OFK_Cuda &&
+ Args.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))),
ShouldLinkDeviceLibs(ShouldLinkDeviceLibs) {
loadMultilibsFromYAML(Args, D);
@@ -709,8 +731,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)
@@ -889,7 +913,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 94f9a26aac39f..035bc3f5b4273 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -53,6 +53,7 @@
#include "llvm/Support/Path.h"
#include "llvm/Support/Process.h"
#include "llvm/Support/YAMLParser.h"
+#include "llvm/Support/raw_ostream.h"
#include "llvm/TargetParser/AArch64TargetParser.h"
#include "llvm/TargetParser/ARMTargetParserCommon.h"
#include "llvm/TargetParser/Host.h"
@@ -952,9 +953,12 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
// 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))
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ if (JA.isOffloading(Action::OFK_Cuda) && !UsesLLVMOffloading) {
getToolChain().AddCudaIncludeArgs(Args, CmdArgs);
- if (JA.isOffloading(Action::OFK_HIP))
+ }
+ if (JA.isOffloading(Action::OFK_HIP) && !UsesLLVMOffloading)
getToolChain().AddHIPIncludeArgs(Args, CmdArgs);
if (JA.isOffloading(Action::OFK_SYCL))
getToolChain().addSYCLIncludeArgs(Args, CmdArgs);
@@ -979,17 +983,35 @@ void Clang::AddPreprocessingOptions(Compi...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/213390
More information about the cfe-commits
mailing list