AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-551244 - Fixes CUDA 13 compilation issues (#1237)

Browse Source
This commit is contained in:
Sourabh U Betigeri
2025-11-04 16:56:49 -08:00
committed by GitHub
parent 1e501dd89a
commit 78e9e47852
1 changed files with 133 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
+133 -1
View File
@@ -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(&params, 0, sizeof(params));
return hipCUResultTohipError(cuCtxCreate(ctx, &params, 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,