SWDEV-292714 - Added support for hipPointerGetAttribute and

hipDrvPointerGetAttributes

Change-Id: I44a0b5c5d9eb5a9a63406b11e8de83e5cb245a59


[ROCm/clr commit: 760f275861]

projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h

@@ -2978,6 +2978,10 @@ typedef struct hip_api_data_s {
 };
 // hipDrvPointerGetAttributes[('unsigned int', 'numAttributes'), ('hipPointer_attribute*', 'attributes'), ('void**', 'data'), ('hipDeviceptr_t', 'ptr')]
 #define INIT_hipDrvPointerGetAttributes_CB_ARGS_DATA(cb_data) { \
+  cb_data.args.hipDrvPointerGetAttributes.numAttributes = (unsigned int)numAttributes; \
+  cb_data.args.hipDrvPointerGetAttributes.attributes = (hipPointer_attribute*)attributes; \
+  cb_data.args.hipDrvPointerGetAttributes.data = (void**)data; \
+  cb_data.args.hipDrvPointerGetAttributes.ptr = (hipDeviceptr_t)ptr; \
 };
 // hipEventCreate[('hipEvent_t*', 'event')]
 #define INIT_hipEventCreate_CB_ARGS_DATA(cb_data) { \
@@ -4249,6 +4253,9 @@ typedef struct hip_api_data_s {
 };
 // hipPointerGetAttribute[('void*', 'data'), ('hipPointer_attribute', 'attribute'), ('hipDeviceptr_t', 'ptr')]
 #define INIT_hipPointerGetAttribute_CB_ARGS_DATA(cb_data) { \
+  cb_data.args.hipPointerGetAttribute.data = (void*)data; \
+  cb_data.args.hipPointerGetAttribute.attribute = (hipPointer_attribute)attribute; \
+  cb_data.args.hipPointerGetAttribute.ptr = (hipDeviceptr_t)ptr; \
 };
 // hipPointerGetAttributes[('hipPointerAttribute_t*', 'attributes'), ('const void*', 'ptr')]
 #define INIT_hipPointerGetAttributes_CB_ARGS_DATA(cb_data) { \

projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h

@@ -362,6 +362,7 @@ typedef struct cudaArray* hipArray_const_t;
 typedef struct cudaFuncAttributes hipFuncAttributes;
 typedef struct cudaLaunchParams hipLaunchParams;
 #define hipFunction_attribute CUfunction_attribute
+#define hipPointer_attribute CUpointer_attribute
 #define hip_Memcpy2D CUDA_MEMCPY2D
 #define HIP_MEMCPY3D CUDA_MEMCPY3D
 #define hipMemcpy3DParms cudaMemcpy3DParms
@@ -421,6 +422,25 @@ typedef struct cudaResourceViewDesc hipResourceViewDesc;
 #define HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT
 #define HIP_FUNC_ATTRIBUTE_MAX CU_FUNC_ATTRIBUTE_MAX
 
