From 6bbcf08fc8de3963decedf41808888e1fa586247 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Mon, 27 Sep 2021 19:08:32 +0000 Subject: [PATCH] SWDEV-292714 - Added support for hipPointerGetAttribute and hipDrvPointerGetAttributes Change-Id: I44a0b5c5d9eb5a9a63406b11e8de83e5cb245a59 [ROCm/clr commit: 760f275861b04f42cc6c5d9f7adfdfe0c7a5a366] --- .../include/hip/amd_detail/hip_prof_str.h | 7 + .../nvidia_detail/nvidia_hip_runtime_api.h | 31 +++ projects/clr/hipamd/src/amdhip.def | 2 + projects/clr/hipamd/src/hip_event_ipc.cpp | 2 +- projects/clr/hipamd/src/hip_hcc.def.in | 2 + projects/clr/hipamd/src/hip_hcc.map.in | 8 + projects/clr/hipamd/src/hip_memory.cpp | 208 ++++++++++++++++++ 7 files changed, 259 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index 5041c92e29..13790bac1b 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/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) { \ diff --git a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index 6e7989e1e2..3cc16415eb 100644 --- a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/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)); } diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index d282a0bbc6..4f6835c078 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -288,6 +288,8 @@ hipStreamIsCapturing hipStreamBeginCapture hipStreamEndCapture hipGraphExecDestroy +hipPointerGetAttribute +hipDrvPointerGetAttributes hipImportExternalSemaphore hipSignalExternalSemaphoresAsync hipWaitExternalSemaphoresAsync diff --git a/projects/clr/hipamd/src/hip_event_ipc.cpp b/projects/clr/hipamd/src/hip_event_ipc.cpp index 720701aac3..4dedb4e506 100644 --- a/projects/clr/hipamd/src/hip_event_ipc.cpp +++ b/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) { diff --git a/projects/clr/hipamd/src/hip_hcc.def.in b/projects/clr/hipamd/src/hip_hcc.def.in index 1d2ab341de..c67d537113 100644 --- a/projects/clr/hipamd/src/hip_hcc.def.in +++ b/projects/clr/hipamd/src/hip_hcc.def.in @@ -288,6 +288,8 @@ hipStreamIsCapturing hipStreamBeginCapture hipStreamEndCapture hipGraphExecDestroy +hipPointerGetAttribute +hipDrvPointerGetAttributes hipImportExternalSemaphore hipSignalExternalSemaphoresAsync hipWaitExternalSemaphoresAsync diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index f02649e6a3..9cea9f30d4 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/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; diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 89b553a0a2..91903a7d7d 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/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(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((static_cast(ptr))->data); + if (!is_valid(dstMemObj)) { + return hipErrorInvalidValue; + } + amd::Image* dstImage = as_amd(dstMemObj)->asImage(); + if (dstImage){ + *reinterpret_cast(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(data[idx]) = + reinterpret_cast(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(data[idx]) = + static_cast(memObj->getHostMem()) + offset; + } else { + // Prepinned memory + *reinterpret_cast(data[idx]) = + static_cast(memObj->getSvmPtr()) + offset; + } + } else { + *reinterpret_cast(data[idx]) = nullptr; + } + } else { // Host Memory + *reinterpret_cast(data[idx]) = static_cast(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(data[idx]) = true; + break; + } + case HIP_POINTER_ATTRIBUTE_BUFFER_ID : { + if (memObj) { + *reinterpret_cast(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(data[idx]) = + ((memObj->getMemFlags() & kManagedAlloc) == kManagedAlloc) ? true : false; + } else { + return hipErrorInvalidValue; + } + break; + } + case HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL : { + if (memObj) { + *reinterpret_cast(data[idx]) = memObj->getUserData().deviceId; + } else { + // for host memory, 100 is returned by default similar to cuda + *reinterpret_cast(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(data[idx]) = + static_cast(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(data[idx]) = + reinterpret_cast(devMem->virtualAddress()); + } + } else { // Host Memory + *reinterpret_cast(data[idx]) = static_cast(ptr); + } + break; + } + case HIP_POINTER_ATTRIBUTE_RANGE_SIZE : { + if (memObj) { + *reinterpret_cast(data[idx]) = memObj->getSize(); + } else { + // for host memory, 10 is returned by default similar to cuda + *reinterpret_cast(data[idx]) = 10; + } + break; + } + case HIP_POINTER_ATTRIBUTE_MAPPED : { + *reinterpret_cast(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(data[idx]) = 0; + break; + } + case HIP_POINTER_ATTRIBUTE_ACCESS_FLAGS : { + if (memObj) { + *reinterpret_cast(data[idx]) = memObj->getMemFlags() >> 16; + } else { + *reinterpret_cast(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(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(&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(attributes), + numAttributes, ptr)); +} + // ================================================================================================ hipError_t hipArrayDestroy(hipArray* array) { HIP_INIT_API(hipArrayDestroy, array);