Skip to content

Commit 13de2c2

Browse files
rwgkleofang
rwgk
and
leofang
authored
CI: run cuda.bindings examples on Linux and Windows (#1517)
* ci: run cuda.bindings examples on Linux * ci: run cuda.bindings examples on Windows * ci: run bindings examples via pytest entrypoint Analysis: - examples were invoked via `python -m pytest` from within `cuda_bindings` so the repo checkout was on sys.path and imports resolved to the source tree - `setuptools_scm` generates `cuda/bindings/_version.py` only in the built wheel, so the source tree lacks this file and `from cuda.bindings._version import __version__` fails during example collection - running `pytest` via the installed entrypoint avoids CWD precedence and keeps imports coming from the installed wheel, which includes the generated version file Change: - switch Linux and Windows example steps to call `pytest` entrypoint * use pathfinder.find_nvidia_header_directory() in cuda_bindings examples KernelHelper * Remove pytest_skipif_cuda_include_not_found() and get_cuda_home() entirely under cuda_bindings/examples/ * Replace py3.13 2-GPU job with 3.14t 2-GPU job * Fix kernel arg lifetimes in isoFDModelling example Keep pointer arrays alive through launches to avoid free-threaded Python misaligned-address failures caused by temporary argument buffers. --------- Co-authored-by: Leo Fang <leof@nvidia.com>
1 parent 070f387 commit 13de2c2

6 files changed

Lines changed: 44 additions & 42 deletions

File tree

‎.github/workflows/test-wheel-linux.yml‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -254,6 +254,16 @@ jobs:
254254
LOCAL_CTK: ${{ matrix.LOCAL_CTK }}
255255
run: run-tests bindings
256256

257+
- name: Run cuda.bindings examples
258+
if: ${{ env.SKIP_CUDA_BINDINGS_TEST == '0' }}
259+
env:
260+
CUDA_VER: ${{ matrix.CUDA_VER }}
261+
LOCAL_CTK: ${{ matrix.LOCAL_CTK }}
262+
run: |
263+
pushd cuda_bindings
264+
${SANITIZER_CMD} pytest -ra -s -vv examples/
265+
popd
266+
257267
- name: Run cuda.core tests
258268
env:
259269
CUDA_VER: ${{ matrix.CUDA_VER }}

‎.github/workflows/test-wheel-windows.yml‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -226,6 +226,17 @@ jobs:
226226
shell: bash --noprofile --norc -xeuo pipefail {0}
227227
run: run-tests bindings
228228

229+
- name: Run cuda.bindings examples
230+
if: ${{ env.SKIP_CUDA_BINDINGS_TEST == '0' }}
231+
env:
232+
CUDA_VER: ${{ matrix.CUDA_VER }}
233+
LOCAL_CTK: ${{ matrix.LOCAL_CTK }}
234+
shell: bash --noprofile --norc -xeuo pipefail {0}
235+
run: |
236+
pushd cuda_bindings
237+
${SANITIZER_CMD} pytest -ra -s -vv examples/
238+
popd
239+
229240
- name: Run cuda.core tests
230241
env:
231242
CUDA_VER: ${{ matrix.CUDA_VER }}

‎ci/test-matrix.yml‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -58,8 +58,8 @@ linux:
5858
# special runners
5959
- { ARCH: 'amd64', PY_VER: '3.13', CUDA_VER: '13.0.2', LOCAL_CTK: '1', GPU: 'h100', GPU_COUNT: '1', DRIVER: 'latest' }
6060
- { ARCH: 'amd64', PY_VER: '3.13', CUDA_VER: '13.1.1', LOCAL_CTK: '1', GPU: 'h100', GPU_COUNT: '1', DRIVER: 'latest' }
61-
- { ARCH: 'amd64', PY_VER: '3.14', CUDA_VER: '13.1.1', LOCAL_CTK: '1', GPU: 'h100', GPU_COUNT: '2', DRIVER: 'latest' }
6261
- { ARCH: 'amd64', PY_VER: '3.14', CUDA_VER: '13.1.1', LOCAL_CTK: '1', GPU: 't4', GPU_COUNT: '2', DRIVER: 'latest' }
62+
- { ARCH: 'amd64', PY_VER: '3.14t', CUDA_VER: '13.1.1', LOCAL_CTK: '1', GPU: 'h100', GPU_COUNT: '2', DRIVER: 'latest' }
6363
nightly: []
6464

6565
windows:

‎cuda_bindings/examples/3_CUDA_Features/globalToShmemAsyncCopy_test.py‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1114,7 +1114,6 @@ def MatrixMultiply(dimsA, dimsB, kernel_number):
11141114

11151115

11161116
def main():
1117-
common.pytest_skipif_cuda_include_not_found()
11181117
common.pytest_skipif_compute_capability_too_low(findCudaDevice(), (7, 0))
11191118

11201119
print("[globalToShmemAsyncCopy] - Starting...")

‎cuda_bindings/examples/common/common.py‎

Lines changed: 10 additions & 28 deletions
Original file line numberDiff line numberDiff line change
@@ -1,33 +1,15 @@
11
# Copyright 2021-2025 NVIDIA Corporation. All rights reserved.
22
# SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE
33

4-
import os
54

65
import numpy as np
76
from common.helper_cuda import checkCudaErrors
7+
from cuda import pathfinder
88
from cuda.bindings import driver as cuda
99
from cuda.bindings import nvrtc
1010
from cuda.bindings import runtime as cudart
1111

1212

13-
def get_cuda_home():
14-
cuda_home = os.getenv("CUDA_HOME")
15-
if cuda_home is None:
16-
cuda_home = os.getenv("CUDA_PATH")
17-
return cuda_home
18-
19-
20-
def pytest_skipif_cuda_include_not_found():
21-
import pytest
22-
23-
cuda_home = get_cuda_home()
24-
if cuda_home is None:
25-
pytest.skip("CUDA_HOME/CUDA_PATH not set")
26-
cuda_include = os.path.join(cuda_home, "include")
27-
if not os.path.exists(cuda_include):
28-
pytest.skip(f"$CUDA_HOME/include does not exist: '{cuda_include}'")
29-
30-
3113
def pytest_skipif_compute_capability_too_low(devID, required_cc_major_minor):
3214
import pytest
3315

@@ -44,16 +26,16 @@ def pytest_skipif_compute_capability_too_low(devID, required_cc_major_minor):
4426

4527
class KernelHelper:
4628
def __init__(self, code, devID):
47-
prog = checkCudaErrors(nvrtc.nvrtcCreateProgram(str.encode(code), b"sourceCode.cu", 0, None, None))
29+
include_dirs = []
30+
for libname in ("cudart", "cccl"):
31+
hdr_dir = pathfinder.find_nvidia_header_directory(libname)
32+
if hdr_dir is None:
33+
import pytest
4834

49-
cuda_home = get_cuda_home()
50-
assert cuda_home is not None
51-
cuda_include = os.path.join(cuda_home, "include")
52-
assert os.path.isdir(cuda_include)
53-
include_dirs = [cuda_include]
54-
cccl_include = os.path.join(cuda_include, "cccl")
55-
if os.path.isdir(cccl_include):
56-
include_dirs.insert(0, cccl_include)
35+
pytest.skip(f'pathfinder.find_nvidia_header_directory("{libname}") returned None')
36+
include_dirs.append(hdr_dir)
37+
38+
prog = checkCudaErrors(nvrtc.nvrtcCreateProgram(str.encode(code), b"sourceCode.cu", 0, None, None))
5739

5840
# Initialize CUDA
5941
checkCudaErrors(cudart.cudaFree(0))

‎cuda_bindings/examples/extra/isoFDModelling_test.py‎

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -310,7 +310,7 @@ def createSource(self, kernel):
310310
freq = np.array(self.params.freqMax, dtype=np.float32)
311311

312312
args = [buf, dt, freq, nt]
313-
args = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
313+
argsp = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
314314
checkCudaErrors(cuda.cuCtxSetCurrent(self.context))
315315
checkCudaErrors(
316316
cuda.cuLaunchKernel(
@@ -323,7 +323,7 @@ def createSource(self, kernel):
323323
1, # block dim
324324
0,
325325
self.streamHalo, # shared mem and stream
326-
args.ctypes.data,
326+
argsp.ctypes.data,
327327
0,
328328
)
329329
) # arguments
@@ -351,7 +351,7 @@ def injectSource(self, kernel, iter):
351351
np_it = np.array(iter, dtype=np.uint32)
352352

353353
args = [wavein + offset_sourceInject, src, np_it]
354-
args = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
354+
argsp = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
355355
checkCudaErrors(
356356
cuda.cuLaunchKernel(
357357
kernel.injectSource,
@@ -363,7 +363,7 @@ def injectSource(self, kernel, iter):
363363
1, # block dim
364364
0,
365365
self.streamHalo, # shared mem and stream
366-
args.ctypes.data,
366+
argsp.ctypes.data,
367367
0,
368368
)
369369
) # arguments
@@ -391,7 +391,7 @@ def createVelocity(self, kernel):
391391
np_stride = np.array(stride, dtype=np.uint32)
392392

393393
args = [vel + offset_velocity, np_dx_dt2, np_nz, np_nx, np_stride]
394-
args = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
394+
argsp = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
395395

396396
checkCudaErrors(cuda.cuCtxSetCurrent(self.context))
397397

@@ -407,7 +407,7 @@ def createVelocity(self, kernel):
407407
1, # block dim
408408
0,
409409
self.streamHalo, # shared mem and stream
410-
args.ctypes.data,
410+
argsp.ctypes.data,
411411
0,
412412
)
413413
) # arguments
@@ -448,7 +448,7 @@ def executeCenter(self, kernel):
448448
np_nx,
449449
np_stride,
450450
]
451-
args = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
451+
argsp = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
452452

