https://github.com/jdoerfert created https://github.com/llvm/llvm-project/pull/94549
Through the new `-foffload-via-llvm` flag, CUDA kernels can now be lowered to the LLVM/Offload API. On the Clang side, this is simply done by using the OpenMP offload toolchain and emitting calls to `llvm*` functions to orchestrate the kernel launch rather than `cuda*` functions. These `llvm*` functions are implemented on top of the existing LLVM/Offload API. As we are about to redefine the Offload API, this wil help us in the design process as a second offload language. We do not support any CUDA APIs yet, however, we could: https://www.osti.gov/servlets/purl/1892137 For proper host execution we need to resurrect/rebase https://tianshilei.me/wp-content/uploads/2021/12/llpp-2021.pdf (which was designed for debugging). ``` ❯❯❯ cat test.cu 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); square<<<1, 1>>>(Ptr); printf("Ptr %p, *Ptr %i\n", Ptr, *Ptr); llvm_omp_target_free_shared(Ptr, DevNo); } ❯❯❯ clang++ test.cu -O3 -o test123 -foffload-via-llvm --offload-arch=native ❯❯❯ llvm-objdump --offloading test123 test123: file format elf64-x86-64 OFFLOADING IMAGE [0]: kind elf arch gfx90a triple amdgcn-amd-amdhsa producer openmp ❯❯❯ LIBOMPTARGET_INFO=16 ./test123 Ptr 0x155448ac8000, *Ptr 7 Ptr 0x155448ac8000, *Ptr 42 ``` >From 36618e65d94ffa3e83464b7d19ff6cd7d5855abf Mon Sep 17 00:00:00 2001 From: Johannes Doerfert <johan...@jdoerfert.de> Date: Wed, 5 Jun 2024 16:51:51 -0700 Subject: [PATCH 1/2] [Offload][NFCI] Initialize the KernelArgsTy to default values --- offload/include/Shared/APITypes.h | 30 +++++++++++++++++------------- 1 file changed, 17 insertions(+), 13 deletions(-) diff --git a/offload/include/Shared/APITypes.h b/offload/include/Shared/APITypes.h index e8fc27785b6c2..fd315c6b992b9 100644 --- a/offload/include/Shared/APITypes.h +++ b/offload/include/Shared/APITypes.h @@ -89,22 +89,26 @@ struct __tgt_async_info { /// This struct contains all of the arguments to a target kernel region launch. struct KernelArgsTy { - uint32_t Version; // Version of this struct for ABI compatibility. - uint32_t NumArgs; // Number of arguments in each input pointer. - void **ArgBasePtrs; // Base pointer of each argument (e.g. a struct). - void **ArgPtrs; // Pointer to the argument data. - int64_t *ArgSizes; // Size of the argument data in bytes. - int64_t *ArgTypes; // Type of the data (e.g. to / from). - void **ArgNames; // Name of the data for debugging, possibly null. - void **ArgMappers; // User-defined mappers, possibly null. - uint64_t Tripcount; // Tripcount for the teams / distribute loop, 0 otherwise. + uint32_t Version = 0; // Version of this struct for ABI compatibility. + uint32_t NumArgs = 0; // Number of arguments in each input pointer. + void **ArgBasePtrs = + nullptr; // Base pointer of each argument (e.g. a struct). + void **ArgPtrs = nullptr; // Pointer to the argument data. + int64_t *ArgSizes = nullptr; // Size of the argument data in bytes. + int64_t *ArgTypes = nullptr; // Type of the data (e.g. to / from). + void **ArgNames = nullptr; // Name of the data for debugging, possibly null. + void **ArgMappers = nullptr; // User-defined mappers, possibly null. + uint64_t Tripcount = + 0; // Tripcount for the teams / distribute loop, 0 otherwise. struct { uint64_t NoWait : 1; // Was this kernel spawned with a `nowait` clause. uint64_t Unused : 63; - } Flags; - uint32_t NumTeams[3]; // The number of teams (for x,y,z dimension). - uint32_t ThreadLimit[3]; // The number of threads (for x,y,z dimension). - uint32_t DynCGroupMem; // Amount of dynamic cgroup memory requested. + } Flags = {0, 0}; + uint32_t NumTeams[3] = {0, 0, + 0}; // The number of teams (for x,y,z dimension). + uint32_t ThreadLimit[3] = {0, 0, + 0}; // The number of threads (for x,y,z dimension). + uint32_t DynCGroupMem = 0; // Amount of dynamic cgroup memory requested. }; static_assert(sizeof(KernelArgsTy().Flags) == sizeof(uint64_t), "Invalid struct size"); >From 3d5c61a78e91ecb379a2bfac71988eaf8e5cd9cd Mon Sep 17 00:00:00 2001 From: Johannes Doerfert <johan...@jdoerfert.de> Date: Mon, 3 Jun 2024 19:52:12 -0700 Subject: [PATCH 2/2] [Offload][CUDA] Allow CUDA kernels to use LLVM/Offload MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Through the new `-foffload-via-llvm` flag, CUDA kernels can now be lowered to the LLVM/Offload API. On the Clang side, this is simply done by using the OpenMP offload toolchain and emitting calls to `llvm*` functions to orchestrate the kernel launch rather than `cuda*` functions. These `llvm*` functions are implemented on top of the existing LLVM/Offload API. As we are about to redefine the Offload API, this wil help us in the design process as a second offload language. We do not support any CUDA APIs yet, however, we could: https://www.osti.gov/servlets/purl/1892137 For proper host execution we need to resurrect/rebase https://tianshilei.me/wp-content/uploads/2021/12/llpp-2021.pdf (which was designed for debugging). ``` ❯❯❯ cat test.cu 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); square<<<1, 1>>>(Ptr); printf("Ptr %p, *Ptr %i\n", Ptr, *Ptr); llvm_omp_target_free_shared(Ptr, DevNo); } ❯❯❯ clang++ test.cu -O3 -o test123 -foffload-via-llvm --offload-arch=native ❯❯❯ llvm-objdump --offloading test123 test123: file format elf64-x86-64 OFFLOADING IMAGE [0]: kind elf arch gfx90a triple amdgcn-amd-amdhsa producer openmp ❯❯❯ LIBOMPTARGET_INFO=16 ./test123 Ptr 0x155448ac8000, *Ptr 7 Ptr 0x155448ac8000, *Ptr 42 ``` --- clang/include/clang/Basic/LangOptions.def | 1 + clang/include/clang/Driver/Options.td | 6 ++ clang/lib/CodeGen/CGCUDANV.cpp | 34 ++++++--- clang/lib/Driver/Driver.cpp | 19 +++-- clang/lib/Driver/ToolChains/Clang.cpp | 26 ++++++- clang/lib/Driver/ToolChains/CommonArgs.cpp | 7 +- clang/lib/Headers/CMakeLists.txt | 3 + .../__clang_openmp_device_functions.h | 6 +- .../Headers/openmp_wrappers/__llvm_offload.h | 31 ++++++++ .../openmp_wrappers/__llvm_offload_device.h | 10 +++ .../openmp_wrappers/__llvm_offload_host.h | 15 ++++ clang/lib/Sema/SemaCUDA.cpp | 3 + clang/test/Driver/cuda-via-liboffload.cu | 23 ++++++ offload/include/Shared/APITypes.h | 5 +- offload/include/omptarget.h | 2 +- offload/plugins-nextgen/amdgpu/src/rtl.cpp | 12 ++- .../common/src/PluginInterface.cpp | 14 +++- offload/src/CMakeLists.txt | 1 + offload/src/KernelLanguage/API.cpp | 76 +++++++++++++++++++ offload/src/exports | 3 + offload/test/offloading/CUDA/basic_launch.cu | 29 +++++++ .../CUDA/basic_launch_blocks_and_threads.cu | 30 ++++++++ .../offloading/CUDA/basic_launch_multi_arg.cu | 33 ++++++++ 23 files changed, 356 insertions(+), 33 deletions(-) create mode 100644 clang/lib/Headers/openmp_wrappers/__llvm_offload.h create mode 100644 clang/lib/Headers/openmp_wrappers/__llvm_offload_device.h create mode 100644 clang/lib/Headers/openmp_wrappers/__llvm_offload_host.h create mode 100644 clang/test/Driver/cuda-via-liboffload.cu create mode 100644 offload/src/KernelLanguage/API.cpp create mode 100644 offload/test/offloading/CUDA/basic_launch.cu create mode 100644 offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu create mode 100644 offload/test/offloading/CUDA/basic_launch_multi_arg.cu diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index 4061451b2150a..8aff98867202e 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -288,6 +288,7 @@ LANGOPT(GPUMaxThreadsPerBlock, 32, 1024, "default max threads per block for kern LANGOPT(GPUDeferDiag, 1, 0, "defer host/device related diagnostic messages for CUDA/HIP") LANGOPT(GPUExcludeWrongSideOverloads, 1, 0, "always exclude wrong side overloads in overloading resolution for CUDA/HIP") LANGOPT(OffloadingNewDriver, 1, 0, "use the new driver for generating offloading code.") +LANGOPT(OffloadViaLLVM, 1, 0, "target LLVM/Offload as portable offloading runtime.") LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device") LANGOPT(SYCLIsHost , 1, 0, "SYCL host compilation") diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index 57f37c5023110..a09d75917ff98 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -1275,6 +1275,12 @@ def no_offload_compress : Flag<["--"], "no-offload-compress">; def offload_compression_level_EQ : Joined<["--"], "offload-compression-level=">, Flags<[HelpHidden]>, HelpText<"Compression level for offload device binaries (HIP only)">; + +defm offload_via_llvm : BoolFOption<"offload-via-llvm", + LangOpts<"OffloadViaLLVM">, DefaultFalse, + PosFlag<SetTrue, [], [ClangOption, CC1Option], "Use">, + NegFlag<SetFalse, [], [ClangOption], "Don't use">, + BothFlags<[], [ClangOption], " LLVM/Offload as portable offloading runtime.">>; } // CUDA options diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp index 43dfbbb90dd52..acbc7e9e43635 100644 --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -19,6 +19,7 @@ #include "clang/Basic/Cuda.h" #include "clang/CodeGen/CodeGenABITypes.h" #include "clang/CodeGen/ConstantInitBuilder.h" +#include "llvm/ADT/StringRef.h" #include "llvm/Frontend/Offloading/Utility.h" #include "llvm/IR/BasicBlock.h" #include "llvm/IR/Constants.h" @@ -36,6 +37,11 @@ constexpr unsigned HIPFatMagic = 0x48495046; // "HIPF" class CGNVCUDARuntime : public CGCUDARuntime { + /// The prefix used for function calls and section names (CUDA, HIP, LLVM) + StringRef Prefix; + /// TODO: We should transition the OpenMP section to LLVM/Offload + StringRef SectionPrefix; + private: llvm::IntegerType *IntTy, *SizeTy; llvm::Type *VoidTy; @@ -191,15 +197,11 @@ class CGNVCUDARuntime : public CGCUDARuntime { } // end anonymous namespace std::string CGNVCUDARuntime::addPrefixToName(StringRef FuncName) const { - if (CGM.getLangOpts().HIP) - return ((Twine("hip") + Twine(FuncName)).str()); - return ((Twine("cuda") + Twine(FuncName)).str()); + return (Prefix + FuncName).str(); } std::string CGNVCUDARuntime::addUnderscoredPrefixToName(StringRef FuncName) const { - if (CGM.getLangOpts().HIP) - return ((Twine("__hip") + Twine(FuncName)).str()); - return ((Twine("__cuda") + Twine(FuncName)).str()); + return ("__" + Prefix + FuncName).str(); } static std::unique_ptr<MangleContext> InitDeviceMC(CodeGenModule &CGM) { @@ -227,6 +229,14 @@ CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM) SizeTy = CGM.SizeTy; VoidTy = CGM.VoidTy; PtrTy = CGM.UnqualPtrTy; + + if (CGM.getLangOpts().OffloadViaLLVM) { + Prefix = "llvm"; + SectionPrefix = "omp"; + } else if (CGM.getLangOpts().HIP) + SectionPrefix = Prefix = "hip"; + else + SectionPrefix = Prefix = "cuda"; } llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const { @@ -305,7 +315,8 @@ void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF, } if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(), CudaFeature::CUDA_USES_NEW_LAUNCH) || - (CGF.getLangOpts().HIP && CGF.getLangOpts().HIPUseNewLaunchAPI)) + (CGF.getLangOpts().HIP && CGF.getLangOpts().HIPUseNewLaunchAPI) || + (CGF.getLangOpts().OffloadViaLLVM)) emitDeviceStubBodyNew(CGF, Args); else emitDeviceStubBodyLegacy(CGF, Args); @@ -1129,8 +1140,9 @@ void CGNVCUDARuntime::transformManagedVars() { // registered. The linker will provide a pointer to this section so we can // register the symbols with the linked device image. void CGNVCUDARuntime::createOffloadingEntries() { - StringRef Section = CGM.getLangOpts().HIP ? "hip_offloading_entries" - : "cuda_offloading_entries"; + SmallVector<char, 32> Out; + StringRef Section = (SectionPrefix + "_offloading_entries").toStringRef(Out); + llvm::Module &M = CGM.getModule(); for (KernelInfo &I : EmittedKernels) llvm::offloading::emitOffloadingEntry( @@ -1199,7 +1211,9 @@ llvm::Function *CGNVCUDARuntime::finalizeModule() { } return nullptr; } - if (CGM.getLangOpts().OffloadingNewDriver && RelocatableDeviceCode) + if (CGM.getLangOpts().OffloadViaLLVM) + createOffloadingEntries(); + else if (CGM.getLangOpts().OffloadingNewDriver && RelocatableDeviceCode) createOffloadingEntries(); else return makeModuleCtorFunction(); diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index f5ea73a04ae5c..815149a49d018 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -792,11 +792,13 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C, }) || C.getInputArgs().hasArg(options::OPT_hip_link) || C.getInputArgs().hasArg(options::OPT_hipstdpar); + bool UseLLVMOffload = C.getInputArgs().hasArg( + options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false); if (IsCuda && IsHIP) { Diag(clang::diag::err_drv_mix_cuda_hip); return; } - if (IsCuda) { + if (IsCuda && !UseLLVMOffload) { const ToolChain *HostTC = C.getSingleOffloadToolChain<Action::OFK_Host>(); const llvm::Triple &HostTriple = HostTC->getTriple(); auto OFK = Action::OFK_Cuda; @@ -818,7 +820,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C, CudaInstallation.WarnIfUnsupportedVersion(); } C.addOffloadDeviceToolChain(CudaTC.get(), OFK); - } else if (IsHIP) { + } else if (IsHIP && !UseLLVMOffload) { if (auto *OMPTargetArg = C.getInputArgs().getLastArg(options::OPT_fopenmp_targets_EQ)) { Diag(clang::diag::err_drv_unsupported_opt_for_language_mode) @@ -842,10 +844,11 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C, // We need to generate an OpenMP toolchain if the user specified targets with // the -fopenmp-targets option or used --offload-arch with OpenMP enabled. bool IsOpenMPOffloading = - C.getInputArgs().hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ, - options::OPT_fno_openmp, false) && - (C.getInputArgs().hasArg(options::OPT_fopenmp_targets_EQ) || - C.getInputArgs().hasArg(options::OPT_offload_arch_EQ)); + ((IsCuda || IsHIP) && UseLLVMOffload) || + (C.getInputArgs().hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ, + options::OPT_fno_openmp, false) && + (C.getInputArgs().hasArg(options::OPT_fopenmp_targets_EQ) || + C.getInputArgs().hasArg(options::OPT_offload_arch_EQ))); if (IsOpenMPOffloading) { // We expect that -fopenmp-targets is always used in conjunction with the // option -fopenmp specifying a valid runtime with offloading support, i.e. @@ -873,7 +876,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C, for (StringRef T : OpenMPTargets->getValues()) OpenMPTriples.insert(T); } else if (C.getInputArgs().hasArg(options::OPT_offload_arch_EQ) && - !IsHIP && !IsCuda) { + ((!IsHIP && !IsCuda) || UseLLVMOffload)) { const ToolChain *HostTC = C.getSingleOffloadToolChain<Action::OFK_Host>(); auto AMDTriple = getHIPOffloadTargetTriple(*this, C.getInputArgs()); auto NVPTXTriple = getNVIDIAOffloadTargetTriple(*this, C.getInputArgs(), @@ -4138,6 +4141,8 @@ void Driver::BuildActions(Compilation &C, DerivedArgList &Args, bool UseNewOffloadingDriver = C.isOffloadingHostKind(Action::OFK_OpenMP) || + Args.hasFlag(options::OPT_foffload_via_llvm, + options::OPT_fno_offload_via_llvm, false) || Args.hasFlag(options::OPT_offload_new_driver, options::OPT_no_offload_new_driver, false); diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 4e1c52462e584..07b7e60caa2b4 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -1125,6 +1125,18 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA, CmdArgs.push_back("__clang_openmp_device_functions.h"); } + if (Args.hasArg(options::OPT_foffload_via_llvm)) { + CmdArgs.push_back("-include"); + SmallString<128> P(D.ResourceDir); + llvm::sys::path::append(P, "include"); + llvm::sys::path::append(P, "openmp_wrappers"); + if (JA.isDeviceOffloading(Action::OFK_OpenMP)) + llvm::sys::path::append(P, "__llvm_offload_device.h"); + else + llvm::sys::path::append(P, "__llvm_offload_host.h"); + CmdArgs.push_back(Args.MakeArgString(P)); + } + // Add -i* options, and automatically translate to // -include-pch/-include-pth for transparent PCH support. It's // wonky, but we include looking for .gch so we can support seamless @@ -6672,11 +6684,17 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, Args.addOptOutFlag(CmdArgs, options::OPT_fopenmp_extensions, options::OPT_fno_openmp_extensions); } - - // Forward the new driver to change offloading code generation. - if (Args.hasFlag(options::OPT_offload_new_driver, - options::OPT_no_offload_new_driver, false)) + // Forward the offload runtime change to code generation, liboffload implies + // new driver. Otherwise, check if we should forward the new driver to change + // offloading code generation. + if (Args.hasFlag(options::OPT_foffload_via_llvm, + options::OPT_fno_offload_via_llvm, false)) { + CmdArgs.push_back("--offload-new-driver"); + CmdArgs.push_back("-foffload-via-llvm"); + } else if (Args.hasFlag(options::OPT_offload_new_driver, + options::OPT_no_offload_new_driver, false)) { CmdArgs.push_back("--offload-new-driver"); + } SanitizeArgs.addArgs(TC, Args, CmdArgs, InputType); diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp b/clang/lib/Driver/ToolChains/CommonArgs.cpp index 71e993119436a..74a69f65f7ad5 100644 --- a/clang/lib/Driver/ToolChains/CommonArgs.cpp +++ b/clang/lib/Driver/ToolChains/CommonArgs.cpp @@ -1144,8 +1144,13 @@ bool tools::addOpenMPRuntime(const Compilation &C, ArgStringList &CmdArgs, bool ForceStaticHostRuntime, bool IsOffloadingHost, bool GompNeedsRT) { if (!Args.hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ, - options::OPT_fno_openmp, false)) + 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"); return false; + } Driver::OpenMPRuntimeKind RTKind = TC.getDriver().getOpenMPRuntime(Args); diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt index d3090e488306f..215dfe207c228 100644 --- a/clang/lib/Headers/CMakeLists.txt +++ b/clang/lib/Headers/CMakeLists.txt @@ -323,6 +323,9 @@ set(openmp_wrapper_files openmp_wrappers/__clang_openmp_device_functions.h openmp_wrappers/complex_cmath.h openmp_wrappers/new + openmp_wrappers/__llvm_offload.h + openmp_wrappers/__llvm_offload_host.h + openmp_wrappers/__llvm_offload_device.h ) set(llvm_libc_wrapper_files diff --git a/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h index d5b6846b03488..889d607557ecb 100644 --- a/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h +++ b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h @@ -10,14 +10,11 @@ #ifndef __CLANG_OPENMP_DEVICE_FUNCTIONS_H__ #define __CLANG_OPENMP_DEVICE_FUNCTIONS_H__ -#ifndef _OPENMP -#error "This file is for OpenMP compilation only." -#endif - #ifdef __cplusplus extern "C" { #endif +#ifdef __NVPTX__ #pragma omp begin declare variant match( \ device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) @@ -34,6 +31,7 @@ extern "C" { #undef __CUDA__ #pragma omp end declare variant +#endif #ifdef __AMDGCN__ #pragma omp begin declare variant match(device = {arch(amdgcn)}) diff --git a/clang/lib/Headers/openmp_wrappers/__llvm_offload.h b/clang/lib/Headers/openmp_wrappers/__llvm_offload.h new file mode 100644 index 0000000000000..d78e3b41f99a5 --- /dev/null +++ b/clang/lib/Headers/openmp_wrappers/__llvm_offload.h @@ -0,0 +1,31 @@ +/*===------ LLVM/Offload helpers for kernel languages (CUDA/HIP) -*- c++ -*-=== + * + * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. + * See https://llvm.org/LICENSE.txt for license information. + * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + * + *===-----------------------------------------------------------------------=== + */ + +#include <stdlib.h> + +#define __host__ __attribute__((host)) +#define __device__ __attribute__((device)) +#define __global__ __attribute__((global)) +#define __shared__ __attribute__((shared)) +#define __constant__ __attribute__((constant)) +#define __managed__ __attribute__((managed)) + +extern "C" { + +typedef struct dim3 { + dim3() {} + dim3(unsigned x) : x(x) {} + unsigned x = 0, y = 0, z = 0; +} dim3; + +// TODO: For some reason the CUDA device compilation requires this declaration +// to be present but it should not. +unsigned __llvmPushCallConfiguration(dim3 gridDim, dim3 blockDim, + size_t sharedMem = 0, void *stream = 0); +} diff --git a/clang/lib/Headers/openmp_wrappers/__llvm_offload_device.h b/clang/lib/Headers/openmp_wrappers/__llvm_offload_device.h new file mode 100644 index 0000000000000..1a813b331515b --- /dev/null +++ b/clang/lib/Headers/openmp_wrappers/__llvm_offload_device.h @@ -0,0 +1,10 @@ +/*===------ LLVM/Offload helpers for kernel languages (CUDA/HIP) -*- c++ -*-=== + * + * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. + * See https://llvm.org/LICENSE.txt for license information. + * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + * + *===-----------------------------------------------------------------------=== + */ + +#include "__llvm_offload.h" diff --git a/clang/lib/Headers/openmp_wrappers/__llvm_offload_host.h b/clang/lib/Headers/openmp_wrappers/__llvm_offload_host.h new file mode 100644 index 0000000000000..160289d169b55 --- /dev/null +++ b/clang/lib/Headers/openmp_wrappers/__llvm_offload_host.h @@ -0,0 +1,15 @@ +/*===------ LLVM/Offload helpers for kernel languages (CUDA/HIP) -*- c++ -*-=== + * + * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. + * See https://llvm.org/LICENSE.txt for license information. + * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + * + *===-----------------------------------------------------------------------=== + */ + +#include "__llvm_offload.h" + +extern "C" { +unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim, + void **args, size_t sharedMem = 0, void *stream = 0); +} diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 80ea43dc5316e..b507e19556363 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -1059,6 +1059,9 @@ void SemaCUDA::inheritTargetAttrs(FunctionDecl *FD, } std::string SemaCUDA::getConfigureFuncName() const { + if (getLangOpts().OffloadViaLLVM) + return "__llvmPushCallConfiguration"; + if (getLangOpts().HIP) return getLangOpts().HIPUseNewLaunchAPI ? "__hipPushCallConfiguration" : "hipConfigureCall"; diff --git a/clang/test/Driver/cuda-via-liboffload.cu b/clang/test/Driver/cuda-via-liboffload.cu new file mode 100644 index 0000000000000..68dc963e906b2 --- /dev/null +++ b/clang/test/Driver/cuda-via-liboffload.cu @@ -0,0 +1,23 @@ +// 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 %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-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" - "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:.+]]" + +// RUN: %clang -### -target x86_64-linux-gnu -ccc-print-bindings --offload-link -foffload-via-llvm %s 2>&1 | FileCheck -check-prefix DEVICE-LINK %s + +// DEVICE-LINK: "x86_64-unknown-linux-gnu" - "Offload::Linker", inputs: ["[[INPUT:.+]]"], output: "a.out" diff --git a/offload/include/Shared/APITypes.h b/offload/include/Shared/APITypes.h index fd315c6b992b9..2c41a6e2ba43c 100644 --- a/offload/include/Shared/APITypes.h +++ b/offload/include/Shared/APITypes.h @@ -102,8 +102,9 @@ struct KernelArgsTy { 0; // Tripcount for the teams / distribute loop, 0 otherwise. struct { uint64_t NoWait : 1; // Was this kernel spawned with a `nowait` clause. - uint64_t Unused : 63; - } Flags = {0, 0}; + uint64_t IsCUDA : 1; // Was this kernel spawned via CUDA. + uint64_t Unused : 62; + } Flags = {0, 0, 0}; uint32_t NumTeams[3] = {0, 0, 0}; // The number of teams (for x,y,z dimension). uint32_t ThreadLimit[3] = {0, 0, diff --git a/offload/include/omptarget.h b/offload/include/omptarget.h index 323dee41630f2..2b6445e9fbe55 100644 --- a/offload/include/omptarget.h +++ b/offload/include/omptarget.h @@ -107,7 +107,7 @@ enum TargetAllocTy : int32_t { inline KernelArgsTy CTorDTorKernelArgs = {1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, - 0, {0,0}, {1, 0, 0}, {1, 0, 0}, 0}; + 0, {0,0,0}, {1, 0, 0}, {1, 0, 0}, 0}; struct DeviceTy; diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp index c6dd954746e4a..bb9c2291c80a8 100644 --- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp +++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp @@ -3268,6 +3268,11 @@ Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice, uint32_t NumThreads, uint64_t NumBlocks, KernelArgsTy &KernelArgs, void *Args, AsyncInfoWrapperTy &AsyncInfoWrapper) const { + if (KernelArgs.Flags.IsCUDA) { + // For CUDA kernels we compute the number of arguments here. + KernelArgs.NumArgs = (ArgsSize - ImplicitArgsSize) / sizeof(void *); + } + const uint32_t KernelArgsSize = KernelArgs.NumArgs * sizeof(void *); if (ArgsSize < KernelArgsSize) @@ -3310,9 +3315,14 @@ Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice, // Copy the explicit arguments. // TODO: We should expose the args memory manager alloc to the common part as // alternative to copying them twice. - if (KernelArgs.NumArgs) + if (KernelArgs.NumArgs && !KernelArgs.Flags.IsCUDA) { std::memcpy(AllArgs, *static_cast<void **>(Args), sizeof(void *) * KernelArgs.NumArgs); + } else { + for (uint32_t I = 0; I < KernelArgs.NumArgs; ++I) + std::memcpy(advanceVoidPtr(AllArgs, sizeof(void *) * I), + static_cast<void **>(Args)[I], sizeof(void *)); + } AMDGPUDeviceTy &AMDGPUDevice = static_cast<AMDGPUDeviceTy &>(GenericDevice); diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp index 913721a15d713..7ecb98a156f7c 100644 --- a/offload/plugins-nextgen/common/src/PluginInterface.cpp +++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp @@ -549,9 +549,17 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs, if (!KernelLaunchEnvOrErr) return KernelLaunchEnvOrErr.takeError(); - void *KernelArgsPtr = - prepareArgs(GenericDevice, ArgPtrs, ArgOffsets, KernelArgs.NumArgs, Args, - Ptrs, *KernelLaunchEnvOrErr); + void *KernelArgsPtr; + + // Kernel languages don't use indirection. + // TODO: In the new API we should have the "source language" encoded. + if (KernelArgs.Flags.IsCUDA) { + KernelArgsPtr = KernelArgs.ArgPtrs; + } else { + KernelArgsPtr = + prepareArgs(GenericDevice, ArgPtrs, ArgOffsets, KernelArgs.NumArgs, + Args, Ptrs, *KernelLaunchEnvOrErr); + } uint32_t NumThreads = getNumThreads(GenericDevice, KernelArgs.ThreadLimit); uint64_t NumBlocks = diff --git a/offload/src/CMakeLists.txt b/offload/src/CMakeLists.txt index efa5cdab33ec9..b442df45deaa5 100644 --- a/offload/src/CMakeLists.txt +++ b/offload/src/CMakeLists.txt @@ -22,6 +22,7 @@ add_llvm_library(omptarget OpenMP/InteropAPI.cpp OpenMP/OMPT/Callback.cpp + KernelLanguage/API.cpp ADDITIONAL_HEADER_DIRS ${LIBOMPTARGET_INCLUDE_DIR} diff --git a/offload/src/KernelLanguage/API.cpp b/offload/src/KernelLanguage/API.cpp new file mode 100644 index 0000000000000..07517ad1fc18a --- /dev/null +++ b/offload/src/KernelLanguage/API.cpp @@ -0,0 +1,76 @@ +//===------ API.cpp - Kernel Language (CUDA/HIP) entry points ----- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +//===----------------------------------------------------------------------===// + + +#include "Shared/APITypes.h" + +#include <cstdio> + +struct dim3 { + unsigned x = 0, y = 0, z = 0; +}; + +struct __omp_kernel_t { + dim3 __grid_size; + dim3 __block_size; + size_t __shared_memory; + + void *__stream; +}; + +static __omp_kernel_t __current_kernel = {}; +#pragma omp threadprivate(__current_kernel); + +extern "C" { + +// TODO: There is little reason we need to keep these names or the way calls are +// issued. For now we do to avoid modifying Clang's CUDA codegen. Unclear when +// we actually need to push/pop configurations. +unsigned __llvmPushCallConfiguration(dim3 __grid_size, dim3 __block_size, + size_t __shared_memory, void *__stream) { + __omp_kernel_t &__kernel = __current_kernel; + __kernel.__grid_size = __grid_size; + __kernel.__block_size = __block_size; + __kernel.__shared_memory = __shared_memory; + __kernel.__stream = __stream; + return 0; +} + +unsigned __llvmPopCallConfiguration(dim3 *__grid_size, dim3 *__block_size, + size_t *__shared_memory, void *__stream) { + __omp_kernel_t &__kernel = __current_kernel; + *__grid_size = __kernel.__grid_size; + *__block_size = __kernel.__block_size; + *__shared_memory = __kernel.__shared_memory; + *((void **)__stream) = __kernel.__stream; + return 0; +} + +int __tgt_target_kernel(void *Loc, int64_t DeviceId, int32_t NumTeams, + int32_t ThreadLimit, const void *HostPtr, + KernelArgsTy *Args); + +unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim, + void **args, size_t sharedMem, void *stream) { + KernelArgsTy Args = {}; + Args.DynCGroupMem = sharedMem; + Args.NumTeams[0] = gridDim.x; + Args.NumTeams[1] = gridDim.y; + Args.NumTeams[2] = gridDim.z; + Args.ThreadLimit[0] = blockDim.x; + Args.ThreadLimit[1] = blockDim.y; + Args.ThreadLimit[2] = blockDim.z; + Args.ArgPtrs = args; + Args.Flags.IsCUDA = true; + int rv = __tgt_target_kernel(nullptr, 0, gridDim.x, + blockDim.x, func, &Args); + return rv; +} +} diff --git a/offload/src/exports b/offload/src/exports index f95544ec8329c..7bdc7d2a531bb 100644 --- a/offload/src/exports +++ b/offload/src/exports @@ -71,6 +71,9 @@ VERS1.0 { __tgt_interop_use; __tgt_interop_destroy; ompt_libomptarget_connect; + __llvmPushCallConfiguration; + __llvmPopCallConfiguration; + llvmLaunchKernel; local: *; }; diff --git a/offload/test/offloading/CUDA/basic_launch.cu b/offload/test/offloading/CUDA/basic_launch.cu new file mode 100644 index 0000000000000..a4955f8565f2d --- /dev/null +++ b/offload/test/offloading/CUDA/basic_launch.cu @@ -0,0 +1,29 @@ +// RUN: %clang++ -foffload-via-llvm --offload-arch=native %s -o %t +// RUN: %t | %fcheck-generic + +// UNSUPPORTED: aarch64-unknown-linux-gnu +// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO +// UNSUPPORTED: x86_64-pc-linux-gnu +// UNSUPPORTED: x86_64-pc-linux-gnu-LTO + + +#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 + square<<<1, 1>>>(Ptr); + printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr); + // CHECK: Ptr [[Ptr]], *Ptr: 42 + llvm_omp_target_free_shared(Ptr, DevNo); +} diff --git a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu new file mode 100644 index 0000000000000..401b9e2d559b1 --- /dev/null +++ b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu @@ -0,0 +1,30 @@ +// RUN: %clang++ -foffload-via-llvm --offload-arch=native %s -o %t +// RUN: %t | %fcheck-generic + +// UNSUPPORTED: aarch64-unknown-linux-gnu +// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO +// UNSUPPORTED: x86_64-pc-linux-gnu +// UNSUPPORTED: x86_64-pc-linux-gnu-LTO + +#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: 7 + square<<<7, 6>>>(Ptr); + printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr); + // CHECK: Ptr [[Ptr]], *Ptr: 42 + llvm_omp_target_free_shared(Ptr, DevNo); +} diff --git a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu new file mode 100644 index 0000000000000..c836292127fea --- /dev/null +++ b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu @@ -0,0 +1,33 @@ +// RUN: %clang++ -foffload-via-llvm --offload-arch=native %s -o %t +// RUN: %t | %fcheck-generic + +// UNSUPPORTED: aarch64-unknown-linux-gnu +// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO +// UNSUPPORTED: x86_64-pc-linux-gnu +// UNSUPPORTED: x86_64-pc-linux-gnu-LTO + +#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, int *Src, short Q, short P) { + *Dst = (Src[0] + Src[1]) * (Q + 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 + square<<<1, 1>>>(Ptr, Src, 3, 4); + printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr); + // CHECK: Ptr [[Ptr]], *Ptr: 42 + llvm_omp_target_free_shared(Ptr, DevNo); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
