[libomptarget][cuda] Call v2 functions explicitly

[JonChesterfield]

Summary

[libomptarget][cuda] Call v2 functions explicitly

rtl.cpp calls functions like cuMemFree that are replaced by a macro
in cuda.h with cuMemFree_v2. This patch changes the source to use
the v2 names consistently.

See also D95104, D95155 for the idea. Alternatives are to use a mixture,
e.g. call the macro names and explictly dlopen the _v2 names, or to keep
the current status where the symbols are replaced by macros in both files

[jdoerfert]

LGTM

[kkwli0]

I am seeing these errors.

/llvm-project/openmp/runtime/src/kmp_runtime.cpp:36:10: fatal error: 'llvm/Support/TimeProfiler.h' file not found
#include "llvm/Support/TimeProfiler.h"
         ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~
1 error generated.
make[5]: *** [openmp/runtime/src/CMakeFiles/omp.dir/kmp_runtime.cpp.o] Error 1
make[5]: *** Waiting for unfinished jobs....
/llvm-project/openmp/libomptarget/plugins/cuda/src/rtl.cpp:469:21: error: use of undeclared identifier 'cuDevicePrimaryCtxRelease_v2'; did you mean 'cuDevicePrimaryCtxRelease'?
        checkResult(cuDevicePrimaryCtxRelease_v2(Device),
                    ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
                    cuDevicePrimaryCtxRelease
/usr/local/cuda/include/cuda.h:3437:18: note: 'cuDevicePrimaryCtxRelease' declared here
CUresult CUDAAPI cuDevicePrimaryCtxRelease(CUdevice dev);
                 ^
/llvm-project/openmp/libomptarget/plugins/cuda/src/rtl.cpp:509:13: error: use of undeclared identifier 'cuDevicePrimaryCtxSetFlags_v2'; did you mean 'cuDevicePrimaryCtxSetFlags'?
      Err = cuDevicePrimaryCtxSetFlags_v2(Device, CU_CTX_SCHED_BLOCKING_SYNC);
            ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~
            cuDevicePrimaryCtxSetFlags
/usr/local/cuda/include/cuda.h:3502:18: note: 'cuDevicePrimaryCtxSetFlags' declared here
CUresult CUDAAPI cuDevicePrimaryCtxSetFlags(CUdevice dev, unsigned int flags);
                 ^
make[4]: *** [openmp/runtime/src/CMakeFiles/omp.dir/all] Error 2
make[4]: *** Waiting for unfinished jobs....

[JonChesterfield]

I think missing TimeProfiler.h is a known problem, independent from this. It can be worked around locally by disabling the profiling.

The error: use of undeclared identifier 'cuDevicePrimaryCtxRelease_v2'; did you mean 'cuDevicePrimaryCtxRelease'? is a hazard though. This means that at least one person (you!) is using a cuda version which is old enough to not use the _v2 appended functions.

We can unblock you by going back to the non-suffixed versions in rtl.cpp and adding the macros back to dynamic_cuda/cuda.h. That will mean the following combinations work:

• directly linking against an old cuda
• directly linking against a new cuda
• dlopen of a new cuda

What version of cuda are you using? It might be useful to know when things changed.

I think we can dlopen cuda, and try to load cuMemcpyDtoHAsync_v2, and then try to load cuMemcpyDtoHAsync on failure. That would get us some degree of robustness to different symbol versions.

[jdoerfert]

In D95274#2520002, @JonChesterfield wrote:

We can unblock you by going back to the non-suffixed versions in rtl.cpp and adding the macros back to dynamic_cuda/cuda.h. That will mean the following combinations work:

Let's do that, we revert this for now.

[JonChesterfield]

Patch at D95367. Can't do a simple revert as we also need to back out e5e448aafa7699c17f78aaffb001b665b607e5ae

[kkwli0]

In D95274#2520002, @JonChesterfield wrote:

What version of cuda are you using? It might be useful to know when things changed.

I think we can dlopen cuda, and try to load cuMemcpyDtoHAsync_v2, and then try to load cuMemcpyDtoHAsync on failure. That would get us some degree of robustness to different symbol versions.

$ cat /usr/local/cuda/version.txt 
CUDA Version 10.2.89

[jdoerfert]

In D95274#2520112, @kkwli0 wrote:

In D95274#2520002, @JonChesterfield wrote:

What version of cuda are you using? It might be useful to know when things changed.

I think we can dlopen cuda, and try to load cuMemcpyDtoHAsync_v2, and then try to load cuMemcpyDtoHAsync on failure. That would get us some degree of robustness to different symbol versions.

$ cat /usr/local/cuda/version.txt 
CUDA Version 10.2.89

This means dyn_cuda won't work for this version (and others I assume). We might need to look for both symbols in dyn_cuda, wdyt Jon?

[JonChesterfield]

In D95274#2520128, @jdoerfert wrote:

This means dyn_cuda won't work for this version (and others I assume). We might need to look for both symbols in dyn_cuda, wdyt Jon?

Based on the macro renaming used to patch application code to work with either, it seems the calling conventions are unchanged.

If so, logic that looks up foo_v2 first, and then tries foo second, and only fails if both are missing seems valid.

The clean way to do that is probably:

• rtl.cpp calls cuMemcpyDtoHAsync
• dynamic_cuda/cuda.h does not have a macro that replaces that with cuMemcpyDtoHAsync_v2
• dynamic_cuda/cuda.cpp implements cuMemcpyDtoHAsync as a call to a function pointer which is initialised with cuMemcpyDtoHAsync_v2 or cuMemcpyDtoHAsync

openmp/libomptarget/plugins/cuda/dynamic_cuda/cuda.h

typedef enum CUfunction_attribute_enum {
   CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK = 0,
 } CUfunction_attribute;
 
 typedef enum CUctx_flags_enum {
   CU_CTX_SCHED_BLOCKING_SYNC = 0x04,
   CU_CTX_SCHED_MASK = 0x07,
 } CUctx_flags;
 
-#define cuMemFree cuMemFree_v2
-#define cuMemAlloc cuMemAlloc_v2
-#define cuMemcpyDtoH cuMemcpyDtoH_v2
-#define cuMemcpyHtoD cuMemcpyHtoD_v2
-#define cuStreamDestroy cuStreamDestroy_v2
-#define cuModuleGetGlobal cuModuleGetGlobal_v2
-#define cuMemcpyDtoHAsync cuMemcpyDtoHAsync_v2
-#define cuMemcpyDtoDAsync cuMemcpyDtoDAsync_v2
-#define cuMemcpyHtoDAsync cuMemcpyHtoDAsync_v2
-#define cuDevicePrimaryCtxRelease cuDevicePrimaryCtxRelease_v2
-#define cuDevicePrimaryCtxSetFlags cuDevicePrimaryCtxSetFlags_v2
-
 CUresult cuCtxGetDevice(CUdevice *);
 CUresult cuDeviceGet(CUdevice *, int);
 CUresult cuDeviceGetAttribute(int *, CUdevice_attribute, CUdevice);
 CUresult cuDeviceGetCount(int *);
 CUresult cuFuncGetAttribute(int *, CUfunction_attribute, CUfunction);
 
 CUresult cuGetErrorString(CUresult, const char **);
 CUresult cuInit(unsigned);
 CUresult cuLaunchKernel(CUfunction, unsigned, unsigned, unsigned, unsigned,
                         unsigned, unsigned, unsigned, CUstream, void **,
                         void **);
 
-CUresult cuMemAlloc(CUdeviceptr *, size_t);
-CUresult cuMemcpyDtoDAsync(CUdeviceptr, CUdeviceptr, size_t, CUstream);
+CUresult cuMemAlloc_v2(CUdeviceptr *, size_t);
+CUresult cuMemcpyDtoDAsync_v2(CUdeviceptr, CUdeviceptr, size_t, CUstream);
 
-CUresult cuMemcpyDtoH(void *, CUdeviceptr, size_t);
-CUresult cuMemcpyDtoHAsync(void *, CUdeviceptr, size_t, CUstream);
-CUresult cuMemcpyHtoD(CUdeviceptr, const void *, size_t);
-CUresult cuMemcpyHtoDAsync(CUdeviceptr, const void *, size_t, CUstream);
+CUresult cuMemcpyDtoH_v2(void *, CUdeviceptr, size_t);
+CUresult cuMemcpyDtoHAsync_v2(void *, CUdeviceptr, size_t, CUstream);
+CUresult cuMemcpyHtoD_v2(CUdeviceptr, const void *, size_t);
+CUresult cuMemcpyHtoDAsync_v2(CUdeviceptr, const void *, size_t, CUstream);
 
-CUresult cuMemFree(CUdeviceptr);
+CUresult cuMemFree_v2(CUdeviceptr);
 CUresult cuModuleGetFunction(CUfunction *, CUmodule, const char *);
-CUresult cuModuleGetGlobal(CUdeviceptr *, size_t *, CUmodule, const char *);
+CUresult cuModuleGetGlobal_v2(CUdeviceptr *, size_t *, CUmodule, const char *);
 
 CUresult cuModuleUnload(CUmodule);
 CUresult cuStreamCreate(CUstream *, unsigned);
-CUresult cuStreamDestroy(CUstream);
+CUresult cuStreamDestroy_v2(CUstream);
 CUresult cuStreamSynchronize(CUstream);
 CUresult cuCtxSetCurrent(CUcontext);
-CUresult cuDevicePrimaryCtxRelease(CUdevice);
+CUresult cuDevicePrimaryCtxRelease_v2(CUdevice);
 CUresult cuDevicePrimaryCtxGetState(CUdevice, unsigned *, int *);
-CUresult cuDevicePrimaryCtxSetFlags(CUdevice, unsigned);
+CUresult cuDevicePrimaryCtxSetFlags_v2(CUdevice, unsigned);
 CUresult cuDevicePrimaryCtxRetain(CUcontext *, CUdevice);
 CUresult cuModuleLoadDataEx(CUmodule *, const void *, unsigned, void *,
                             void **);
 
 CUresult cuDeviceCanAccessPeer(int *, CUdevice, CUdevice);
 CUresult cuCtxEnablePeerAccess(CUcontext, unsigned);
 CUresult cuMemcpyPeerAsync(CUdeviceptr, CUcontext, CUdeviceptr, CUcontext,
                            size_t, CUstream);
 
 #endif

openmp/libomptarget/plugins/cuda/src/rtl.cpp

bool checkResult(CUresult Err, const char *ErrMsg) {
 
   REPORT("%s", ErrMsg);
   CUDA_ERR_STRING(Err);
   return false;
 }
 
 int memcpyDtoD(const void *SrcPtr, void *DstPtr, int64_t Size,
                CUstream Stream) {
-  CUresult Err =
-      cuMemcpyDtoDAsync((CUdeviceptr)DstPtr, (CUdeviceptr)SrcPtr, Size, Stream);
+  CUresult Err = cuMemcpyDtoDAsync_v2((CUdeviceptr)DstPtr, (CUdeviceptr)SrcPtr,
+                                      Size, Stream);
 
   if (Err != CUDA_SUCCESS) {
     REPORT("Error when copying data from device to device. Pointers: src "
            "= " DPxMOD ", dst = " DPxMOD ", size = %" PRId64 "\n",
            DPxPTR(SrcPtr), DPxPTR(DstPtr), Size);
     CUDA_ERR_STRING(Err);
     return OFFLOAD_FAIL;
   }
public:
   ~StreamManagerTy() {
     // Destroy streams
     for (int I = 0; I < NumberOfDevices; ++I) {
       checkResult(cuCtxSetCurrent(DeviceData[I].Context),
                   "Error returned from cuCtxSetCurrent\n");
 
       for (CUstream &S : StreamPool[I]) {
         if (S)
-          checkResult(cuStreamDestroy(S),
-                      "Error returned from cuStreamDestroy\n");
+          checkResult(cuStreamDestroy_v2(S),
+                      "Error returned from cuStreamDestroy_v2\n");
       }
     }
   }
 
   // Get a CUstream from pool. Per-device next stream id always points to the
   // next available CUstream. That means, CUstreams [0, id-1] have been
   // assigned, and [id,] are still available. If there is no CUstream left, we
   // will ask more CUstreams from CUDA RT. Each time a CUstream is assigned,
void *allocate(size_t Size, void *) override {
       if (Size == 0)
         return nullptr;
 
       CUresult Err = cuCtxSetCurrent(DeviceData[DeviceId].Context);
       if (!checkResult(Err, "Error returned from cuCtxSetCurrent\n"))
         return nullptr;
 
       CUdeviceptr DevicePtr;
-      Err = cuMemAlloc(&DevicePtr, Size);
-      if (!checkResult(Err, "Error returned from cuMemAlloc\n"))
+      Err = cuMemAlloc_v2(&DevicePtr, Size);
+      if (!checkResult(Err, "Error returned from cuMemAlloc_v2\n"))
         return nullptr;
 
       return (void *)DevicePtr;
     }
 
     int free(void *TgtPtr) override {
       CUresult Err = cuCtxSetCurrent(DeviceData[DeviceId].Context);
       if (!checkResult(Err, "Error returned from cuCtxSetCurrent\n"))
         return OFFLOAD_FAIL;
 
-      Err = cuMemFree((CUdeviceptr)TgtPtr);
-      if (!checkResult(Err, "Error returned from cuMemFree\n"))
+      Err = cuMemFree_v2((CUdeviceptr)TgtPtr);
+      if (!checkResult(Err, "Error returned from cuMemFree_v2\n"))
         return OFFLOAD_FAIL;
 
       return OFFLOAD_SUCCESS;
     }
   };
 
   /// A vector of device allocators
   std::vector<CUDADeviceAllocatorTy> DeviceAllocators;
~DeviceRTLTy() {
     for (DeviceDataTy &D : DeviceData) {
       // Destroy context
       if (D.Context) {
         checkResult(cuCtxSetCurrent(D.Context),
                     "Error returned from cuCtxSetCurrent\n");
         CUdevice Device;
         checkResult(cuCtxGetDevice(&Device),
                     "Error returned from cuCtxGetDevice\n");
-        checkResult(cuDevicePrimaryCtxRelease(Device),
-                    "Error returned from cuDevicePrimaryCtxRelease\n");
+        checkResult(cuDevicePrimaryCtxRelease_v2(Device),
+                    "Error returned from cuDevicePrimaryCtxRelease_v2\n");
       }
     }
   }
 
   // Check whether a given DeviceId is valid
   bool isValidDeviceId(const int DeviceId) const {
     return DeviceId >= 0 && DeviceId < NumberOfDevices;
   }
int initDevice(const int DeviceId) {
     if (FormerPrimaryCtxIsActive) {
       DP("The primary context is active, no change to its flags\n");
       if ((FormerPrimaryCtxFlags & CU_CTX_SCHED_MASK) !=
           CU_CTX_SCHED_BLOCKING_SYNC)
         DP("Warning the current flags are not CU_CTX_SCHED_BLOCKING_SYNC\n");
     } else {
       DP("The primary context is inactive, set its flags to "
          "CU_CTX_SCHED_BLOCKING_SYNC\n");
-      Err = cuDevicePrimaryCtxSetFlags(Device, CU_CTX_SCHED_BLOCKING_SYNC);
-      if (!checkResult(Err, "Error returned from cuDevicePrimaryCtxSetFlags\n"))
+      Err = cuDevicePrimaryCtxSetFlags_v2(Device, CU_CTX_SCHED_BLOCKING_SYNC);
+      if (!checkResult(Err,
+                       "Error returned from cuDevicePrimaryCtxSetFlags_v2\n"))
         return OFFLOAD_FAIL;
     }
 
     // Retain the per device primary context and save it to use whenever this
     // device is selected.
     Err = cuDevicePrimaryCtxRetain(&DeviceData[DeviceId].Context, Device);
     if (!checkResult(Err, "Error returned from cuDevicePrimaryCtxRetain\n"))
       return OFFLOAD_FAIL;
for (const __tgt_offload_entry *E = HostBegin; E != HostEnd; ++E) {
         DP("Invalid binary: host entry '<null>' (size = %zd)...\n", E->size);
         return nullptr;
       }
 
       if (E->size) {
         __tgt_offload_entry Entry = *E;
         CUdeviceptr CUPtr;
         size_t CUSize;
-        Err = cuModuleGetGlobal(&CUPtr, &CUSize, Module, E->name);
+        Err = cuModuleGetGlobal_v2(&CUPtr, &CUSize, Module, E->name);
         // We keep this style here because we need the name
         if (Err != CUDA_SUCCESS) {
           REPORT("Loading global '%s' Failed\n", E->name);
           CUDA_ERR_STRING(Err);
           return nullptr;
         }
 
         if (CUSize != E->size) {
for (const __tgt_offload_entry *E = HostBegin; E != HostEnd; ++E) {
         // the below condition should be changed to explicitly
         // check for to and link variables types:
         // (RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && (e->flags &
         // OMP_DECLARE_TARGET_LINK || e->flags == OMP_DECLARE_TARGET_TO))
         if (RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
           // If unified memory is present any target link or to variables
           // can access host addresses directly. There is no longer a
           // need for device copies.
-          cuMemcpyHtoD(CUPtr, E->addr, sizeof(void *));
+          cuMemcpyHtoD_v2(CUPtr, E->addr, sizeof(void *));
           DP("Copy linked variable host address (" DPxMOD
              ") to device address (" DPxMOD ")\n",
              DPxPTR(*((void **)E->addr)), DPxPTR(CUPtr));
         }
 
         addOffloadEntry(DeviceId, Entry);
 
         continue;
for (const __tgt_offload_entry *E = HostBegin; E != HostEnd; ++E) {
       // default value GENERIC (in case symbol is missing from cubin file)
       int8_t ExecModeVal = ExecutionModeType::GENERIC;
       std::string ExecModeNameStr(E->name);
       ExecModeNameStr += "_exec_mode";
       const char *ExecModeName = ExecModeNameStr.c_str();
 
       CUdeviceptr ExecModePtr;
       size_t CUSize;
-      Err = cuModuleGetGlobal(&ExecModePtr, &CUSize, Module, ExecModeName);
+      Err = cuModuleGetGlobal_v2(&ExecModePtr, &CUSize, Module, ExecModeName);
       if (Err == CUDA_SUCCESS) {
         if (CUSize != sizeof(int8_t)) {
           DP("Loading global exec_mode '%s' - size mismatch (%zd != %zd)\n",
              ExecModeName, CUSize, sizeof(int8_t));
           return nullptr;
         }
 
-        Err = cuMemcpyDtoH(&ExecModeVal, ExecModePtr, CUSize);
+        Err = cuMemcpyDtoH_v2(&ExecModeVal, ExecModePtr, CUSize);
         if (Err != CUDA_SUCCESS) {
           REPORT("Error when copying data from device to host. Pointers: "
                  "host = " DPxMOD ", device = " DPxMOD ", size = %zd\n",
                  DPxPTR(&ExecModeVal), DPxPTR(ExecModePtr), CUSize);
           CUDA_ERR_STRING(Err);
           return nullptr;
         }
 
#ifdef OMPTARGET_DEBUG
       if (const char *EnvStr = getenv("LIBOMPTARGET_DEVICE_RTL_DEBUG"))
         DeviceEnv.debug_level = std::stoi(EnvStr);
 #endif
 
       const char *DeviceEnvName = "omptarget_device_environment";
       CUdeviceptr DeviceEnvPtr;
       size_t CUSize;
 
-      Err = cuModuleGetGlobal(&DeviceEnvPtr, &CUSize, Module, DeviceEnvName);
+      Err = cuModuleGetGlobal_v2(&DeviceEnvPtr, &CUSize, Module, DeviceEnvName);
       if (Err == CUDA_SUCCESS) {
         if (CUSize != sizeof(DeviceEnv)) {
           REPORT(
               "Global device_environment '%s' - size mismatch (%zu != %zu)\n",
               DeviceEnvName, CUSize, sizeof(int32_t));
           CUDA_ERR_STRING(Err);
           return nullptr;
         }
 
-        Err = cuMemcpyHtoD(DeviceEnvPtr, &DeviceEnv, CUSize);
+        Err = cuMemcpyHtoD_v2(DeviceEnvPtr, &DeviceEnv, CUSize);
         if (Err != CUDA_SUCCESS) {
           REPORT("Error when copying data from host to device. Pointers: "
                  "host = " DPxMOD ", device = " DPxMOD ", size = %zu\n",
                  DPxPTR(&DeviceEnv), DPxPTR(DeviceEnvPtr), CUSize);
           CUDA_ERR_STRING(Err);
           return nullptr;
         }
 
int dataSubmit(const int DeviceId, const void *TgtPtr, const void *HstPtr,
     assert(AsyncInfoPtr && "AsyncInfoPtr is nullptr");
 
     CUresult Err = cuCtxSetCurrent(DeviceData[DeviceId].Context);
     if (!checkResult(Err, "Error returned from cuCtxSetCurrent\n"))
       return OFFLOAD_FAIL;
 
     CUstream Stream = getStream(DeviceId, AsyncInfoPtr);
 
-    Err = cuMemcpyHtoDAsync((CUdeviceptr)TgtPtr, HstPtr, Size, Stream);
+    Err = cuMemcpyHtoDAsync_v2((CUdeviceptr)TgtPtr, HstPtr, Size, Stream);
     if (Err != CUDA_SUCCESS) {
       REPORT("Error when copying data from host to device. Pointers: host "
              "= " DPxMOD ", device = " DPxMOD ", size = %" PRId64 "\n",
              DPxPTR(HstPtr), DPxPTR(TgtPtr), Size);
       CUDA_ERR_STRING(Err);
       return OFFLOAD_FAIL;
     }
 
     return OFFLOAD_SUCCESS;
   }
 
   int dataRetrieve(const int DeviceId, void *HstPtr, const void *TgtPtr,
                    const int64_t Size, __tgt_async_info *AsyncInfoPtr) const {
     assert(AsyncInfoPtr && "AsyncInfoPtr is nullptr");
 
     CUresult Err = cuCtxSetCurrent(DeviceData[DeviceId].Context);
     if (!checkResult(Err, "Error returned from cuCtxSetCurrent\n"))
       return OFFLOAD_FAIL;
 
     CUstream Stream = getStream(DeviceId, AsyncInfoPtr);
 
-    Err = cuMemcpyDtoHAsync(HstPtr, (CUdeviceptr)TgtPtr, Size, Stream);
+    Err = cuMemcpyDtoHAsync_v2(HstPtr, (CUdeviceptr)TgtPtr, Size, Stream);
     if (Err != CUDA_SUCCESS) {
       REPORT("Error when copying data from device to host. Pointers: host "
              "= " DPxMOD ", device = " DPxMOD ", size = %" PRId64 "\n",
              DPxPTR(HstPtr), DPxPTR(TgtPtr), Size);
       CUDA_ERR_STRING(Err);
       return OFFLOAD_FAIL;
     }