Skip to content

Commit 655bc2d

Browse files
Andy-JostAndy-Jost
committed
Merge remote-tracking branch 'origin/main' into ajost/cuda-core-1.2.2-docs
2 parents edbd469 + d9feb99 commit 655bc2d

32 files changed

Lines changed: 2425 additions & 212 deletions

‎cuda_bindings/cuda/bindings/driver.pyx‎

Lines changed: 2 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
# SPDX-License-Identifier: Apache-2.0
33

44
# This code was automatically generated with version 13.4.1. Do not modify it directly.
5-
# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=175f5b5d51fbe64e7d8f5eca92e8b557cab5dea406091f66e10a4d19400dcefa
5+
# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=548b03ca3d3021957ab650d666fcd33a16d9af80a5004254c2733f807836679b
66
from typing import Any, Optional
77
import cython
88
import ctypes
@@ -43815,19 +43815,6 @@ def cuLaunchCooperativeKernel(f, unsigned int gridDimX, unsigned int gridDimY, u
4381543815
err = cydriver.cuLaunchCooperativeKernel(cyf, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, cyhStream, cykernelParams_ptr)
4381643816
return (_CUresult(err),)
4381743817

43818-
ctypedef struct cuHostCallbackData_st:
43819-
cydriver.CUhostFn callback
43820-
void *userData
43821-
43822-
ctypedef cuHostCallbackData_st cuHostCallbackData
43823-
43824-
@cython.show_performance_hints(False)
43825-
cdef void cuHostCallbackWrapper(void *data) nogil:
43826-
cdef cuHostCallbackData *cbData = <cuHostCallbackData *>data
43827-
with gil:
43828-
cbData.callback(cbData.userData)
43829-
free(cbData)
43830-
4383143818
@cython.embedsignature(True)
4383243819
def cuLaunchHostFunc(hStream, fn, userData):
4383343820
""" Enqueues a host function call in a stream.
@@ -43907,18 +43894,8 @@ def cuLaunchHostFunc(hStream, fn, userData):
4390743894
cyhStream = <cydriver.CUstream><void_ptr>phStream
4390843895
cdef _HelperInputVoidPtrStruct cyuserDataHelper
4390943896
cdef void* cyuserData = _helper_input_void_ptr(userData, &cyuserDataHelper)
43910-
43911-
cdef cuHostCallbackData *cbData = NULL
43912-
cbData = <cuHostCallbackData *>malloc(sizeof(cbData[0]))
43913-
if cbData == NULL:
43914-
return (CUresult.CUDA_ERROR_OUT_OF_MEMORY,)
43915-
cbData.callback = cyfn
43916-
cbData.userData = cyuserData
43917-
4391843897
with nogil:
43919-
err = cydriver.cuLaunchHostFunc(cyhStream, <cydriver.CUhostFn>cuHostCallbackWrapper, <void *>cbData)
43920-
if err != cydriver.CUDA_SUCCESS:
43921-
free(cbData)
43898+
err = cydriver.cuLaunchHostFunc(cyhStream, cyfn, cyuserData)
4392243899
_helper_input_void_ptr_free(&cyuserDataHelper)
4392343900
return (_CUresult(err),)
4392443901

‎cuda_bindings/cuda/bindings/runtime.pyx‎

Lines changed: 2 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
# SPDX-License-Identifier: Apache-2.0
33

44
# This code was automatically generated with version 13.4.1. Do not modify it directly.
5-
# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=4c952a0aa0dcb7a2948769d3eee64d0ef217ca1a8bd72873d99d94e9a34f2a65
5+
# CYTHON-BINDINGS-GENERATED-DO-NOT-MODIFY-THIS-FILE: format=1; content-sha256=ea95eb12334cbd4c1bf3fc8e890fb6f9493063b12faea853bb3d8c6288934182
66
from typing import Any, Optional
77
import cython
88
import ctypes
@@ -25014,19 +25014,6 @@ def cudaFuncGetParamCount(func):
2501425014
return (_cudaError_t(err), None)
2501525015
return (_cudaError_t_SUCCESS, paramCount)
2501625016

25017-
ctypedef struct cudaStreamHostCallbackData_st:
25018-
cyruntime.cudaHostFn_t callback
25019-
void *userData
25020-
25021-
ctypedef cudaStreamHostCallbackData_st cudaStreamHostCallbackData
25022-
25023-
@cython.show_performance_hints(False)
25024-
cdef void cudaStreamRtHostCallbackWrapper(void *data) nogil:
25025-
cdef cudaStreamHostCallbackData *cbData = <cudaStreamHostCallbackData *>data
25026-
with gil:
25027-
cbData.callback(cbData.userData)
25028-
free(cbData)
25029-
2503025017
@cython.embedsignature(True)
2503125018
def cudaLaunchHostFunc(stream, fn, userData):
2503225019
""" Enqueues a host function call in a stream.
@@ -25107,20 +25094,9 @@ def cudaLaunchHostFunc(stream, fn, userData):
2510725094
cystream = <cyruntime.cudaStream_t><void_ptr>pstream
2510825095
cdef _HelperInputVoidPtrStruct cyuserDataHelper
2510925096
cdef void* cyuserData = _helper_input_void_ptr(userData, &cyuserDataHelper)
25110-
25111-
cdef cudaStreamHostCallbackData *cbData = NULL
25112-
cbData = <cudaStreamHostCallbackData *>malloc(sizeof(cbData[0]))
25113-
if cbData == NULL:
25114-
return (cudaError_t.cudaErrorMemoryAllocation,)
25115-
cbData.callback = cyfn
25116-
cbData.userData = cyuserData
25117-
2511825097
with nogil:
25119-
err = cyruntime.cudaLaunchHostFunc(cystream, <cyruntime.cudaHostFn_t>cudaStreamRtHostCallbackWrapper, <void *>cbData)
25120-
if err != cyruntime.cudaSuccess:
25121-
free(cbData)
25098+
err = cyruntime.cudaLaunchHostFunc(cystream, cyfn, cyuserData)
2512225099
_helper_input_void_ptr_free(&cyuserDataHelper)
25123-
2512425100
return (_cudaError_t(err),)
2512525101

2512625102
@cython.embedsignature(True)

‎cuda_bindings/docs/source/overview.rst‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -571,3 +571,25 @@ For ctypes, we leverage the special handling of ``None`` type since each Python
571571
None,
572572
)
573573
kernelArgs = (kernelValues, kernelTypes)
574+
575+
.. _host-function-lifetime:
576+
577+
Host functions
578+
--------------
579+
580+
:func:`cuda.bindings.driver.cuLaunchHostFunc` and
581+
:func:`cuda.bindings.runtime.cudaLaunchHostFunc` enqueue a host function on a
582+
stream. The caller owns the memory behind ``fn`` and ``userData``; the bindings
583+
do not keep the corresponding Python objects alive.
584+
585+
Keep strong references to the callback object (such as a ``ctypes.CFUNCTYPE``
586+
instance) and the storage backing ``userData`` until the host function has
587+
completed, as confirmed by stream synchronization. A ``CUhostFn`` or
588+
``cudaHostFn_t`` wrapper alone does not retain the underlying ctypes callback.
589+
590+
When a host function is captured into a CUDA graph, it runs on every launch of
591+
that graph. Keep both references alive until the graph and all graph executables
592+
containing that host node have been destroyed and all pending launches have
593+
completed. Releasing either object after the first invocation can cause later
594+
graph launches to access freed memory. The same lifetime requirements apply to
595+
the ``_v2`` variants of these APIs.