+//Pointer Attributes
+#define HIP_POINTER_ATTRIBUTE_CONTEXT           CU_POINTER_ATTRIBUTE_CONTEXT
+#define HIP_POINTER_ATTRIBUTE_MEMORY_TYPE       CU_POINTER_ATTRIBUTE_MEMORY_TYPE
+#define HIP_POINTER_ATTRIBUTE_DEVICE_POINTER    CU_POINTER_ATTRIBUTE_DEVICE_POINTER
+#define HIP_POINTER_ATTRIBUTE_HOST_POINTER      CU_POINTER_ATTRIBUTE_HOST_POINTER
+#define HIP_POINTER_ATTRIBUTE_P2P_TOKENS        CU_POINTER_ATTRIBUTE_P2P_TOKENS
+#define HIP_POINTER_ATTRIBUTE_SYNC_MEMOPS       CU_POINTER_ATTRIBUTE_SYNC_MEMOPS
+#define HIP_POINTER_ATTRIBUTE_BUFFER_ID         CU_POINTER_ATTRIBUTE_BUFFER_ID
+#define HIP_POINTER_ATTRIBUTE_IS_MANAGED        CU_POINTER_ATTRIBUTE_IS_MANAGED
+#define HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL    CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL
+#define HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE  CU_POINTER_ATTRIBUTE_IS_LEGACY_CUDA_IPC_CAPABLE
+#define HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR  CU_POINTER_ATTRIBUTE_RANGE_START_ADDR
+#define HIP_POINTER_ATTRIBUTE_RANGE_SIZE        CU_POINTER_ATTRIBUTE_RANGE_SIZE
+#define HIP_POINTER_ATTRIBUTE_MAPPED            CU_POINTER_ATTRIBUTE_MAPPED
+#define HIP_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES CU_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES
+#define HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE CU_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE
+#define HIP_POINTER_ATTRIBUTE_ACCESS_FLAGS      CU_POINTER_ATTRIBUTE_ACCESS_FLAGS
+#define HIP_POINTER_ATTRIBUTE_MEMPOOL_HANDLE    CU_POINTER_ATTRIBUTE_MEMPOOL_HANDLE
+
 #if CUDA_VERSION >= 9000
 #define __shfl(...)      __shfl_sync(0xffffffff, __VA_ARGS__)
 #define __shfl_up(...)   __shfl_up_sync(0xffffffff, __VA_ARGS__)
@@ -1731,6 +1751,17 @@ inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attribut
     return err;
 }
 
+inline static hipError_t hipPointerGetAttribute(void* data, hipPointer_attribute attribute,
+                                                hipDeviceptr_t ptr) {
+    return hipCUResultTohipError(cuPointerGetAttribute(data, attribute, ptr));
+}
+
+inline static hipError_t hipDrvPointerGetAttributes(unsigned int numAttributes,
+                                                    hipPointer_attribute* attributes,
+                                                    void** data, hipDeviceptr_t ptr) {
+    return hipCUResultTohipError(cuPointerGetAttributes(numAttributes, attributes, data, ptr));
+}
+
 inline static hipError_t hipMemGetInfo(size_t* free, size_t* total) {
     return hipCUDAErrorTohipError(cudaMemGetInfo(free, total));
 }

projects/clr/hipamd/src/amdhip.def

@@ -288,6 +288,8 @@ hipStreamIsCapturing
 hipStreamBeginCapture
 hipStreamEndCapture
 hipGraphExecDestroy
+hipPointerGetAttribute
+hipDrvPointerGetAttributes
 hipImportExternalSemaphore
 hipSignalExternalSemaphoresAsync
 hipWaitExternalSemaphoresAsync

projects/clr/hipamd/src/hip_event_ipc.cpp

@@ -226,7 +226,7 @@ hipError_t hipIpcGetEventHandle(hipIpcEventHandle_t* handle, hipEvent_t event) {
 }
 
 hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle) {
-  HIP_INIT_API(NONE, event, handle);
+  HIP_INIT_API(hipIpcOpenEventHandle, event, handle);
 #if !defined(_MSC_VER)
   hipError_t hip_err = hipSuccess;
   if (event == nullptr) {

projects/clr/hipamd/src/hip_hcc.def.in

@@ -288,6 +288,8 @@ hipStreamIsCapturing
 hipStreamBeginCapture
 hipStreamEndCapture
 hipGraphExecDestroy
+hipPointerGetAttribute
+hipDrvPointerGetAttributes
 hipImportExternalSemaphore
 hipSignalExternalSemaphoresAsync
 hipWaitExternalSemaphoresAsync

projects/clr/hipamd/src/hip_hcc.map.in

@@ -387,3 +387,11 @@ global:
 local:
     *;
 } hip_4.4;
+
+hip_5.0 {
+global:
+    hipPointerGetAttribute;
+    hipDrvPointerGetAttributes;
+local:
+    *;
+} hip_4.5;

projects/clr/hipamd/src/hip_memory.cpp

@@ -2651,6 +2651,214 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void
   HIP_RETURN(hipErrorInvalidValue);
 }
 
