Skip to content

Update identify_stream_usage CUDA runtime hooks to CUDA 13 #19807

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
rapids-bot merged 2 commits into from
Aug 27, 2025
+41 −0
Merged

Update identify_stream_usage CUDA runtime hooks to CUDA 13 #19807

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
a7a823d
Update identify_stream_usage CUDA runtime hooks to CUDA 13
robertmaynard Aug 27, 2025
05ef44a
Improve comment
robertmaynard Aug 27, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
41 changes: 41 additions & 0 deletions cpp/tests/utilities/identify_stream_usage.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -209,6 +209,40 @@ DEFINE_OVERLOAD(cudaLaunchKernel,
size_t sharedMem,
cudaStream_t stream),
ARG(func, gridDim, blockDim, args, sharedMem, stream));

#if CUDART_VERSION >= 13000
// We need to define the __cudaLaunchKernel ABI as
// it isn't part of cuda_runtime.h when compiling as a C++ source
Comment on lines +214 to +215
Copy link
Contributor

@vuule vuule Aug 27, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

do we know why this is the case in CUDA 13?

Copy link
Contributor

@davidwendt davidwendt Aug 27, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

They changed their API.
My understanding is that the <<< >>> pattern changed to map to a different API which we need to now override.
The internal __cudaLaunchKernel declaration is hidden under ifdef __CUDACC__ in the header files in 13 which makes sense given how it is intended to be used. We just need to hook it here (without calling it) so the forward references are required in this .cpp file.

extern "C" cudaError_t CUDARTAPI __cudaLaunchKernel(cudaKernel_t kernel,
dim3 gridDim,
dim3 blockDim,
void** args,
size_t sharedMem,
cudaStream_t stream);
extern "C" cudaError_t CUDARTAPI __cudaLaunchKernel_ptsz(cudaKernel_t kernel,
dim3 gridDim,
dim3 blockDim,
void** args,
size_t sharedMem,
cudaStream_t stream);
DEFINE_OVERLOAD(__cudaLaunchKernel,
ARG(cudaKernel_t kernel,
dim3 gridDim,
dim3 blockDim,
void** args,
size_t sharedMem,
cudaStream_t stream),
ARG(kernel, gridDim, blockDim, args, sharedMem, stream));
DEFINE_OVERLOAD(__cudaLaunchKernel_ptsz,
ARG(cudaKernel_t kernel,
dim3 gridDim,
dim3 blockDim,
void** args,
size_t sharedMem,
cudaStream_t stream),
ARG(kernel, gridDim, blockDim, args, sharedMem, stream));
#endif

DEFINE_OVERLOAD(cudaLaunchCooperativeKernel,
ARG(void const* func,
dim3 gridDim,
Expand All @@ -223,9 +257,16 @@ DEFINE_OVERLOAD(cudaLaunchHostFunc,

// Memory transfer APIS:
// https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY
#if CUDART_VERSION >= 13000
DEFINE_OVERLOAD(
cudaMemPrefetchAsync,
ARG(void const* devPtr, size_t count, cudaMemLocation loc, int flags, cudaStream_t stream),
ARG(devPtr, count, loc, flags, stream));
#else
DEFINE_OVERLOAD(cudaMemPrefetchAsync,
ARG(void const* devPtr, size_t count, int dstDevice, cudaStream_t stream),
ARG(devPtr, count, dstDevice, stream));
#endif
DEFINE_OVERLOAD(cudaMemcpy2DAsync,
ARG(void* dst,
size_t dpitch,
Expand Down