‎cuda_bindings/docs/source/release/13.5.0-notes.rst‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -323,6 +323,14 @@ Bugfixes
323323
- A number of incorrect docstrings have been corrected.
324324
- Setting a ``cuda_bindings`` class instance to itself (``obj = cuda_bindings.driver.Class(); obj[0] = obj``) will no longer crash or use garbage values.
325325
- Fix a potential use-after-free error in the destructors of classes in ``driver`` and ``runtime``.
326+
- ``driver.cuLaunchHostFunc`` and ``runtime.cudaLaunchHostFunc`` now pass the host
327+
function and its user data to CUDA unchanged. They used to route the call through
328+
an internal trampoline that freed its state after the first invocation, so a host
329+
function captured into a CUDA graph crashed on the second launch of that graph.
330+
(`Issue #790 <https://github.com/NVIDIA/cuda-python/issues/790>`_)
331+
Callers must retain the callback object and user data storage for the lifetime
332+
of the graph and its executables, including pending launches; see
333+
:ref:`host-function-lifetime`.
326334

327335
Prerelease feature
328336
------------------

‎cuda_bindings/tests/test_cuda.py‎

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -691,6 +691,48 @@ def test_stream_capture():
691691
pass
692692

693693

694+
@pytest.mark.agent_authored(model="claude-fable-5-1")
695+
def test_cuda_launch_host_func_graph_replay():
696+
# A host function captured into a graph runs once per graph launch. The
697+
# bindings must hand the function pointer and the user data to CUDA
698+
# unchanged so that every launch reaches the same function (issue #790).
699+
counter = ctypes.c_int(0)
700+
701+
@ctypes.CFUNCTYPE(None, ctypes.c_void_p)
702+
def host_fn(userData):
703+
ctypes.c_int.from_address(userData).value += 1
704+
705+
callback = cuda.CUhostFn(_ptr=ctypes.addressof(host_fn))
706+
err, stream = cuda.cuStreamCreate(0)
707+
assert err == cuda.CUresult.CUDA_SUCCESS
708+
try:
709+
(err,) = cuda.cuStreamBeginCapture(stream, cuda.CUstreamCaptureMode.CU_STREAM_CAPTURE_MODE_GLOBAL)
710+
assert err == cuda.CUresult.CUDA_SUCCESS
711+
(err,) = cuda.cuLaunchHostFunc(stream, callback, ctypes.addressof(counter))
712+
assert err == cuda.CUresult.CUDA_SUCCESS
713+
err, graph = cuda.cuStreamEndCapture(stream)
714+
assert err == cuda.CUresult.CUDA_SUCCESS
715+
try:
716+
err, graphExec = cuda.cuGraphInstantiate(graph, 0)
717+
assert err == cuda.CUresult.CUDA_SUCCESS
718+
try:
719+
for launch in range(1, 6):
720+
(err,) = cuda.cuGraphLaunch(graphExec, stream)
721+
assert err == cuda.CUresult.CUDA_SUCCESS
722+
(err,) = cuda.cuStreamSynchronize(stream)
723+
assert err == cuda.CUresult.CUDA_SUCCESS
724+
assert counter.value == launch
725+
finally:
726+
(err,) = cuda.cuGraphExecDestroy(graphExec)
727+
assert err == cuda.CUresult.CUDA_SUCCESS
728+
finally:
729+
(err,) = cuda.cuGraphDestroy(graph)
730+
assert err == cuda.CUresult.CUDA_SUCCESS
731+
finally:
732+
(err,) = cuda.cuStreamDestroy(stream)
733+
assert err == cuda.CUresult.CUDA_SUCCESS
734+
735+
694736
def test_profiler():
695737
(err,) = cuda.cuProfilerStart()
696738
assert err == cuda.CUresult.CUDA_SUCCESS

‎cuda_bindings/tests/test_cudart.py‎

Lines changed: 65 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -1395,33 +1395,32 @@ def test_cudart_eglFrame():
13951395
assert int(frame.frame.pPitch[2].ysize) == 0
13961396

13971397

1398-
def cudart_func_stream_callback(use_host_api):
1398+
@pytest.mark.parametrize("use_host_api", [False, True], ids=["stream_callback", "host_function"])
1399+
@pytest.mark.agent_authored(model="gpt-6")
1400+
def test_cudart_func_callback(use_host_api):
13991401
class testStruct(ctypes.Structure):
14001402
_fields_ = [
14011403
("a", ctypes.c_int),
14021404
("b", ctypes.c_int),
14031405
("c", ctypes.c_int),
14041406
]
14051407

1408+
calls = []
1409+
stream_results = []
1410+
14061411
def task_callback_host(userData):
14071412
data = testStruct.from_address(userData)
1408-
assert data.a == 1
1409-
assert data.b == 2
1410-
assert data.c == 3
1411-
return 0
1413+
calls.append((userData, data.a, data.b, data.c))
14121414

14131415
def task_callback_stream(stream, status, userData):
1414-
data = testStruct.from_address(userData)
1415-
assert data.a == 1
1416-
assert data.b == 2
1417-
assert data.c == 3
1418-
return 0
1416+
stream_results.append((stream, status))
1417+
task_callback_host(userData)
14191418

14201419
if use_host_api:
1421-
callback_type = ctypes.PYFUNCTYPE(ctypes.c_int, ctypes.c_void_p)
1420+
callback_type = ctypes.CFUNCTYPE(None, ctypes.c_void_p)
14221421
target_task = task_callback_host
14231422
else:
1424-
callback_type = ctypes.PYFUNCTYPE(ctypes.c_int, ctypes.c_void_p, ctypes.c_int, ctypes.c_void_p)
1423+
callback_type = ctypes.CFUNCTYPE(None, ctypes.c_void_p, ctypes.c_int, ctypes.c_void_p)
14251424
target_task = task_callback_stream
14261425

14271426
# Construct ctype data
@@ -1437,19 +1436,64 @@ def task_callback_stream(stream, status, userData):
14371436
# Run
14381437
err, stream = cudart.cudaStreamCreate()
14391438
assertSuccess(err)
1440-
if use_host_api:
1441-
(err,) = cudart.cudaLaunchHostFunc(stream, callback, ctypes.addressof(c_data))
1439+
try:
1440+
if use_host_api:
1441+
(err,) = cudart.cudaLaunchHostFunc(stream, callback, ctypes.addressof(c_data))
1442+
else:
1443+
(err,) = cudart.cudaStreamAddCallback(stream, callback, ctypes.addressof(c_data), 0)
14421444
assertSuccess(err)
1443-
else:
1444-
(err,) = cudart.cudaStreamAddCallback(stream, callback, ctypes.addressof(c_data), 0)
1445+
(err,) = cudart.cudaStreamSynchronize(stream)
14451446
assertSuccess(err)
1446-
(err,) = cudart.cudaDeviceSynchronize()
1447-
assertSuccess(err)
14481447

1448+
# ctypes swallows callback exceptions, so assert on the test thread.
1449+
assert calls == [(ctypes.addressof(c_data), 1, 2, 3)]
1450+
if not use_host_api:
1451+
assert stream_results == [(int(stream), int(cudart.cudaError_t.cudaSuccess))]
1452+
finally:
1453+
(err,) = cudart.cudaStreamDestroy(stream)
1454+
assertSuccess(err)
1455+
1456+
1457+
@pytest.mark.agent_authored(model="claude-fable-5-1")
1458+
def test_cudart_host_func_graph_replay():
1459+
# A host function captured into a graph runs once per graph launch. The
1460+
# bindings must hand the function pointer and the user data to CUDA
1461+
# unchanged so that every launch reaches the same function (issue #790).
1462+
counter = ctypes.c_int(0)
1463+
1464+
@ctypes.CFUNCTYPE(None, ctypes.c_void_p)
1465+
def host_fn(userData):
1466+
ctypes.c_int.from_address(userData).value += 1
14491467

1450-
def test_cudart_func_callback():
1451-
cudart_func_stream_callback(use_host_api=False)
1452-
cudart_func_stream_callback(use_host_api=True)
1468+
callback = cudart.cudaHostFn_t(_ptr=ctypes.addressof(host_fn))
1469+
err, stream = cudart.cudaStreamCreate()
1470+
assertSuccess(err)
1471+
try:
1472+
(err,) = cudart.cudaStreamBeginCapture(stream, cudart.cudaStreamCaptureMode.cudaStreamCaptureModeGlobal)
1473+
assertSuccess(err)
1474+
(err,) = cudart.cudaLaunchHostFunc(stream, callback, ctypes.addressof(counter))
1475+
assertSuccess(err)
1476+
err, graph = cudart.cudaStreamEndCapture(stream)
1477+
assertSuccess(err)
1478+
try:
1479+
err, graphExec = cudart.cudaGraphInstantiate(graph, 0)
1480+
assertSuccess(err)
1481+
try:
1482+
for launch in range(1, 6):
1483+
(err,) = cudart.cudaGraphLaunch(graphExec, stream)
1484+
assertSuccess(err)
1485+
(err,) = cudart.cudaStreamSynchronize(stream)
1486+
assertSuccess(err)
1487+
assert counter.value == launch
1488+
finally:
1489+
(err,) = cudart.cudaGraphExecDestroy(graphExec)
1490+
assertSuccess(err)
1491+
finally:
1492+
(err,) = cudart.cudaGraphDestroy(graph)
1493+
assertSuccess(err)
1494+
finally:
1495+
(err,) = cudart.cudaStreamDestroy(stream)
1496+
assertSuccess(err)
14531497

14541498

14551499
@pytest.mark.skipif(

0 commit comments

Comments
 (0)