+// ================================================================================================
+hipError_t ihipPointerGetAttributes(void** data, int* attributes,
+                                    int numAttributes, hipDeviceptr_t ptr) {
+
+  size_t offset = 0;
+  amd::Memory* memObj = getMemoryObject(ptr, offset);
+  constexpr uint32_t kManagedAlloc = (CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_ALLOC_HOST_PTR);
+
+  hipError_t status = hipSuccess;
+
+  uint32_t idx = 0;
+  for (int i = 0; i < numAttributes; i++) {
+    switch (attributes[i]) {
+      case HIP_POINTER_ATTRIBUTE_CONTEXT : {
+        status = hipErrorNotSupported;
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_MEMORY_TYPE : {
+        if (memObj) { // checks for host type or device type
+          *reinterpret_cast<uint32_t*>(data[idx]) =
+          ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) &
+            memObj->getMemFlags())? hipMemoryTypeHost : hipMemoryTypeDevice;
+        } else { // checks for array type
+          cl_mem dstMemObj = reinterpret_cast<cl_mem>((static_cast<hipArray*>(ptr))->data);
+          if (!is_valid(dstMemObj)) {
+            return hipErrorInvalidValue;
+          }
+          amd::Image* dstImage = as_amd(dstMemObj)->asImage();
+          if (dstImage){
+            *reinterpret_cast<uint32_t*>(data[idx]) = hipMemoryTypeArray;
+          } else {
+            return hipErrorInvalidValue;
+          }
+        }
+       break;
+      }
+      case HIP_POINTER_ATTRIBUTE_DEVICE_POINTER : {
+        if (memObj) {
+          device::Memory* devMem = memObj->getDeviceMemory(*hip::getCurrentDevice()->devices()[0]);
+
+          //getDeviceMemory can fail, hence validate the sanity of the mem obtained
+          if (nullptr == devMem) {
+            DevLogPrintfError("getDeviceMemory for ptr failed : %p \n", ptr);
+            return hipErrorMemoryAllocation;
+          }
+          *reinterpret_cast<hipDeviceptr_t*>(data[idx]) =
+            reinterpret_cast<hipDeviceptr_t*>(devMem->virtualAddress() + offset);
+        } else {
+          return hipErrorInvalidValue;
+        }
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_HOST_POINTER : {
+        if (memObj) {
+          if ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) & memObj->getMemFlags()) {
+            if (memObj->getHostMem() != nullptr) {
+              // Registered memory
+              *reinterpret_cast<char**>(data[idx]) =
+                static_cast<char*>(memObj->getHostMem()) + offset;
+            } else {
+              // Prepinned memory
+              *reinterpret_cast<char**>(data[idx]) =
+                static_cast<char*>(memObj->getSvmPtr()) + offset;
+            }
+          } else {
+            *reinterpret_cast<char**>(data[idx]) = nullptr;
+          }
+        } else { // Host Memory
+          *reinterpret_cast<char**>(data[idx]) = static_cast<char*>(ptr);
+        }
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_P2P_TOKENS : {
+        // Currently not supported, deprecated in cuda as well
+        status = hipErrorNotSupported;
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_SYNC_MEMOPS : {
+        // This attribute is ideally used in hipPointerSetAttribute, defaults to true
+        *reinterpret_cast<bool*>(data[idx]) = true;
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_BUFFER_ID : {
+        if (memObj) {
+          *reinterpret_cast<uint32_t*>(data[idx]) = memObj->getUniqueId();
+        } else { // ptr passed must be allocated using HIP memory allocation API
+          return hipErrorInvalidValue;
+        }
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_IS_MANAGED : {
+        if (memObj) {
+          *reinterpret_cast<bool*>(data[idx]) =
+               ((memObj->getMemFlags() & kManagedAlloc) == kManagedAlloc) ? true : false;
+        } else {
+          return hipErrorInvalidValue;
+        }
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL : {
+        if (memObj) {
+          *reinterpret_cast<int*>(data[idx]) = memObj->getUserData().deviceId;
+        } else {
+          // for host memory, 100 is returned by default similar to cuda
+          *reinterpret_cast<int*>(data[idx]) = 100;
+        }
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE : {
+        // TODO: Unclear what to be done for this attribute
+        status = hipErrorNotSupported;
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR : {
+        if (memObj) {
+          if (memObj->getHostMem() != nullptr) {
+            *reinterpret_cast<hipDeviceptr_t*>(data[idx]) =
+              static_cast<char*>(memObj->getHostMem());
+          } else {
+            device::Memory* devMem =
+              memObj->getDeviceMemory(*hip::getCurrentDevice()->devices()[0]);
+
+            //getDeviceMemory can fail, hence validate the sanity of the mem obtained
+            if (nullptr == devMem) {
+              DevLogPrintfError("getDeviceMemory for ptr failed : %p \n", ptr);
+              return hipErrorMemoryAllocation;
+            }
+            *reinterpret_cast<hipDeviceptr_t*>(data[idx]) =
+                 reinterpret_cast<char*>(devMem->virtualAddress());
+          }
+        } else { // Host Memory
+          *reinterpret_cast<char**>(data[idx]) = static_cast<char*>(ptr);
+        }
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_RANGE_SIZE : {
+        if (memObj) {
+          *reinterpret_cast<uint32_t*>(data[idx]) = memObj->getSize();
+        } else {
+          // for host memory, 10 is returned by default similar to cuda
+          *reinterpret_cast<uint32_t*>(data[idx]) = 10;
+        }
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_MAPPED : {
+        *reinterpret_cast<bool*>(data[idx]) = (memObj) ? true : false;
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES : {
+        // hipMemAllocationHandleType is not yet supported
+        status = hipErrorNotSupported;
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE : {
+        // GPUDirect RDMA API is not yet supported, hence returning 0
+        LogPrintfWarning("attribute %d is not supported, defaults to 0", attributes[i]);
+        *reinterpret_cast<bool*>(data[idx]) = 0;
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_ACCESS_FLAGS : {
+        if (memObj) {
+          *reinterpret_cast<uint32_t*>(data[idx]) = memObj->getMemFlags() >> 16;
+        } else {
+          *reinterpret_cast<uint32_t*>(data[idx]) = 0;
+        }
+        break;
+      }
+      case HIP_POINTER_ATTRIBUTE_MEMPOOL_HANDLE : {
+        // allocations from mempool are not yet supported, hence returning 0
+        LogPrintfWarning("attribute %d is not supported, defaults to 0", attributes[i]);
+        *reinterpret_cast<bool*>(data[idx]) = 0;
+        break;
+      }
+      default: {
+        LogPrintfError("Invalid attribute: %d ", attributes[i]);
+        status = hipErrorInvalidValue;
+        break;
+      }
+    }
+  ++idx;
+  }
+  return status;
+}
+
+// ================================================================================================
+hipError_t hipPointerGetAttribute(void* data, hipPointer_attribute attribute, hipDeviceptr_t ptr) {
+  HIP_INIT_API(hipPointerGetAttribute, data, attribute, ptr);
+
+  if (ptr == nullptr || data == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  HIP_RETURN(ihipPointerGetAttributes(&data, reinterpret_cast<int*>(&attribute), 1, ptr));
+}
+
+// ================================================================================================
+hipError_t hipDrvPointerGetAttributes(unsigned int numAttributes, hipPointer_attribute* attributes,
+                                      void** data, hipDeviceptr_t ptr) {
+   HIP_INIT_API(hipDrvPointerGetAttributes, numAttributes, attributes, data, ptr);
+
+   if (numAttributes == 0 || attributes == nullptr || data == nullptr || ptr == nullptr) {
+     HIP_RETURN(hipErrorInvalidValue);
+   }
+
+   HIP_RETURN(ihipPointerGetAttributes(data, reinterpret_cast<int*>(attributes),
+                                       numAttributes, ptr));
+}
+
 // ================================================================================================
 hipError_t hipArrayDestroy(hipArray* array) {
   HIP_INIT_API(hipArrayDestroy, array);