From 78e9e47852198171cb84168c7aec5b57e65fe9c7 Mon Sep 17 00:00:00 2001 From: Sourabh U Betigeri Date: Tue, 4 Nov 2025 16:56:49 -0800 Subject: [PATCH] SWDEV-551244 - Fixes CUDA 13 compilation issues (#1237) --- .../nvidia_detail/nvidia_hip_runtime_api.h | 134 +++++++++++++++++- 1 file changed, 133 insertions(+), 1 deletion(-) 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 6828d11252..1fc60e8680 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 @@ -1931,6 +1931,7 @@ typedef enum cudaMemLocationType hipMemLocationType; #define hipMemLocationTypeInvalid cudaMemLocationTypeInvalid #define hipMemLocationTypeDevice cudaMemLocationTypeDevice #define hipMemLocationTypeHost cudaMemLocationTypeHost +#define hipMemLocationTypeHostNuma cudaMemLocationTypeHostNuma #define hipMemHandleTypeNone cudaMemHandleTypeNone #define hipMemHandleTypePosixFileDescriptor cudaMemHandleTypePosixFileDescriptor #define hipMemHandleTypeWin32 cudaMemHandleTypeWin32 @@ -2048,13 +2049,32 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla inline static hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice, int device) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ uses cudaMemLocation instead of int device + cudaMemLocation location; + location.type = cudaMemLocationTypeDevice; + location.id = device; + return hipCUDAErrorTohipError( + cudaMemAdvise(dev_ptr, count, hipMemoryAdviseTocudaMemoryAdvise(advice), location)); +#else + // CUDA < 13 uses int device directly return hipCUDAErrorTohipError( cudaMemAdvise(dev_ptr, count, hipMemoryAdviseTocudaMemoryAdvise(advice), device)); +#endif } inline static hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device, hipStream_t stream __dparm(0)) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ uses cudaMemLocation and flags parameter + cudaMemLocation location; + location.type = cudaMemLocationTypeDevice; + location.id = device; + return hipCUDAErrorTohipError(cudaMemPrefetchAsync(dev_ptr, count, location, 0, stream)); +#else + // CUDA < 13 uses int device directly return hipCUDAErrorTohipError(cudaMemPrefetchAsync(dev_ptr, count, device, stream)); +#endif } inline static hipError_t hipMemPrefetchAsync_v2(const void* dev_ptr, size_t count, @@ -2195,14 +2215,20 @@ inline static hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* pro cdprop.regsPerBlock = prop->regsPerBlock; cdprop.warpSize = prop->warpSize; cdprop.maxThreadsPerBlock = prop->maxThreadsPerBlock; +#if CUDA_VERSION < 13000 cdprop.clockRate = prop->clockRate; +#endif cdprop.totalConstMem = prop->totalConstMem; cdprop.multiProcessorCount = prop->multiProcessorCount; cdprop.l2CacheSize = prop->l2CacheSize; cdprop.maxThreadsPerMultiProcessor = prop->maxThreadsPerMultiProcessor; +#if CUDA_VERSION < 13000 cdprop.computeMode = prop->computeMode; +#endif cdprop.canMapHostMemory = prop->canMapHostMemory; +#if CUDA_VERSION < 13000 cdprop.memoryClockRate = prop->memoryClockRate; +#endif cdprop.memoryBusWidth = prop->memoryBusWidth; return hipCUDAErrorTohipError(cudaChooseDevice(device, &cdprop)); } @@ -2407,14 +2433,33 @@ inline static hipError_t hipMemcpy2DToArrayAsync(hipArray_t dst, size_t wOffset, 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) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ signature: failIdx removed, const qualifiers added + if (failIdx != nullptr) { + *failIdx = 0; + } + return hipCUDAErrorTohipError( + cudaMemcpyBatchAsync((void *const *)dsts, (const void *const *)srcs, (const size_t *)sizes, count, attrs, attrsIdxs, numAttrs, stream)); +#else + // CUDA < 13 signature: failIdx supported, no const qualifiers return hipCUDAErrorTohipError( cudaMemcpyBatchAsync(dsts, srcs, sizes, count, attrs, attrsIdxs, numAttrs, failIdx, stream)); +#endif } inline static hipError_t hipMemcpy3DBatchAsync(size_t numOps, hipMemcpy3DBatchOp* opList, size_t* failIdx, unsigned long long flags, hipStream_t stream) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ signature: failIdx removed + if (failIdx != nullptr) { + *failIdx = 0; + } + return hipCUDAErrorTohipError(cudaMemcpy3DBatchAsync(numOps, opList, flags, stream)); +#else + // CUDA < 13 signature: failIdx supported return hipCUDAErrorTohipError(cudaMemcpy3DBatchAsync(numOps, opList, failIdx, flags, stream)); +#endif } inline static hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms* p) { return hipCUDAErrorTohipError(cudaMemcpy3DPeer(p)); @@ -2634,21 +2679,31 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev p_prop->maxGridSize[0] = cdprop.maxGridSize[0]; p_prop->maxGridSize[1] = cdprop.maxGridSize[1]; p_prop->maxGridSize[2] = cdprop.maxGridSize[2]; +#if CUDA_VERSION < 13000 p_prop->clockRate = cdprop.clockRate; +#endif p_prop->totalConstMem = cdprop.totalConstMem; p_prop->major = cdprop.major; p_prop->minor = cdprop.minor; p_prop->textureAlignment = cdprop.textureAlignment; p_prop->texturePitchAlignment = cdprop.texturePitchAlignment; +#if CUDA_VERSION < 13000 p_prop->deviceOverlap = cdprop.deviceOverlap; +#endif p_prop->multiProcessorCount = cdprop.multiProcessorCount; +#if CUDA_VERSION < 13000 p_prop->kernelExecTimeoutEnabled = cdprop.kernelExecTimeoutEnabled; +#endif p_prop->integrated = cdprop.integrated; p_prop->canMapHostMemory = cdprop.canMapHostMemory; +#if CUDA_VERSION < 13000 p_prop->computeMode = cdprop.computeMode; +#endif p_prop->maxTexture1D = cdprop.maxTexture1D; p_prop->maxTexture1DMipmap = cdprop.maxTexture1DMipmap; +#if CUDA_VERSION < 13000 p_prop->maxTexture1DLinear = cdprop.maxTexture1DLinear; +#endif p_prop->maxTexture2D[0] = cdprop.maxTexture2D[0]; p_prop->maxTexture2D[1] = cdprop.maxTexture2D[1]; p_prop->maxTexture2DMipmap[0] = cdprop.maxTexture2DMipmap[0]; @@ -2695,7 +2750,9 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev p_prop->tccDriver = cdprop.tccDriver; p_prop->asyncEngineCount = cdprop.asyncEngineCount; p_prop->unifiedAddressing = cdprop.unifiedAddressing; +#if CUDA_VERSION < 13000 p_prop->memoryClockRate = cdprop.memoryClockRate; +#endif p_prop->memoryBusWidth = cdprop.memoryBusWidth; p_prop->l2CacheSize = cdprop.l2CacheSize; p_prop->maxThreadsPerMultiProcessor = cdprop.maxThreadsPerMultiProcessor; @@ -2708,13 +2765,17 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev p_prop->isMultiGpuBoard = cdprop.isMultiGpuBoard; p_prop->multiGpuBoardGroupID = cdprop.multiGpuBoardGroupID; p_prop->hostNativeAtomicSupported = cdprop.hostNativeAtomicSupported; +#if CUDA_VERSION < 13000 p_prop->singleToDoublePrecisionPerfRatio = cdprop.singleToDoublePrecisionPerfRatio; +#endif p_prop->pageableMemoryAccess = cdprop.pageableMemoryAccess; p_prop->concurrentManagedAccess = cdprop.concurrentManagedAccess; p_prop->computePreemptionSupported = cdprop.computePreemptionSupported; p_prop->canUseHostPointerForRegisteredMem = cdprop.canUseHostPointerForRegisteredMem; p_prop->cooperativeLaunch = cdprop.cooperativeLaunch; +#if CUDA_VERSION < 13000 p_prop->cooperativeMultiDeviceLaunch = cdprop.cooperativeMultiDeviceLaunch; +#endif p_prop->sharedMemPerBlockOptin = cdprop.sharedMemPerBlockOptin; p_prop->pageableMemoryAccessUsesHostPageTables = cdprop.pageableMemoryAccessUsesHostPageTables; p_prop->directManagedMemAccessFromHost = cdprop.directManagedMemAccessFromHost; @@ -2869,7 +2930,12 @@ inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t att cdattr = cudaDevAttrCooperativeLaunch; break; case hipDeviceAttributeCooperativeMultiDeviceLaunch: +#if CUDA_VERSION < 13000 && defined(cudaDevAttrCooperativeMultiDeviceLaunch) cdattr = cudaDevAttrCooperativeMultiDeviceLaunch; +#else + // cudaDevAttrCooperativeMultiDeviceLaunch removed in CUDA 13+ + return hipErrorInvalidValue; +#endif break; case hipDeviceAttributeHostRegisterSupported: cdattr = cudaDevAttrHostRegisterSupported; @@ -3447,7 +3513,14 @@ inline static hipError_t hipEventQuery(hipEvent_t event) { } inline static hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ uses a different signature with CUctxCreateParams + CUctxCreateParams params; + memset(¶ms, 0, sizeof(params)); + return hipCUResultTohipError(cuCtxCreate(ctx, ¶ms, flags, device)); +#else return hipCUResultTohipError(cuCtxCreate(ctx, flags, device)); +#endif } inline static hipError_t hipCtxDestroy(hipCtx_t ctx) { @@ -3832,8 +3905,14 @@ inline static hipError_t hipModuleLaunchCooperativeKernel( inline static hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags) { +#if CUDA_VERSION < 13000 + // cudaLaunchCooperativeKernelMultiDevice available in CUDA < 13 return hipCUDAErrorTohipError( cudaLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags)); +#else + // cudaLaunchCooperativeKernelMultiDevice removed in CUDA 13+ + return hipErrorNotSupported; +#endif } inline static hipError_t hipModuleLaunchCooperativeKernelMultiDevice( @@ -4275,7 +4354,11 @@ inline static hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGra return hipCUDAErrorTohipError(cudaStreamBeginCaptureToGraph( stream, graph, dependencies, dependencyData, numDependencies, mode)); } +#endif +#if CUDA_VERSION >= CUDA_12030 && CUDA_VERSION < 13000 +// Note: cudaGraphNodeGetDependentNodes_v2 only exists in CUDA 12.3-12.x +// In CUDA 13+, the main cudaGraphNodeGetDependentNodes function signature was updated inline static hipError_t hipGraphNodeGetDependentNodes_v2(hipGraphNode_t node, hipGraphNode_t* pDependentNodes, hipGraphEdgeData* edgeData, @@ -4451,7 +4534,13 @@ inline static hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphEx inline static hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from, const hipGraphNode_t* to, size_t numDependencies) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ signature update: edgeData is optional array of edge data. + // If NULL, default (zeroed) edge data is assumed. + return hipCUDAErrorTohipError(cudaGraphAddDependencies(graph, from, to, NULL, numDependencies)); +#else return hipCUDAErrorTohipError(cudaGraphAddDependencies(graph, from, to, numDependencies)); +#endif } inline static hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, @@ -4568,26 +4657,51 @@ inline static hipError_t hipGraphExecBatchMemOpNodeSetParams( inline static hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* from, const hipGraphNode_t* to, size_t numDependencies) { +#if CUDA_VERSION >= 13000 +// CUDA 13+ signature update:edgeData is optional array of edge data. +// If NULL, edge data is assumed to be default (zeroed). + return hipCUDAErrorTohipError(cudaGraphRemoveDependencies(graph, from, to, NULL, numDependencies)); +#else return hipCUDAErrorTohipError(cudaGraphRemoveDependencies(graph, from, to, numDependencies)); +#endif } inline static hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, hipGraphNode_t* to, size_t* numEdges) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ signature update: edgeData is optional location to return edge data + return hipCUDAErrorTohipError(cudaGraphGetEdges(graph, from, to, NULL, numEdges)); +#else return hipCUDAErrorTohipError(cudaGraphGetEdges(graph, from, to, numEdges)); +#endif } inline static hipError_t hipGraphNodeGetDependencies(hipGraphNode_t node, hipGraphNode_t* pDependencies, size_t* pNumDependencies) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ signature update: + // edgeData is optional array to return edge data for each dependency + return hipCUDAErrorTohipError( + cudaGraphNodeGetDependencies(node, pDependencies, NULL, pNumDependencies)); +#else return hipCUDAErrorTohipError( cudaGraphNodeGetDependencies(node, pDependencies, pNumDependencies)); +#endif } inline static hipError_t hipGraphNodeGetDependentNodes(hipGraphNode_t node, hipGraphNode_t* pDependentNodes, size_t* pNumDependentNodes) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ signature update: + // edgeData is optional pointer to return edge data for dependent nodes + return hipCUDAErrorTohipError( + cudaGraphNodeGetDependentNodes(node, pDependentNodes, NULL, pNumDependentNodes)); +#else return hipCUDAErrorTohipError( cudaGraphNodeGetDependentNodes(node, pDependentNodes, pNumDependentNodes)); +#endif } inline static hipError_t hipGraphNodeGetType(hipGraphNode_t node, hipGraphNodeType* pType) { @@ -4632,7 +4746,9 @@ inline static hipError_t hipStreamGetCaptureInfo(hipStream_t stream, return hipCUDAErrorTohipError(cudaStreamGetCaptureInfo(stream, pCaptureStatus, pId)); } -#if CUDA_VERSION >= CUDA_11030 || defined(__CUDA_API_VERSION_INTERNAL) +#if CUDA_VERSION >= CUDA_11030 && CUDA_VERSION < 13000 +// Note: cuStreamGetCaptureInfo_v2 only exists in CUDA 11.3-12.x +// In CUDA 13+, it was superseded by v3 inline static hipError_t hipStreamGetCaptureInfo_v2( hipStream_t stream, hipStreamCaptureStatus* captureStatus_out, unsigned long long* id_out __dparm(0), hipGraph_t* graph_out __dparm(0), @@ -4653,8 +4769,16 @@ inline static hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies, size_t numDependencies, unsigned int flags __dparm(0)) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ signature update: + // dependencyData is optional array of data associated with each dependency. + // If NULL, default (zeroed) data is assumed. + return hipCUDAErrorTohipError( + cudaStreamUpdateCaptureDependencies(stream, dependencies, NULL, numDependencies, flags)); +#else return hipCUDAErrorTohipError( cudaStreamUpdateCaptureDependencies(stream, dependencies, numDependencies, flags)); +#endif } #endif @@ -4845,8 +4969,16 @@ inline static hipError_t hipGraphExternalSemaphoresSignalNodeGetParams( inline static hipError_t hipGraphAddNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, const hipGraphNode_t* pDependencies, size_t numDependencies, hipGraphNodeParams* nodeParams) { +#if CUDA_VERSION >= 13000 + // CUDA 13+ signature update: + // dependencyData is optional edge data for the dependencies. + // If NULL, the data is assumed to be default (zeroed) for all dependencies. + return hipCUDAErrorTohipError( + cudaGraphAddNode(pGraphNode, graph, pDependencies, NULL, numDependencies, nodeParams)); +#else return hipCUDAErrorTohipError( cudaGraphAddNode(pGraphNode, graph, pDependencies, numDependencies, nodeParams)); +#endif } inline static hipError_t hipGraphExecNodeSetParams(hipGraphExec_t graphExec, hipGraphNode_t node,