Openmp memory interop

[Waqar-ukaea]

I wanted to open a PR to keep track of where I'm at in terms of getting OpenMP memory interop with CUDA (and eventually GPRT/Vulkan) working. For now this is going to be an internal PR to my fork since its a lot of experimenting so far.

It's a PR onto the batch-query-api branch that I have also been working on as the idea is probably to eventually make use of some of those batch api features to demonstrate functionality of the interop.

I have started to define a super thin wrapper to act as the "interop" (which I've called OpenMP Memory Interop - OMI) layer between OpenMP and CUDA. The idea being that in theory, this wrapper could be extended to support multiple different "backends" in GPU runtime contexts (HIP/CUDA). All it really needs to do is expose a few methods for setting up and freeing device memory. And the idea is that both OpenMP and (eventually GPRT through Vulkan's external memory extension) can take the device pointers returned by the interop layer OMI and use them to work with device data.

[Waqar-ukaea]

Right now, I am trying to compile on my local machine which contains an RTX 2000 ADA and CUDA 12.6 installed. But I am running into some compilation errors whilst using LLVM/Clang 19.1.0 built with OpenMP offload support. I get the following error when attempting to compile my application which contains OpenMP target offload regions and the CUDA memory calls:

[ 99%] Building CXX object tools/CMakeFiles/omp-cuda-gprt-interop-test.dir/omp_cuda_gprt_interop_test.cpp.o
ptxas /tmp/omp_cuda_gprt_interop_test-4cbf6f.s, line 324; error   : Call has wrong number of parameters
ptxas /tmp/omp_cuda_gprt_interop_test-4cbf6f.s, line 324; error   : Type of argument does not match formal parameter '__kmpc_target_init_param_1'
ptxas /tmp/omp_cuda_gprt_interop_test-4cbf6f.s, line 324; error   : Alignment of argument does not match formal parameter '__kmpc_target_init_param_1'
ptxas /tmp/omp_cuda_gprt_interop_test-4cbf6f.s, line 324; error   : Call has wrong number of parameters
ptxas /tmp/omp_cuda_gprt_interop_test-4cbf6f.s, line 384; error   : Call has wrong number of parameters
ptxas /tmp/omp_cuda_gprt_interop_test-4cbf6f.s, line 384; error   : Call has wrong number of parameters
ptxas fatal   : Ptx assembly aborted due to errors
clang++: error: ptxas command failed with exit code 255 (use -v to see invocation)
clang version 19.1.0 (https://github.com/llvm/llvm-project.git a4bf6cd7cfb1a1421ba92bca9d017b49936c55e4)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/waqar/compilers/llvm-dev/bin
Build config: +assertions
clang++: note: diagnostic msg: 
********************

PLEASE ATTACH THE FOLLOWING FILES TO THE BUG REPORT:
Preprocessed source(s) and associated run script(s) are located at:
clang++: note: diagnostic msg: /tmp/omp_cuda_gprt_interop_test-8bc1d4.cpp
clang++: note: diagnostic msg: /tmp/omp_cuda_gprt_interop_test-ec947b.cpp
clang++: note: diagnostic msg: /tmp/omp_cuda_gprt_interop_test-8bc1d4.sh
clang++: note: diagnostic msg: 

********************
make[2]: *** [tools/CMakeFiles/omp-cuda-gprt-interop-test.dir/build.make:79: tools/CMakeFiles/omp-cuda-gprt-interop-test.dir/omp_cuda_gprt_interop_test.cpp.o] Error 1
make[1]: *** [CMakeFiles/Makefile2:3848: tools/CMakeFiles/omp-cuda-gprt-interop-test.dir/all] Error 2
make: *** [Makefile:166: all] Error 2

[Waqar-ukaea]

Switched to LLVM-18.1.8 and my simple OpenMP-CUDA test application looks to be working. Though, the fresh LLVM build/install may have been the real fix tbh.

Where dPtr is a pointer returned from cudamemalloc and the data was created on host before being memcopied to device:

  #pragma omp target teams distribute parallel for is_device_ptr(dPtr)
    for (int i = 0; i < N; ++i) {
      dPtr[i] = i*2.0f;
      if (i == 0) printf("dPtr[%d] = %f \n", i, dPtr[i]);
      if (i == 512) printf("dPtr[%d] = %f \n", i, dPtr[i]);
      if (i == 1023) printf("dPtr[%d] = %f \n", i, dPtr[i]);      
    }

Prints the following:

Successfully created CUDA device pointer: 0x7f2863a00000
Succesfully called cuda memcpy via OMI interop layer
dPtr[512] = 1024.000000 
dPtr[0] = 0.000000 
dPtr[1023] = 2046.000000 

So we're on a pretty good path here where OpenMP and CUDA are interacting completely fine with each other. Some extra work to fully abstract/encapsulate the GPU runtime backend (and extend for HIP) is still required. But i've confirmed that I can get the easy part of the handshake (OpenMP --> vendor specific memory) to work.

Next step is to bring in Vulkan

[pshriwise]

Good stuff here @Waqar-ukaea! I'll pull this and do some testing myself.

I like the name OMI :) Maybe eventually VOMI (Vulkan-OpenMP Memory Interface).