[libomptarget] Add support for target allocators to dynamic cuda RTL

[JonChesterfield]

Summary

[libomptarget] Add support for target allocators to dynamic cuda RTL

Follow on to D102000 which introduced new calls into libcuda. This patch adds
the corresponding entry points to dynamic_cuda, fixing the build for systems
that do not have the cuda toolkit installed.

Function types and enum from https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MEM.html

[pdhaliwal]

LGTM. Thanks for fixing this!

[JonChesterfield]

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

52

Follows the above pattern but I don't see why we need the name_enum identifiers. Maybe one for a later cleanup.

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;
 
+typedef enum CUmemAttach_flags_enum {
+  CU_MEM_ATTACH_GLOBAL = 0x1,
+  CU_MEM_ATTACH_HOST = 0x2,
+  CU_MEM_ATTACH_SINGLE = 0x4,
+} CUmemAttach_flags;
+
 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 cuMemAllocHost(void **, size_t);
+CUresult cuMemAllocManaged(CUdeviceptr *, size_t, unsigned int);
 
+CUresult cuMemcpyDtoDAsync(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 cuMemFree(CUdeviceptr);
+CUresult cuMemFreeHost(void *);
+
 CUresult cuModuleGetFunction(CUfunction *, CUmodule, const char *);
 CUresult cuModuleGetGlobal(CUdeviceptr *, size_t *, CUmodule, const char *);
 
 CUresult cuModuleUnload(CUmodule);
 CUresult cuStreamCreate(CUstream *, unsigned);
 CUresult cuStreamDestroy(CUstream);
 CUresult cuStreamSynchronize(CUstream);
 CUresult cuCtxSetCurrent(CUcontext);

openmp/libomptarget/plugins/cuda/dynamic_cuda/cuda.cpp

 //===--- cuda/dynamic_cuda/cuda.pp ------------------------------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Implement subset of cuda api by calling into cuda library via dlopen
 // Does the dlopen/dlsym calls as part of the call to cuInit
 //
 //===----------------------------------------------------------------------===//
 
 #include "cuda.h"
 #include "Debug.h"
 #include "dlwrap.h"
 
 #include <string>
 #include <unordered_map>
 
 #include <dlfcn.h>
 
 DLWRAP_INTERNAL(cuInit, 1);
 
 DLWRAP(cuCtxGetDevice, 1);
 DLWRAP(cuDeviceGet, 2);
 DLWRAP(cuDeviceGetAttribute, 3);
 DLWRAP(cuDeviceGetCount, 1);
 DLWRAP(cuFuncGetAttribute, 3);
 
 DLWRAP(cuGetErrorString, 2);
 DLWRAP(cuLaunchKernel, 11);
 
 DLWRAP(cuMemAlloc, 2);
-DLWRAP(cuMemcpyDtoDAsync, 4);
+DLWRAP(cuMemAllocHost, 2);
+DLWRAP(cuMemAllocManaged, 3);
 
+DLWRAP(cuMemcpyDtoDAsync, 4);
 DLWRAP(cuMemcpyDtoH, 3);
 DLWRAP(cuMemcpyDtoHAsync, 4);
 DLWRAP(cuMemcpyHtoD, 3);
 DLWRAP(cuMemcpyHtoDAsync, 4);
 
 DLWRAP(cuMemFree, 1);
+DLWRAP(cuMemFreeHost, 1);
 DLWRAP(cuModuleGetFunction, 3);
 DLWRAP(cuModuleGetGlobal, 4);
 
 DLWRAP(cuModuleUnload, 1);
 DLWRAP(cuStreamCreate, 2);
 DLWRAP(cuStreamDestroy, 1);
 DLWRAP(cuStreamSynchronize, 1);
 DLWRAP(cuCtxSetCurrent, 1);