[all-commits] [llvm/llvm-project] 80525d: [Offload][CUDA] Allow CUDA kernels to use LLVM/Off...

[Johannes Doerfert via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 80525dfcde5bf8aae6ab6b0810124ba502de6096
      https://github.com/llvm/llvm-project/commit/80525dfcde5bf8aae6ab6b0810124ba502de6096
  Author: Johannes Doerfert <johannes at jdoerfert.de>
  Date:   2024-08-12 (Mon, 12 Aug 2024)

  Changed paths:
    M clang/include/clang/Basic/LangOptions.def
    M clang/include/clang/Driver/Options.td
    M clang/lib/CodeGen/CGCUDANV.cpp
    M clang/lib/CodeGen/CodeGenFunction.cpp
    M clang/lib/Driver/Driver.cpp
    M clang/lib/Driver/ToolChains/Clang.cpp
    M clang/lib/Driver/ToolChains/CommonArgs.cpp
    M clang/lib/Driver/ToolChains/Cuda.cpp
    M clang/lib/Headers/CMakeLists.txt
    A clang/lib/Headers/llvm_offload_wrappers/__llvm_offload.h
    A clang/lib/Headers/llvm_offload_wrappers/__llvm_offload_device.h
    A clang/lib/Headers/llvm_offload_wrappers/__llvm_offload_host.h
    M clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
    M clang/lib/Sema/SemaCUDA.cpp
    M clang/test/CodeGenCUDA/Inputs/cuda.h
    A clang/test/CodeGenCUDA/offload_via_llvm.cu
    A clang/test/Driver/cuda-via-liboffload.cu
    M offload/include/Shared/APITypes.h
    M offload/include/omptarget.h
    M offload/plugins-nextgen/common/src/PluginInterface.cpp
    M offload/src/CMakeLists.txt
    A offload/src/KernelLanguage/API.cpp
    M offload/src/exports
    M offload/test/lit.cfg
    A offload/test/offloading/CUDA/basic_launch.cu
    A offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
    A offload/test/offloading/CUDA/basic_launch_multi_arg.cu
    A offload/test/offloading/CUDA/kernel_tu.cu.inc
    A offload/test/offloading/CUDA/launch_tu.cu

  Log Message:
  -----------
  [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
```



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications