asynchronous offload

[lucaparisi91]

• map dependencies
• kernel dependencies ( openmp tasks ? )
• openmp detach functionality

Compiler	Support	Notes
nvhpc 24.5	Yes	No overlapping with data transfers
clang 18.1.8	Yes
cce 16.0.1	Yes	Requires multiple CPU threads for concurrent launches

[lucaparisi91]

Resources:

• https://www.youtube.com/watch?v=XTFCyKW4Z54
• https://www.youtube.com/watch?v=TewAwn_ojgU
• https://www.youtube.com/watch?v=qjkTL2CT37w
• https://www.openmp.org/wp-content/uploads/2021_03_26_Leveraging_Implicit_CUDA_Streams_and_Asynchronous_OpenMP_offload_Features_in_LLVM.pdf

[lucaparisi91]

Async Map

#pragma omp map(to: a[:N] ) nowait

In OpenMP 5.0 a task can be detached, even if not completed. Needs to use API to signal the completion of the task.

omp_event_handle_t *event;

#pragma omp task A 
{
}
#pragma omp task B detach(event)
{
    do_stuff()
    hipStreamAddCallback(stream,callback,&hip_event,0)

}

void callback(hipStream_t stream, hipError_t status, void * cb_dat){
omp_fullfill_event( *(omp_event_handle_t *) cb_data );
}
#pragma omp taskwait;

[lucaparisi91]

The main_single_transfer.cpp file demonstrates using tasking and openmp offload.
It runs in parallel with both nnhpc 24.5 and clang 18.1.8. Below a screenshot for nvidia nvhpc 24.5.

[lucaparisi91]

The main_multiple_transfer.cpp file demonstrate overlapping computation and memory transfer.
This kind-of works with clang 18.1.8 . See the image below. Each task creates a kernel in a new gpu thread.

The nvidia compiler also creates different streams , however there does not seem to be any overlap between the running kernels or with the memory transfer, as per the screenshot below.

[lucaparisi91]

On A2, with cce 16.0.1 , main_single_transfer.cpp runs in serial. However they do run concurrently when using multiple threads. Below a screenshot for 4 threads with main_single_transfer_multiple_threads.cpp. The number of threads is equal to the number of compilers.