From 4fd3f70f51718cd464067f647434fb5aada99e74 Mon Sep 17 00:00:00 2001 From: Sourabh Betigeri Date: Wed, 11 Sep 2024 19:04:02 -0700 Subject: [PATCH] SWDEV-484578 SWDEV-484575 SWDEV-484573 SWDEV-483324 SWDEV-483323 - Adds nvidia mappings For hipStreamBatchMemOp, hipGraphAddBatchMemOpNode, hipGraphBatchMemOpNodeSetParams, hipGraphBatchMemOpNodeGetParams, hipGraphExecBatchMemOpNodeSetParams APIs Change-Id: I54a42fe070ee4a9e0ae41718b79f96aed883db5e --- .../nvidia_detail/nvidia_hip_runtime_api.h | 178 ++++++++++++++++++ 1 file changed, 178 insertions(+) diff --git a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index 61cb1703fe..237642174f 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -519,6 +519,146 @@ static inline void hipResourceDesTocudaResourceDes(CUDA_RESOURCE_DESC* a, const a->flags = p->flags; } + +/** Operations for hipStreamBatchMemOp*/ +typedef enum hipStreamBatchMemOpType { + hipStreamMemOpWaitValue32 = 0x1, + hipStreamMemOpWriteValue32 = 0x2, + hipStreamMemOpWaitValue64 = 0x4, + hipStreamMemOpWriteValue64 = 0x5, + hipStreamMemOpBarrier = 0x6, ///< Currently not supported + hipStreamMemOpFlushRemoteWrites = 0x3 ///< Currently not supported +} hipStreamBatchMemOpType; + + + +inline static CUstreamBatchMemOpType hipStreamBatchMemOpType_enumToCUstreamBatchMemOpType( + hipStreamBatchMemOpType memOpType) { + switch (memOpType) { + case hipStreamMemOpWaitValue32: + return CU_STREAM_MEM_OP_WAIT_VALUE_32; + case hipStreamMemOpWriteValue32: + return CU_STREAM_MEM_OP_WRITE_VALUE_32; + case hipStreamMemOpWaitValue64: + return CU_STREAM_MEM_OP_WAIT_VALUE_64; + case hipStreamMemOpWriteValue64: + return CU_STREAM_MEM_OP_WRITE_VALUE_64; + case hipStreamMemOpBarrier: + return CU_STREAM_MEM_OP_BARRIER; + case hipStreamMemOpFlushRemoteWrites: + return CU_STREAM_MEM_OP_FLUSH_REMOTE_WRITES; + default: + return CU_STREAM_MEM_OP_WAIT_VALUE_32; + } +} + +typedef union hipStreamBatchMemOpParams_union { + hipStreamBatchMemOpType operation; + struct hipStreamMemOpWaitValueParams_t{ + hipStreamBatchMemOpType operation; + hipDeviceptr_t address; + union { + uint32_t value; + uint64_t value64; + }; + unsigned int flags; + hipDeviceptr_t alias; ///< Not valid for AMD backend. Initial value is unimportant + } waitValue; + struct hipStreamMemOpWriteValueParams_t{ + hipStreamBatchMemOpType operation; + hipDeviceptr_t address; + union { + uint32_t value; + uint64_t value64; + }; + unsigned int flags; + hipDeviceptr_t alias; ///< Not valid for AMD backend. Initial value is unimportant + } writeValue; + struct hipStreamMemOpFlushRemoteWritesParams_t{ + hipStreamBatchMemOpType operation; + unsigned int flags; + } flushRemoteWrites; ///< Currently not supported on AMD + struct hipStreamMemOpMemoryBarrierParams_t{ + hipStreamBatchMemOpType operation; + unsigned int flags; + } memoryBarrier; ///< Currently not supported on AMD + uint64_t pad[6]; +} hipStreamBatchMemOpParams; +// hipStreamBatchMemOpType + +typedef struct hipBatchMemOpNodeParams { + hipCtx_t ctx; + unsigned int count; + hipStreamBatchMemOpParams* paramArray; + unsigned int flags; +} hipBatchMemOpNodeParams; + +#define hipStreamBatchMemOpType CUstreamBatchMemOpType + +static inline void hipBatchMemOpParamsTocudaBatchMemOpParams(CUstreamBatchMemOpParams* a, + const hipStreamBatchMemOpParams* p, + unsigned int count) { + for (unsigned int i = 0; i < count; i++) { + switch (p[i].operation) { + case hipStreamMemOpWaitValue32: + a[i].operation = CU_STREAM_MEM_OP_WAIT_VALUE_32; + break; + case hipStreamMemOpWriteValue32: + a[i].operation = CU_STREAM_MEM_OP_WRITE_VALUE_32; + break; + case hipStreamMemOpWaitValue64: + a[i].operation = CU_STREAM_MEM_OP_WAIT_VALUE_64; + break; + case hipStreamMemOpWriteValue64: + a[i].operation = CU_STREAM_MEM_OP_WRITE_VALUE_64; + break; + case hipStreamMemOpBarrier: + a[i].operation = CU_STREAM_MEM_OP_BARRIER; + break; + case hipStreamMemOpFlushRemoteWrites: + a[i].operation = CU_STREAM_MEM_OP_FLUSH_REMOTE_WRITES; + break; + default: + a[i].operation = CU_STREAM_MEM_OP_WAIT_VALUE_32; + break; + } + a[i].waitValue.operation = hipStreamBatchMemOpType_enumToCUstreamBatchMemOpType( + p[i].waitValue.operation); + a[i].waitValue.address = p[i].waitValue.address; + a[i].waitValue.value = static_cast(p[i].waitValue.value); + a[i].waitValue.value64 = static_cast(p[i].waitValue.value64); + a[i].waitValue.flags = p[i].waitValue.flags; + a[i].waitValue.alias = (CUdeviceptr)p[i].waitValue.alias; + + a[i].writeValue.operation = hipStreamBatchMemOpType_enumToCUstreamBatchMemOpType( + p[i].waitValue.operation); + a[i].writeValue.address = p[i].writeValue.address; + a[i].writeValue.value = static_cast(p[i].writeValue.value); + a[i].writeValue.value64 = static_cast(p[i].writeValue.value64); + a[i].writeValue.flags = p[i].writeValue.flags; + a[i].writeValue.alias = (CUdeviceptr)p[i].writeValue.alias; + + a[i].flushRemoteWrites.operation = hipStreamBatchMemOpType_enumToCUstreamBatchMemOpType( + p[i].flushRemoteWrites.operation); + a[i].flushRemoteWrites.flags = p[i].flushRemoteWrites.flags; + + a[i].memoryBarrier.operation = hipStreamBatchMemOpType_enumToCUstreamBatchMemOpType( + p[i].memoryBarrier.operation); + a[i].memoryBarrier.flags = p[i].memoryBarrier.flags; + } +} + +static inline void hipBatchMemOpNodeParamsTocudaBatchMemOpNodeParams( + CUDA_BATCH_MEM_OP_NODE_PARAMS* a, + const hipBatchMemOpNodeParams* p) { + CUstreamBatchMemOpParams cuParamArray[p->count]; + hipBatchMemOpParamsTocudaBatchMemOpParams(cuParamArray, p->paramArray, p->count); + a->ctx = (CUcontext)p->ctx; + a->count = p->count; + a->paramArray = cuParamArray; + a->flags = p->flags; +} + typedef struct hip_Memcpy2D { size_t srcXInBytes; size_t srcY; @@ -4016,6 +4156,44 @@ inline static hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, int static_cast(value), flags)); } +inline static hipError_t hipStreamBatchMemOp(hipStream_t stream, unsigned int count, + hipStreamBatchMemOpParams* paramArray, + unsigned int flags) { + CUstreamBatchMemOpParams cuParamArray[count]; + hipBatchMemOpParamsTocudaBatchMemOpParams(cuParamArray, paramArray, count); + return hipCUResultTohipError(cuStreamBatchMemOp(stream, count, cuParamArray, flags)); +} + +inline static hipError_t hipGraphAddBatchMemOpNode(hipGraphNode_t *phGraphNode, hipGraph_t hGraph, + const hipGraphNode_t *dependencies, + size_t numDependencies, + const hipBatchMemOpNodeParams* nodeParams) { + CUDA_BATCH_MEM_OP_NODE_PARAMS cuBatchMemOpNodeParams; + hipBatchMemOpNodeParamsTocudaBatchMemOpNodeParams(&cuBatchMemOpNodeParams, nodeParams); + return hipCUDAErrorTohipError(cuGraphAddBatchMemOpNode(phGraphNode, hGraph, dependencies, + numDependencies, + (const CUDA_BATCH_MEM_OP_NODE_PARAMS*)&cuBatchMemOpNodeParams)); +} + +inline static hipError_t hipGraphBatchMemOpNodeGetParams(hipGraphNode_t hNode, + hipBatchMemOpNodeParams* nodeParams_out) { + return hipCUDAErrorTohipError(cuGraphBatchMemOpNodeGetParams(hNode, nodeParams_out)); +} + +inline static hipError_t hipGraphBatchMemOpNodeSetParams(hipGraphNode_t hNode, + hipBatchMemOpNodeParams* nodeParams) { + return hipCUDAErrorTohipError(cuGraphBatchMemOpNodeSetParams (hNode, + (const CUDA_BATCH_MEM_OP_NODE_PARAMS*)nodeParams)); +} + +inline static hipError_t hipGraphExecBatchMemOpNodeSetParams(hipGraphExec_t hGraphExec, + hipGraphNode_t hNode, + const hipBatchMemOpNodeParams* nodeParams) { + return hipCUDAErrorTohipError(cuGraphExecBatchMemOpNodeSetParams(hGraphExec, hNode, + (const CUDA_BATCH_MEM_OP_NODE_PARAMS*)nodeParams)); +} + + inline static hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* from, const hipGraphNode_t* to, size_t numDependencies) {