https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/212373
>From 175a0cc400ca80efd3556ea4f7642ffd1c056e31 Mon Sep 17 00:00:00 2001 From: Sophia Herrmann <[email protected]> Date: Mon, 27 Jul 2026 15:48:11 -0700 Subject: [PATCH] decouple front/back end Co-authored-by: Johannes Doerfert <[email protected]> Co-authored-by: Jonas Greifenhain <[email protected]> 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 | 19 +++- clang/lib/Driver/ToolChains/Clang.cpp | 106 +++++++++++------- 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 ++- .../linker-wrapper-image.c | 38 ++++--- .../ClangLinkerWrapper.cpp | 13 ++- .../Frontend/Offloading/OffloadWrapper.cpp | 4 +- .../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 | 7 +- .../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 | 40 +++++++ 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 | 40 +++++++ 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 ++++++++ 46 files changed, 1139 insertions(+), 286 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 d4719f37e5b4d..3ec91d15f2c88 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 && @@ -7113,7 +7125,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 3d8d1e493570a..9f701b47b8065 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 @@ -723,7 +727,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); @@ -733,8 +741,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 +923,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 d2e22920aa432..77dae5913a762 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -52,6 +52,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" @@ -915,48 +916,64 @@ 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) { + CmdArgs.append({"-include", "__clang_gpu_device_functions.h"}); + if (JA.isOffloading(Action::OFK_Cuda)) { + SmallString<128> OffloadCudaInclude(D.Dir); + llvm::sys::path::append(OffloadCudaInclude, "..", "include", "offload", + "cuda"); + CmdArgs.append({"-internal-isystem", + Args.MakeArgString(OffloadCudaInclude), "-include"}); + CmdArgs.push_back("cuda_runtime.h"); + } + if (JA.isOffloading(Action::OFK_HIP) && + !Args.hasArg(options::OPT_nohipwrapperinc)) { + SmallString<128> OffloadHIPInclude(D.Dir); + llvm::sys::path::append(OffloadHIPInclude, "..", "include", "offload"); + CmdArgs.append({"-internal-isystem", + Args.MakeArgString(OffloadHIPInclude), "-include"}); + CmdArgs.push_back("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 @@ -5164,6 +5181,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 = @@ -5282,7 +5301,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 *>( @@ -8289,12 +8308,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/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c index d05e9d54a108a..083d3340f6f81 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c @@ -94,7 +94,7 @@ // CUDA-NEXT: br i1 %1, label %while.entry, label %while.end // // CUDA: while.entry: -// CUDA-NEXT: %entry1 = phi ptr [ @__start_llvm_offload_entries, %entry ], [ %16, %if.end ] +// CUDA-NEXT: %entry1 = phi ptr [ @__start_llvm_offload_entries, %entry ], [ %17, %if.end ] // CUDA-NEXT: %2 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 4 // CUDA-NEXT: %addr = load ptr, ptr %2, align 8 // CUDA-NEXT: %3 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 8 @@ -117,15 +117,16 @@ // CUDA-NEXT: %constant = lshr i32 %11, 4 // CUDA-NEXT: %12 = and i32 %flags, 32 // CUDA-NEXT: %normalized = lshr i32 %12, 5 -// CUDA-NEXT: %13 = icmp eq i16 %kind, 2 -// CUDA-NEXT: br i1 %13, label %if.kind, label %if.end +// CUDA-NEXT: %13 = and i16 %kind, 2 +// CUDA-NEXT: %14 = icmp ne i16 %13, 0 +// CUDA-NEXT: br i1 %14, label %if.kind, label %if.end // // CUDA: if.kind: -// CUDA-NEXT: %14 = icmp eq i64 %size, 0 -// CUDA-NEXT: br i1 %14, label %if.then, label %if.else +// CUDA-NEXT: %15 = icmp eq i64 %size, 0 +// CUDA-NEXT: br i1 %15, label %if.then, label %if.else // // CUDA: if.then: -// CUDA-NEXT: %15 = call i32 @__cudaRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null) +// CUDA-NEXT: %16 = call i32 @__cudaRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null) // CUDA-NEXT: br label %if.end // // CUDA: if.else: @@ -151,9 +152,9 @@ // CUDA-NEXT: br label %if.end // // CUDA: if.end: -// CUDA-NEXT: %16 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1 -// CUDA-NEXT: %17 = icmp eq ptr %16, @__stop_llvm_offload_entries -// CUDA-NEXT: br i1 %17, label %while.end, label %while.entry +// CUDA-NEXT: %17 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1 +// CUDA-NEXT: %18 = icmp eq ptr %17, @__stop_llvm_offload_entries +// CUDA-NEXT: br i1 %18, label %while.end, label %while.entry // // CUDA: while.end: // CUDA-NEXT: ret void @@ -236,7 +237,7 @@ // HIP-NEXT: br i1 %1, label %while.entry, label %while.end // // HIP: while.entry: -// HIP-NEXT: %entry1 = phi ptr [ @{{.*offload_entries.*}}, %entry ], [ %16, %if.end ] +// HIP-NEXT: %entry1 = phi ptr [ @{{.*offload_entries.*}}, %entry ], [ %17, %if.end ] // HIP-NEXT: %2 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 4 // HIP-NEXT: %addr = load ptr, ptr %2, align 8 // HIP-NEXT: %3 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 8 @@ -259,15 +260,16 @@ // HIP-NEXT: %constant = lshr i32 %11, 4 // HIP-NEXT: %12 = and i32 %flags, 32 // HIP-NEXT: %normalized = lshr i32 %12, 5 -// HIP-NEXT: %13 = icmp eq i16 %kind, 4 -// HIP-NEXT: br i1 %13, label %if.kind, label %if.end +// HIP-NEXT: %13 = and i16 %kind, 4 +// HIP-NEXT: %14 = icmp ne i16 %13, 0 +// HIP-NEXT: br i1 %14, label %if.kind, label %if.end // // HIP: if.kind: -// HIP-NEXT: %14 = icmp eq i64 %size, 0 -// HIP-NEXT: br i1 %14, label %if.then, label %if.else +// HIP-NEXT: %15 = icmp eq i64 %size, 0 +// HIP-NEXT: br i1 %15, label %if.then, label %if.else // // HIP: if.then: -// HIP-NEXT: %15 = call i32 @__hipRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null) +// HIP-NEXT: %16 = call i32 @__hipRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null) // HIP-NEXT: br label %if.end // // HIP: if.else: @@ -295,9 +297,9 @@ // HIP-NEXT: br label %if.end // // HIP: if.end: -// HIP-NEXT: %16 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1 -// HIP-NEXT: %17 = icmp eq ptr %16, @{{.*offload_entries.*}} -// HIP-NEXT: br i1 %17, label %while.end, label %while.entry +// HIP-NEXT: %17 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1 +// HIP-NEXT: %18 = icmp eq ptr %17, @{{.*offload_entries.*}} +// HIP-NEXT: br i1 %18, label %while.end, label %while.entry // // HIP: while.end: // HIP-NEXT: ret void 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/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp index ded603a1e00e3..037b81a7c42fb 100644 --- a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp +++ b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp @@ -485,10 +485,12 @@ Function *createRegisterGlobalsFunction(Module &M, bool IsHIP, llvm::offloading::OffloadGlobalNormalized)); auto *Normalized = Builder.CreateLShr( NormalizedBit, ConstantInt::get(Type::getInt32Ty(C), 5), "normalized"); - auto *KindCond = Builder.CreateICmpEQ( + auto *KindAnd = Builder.CreateAnd( Kind, ConstantInt::get(Type::getInt16Ty(C), IsHIP ? object::OffloadKind::OFK_HIP : object::OffloadKind::OFK_Cuda)); + auto *KindCond = + Builder.CreateICmpNE(KindAnd, ConstantInt::get(Type::getInt16Ty(C), 0)); Builder.CreateCondBr(KindCond, IfKindBB, IfEndBB); Builder.SetInsertPoint(IfKindBB); auto *FnCond = Builder.CreateICmpEQ( 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 81442f9f2c507..825cd4259b7fd 100644 --- a/offload/languages/kernel/CMakeLists.txt +++ b/offload/languages/kernel/CMakeLists.txt @@ -75,4 +75,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..0d1a64fe464f8 100644 --- a/offload/languages/kernel/include/Types.h +++ b/offload/languages/kernel/include/Types.h @@ -16,7 +16,12 @@ struct uint3 { unsigned x = 0, y = 0, z = 0; }; -using dim3 = uint3; +struct dim3 { + dim3() = default; + dim3(unsigned X, unsigned Y = 1, unsigned Z = 1) : x(X), y(Y), z(Z) {} + + unsigned x = 1, y = 1, z = 1; +}; 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..f23eac7604582 --- /dev/null +++ b/offload/test/offloading/CUDA/host_alloc.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) { + int *HostAllocPtr = nullptr; + if (cudaHostAlloc(&HostAllocPtr, sizeof(int), cudaHostAllocDefault) != + cudaSuccess) + return 1; + + *HostAllocPtr = 17; + printf("cudaHostAlloc value: %d\n", *HostAllocPtr); + // CHECK: cudaHostAlloc value: 17 + + if (cudaFreeHost(HostAllocPtr) != cudaSuccess) + return 1; + + int *MallocHostPtr = nullptr; + if (cudaMallocHost(&MallocHostPtr, sizeof(int)) != cudaSuccess) + return 1; + + *MallocHostPtr = 23; + printf("cudaMallocHost value: %d\n", *MallocHostPtr); + // CHECK: cudaMallocHost value: 23 + + 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..8b067b39f2f81 --- /dev/null +++ b/offload/test/offloading/HIP/host_alloc.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) { + int *HostAllocPtr = nullptr; + if (hipHostAlloc(&HostAllocPtr, sizeof(int), hipHostAllocDefault) != + hipSuccess) + return 1; + + *HostAllocPtr = 17; + printf("hipHostAlloc value: %d\n", *HostAllocPtr); + // CHECK: hipHostAlloc value: 17 + + if (hipFreeHost(HostAllocPtr) != hipSuccess) + return 1; + + int *MallocHostPtr = nullptr; + if (hipMallocHost(&MallocHostPtr, sizeof(int)) != hipSuccess) + return 1; + + *MallocHostPtr = 23; + printf("hipMallocHost value: %d\n", *MallocHostPtr); + // CHECK: hipMallocHost value: 23 + + 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; +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
