From 607ffbef03ffb8d717f957459c78a47bab29f5ce Mon Sep 17 00:00:00 2001 From: "Manocha, Rahul" Date: Fri, 11 Jul 2025 14:27:46 -0700 Subject: [PATCH 01/11] SWDEV-533787 - Add cuda mapping for hipMemGetHandleForAddressRange (#14) Co-authored-by: Rahul Manocha --- .../nvidia_detail/nvidia_hip_runtime_api.h | 19 +++++++++++++++++++ 1 file changed, 19 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 d9774dddfb..65d8dbf900 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -300,6 +300,12 @@ typedef enum cudaMemRangeAttribute hipMemRangeAttribute; #define hipMemRangeAttributeAccessedBy cudaMemRangeAttributeAccessedBy #define hipMemRangeAttributeLastPrefetchLocation cudaMemRangeAttributeLastPrefetchLocation +#if CUDA_VERSION >= CUDA_12000 +typedef enum CUmemRangeHandleType_enum hipMemRangeHandleType; +#define hipMemRangeHandleTypeDmaBufFd CU_MEM_RANGE_HANDLE_TYPE_DMA_BUF_FD +#define hipMemRangeHandleTypeMax CU_MEM_RANGE_HANDLE_TYPE_MAX +#endif + #define hipSurfaceBoundaryMode cudaSurfaceBoundaryMode #define hipBoundaryModeZero cudaBoundaryModeZero #define hipBoundaryModeTrap cudaBoundaryModeTrap @@ -2998,6 +3004,16 @@ inline static hipError_t hipMemUnmap(hipDeviceptr_t ptr, size_t size) { } #endif // CUDA_VERSION >= CUDA_10020 +#if CUDA_VERSION >= CUDA_12000 +inline static hipError_t hipMemGetHandleForAddressRange(void* handle, hipDeviceptr_t dptr, + size_t size, + hipMemRangeHandleType handleType, + unsigned long long flags) { + return hipCUResultTohipError(cuMemGetHandleForAddressRange(handle, dptr, size, handleType, + flags)); +} +#endif + inline static hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, const void* func, int blockSize, @@ -4830,6 +4846,9 @@ inline static hipError_t hipMemcpy2DArrayToArray(hipArray_t dst, size_t wOffsetD inline static hipError_t hipSetValidDevices(int* device_arr, int len) { return hipCUDAErrorTohipError(cudaSetValidDevices(device_arr, len)); } + + + #endif //__CUDACC__ #endif // HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_RUNTIME_API_H From f638aa91ebcd373ef51e0538aa90f7051388fa83 Mon Sep 17 00:00:00 2001 From: "Stojiljkovic, Vladana" Date: Mon, 14 Jul 2025 11:02:37 +0200 Subject: [PATCH 02/11] SWDEV-541362 - Add nullptr check for memsetParams (#22) --- .../nvidia_detail/nvidia_hip_runtime_api.h | 42 ++++++++++++------- 1 file changed, 28 insertions(+), 14 deletions(-) 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 65d8dbf900..06ba5cf7f9 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -4724,12 +4724,19 @@ inline static void hipMemsetParamsToCUDAMemsetNodeParams(CUDA_MEMSET_NODE_PARAMS } inline static hipError_t hipDrvGraphAddMemsetNode(hipGraphNode_t* phGraphNode, hipGraph_t hGraph, - const hipGraphNode_t* dependencies, size_t numDependencies, - const hipMemsetParams* memsetParams, hipCtx_t ctx) { - CUDA_MEMSET_NODE_PARAMS cuMemsetParams; - hipMemsetParamsToCUDAMemsetNodeParams(&cuMemsetParams, memsetParams); - return hipCUResultTohipError(cuGraphAddMemsetNode(phGraphNode, hGraph, dependencies, numDependencies, - &cuMemsetParams, ctx)); + const hipGraphNode_t* dependencies, + size_t numDependencies, + const hipMemsetParams* memsetParams, hipCtx_t ctx) +{ + if (memsetParams == nullptr) { + return hipCUResultTohipError( + cuGraphAddMemsetNode(phGraphNode, hGraph, dependencies, numDependencies, nullptr, ctx)); + } else { + CUDA_MEMSET_NODE_PARAMS cuMemsetParams; + hipMemsetParamsToCUDAMemsetNodeParams(&cuMemsetParams, memsetParams); + return hipCUResultTohipError(cuGraphAddMemsetNode(phGraphNode, hGraph, dependencies, + numDependencies, &cuMemsetParams, ctx)); + } } inline static hipError_t hipDrvGraphAddMemcpyNode(hipGraphNode_t* phGraphNode, hipGraph_t hGraph, @@ -4787,14 +4794,21 @@ inline static hipError_t hipDrvGraphExecMemcpyNodeSetParams(hipGraphExec_t hGrap } } -inline static hipError_t hipDrvGraphExecMemsetNodeSetParams( - hipGraphExec_t hGraphExec, hipGraphNode_t hNode, const hipMemsetParams* memsetParams, - hipCtx_t ctx) { - CUDA_MEMSET_NODE_PARAMS cuMemsetParams; - hipMemsetParamsToCUDAMemsetNodeParams(&cuMemsetParams, memsetParams); - return hipCUResultTohipError( - cuGraphExecMemsetNodeSetParams(hGraphExec, hNode, &cuMemsetParams, ctx)); - } +inline static hipError_t hipDrvGraphExecMemsetNodeSetParams(hipGraphExec_t hGraphExec, + hipGraphNode_t hNode, + const hipMemsetParams* memsetParams, + hipCtx_t ctx) +{ + if (memsetParams == nullptr) { + return hipCUResultTohipError( + cuGraphExecMemsetNodeSetParams(hGraphExec, hNode, nullptr, ctx)); + } else { + CUDA_MEMSET_NODE_PARAMS cuMemsetParams; + hipMemsetParamsToCUDAMemsetNodeParams(&cuMemsetParams, memsetParams); + return hipCUResultTohipError( + cuGraphExecMemsetNodeSetParams(hGraphExec, hNode, &cuMemsetParams, ctx)); + } +} #endif #if CUDA_VERSION >= CUDA_11040 From 98ef2f64b55a11c2d2ba8352bfc7b76a87fac701 Mon Sep 17 00:00:00 2001 From: "Dittakavi, Satyanvesh" Date: Tue, 12 Aug 2025 16:30:30 +0530 Subject: [PATCH 03/11] SWDEV-545947 - Add Nvidia mapping for hipSetValidDevices (#29) --- hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h | 4 ++++ 1 file changed, 4 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 06ba5cf7f9..ee4ed49e41 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -2097,6 +2097,10 @@ inline static hipError_t hipSetDevice(int device) { return hipCUDAErrorTohipError(cudaSetDevice(device)); } +inline static hipError_t hipSetValidDevices(int* device_arr, int len) { + return hipCUDAErrorTohipError(cudaSetValidDevices(device_arr, len)); +} + inline static hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop) { if (prop == NULL) { From f1838947320e273ab34bf72fea41f3ca531b74c9 Mon Sep 17 00:00:00 2001 From: "GunaShekar, Ajay" Date: Wed, 13 Aug 2025 20:28:35 -0700 Subject: [PATCH 04/11] SWDEV-546179 - nvidia impl hipModuleGetFunctionCount (#25) --- hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h | 4 ++++ 1 file changed, 4 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 ee4ed49e41..04a3ba96f3 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -3455,6 +3455,10 @@ inline static hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule return hipCUResultTohipError(cuModuleGetFunction(function, module, kname)); } +inline static hipError_t hipModuleGetFunctionCount(unsigned int* count, hipModule_t mod) { + return hipCUResultTohipError(cuModuleGetFunctionCount(count, mod)); +} + inline static hipError_t hipModuleGetTexRef(hipTexRef* pTexRef, hipModule_t hmod, const char* name){ return hipCUResultTohipError(cuModuleGetTexRef(pTexRef, hmod, name)); } From 5107c35166a0f2e2a887019d7f6cb5f8f7b233e0 Mon Sep 17 00:00:00 2001 From: "GunaShekar, Ajay" Date: Fri, 15 Aug 2025 12:37:44 -0700 Subject: [PATCH 05/11] SWDEV-546178 - hipModuleLoadFatBinary nvidia impl (#26) --- hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h | 4 ++++ 1 file changed, 4 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 04a3ba96f3..edf29bdc4d 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -3446,6 +3446,10 @@ inline static hipError_t hipModuleLoad(hipModule_t* module, const char* fname) { return hipCUResultTohipError(cuModuleLoad(module, fname)); } +inline static hipError_t hipModuleLoadFatBinary(hipModule_t* module, const void* fatbin) { + return hipCUResultTohipError(cuModuleLoadFatBinary(module, fatbin)); +} + inline static hipError_t hipModuleUnload(hipModule_t hmod) { return hipCUResultTohipError(cuModuleUnload(hmod)); } From 91a7a8d7674acba10edfe1dc23b2d302ddfbd1b9 Mon Sep 17 00:00:00 2001 From: "Manocha, Rahul" Date: Fri, 15 Aug 2025 13:21:52 -0700 Subject: [PATCH 06/11] SWDEV-545952 - cuda mappings for StreamGet/Set Attr (#32) Co-authored-by: Rahul Manocha --- .../nvidia_detail/nvidia_hip_runtime_api.h | 32 +++++++++++++++++++ 1 file changed, 32 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 edf29bdc4d..2be48e3530 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -28,6 +28,7 @@ THE SOFTWARE. #include #include +#include #include #define CUDA_9000 9000 @@ -1818,6 +1819,26 @@ typedef enum cudaKernelNodeAttrID hipKernelNodeAttrID; #define hipKernelNodeAttributeCooperative cudaKernelNodeAttributeCooperative #define hipKernelNodeAttributePriority cudaKernelNodeAttributePriority +#if CUDA_VERSION >= CUDA_12000 +typedef union cudaStreamAttrValue hipStreamAttrValue; +typedef enum cudaStreamAttrID hipStreamAttrID; +#define hipStreamAttributeAccessPolicyWindow cudaStreamAttributeAccessPolicyWindow +#define hipStreamAttributeSynchronizationPolicy cudaStreamAttributeSynchronizationPolicy +#define hipStreamAttributeMemSyncDomainMap cudaStreamAttributeMemSyncDomainMap +#define hipStreamAttributeMemSyncDomain cudaStreamAttributeMemSyncDomain +#define hipStreamAttributePriority cudaStreamAttributePriority + +typedef enum cudaSynchronizationPolicy hipSynchronizationPolicy; +#define hipSyncPolicyAuto cudaSyncPolicyAuto +#define hipSyncPolicySpin cudaSyncPolicySpin +#define hipSyncPolicyYield cudaSyncPolicyYield +#define hipSyncPolicyBlockingSync cudaSyncPolicyBlockingSync + +typedef enum cudaLaunchMemSyncDomain hipLaunchMemSyncDomain; +#define hipLaunchMemSyncDomainDefault cudaLaunchMemSyncDomainDefault +#define hipLaunchMemSyncDomainRemote cudaLaunchMemSyncDomainRemote +#endif + #if CUDA_VERSION >= CUDA_12000 typedef enum cudaGraphInstantiateResult hipGraphInstantiateResult; #define hipGraphInstantiateSuccess cudaGraphInstantiateSuccess @@ -3215,6 +3236,16 @@ inline static hipError_t hipStreamGetDevice(hipStream_t stream, hipDevice_t* dev return hipCUResultTohipError(cuCtxPopCurrent(&context)); } +inline static hipError_t hipStreamSetAttribute(hipStream_t stream, hipStreamAttrID attr, + const hipStreamAttrValue* value) { + return hipCUDAErrorTohipError(cudaStreamSetAttribute(stream, attr, value)); +} + +inline static hipError_t hipStreamGetAttribute(hipStream_t stream, hipStreamAttrID attr, + hipStreamAttrValue* value_out) { + return hipCUDAErrorTohipError(cudaStreamGetAttribute(stream, attr, value_out)); +} + inline static hipError_t hipDriverGetVersion(int* driverVersion) { return hipCUDAErrorTohipError(cudaDriverGetVersion(driverVersion)); } @@ -4875,6 +4906,7 @@ inline static hipError_t hipSetValidDevices(int* device_arr, int len) { + #endif //__CUDACC__ #endif // HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_RUNTIME_API_H From 60ecc06a1c8e7673a966c32de4199e286ad76f4a Mon Sep 17 00:00:00 2001 From: "Manocha, Rahul" Date: Fri, 15 Aug 2025 13:22:27 -0700 Subject: [PATCH 07/11] SWDEV-546295 - cuda mappings for hipMemcpy3DBatchAsync and hipMemcpyPeer (#33) Co-authored-by: Rahul Manocha --- .../nvidia_detail/nvidia_hip_runtime_api.h | 43 +++++++++++++++++++ 1 file changed, 43 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 2be48e3530..36c2ee9efd 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -872,6 +872,29 @@ static inline void hipMemcpy2DTocudaMemcpy2D(CUDA_MEMCPY2D* a, const hip_Memcpy2 a->Height = (size_t)p->Height; } +#if CUDA_VERSION >= CUDA_12020 +typedef enum cudaMemcpyFlags hipMemcpyFlags; +#define hipMemcpyFlagDefault cudaMemcpyFlagDefault +#define hipMemcpyFlagPreferOverlapWithCompute cudaMemcpyFlagPreferOverlapWithCompute + +typedef enum cudaMemcpySrcAccessOrder hipMemcpySrcAccessOrder; +#define hipMemcpySrcAccessOrderInvalid cudaMemcpySrcAccessOrderInvalid +#define hipMemcpySrcAccessOrderStream cudaMemcpySrcAccessOrderStream +#define hipMemcpySrcAccessOrderDuringApiCall cudaMemcpySrcAccessOrderDuringApiCall +#define hipMemcpySrcAccessOrderAny cudaMemcpySrcAccessOrderAny +#define hipMemcpySrcAccessOrderMax cudaMemcpySrcAccessOrderMax + +typedef struct cudaMemcpyAttributes hipMemcpyAttributes; +typedef enum cudaMemcpy3DOperandType hipMemcpy3DOperandType; +#define hipMemcpyOperandTypePointer cudaMemcpyOperandTypePointer +#define hipMemcpyOperandTypeArray cudaMemcpyOperandTypeArray +#define hipMemcpyOperandTypeMax cudaMemcpyOperandTypeMax + +typedef struct cudaOffset3D hipOffset3D; +typedef struct cudaMemcpy3DOperand hipMemcpy3DOperand; +typedef struct cudaMemcpy3DBatchOp hipMemcpy3DBatchOp; +typedef struct cudaMemcpy3DPeerParms hipMemcpy3DPeerParms; +#endif #define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault @@ -2354,6 +2377,26 @@ inline static hipError_t hipMemcpy2DToArrayAsync(hipArray_t dst, size_t wOffset, stream)); } +inline static hipError_t hipMemcpyBatchAsync(void **dsts, void **srcs, size_t *sizes, size_t count, + hipMemcpyAttributes *attrs, size_t *attrsIdxs, + size_t numAttrs, size_t *failIdx, hipStream_t stream) { + return hipCUDAErrorTohipError(cudaMemcpyBatchAsync(dsts, srcs, sizes, count, attrs, attrsIdxs, + numAttrs, failIdx, stream)); +} + +inline static hipError_t hipMemcpy3DBatchAsync(size_t numOps, hipMemcpy3DBatchOp *opList, + size_t *failIdx, unsigned long long flags, + hipStream_t stream) { + return hipCUDAErrorTohipError(cudaMemcpy3DBatchAsync(numOps, opList, failIdx, flags, stream)); +} +inline static hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms *p) { + return hipCUDAErrorTohipError(cudaMemcpy3DPeer(p)); +} +inline static hipError_t hipMemcpy3DPeerAsync(hipMemcpy3DPeerParms *p, hipStream_t stream) { + return hipCUDAErrorTohipError(cudaMemcpy3DPeerAsync(p, stream)); +} + + __HIP_DEPRECATED inline static hipError_t hipMemcpyToArray(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind) { From a7580daf57cfe886d95c9523d0111b58db499dd5 Mon Sep 17 00:00:00 2001 From: "Manocha, Rahul" Date: Fri, 15 Aug 2025 13:22:56 -0700 Subject: [PATCH 08/11] SWDEV-546328 - cuda mappings for hipMemsetD2D8 (#34) Co-authored-by: Rahul Manocha --- .../nvidia_detail/nvidia_hip_runtime_api.h | 33 +++++++++++++++++++ 1 file changed, 33 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 36c2ee9efd..01dc583edc 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -2548,6 +2548,39 @@ inline static hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int valu return hipCUDAErrorTohipError(cudaMemset3DAsync(pitchedDevPtr, value, extent, stream)); } +inline static hipError_t hipMemsetD2D8(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, + size_t width, size_t height) { + return hipCUResultTohipError(cuMemsetD2D8(dst, dstPitch, value, width, height)); +} +inline static hipError_t hipMemsetD2D8Async(hipDeviceptr_t dst, size_t dstPitch, + unsigned char value, size_t width, size_t height, + hipStream_t stream) { + return hipCUResultTohipError(cuMemsetD2D8Async(dst, dstPitch, value, width, height, stream)); +} + +inline static hipError_t hipMemsetD2D16(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, + size_t width, size_t height) { + return hipCUResultTohipError(cuMemsetD2D16(dst, dstPitch, value, width, height)); +} + +inline static hipError_t hipMemsetD2D16Async(hipDeviceptr_t dst, size_t dstPitch, + unsigned short value, size_t width, size_t height, + hipStream_t stream) { + return hipCUResultTohipError(cuMemsetD2D16Async(dst, dstPitch, value, width, height, + stream)); +} + +inline static hipError_t hipMemsetD2D32(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, + size_t width, size_t height) { + return hipCUResultTohipError(cuMemsetD2D32(dst, dstPitch, value, width, height)); +} + +inline static hipError_t hipMemsetD2D32Async(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, + size_t width, size_t height, hipStream_t stream) { + return hipCUResultTohipError(cuMemsetD2D32Async(dst, dstPitch, value, width, height, + stream)); +} + inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int device) { if (p_prop == NULL) { return hipErrorInvalidValue; From 87f262a56faea6733a4cfd62c093091993b35ac7 Mon Sep 17 00:00:00 2001 From: "Li, Todd tiantuo" Date: Fri, 15 Aug 2025 22:38:34 -0700 Subject: [PATCH 09/11] SWDEV-546285 - add hipGetDriverEntryPoint (#35) --- .../hip/nvidia_detail/nvidia_hip_runtime_api.h | 16 ++++++++++++++++ 1 file changed, 16 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 01dc583edc..047538fab1 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -1000,6 +1000,15 @@ typedef enum cudaGraphInstantiateFlags hipGraphInstantiateFlags; #define hipGraphInstantiateFlagDeviceLaunch cudaGraphInstantiateFlagDeviceLaunch #define hipGraphInstantiateFlagUseNodePriority cudaGraphInstantiateFlagUseNodePriority +typedef enum cudaDriverEntryPointQueryResult hipDriverEntryPointQueryResult; +#define hipDriverEntryPointSuccess cudaDriverEntryPointSuccess +#define hipDriverEntryPointSymbolNotFound cudaDriverEntryPointSymbolNotFound +#define hipDriverEntryPointVersionNotSufficent cudaDriverEntryPointVersionNotSufficent + +#define hipEnableDefault cudaEnableDefault +#define hipEnableLegacyStream cudaEnableLegacyStream +#define hipEnablePerThreadDefaultStream cudaEnablePerThreadDefaultStream + inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) { switch (cuError) { case cudaSuccess: @@ -3722,6 +3731,13 @@ inline static hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_ return hipCUDAErrorTohipError(cudaGetChannelDesc(desc,array)); } +inline static hipError_t hipGetDriverEntryPoint(const char* symbol, void** funcPtr, + unsigned long long flags, + hipDriverEntryPointQueryResult* driverStatus) +{ + return hipCUDAErrorTohipError(cudaGetDriverEntryPoint(symbol, funcPtr, flags, driverStatus)); +} + inline static hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDim, void** kernelParams, unsigned int sharedMemBytes, hipStream_t stream) { From bf0302a4944319450e1a26646d82581ff08a327b Mon Sep 17 00:00:00 2001 From: "Betigeri, Sourabh" Date: Fri, 15 Aug 2025 22:44:41 -0700 Subject: [PATCH 10/11] =?UTF-8?q?SWDEV-546293=20-=20Adds=20nvidia=20mappin?= =?UTF-8?q?gs=20to=20hipMemPrefetchAsync=5Fvs=20and=20hip=E2=80=A6=20(#36)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit SWDEV-546293 - Adds nvidia mappings to hipMemPrefetchAsync_vs and hipMemAdvise_v2 --- .../nvidia_detail/nvidia_hip_runtime_api.h | 21 +++++++++++++++++++ 1 file changed, 21 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 047538fab1..5effd3fa70 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -2050,6 +2050,27 @@ inline static hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, return hipCUDAErrorTohipError(cudaMemPrefetchAsync(dev_ptr, count, device, stream)); } +inline static hipError_t hipMemPrefetchAsync_v2(const void *dev_ptr, size_t count, + hipMemLocation location, unsigned int flags, + hipStream_t stream __dparm(0)) { +#if CUDA_VERSION >= 13000 + return hipCUDAErrorTohipError(cudaMemPrefetchAsync(dev_ptr, count, location, flags, stream)); +#else + return hipCUDAErrorTohipError(cudaMemPrefetchAsync_v2(dev_ptr, count, location, flags, + stream)); +#endif +} + +inline static hipError_t hipMemAdvise_v2(const void* dev_ptr, size_t count, + hipMemoryAdvise advice, hipMemLocation location) { +#if CUDA_VERSION >= 13000 + return hipCUDAErrorTohipError(cudaMemAdvise(dev_ptr, count, advice , location)); +#else + return hipCUDAErrorTohipError(cudaMemAdvise_v2(dev_ptr, count, advice , location)); +#endif +} + + inline static hipError_t hipMemRangeGetAttribute(void* data, size_t data_size, hipMemRangeAttribute attribute, const void* dev_ptr, size_t count) { From ed7fae51d884a17cd3bd96c7973d1f1f78b5f3a5 Mon Sep 17 00:00:00 2001 From: "Manocha, Rahul" Date: Mon, 18 Aug 2025 12:09:09 -0700 Subject: [PATCH 11/11] Fix typo in hipBatchMemOpParamsTocudaBatchMemOpParams (#38) Co-authored-by: Rahul Manocha --- hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 5effd3fa70..a4dae25467 100644 --- a/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -672,7 +672,7 @@ static inline void hipBatchMemOpParamsTocudaBatchMemOpParams(CUstreamBatchMemOpP a[i].writeValue.alias = (CUdeviceptr)(p[i].writeValue.alias); } else if (p[i].memoryBarrier.operation == hipStreamMemOpBarrier) { - a[i].memoryBarrier.operation == CU_STREAM_MEM_OP_BARRIER; + a[i].memoryBarrier.operation = CU_STREAM_MEM_OP_BARRIER; a[i].memoryBarrier.flags = p[i].memoryBarrier.flags; } else if (p[i].flushRemoteWrites.operation == hipStreamMemOpFlushRemoteWrites) {