The cuda module operates either in CUBIN or PTX-JIT mode. These modes affect the output kernel executable format, as well as the actions taking place in the compiler and runtime parts of OMPi. 

Both modes have their positive and negative aspects. CUBIN mode removes many runtime overheads but adds a notable delay during kernel compilation. On the other hand, PTX-JIT mode leads to smaller compilation times as PTX file production is much lighter than binary production, however it causes extra overheads during runtime. These overheads are related to the compilation and linking of the PTX file with the appropriate device libraries; they can be, however, minimized by utilizing CUDA JIT cache, a global disk cache for loaded PTX files. 

To make sure the cache is enabled, set the corresponding environment variable:

export CUDA_CACHE_DISABLE=0

OMPi by default produces executable CUBIN files. To enable the PTX-JIT mode, configure OMPi with the --enable-cuda-jit flag:

./configure --prefix=‹install-dir› --enable-cuda-jit LDFLAGS=-rdynamic

The ompicc --devvinfo command should give an output similar to the following (here a Tesla P40 GPU was identified by OMPi):

MODULE [cuda]:
------
OMPi CUDA device module.
Available devices : 1

device id < 1 > { 
 name: Tesla P40 (SM v6.1)
 30 multiprocessors
 128 cores per multiprocessor
 3840 cores in total
 1024 maximum thread block size
 23469184 Kbytes of device global memory
}
------

Here is a sample application that can be used to verify the successful deployment of the cuda module:

#include <stdio.h>
#include <omp.h>

int main(void)
{
    /* 
     * Expected result (assuming that the GPU device ID is 1):
     * Running on CUDA device 
     */
    #pragma omp target device(1)
    {
        if (omp_is_initial_device()) 
            printf("Running on host\n");    
        else 
            printf("Running on CUDA device\n"); 
    }
    return 0;
}

When an application contains several target regions, compilation times may become substantial. One can generate multiple compilation jobs that handle the compilation of different kernels in parallel, by using the -j<number> device option. For example to have 4 parallel kernel compilation jobs, use ompicc as follows:

ompicc --devopt -j4 app.c