453453
# do center propagation from 2 * fd_order to nz - 2 * fd_order
454454
checkCudaErrors(
@@ -462,7 +462,7 @@ def executeCenter(self, kernel):
462462
1, # block dim
463463
0,
464464
self.streamCenter, # shared mem and stream
465-
args.ctypes.data,
465+
argsp.ctypes.data,
466466
0,
467467
)
468468
) # arguments
@@ -503,7 +503,7 @@ def executeHalo(self, kernel):
503503
np_nx,
504504
np_stride,
505505
]
506-
args = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
506+
argsp = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
507507

508508
# do halo up
509509
checkCudaErrors(
@@ -517,7 +517,7 @@ def executeHalo(self, kernel):
517517
1, # block dim
518518
0,
519519
self.streamHalo, # shared mem and stream
520-
args.ctypes.data,
520+
argsp.ctypes.data,
521521
0,
522522
)
523523
) # arguments
@@ -541,7 +541,7 @@ def executeHalo(self, kernel):
541541
np_nx,
542542
np_stride,
543543
]
544-
args = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
544+
argsp = np.array([arg.ctypes.data for arg in args], dtype=np.uint64)
545545
checkCudaErrors(
546546
cuda.cuLaunchKernel(
547547
kernel.fdPropag,
@@ -553,7 +553,7 @@ def executeHalo(self, kernel):
553553
1, # block dim
554554
0,
555555
self.streamHalo, # shared mem and stream
556-
args.ctypes.data,
556+
argsp.ctypes.data,
557557
0,
558558
)
559559
) # arguments

0 commit comments

Comments
 (0)