diff --git a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
index 7a3c37076b..a4dae25467 100644
--- a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
+++ b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
@@ -28,6 +28,7 @@ THE SOFTWARE.
 #include <cuda_profiler_api.h>
 #include <cuda_fp16.h>
 
+#include <driver_types.h>
 #include <stdio.h>
 
 #define CUDA_9000 9000
@@ -300,6 +301,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
@@ -865,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
@@ -970,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:
@@ -1812,6 +1851,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
@@ -1991,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) {
@@ -2091,6 +2171,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) {
@@ -2323,6 +2407,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) {
@@ -2474,6 +2578,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;
@@ -2998,6 +3135,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,
@@ -3195,6 +3342,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));
 }
@@ -3426,6 +3583,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));
 }
@@ -3435,6 +3596,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));
 }
@@ -3587,6 +3752,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) {
@@ -4708,12 +4880,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,
@@ -4771,14 +4950,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
@@ -4830,6 +5016,10 @@ 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