llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-x86 @llvm/pr-subscribers-offload Author: Sophia Herrmann (jellytabby) <details> <summary>Changes</summary> Continuing the work to cover cuda/hip runtime in LLVMOffloadKernel, this PR introduces the `GetErrorName` and `GetErrorString` functions for pretty printing LLVMOffload errors, in a separate `LanguageUtils.cpp` Additionally moves out helper functions ` convertResult` and `getQueueFromStream` into the utils file too. This PR depends on #<!-- -->211694 and #<!-- -->212373 but because I do not have commit access I cannot stack the PR. For review only consider the LAST ONE commit. --- Patch is 150.32 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/212887.diff 70 Files Affected: - (modified) clang/include/clang/Driver/CommonArgs.h (+6) - (modified) clang/lib/CodeGen/CGCUDANV.cpp (+38-29) - (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 (+1-1) - (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 (+16) - (added) offload/languages/hip/CMakeLists.txt (+1) - (added) offload/languages/hip/src/hip_runtime.cpp (+17) - (added) offload/languages/include/cuda/cuda_runtime.h (+22) - (added) offload/languages/include/hip/hip_runtime.h (+59) - (added) offload/languages/include/kernel/DefineLanguageNames.inc (+46) - (added) offload/languages/include/kernel/LanguageRuntime.h (+187) - (added) offload/languages/include/kernel/UndefineLanguageNames.inc (+44) - (added) offload/languages/kernel/CMakeLists.txt (+54) - (added) offload/languages/kernel/exports (+11) - (added) offload/languages/kernel/include/LanguageAliases.h (+40) - (added) offload/languages/kernel/include/LanguageLaunch.h (+50) - (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/include/Types.h (+28) - (added) offload/languages/kernel/src/LanguageCommon.cpp (+18) - (added) offload/languages/kernel/src/LanguageLaunch.cpp (+96) - (added) offload/languages/kernel/src/LanguageRegistration.cpp (+130) - (added) offload/languages/kernel/src/LanguageRuntime.cpp (+170) - (added) offload/languages/kernel/src/LanguageUtils.cpp (+65) - (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 (+61) - (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 (+61) - (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..0ea3ed36fae83 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); @@ -435,7 +445,9 @@ void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF, else if (CGF.getLangOpts().CUDA) KernelLaunchAPI = KernelLaunchAPI + "_ptsz"; } - auto LaunchKernelName = addPrefixToName(KernelLaunchAPI); + /// Use __llvmLaunchKernel for LLVMOffload. + auto LaunchKernelName = UsesLLVMOffloading ? "__llvm" + KernelLaunchAPI + : addPrefixToName(KernelLaunchAPI); const IdentifierInfo &cudaLaunchKernelII = CGM.getContext().Idents.get(LaunchKernelName); FunctionDecl *cudaLaunchKernelFD = nullptr; @@ -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 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.isOffloadi... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/212887 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
