aboutsummaryrefslogtreecommitdiff
path: root/clang/lib/CodeGen/CodeGenFunction.cpp
diff options
context:
space:
mode:
authorJohannes Doerfert <johannes@jdoerfert.de>2024-08-12 17:44:58 -0700
committerGitHub <noreply@github.com>2024-08-12 17:44:58 -0700
commit80525dfcde5bf8aae6ab6b0810124ba502de6096 (patch)
treea580f66b88063a1b8570e70fbdf2e276aa9ab69a /clang/lib/CodeGen/CodeGenFunction.cpp
parent887f7002b65f7376c7a5004535bd08c95bdaa8f8 (diff)
downloadllvm-80525dfcde5bf8aae6ab6b0810124ba502de6096.zip
llvm-80525dfcde5bf8aae6ab6b0810124ba502de6096.tar.gz
llvm-80525dfcde5bf8aae6ab6b0810124ba502de6096.tar.bz2
[Offload][CUDA] Allow CUDA kernels to use LLVM/Offload (#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 ```
Diffstat (limited to 'clang/lib/CodeGen/CodeGenFunction.cpp')
-rw-r--r--clang/lib/CodeGen/CodeGenFunction.cpp3
1 files changed, 2 insertions, 1 deletions
diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index 2b2e23f..7adc5a4 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -1013,7 +1013,8 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy,
}
if (FD && (getLangOpts().OpenCL ||
- (getLangOpts().HIP && getLangOpts().CUDAIsDevice))) {
+ ((getLangOpts().HIP || getLangOpts().OffloadViaLLVM) &&
+ getLangOpts().CUDAIsDevice))) {
// Add metadata for a kernel function.
EmitKernelMetadata(FD, Fn);
}