From 6f4423482f9c36a02f018d553e21f9dcfd1f5cc3 Mon Sep 17 00:00:00 2001 From: Nader Al Awar Date: Wed, 15 Jul 2026 18:58:28 -0500 Subject: [PATCH 1/4] add pytest run parallel --- ci/test_cuda_compute_minimal_python.sh | 22 ++++++++++++++++++- python/cuda_cccl/tests/compute/conftest.py | 22 ++++++++++++++++--- .../cuda_cccl/tests/compute/test_no_numba.py | 18 +++++++-------- 3 files changed, 49 insertions(+), 13 deletions(-) diff --git a/ci/test_cuda_compute_minimal_python.sh b/ci/test_cuda_compute_minimal_python.sh index f0ec1d85760..3eef6324ed8 100755 --- a/ci/test_cuda_compute_minimal_python.sh +++ b/ci/test_cuda_compute_minimal_python.sh @@ -30,7 +30,7 @@ fi # full cu* extras because those pull in numba/numba-cuda. CUDA_CCCL_WHEEL_PATH="$(ls "${wheelhouse_dir}"/cuda_cccl-*.whl)" python -m pip install "${CUDA_CCCL_WHEEL_PATH}[minimal-cu${cuda_major_version}]" -python -m pip install pytest pytest-xdist +python -m pip install pytest pytest-xdist pytest-run-parallel cd "${repo_root}/python/cuda_cccl/tests/" python -m pytest -n 6 -v compute/test_no_numba.py @@ -44,4 +44,24 @@ if [[ "${py_version}" == "3.14t" ]]; then compute/test_free_threading_stress.py \ compute/test_multi_cc_serialization.py::test_aot_build_result_load_failure_is_shared_and_retryable \ compute/test_multi_cc_serialization.py::test_aot_serialization_waits_for_canonical_first_load + + # Broad thread-safety sweep (pytest-run-parallel): re-run the numba-free + # functional suite with each test executed concurrently across threads + # (barrier-synchronized start), stressing the process-wide build cache, + # single-flight coordination, and the Cython bindings from many threads at + # once. Complements test_free_threading_stress.py above, which targets specific + # shared-object scenarios by hand. -n 0 so the threads share one interpreter. + # + # --parallel-threads=2 matches CuPy's free-threading CI (the closest GPU + # precedent); a small fixed count bounds GPU-memory pressure from concurrent + # kernels and stays reproducible across runners, unlike =auto (the runner's + # logical-core count). + # + # Fail fast if the interpreter is not actually GIL-free (wrong build / + # PYTHON_GIL=1): pytest-run-parallel does NOT catch a GIL that is enabled from + # the start -- it would run threads GIL-serialized and pass vacuously. (A GIL + # *re-enabled mid-run* by a non-free-threaded import IS caught by the plugin, + # which is why we do not pass --ignore-gil-enabled.) + python -c "import sys; assert not sys._is_gil_enabled(), 'GIL is enabled; parallel sweep has no signal'" + python -m pytest -n 0 -v --parallel-threads=2 compute/test_no_numba.py fi diff --git a/python/cuda_cccl/tests/compute/conftest.py b/python/cuda_cccl/tests/compute/conftest.py index 57c73ed7de8..d195c46d6cc 100644 --- a/python/cuda_cccl/tests/compute/conftest.py +++ b/python/cuda_cccl/tests/compute/conftest.py @@ -88,14 +88,21 @@ def cuda_stream() -> Generator[Stream, None, None]: @pytest.fixture(scope="function", autouse=True) -def verify_sass(request, monkeypatch): +def verify_sass(request): if request.node.get_closest_marker("no_verify_sass"): return if not check_ldl_stl_in_sass: - print("not checking sass") return + # Pull monkeypatch dynamically rather than as a fixture parameter so this + # autouse fixture does not add monkeypatch to every test's static fixture + # closure. pytest-run-parallel treats monkeypatch as thread-unsafe based on + # that closure, so a parameter here would serialize the entire free-threaded + # parallel sweep -- even though this fixture only patches on the opt-in + # SASS-check path (check_ldl_stl_in_sass, off by default and in CI). + monkeypatch = request.getfixturevalue("monkeypatch") + import cuda.compute._cccl_interop monkeypatch.setattr( @@ -124,8 +131,17 @@ def pytest_collection_modifyitems(config, items): serialization_skip = pytest.mark.skip( reason="serialization not supported on v2 (HostJIT) backend" ) + # Under the pytest-run-parallel sweep (--parallel-threads > 1) skip the + # blanket raise_on_numba_import injection: it monkeypatches + # builtins.__import__, which pytest-run-parallel treats as thread-unsafe and + # would serialize every no_numba test, neutering the sweep. + # + # getoption returns the int 1 from the argparse default but a *string* for + # CLI-passed values ("2", "auto", "1"), so normalize to str before comparing + # -- a bare `> 1` would raise TypeError on the "2" string. + running_parallel = str(config.getoption("parallel_threads", 1)) != "1" for item in items: - if item.get_closest_marker("no_numba"): + if item.get_closest_marker("no_numba") and not running_parallel: if "raise_on_numba_import" not in item.fixturenames: item.fixturenames.append("raise_on_numba_import") if USING_V2 and item.get_closest_marker("serialization"): diff --git a/python/cuda_cccl/tests/compute/test_no_numba.py b/python/cuda_cccl/tests/compute/test_no_numba.py index 264cf117aa0..da605d0dab2 100644 --- a/python/cuda_cccl/tests/compute/test_no_numba.py +++ b/python/cuda_cccl/tests/compute/test_no_numba.py @@ -103,7 +103,11 @@ def _raw_negate_i16_op() -> RawOp: return _raw_op(source, "no_numba_negate_i16") -def test_import_numba_raises(): +def test_import_numba_raises(raise_on_numba_import): + # Request the guard explicitly rather than relying on the no_numba + # auto-injection (skipped under the parallel sweep, see conftest + # pytest_collection_modifyitems): this test exercises the guard directly, and + # the monkeypatch it depends on keeps it single-threaded under the sweep. with pytest.raises( ImportError, match="This test is marked 'no_numba' but attempted to import it" ): @@ -205,9 +209,8 @@ def test_binary_search_explicit_opkind_less(search, side): np.testing.assert_array_equal(d_out.copy_to_host(), expected) -def test_segmented_reduce_well_known_plus(monkeypatch): - monkeypatch.setattr(cuda.compute._cccl_interop, "_check_sass", False) - +@pytest.mark.no_verify_sass +def test_segmented_reduce_well_known_plus(): h_input = np.asarray([1, 2, 3, 4, 5, 6, 7, 8], dtype=np.uint32) h_starts = np.asarray([0, 3, 5], dtype=np.int32) h_ends = np.asarray([3, 5, 8], dtype=np.int32) @@ -288,11 +291,8 @@ def test_segmented_sort_keys(): np.testing.assert_array_equal(d_output.copy_to_host(), expected) -def test_unique_by_key_well_known_equal_to(monkeypatch): - cc_major, _ = cuda.compute._cccl_interop.CudaDevice().compute_capability - if cc_major >= 9: - monkeypatch.setattr(cuda.compute._cccl_interop, "_check_sass", False) - +@pytest.mark.no_verify_sass +def test_unique_by_key_well_known_equal_to(): h_keys = np.asarray([1, 1, 2, 2, 2, 3, 4, 4], dtype=np.int16) h_values = np.asarray([10, 11, 20, 21, 22, 30, 40, 41], dtype=np.int8) d_keys = DeviceArray.from_numpy(h_keys) From ec4557e8d3975c1226b513bf963191bf4f1bfa75 Mon Sep 17 00:00:00 2001 From: Nader Al Awar Date: Wed, 15 Jul 2026 19:19:42 -0500 Subject: [PATCH 2/4] address review --- ci/test_cuda_compute_minimal_python.sh | 6 +++++- python/cuda_cccl/tests/compute/conftest.py | 10 +++++++--- 2 files changed, 12 insertions(+), 4 deletions(-) diff --git a/ci/test_cuda_compute_minimal_python.sh b/ci/test_cuda_compute_minimal_python.sh index 3eef6324ed8..df3d8f8a5b6 100755 --- a/ci/test_cuda_compute_minimal_python.sh +++ b/ci/test_cuda_compute_minimal_python.sh @@ -30,7 +30,7 @@ fi # full cu* extras because those pull in numba/numba-cuda. CUDA_CCCL_WHEEL_PATH="$(ls "${wheelhouse_dir}"/cuda_cccl-*.whl)" python -m pip install "${CUDA_CCCL_WHEEL_PATH}[minimal-cu${cuda_major_version}]" -python -m pip install pytest pytest-xdist pytest-run-parallel +python -m pip install pytest pytest-xdist cd "${repo_root}/python/cuda_cccl/tests/" python -m pytest -n 6 -v compute/test_no_numba.py @@ -57,6 +57,10 @@ if [[ "${py_version}" == "3.14t" ]]; then # kernels and stays reproducible across runners, unlike =auto (the runner's # logical-core count). # + # pytest-run-parallel is only used by this sweep, so install it on the 3.14t + # path rather than for every minimal (e.g. non-free-threaded 3.14) run. + python -m pip install pytest-run-parallel + # Fail fast if the interpreter is not actually GIL-free (wrong build / # PYTHON_GIL=1): pytest-run-parallel does NOT catch a GIL that is enabled from # the start -- it would run threads GIL-serialized and pass vacuously. (A GIL diff --git a/python/cuda_cccl/tests/compute/conftest.py b/python/cuda_cccl/tests/compute/conftest.py index d195c46d6cc..40b5a8d581c 100644 --- a/python/cuda_cccl/tests/compute/conftest.py +++ b/python/cuda_cccl/tests/compute/conftest.py @@ -137,9 +137,13 @@ def pytest_collection_modifyitems(config, items): # would serialize every no_numba test, neutering the sweep. # # getoption returns the int 1 from the argparse default but a *string* for - # CLI-passed values ("2", "auto", "1"), so normalize to str before comparing - # -- a bare `> 1` would raise TypeError on the "2" string. - running_parallel = str(config.getoption("parallel_threads", 1)) != "1" + # CLI-passed values ("2", "auto"). Only bypass for an explicit numeric thread + # count > 1 (what CI passes): "auto" is left to inject the guard so a run the + # plugin executes single-threaded (e.g. "auto" resolving to 1 logical CPU) is + # not mistaken for a parallel run. (A bare `> 1` would also TypeError on the + # "2" string.) + parallel_threads = str(config.getoption("parallel_threads", 1)) + running_parallel = parallel_threads.isdigit() and int(parallel_threads) > 1 for item in items: if item.get_closest_marker("no_numba") and not running_parallel: if "raise_on_numba_import" not in item.fixturenames: From 93a71560479cad3fe3b38d48228dd4f95f207ea0 Mon Sep 17 00:00:00 2001 From: Nader Al Awar Date: Thu, 16 Jul 2026 12:47:21 -0500 Subject: [PATCH 3/4] Add initial TSan CI job --- c/parallel/CMakeLists.txt | 23 +++++ ci/build_cuda_cccl_python.sh | 10 ++ ci/build_cuda_cccl_python_tsan.sh | 12 +++ ci/build_cuda_cccl_wheel.sh | 9 ++ ci/matrix.yaml | 15 +++ ci/test_cuda_compute_minimal_python_tsan.sh | 106 ++++++++++++++++++++ ci/util/workflow/get_wheel_artifact_name.sh | 5 + 7 files changed, 180 insertions(+) create mode 100755 ci/build_cuda_cccl_python_tsan.sh create mode 100755 ci/test_cuda_compute_minimal_python_tsan.sh diff --git a/c/parallel/CMakeLists.txt b/c/parallel/CMakeLists.txt index 7486dd57064..d6974579f4a 100644 --- a/c/parallel/CMakeLists.txt +++ b/c/parallel/CMakeLists.txt @@ -8,6 +8,17 @@ option( "Build cccl.c.parallel standalone headers." OFF ) +# Instrument cccl.c.parallel's host code with ThreadSanitizer for the +# free-threaded (3.14t) TSan CI lane. TSan is a host-only tool, so CUDA device +# code is never instrumented -- only the host portion of each .cu (via +# -Xcompiler). The libtsan runtime is resolved at load time (the TSan lane +# LD_PRELOADs it) rather than linked-and-bundled, so the wheel build must +# --exclude libtsan from auditwheel repair. +option( + CCCL_C_PARALLEL_SANITIZE_THREAD + "Build cccl.c.parallel host code with ThreadSanitizer (-fsanitize=thread)." + OFF +) # FIXME Ideally this would be handled by presets and install rules, but for now # consumers may override this to control the target location of cccl.c.parallel. @@ -45,6 +56,18 @@ if (CCCL_C_PARALLEL_LIBRARY_OUTPUT_DIRECTORY) ) endif() +if (CCCL_C_PARALLEL_SANITIZE_THREAD) + # Host-only: -fsanitize=thread on C++ TUs, and forwarded to the host compiler + # for the host side of each .cu. -g keeps file:line frames in TSan reports. + target_compile_options( + cccl.c.parallel + PRIVATE + $<$:-fsanitize=thread;-g> + $<$:-Xcompiler=-fsanitize=thread;-g> + ) + target_link_options(cccl.c.parallel PRIVATE -fsanitize=thread) +endif() + cccl_get_cub() cccl_get_cudatoolkit() cccl_get_thrust() diff --git a/ci/build_cuda_cccl_python.sh b/ci/build_cuda_cccl_python.sh index 7ac853a427d..a5994942121 100755 --- a/ci/build_cuda_cccl_python.sh +++ b/ci/build_cuda_cccl_python.sh @@ -90,6 +90,7 @@ for ctk in 12 13; do --env "GITHUB_RUN_ID=${GITHUB_RUN_ID:-}" \ --env "JOB_ID=${JOB_ID:-}" \ --env "CCCL_PYTHON_USE_V2=${CCCL_PYTHON_USE_V2:-}" \ + --env "CCCL_C_PARALLEL_SANITIZE_THREAD=${CCCL_C_PARALLEL_SANITIZE_THREAD:-}" \ --env "CCACHE_DIR=/root/.ccache" \ --env "CPM_SOURCE_CACHE=/root/.cpm-cache" \ "$image" \ @@ -133,6 +134,14 @@ echo "Found CUDA 13 wheel: $cu13_wheel" # Merge the wheels python python/cuda_cccl/merge_cuda_wheels.py "$cu12_wheel" "$cu13_wheel" --output-dir wheelhouse_merged +# A ThreadSanitizer wheel links libtsan; keep it external (excluded) so it is +# NOT bundled -- the TSan test lane LD_PRELOADs the runner's matching libtsan +# instead. Harmless for normal builds (the .so has no libtsan dependency). +tsan_exclude=() +if [[ "${CCCL_C_PARALLEL_SANITIZE_THREAD:-}" =~ ^(1|true|TRUE|on|ON)$ ]]; then + tsan_exclude=(--exclude 'libtsan.so.2') +fi + # Install auditwheel and repair the merged wheel python -m pip install patchelf auditwheel for wheel in wheelhouse_merged/cuda_cccl-*.whl; do @@ -145,6 +154,7 @@ for wheel in wheelhouse_merged/cuda_cccl-*.whl; do --exclude 'libcudart.so.12' \ --exclude 'libcudart.so.13' \ --exclude 'libcuda.so.1' \ + "${tsan_exclude[@]}" \ "$wheel" \ --wheel-dir wheelhouse_final done diff --git a/ci/build_cuda_cccl_python_tsan.sh b/ci/build_cuda_cccl_python_tsan.sh new file mode 100755 index 00000000000..d250723e0f0 --- /dev/null +++ b/ci/build_cuda_cccl_python_tsan.sh @@ -0,0 +1,12 @@ +#!/usr/bin/env bash +# Thin wrapper around build_cuda_cccl_python.sh that builds the cuda_cccl wheel +# with ThreadSanitizer instrumentation on the c.parallel (v1) host code, for the +# free-threaded (3.14t) TSan nightly lane. The shared build script honors +# CCCL_C_PARALLEL_SANITIZE_THREAD (passes -DCCCL_C_PARALLEL_SANITIZE_THREAD=ON to +# the wheel build and --exclude libtsan from the auditwheel repair). +set -euo pipefail + +ci_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" + +export CCCL_C_PARALLEL_SANITIZE_THREAD=1 +exec "$ci_dir/build_cuda_cccl_python.sh" "$@" diff --git a/ci/build_cuda_cccl_wheel.sh b/ci/build_cuda_cccl_wheel.sh index 4b95dfda6d5..6bfd0f9594c 100755 --- a/ci/build_cuda_cccl_wheel.sh +++ b/ci/build_cuda_cccl_wheel.sh @@ -91,6 +91,15 @@ if [[ "${CCCL_PYTHON_USE_V2:-}" =~ ^(1|true|TRUE|on|ON)$ ]]; then python -m pip install --upgrade 'cmake>=3.27' fi +# When CCCL_C_PARALLEL_SANITIZE_THREAD is set (=1/true/on), instrument the +# c.parallel (v1) host code with ThreadSanitizer for the free-threaded TSan +# nightly lane. Host-only; the libtsan runtime stays external (the shared +# build_cuda_cccl_python.sh --excludes it from auditwheel). +if [[ "${CCCL_C_PARALLEL_SANITIZE_THREAD:-}" =~ ^(1|true|TRUE|on|ON)$ ]]; then + export CMAKE_ARGS="${CMAKE_ARGS:-} -DCCCL_C_PARALLEL_SANITIZE_THREAD=ON" + echo "Building wheel with ThreadSanitizer-instrumented c.parallel: CMAKE_ARGS=${CMAKE_ARGS}" +fi + # Build the wheel python -m pip wheel --no-deps --verbose --wheel-dir dist . diff --git a/ci/matrix.yaml b/ci/matrix.yaml index 259bba3c3cd..e979afae580 100644 --- a/ci/matrix.yaml +++ b/ci/matrix.yaml @@ -238,6 +238,10 @@ workflows: # so FT regressions (e.g. from dependency bumps) surface between PRs. - {jobs: ['test_py_compute_minimal'], project: 'python', ctk: ['12.X', '13.0', '13.X'], py_version: '3.14t', gpu: 'l4', cxx: 'gcc13'} - {jobs: ['test_py_compute_minimal'], project: 'python_v2', ctk: '13.X', py_version: '3.14t', gpu: 'l4', cxx: 'gcc13'} + # Free-threaded (3.14t) ThreadSanitizer lane: builds c.parallel v1 host code + # with -fsanitize=thread and runs the FT stress + parallel sweep under the + # TSan runtime. + - {jobs: ['test_py_compute_minimal'], project: 'python_tsan', ctk: '13.X', py_version: '3.14t', gpu: 'l4', cxx: 'gcc13'} # CCCL packaging: - {jobs: ['test'], project: 'packaging', ctk: '12.0', cxx: ['gcc10', 'clang14'], gpu: 'rtx2080', args: '-min-cmake'} - {jobs: ['test'], project: 'packaging', ctk: '12.X', cxx: ['gcc10', 'clang14'], gpu: 'rtx2080'} @@ -624,6 +628,17 @@ projects: job_map: build: ['build_py_wheel'] test: ['test_py_par', 'test_py_examples'] + python_tsan: + name: "Python (cuda.compute free-threaded ThreadSanitizer)" + # Nightly-only lane. The producer build_py_wheel runs + # build_cuda_cccl_python_tsan.sh (c.parallel v1 host code built with + # -fsanitize=thread); test_py_compute_minimal runs + # test_cuda_compute_minimal_python_tsan.sh (the FT stress + sweep under the + # TSan runtime). Deliberately absent from the `python-wheels` publish + # workflow, so these instrumented wheels never reach PyPI. + job_map: + build: ['build_py_wheel'] + test: ['test_py_compute_minimal'] cccl_c_parallel: name: 'CCCL C Parallel' stds: [20] diff --git a/ci/test_cuda_compute_minimal_python_tsan.sh b/ci/test_cuda_compute_minimal_python_tsan.sh new file mode 100755 index 00000000000..b29658e1e92 --- /dev/null +++ b/ci/test_cuda_compute_minimal_python_tsan.sh @@ -0,0 +1,106 @@ +#!/usr/bin/env bash + +# ThreadSanitizer variant of the minimal free-threaded cuda.compute lane. +# +# Installs the TSan-instrumented cuda_cccl wheel (produced by the `python_tsan` +# build, which compiles c.parallel v1 host code with -fsanitize=thread) and runs +# the free-threaded stress tests + the pytest-run-parallel sweep under the TSan +# runtime. A real data race in c.parallel (e.g. build-owned state mutated at +# launch and shared across threads) fails the nightly here instead of silently +# corrupting results under some interleaving. +# +# Only meaningful on a free-threaded (3.14t) interpreter; the GIL would serialize +# the threads and hide the very races we are looking for. + +set -euo pipefail + +ci_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +repo_root="$(cd "$ci_dir/.." && pwd)" +source "$ci_dir/pyenv_helper.sh" + +# Parse common arguments +source "$ci_dir/util/python/common_arg_parser.sh" +parse_python_args "$@" +require_py_version "Usage: $0 -py-version " + +if [[ "${py_version}" != *t ]]; then + echo "ERROR: the TSan lane requires a free-threaded (…t) interpreter; got '${py_version}'." >&2 + echo "On a GIL interpreter the sweep serializes and TSan has no signal." >&2 + exit 1 +fi + +# Instrument c.parallel when this script builds the wheel itself (local runs). In +# CI the wheel is the pre-built `python_tsan` artifact, already instrumented; the +# export is harmless there. +export CCCL_C_PARALLEL_SANITIZE_THREAD=1 + +cuda_major_version=$(nvcc --version | grep release | awk '{print $6}' | tr -d ',' | cut -d '.' -f 1 | cut -d 'V' -f 2) + +# Setup Python environment +setup_python_env "${py_version}" + +# Fetch or build the TSan-instrumented cuda_cccl wheel. Under project +# `python_tsan`, get_wheel_artifact_name.sh resolves to the distinct `-tsan` +# artifact, so this never grabs an uninstrumented wheel. +if [[ -n "${GITHUB_ACTIONS:-}" ]]; then + wheel_artifact_name=$("$ci_dir/util/workflow/get_wheel_artifact_name.sh") + "$ci_dir/util/artifacts/download.sh" "${wheel_artifact_name}" "${repo_root}/" + wheelhouse_dir="${repo_root}/wheelhouse" +else + "$ci_dir/build_cuda_cccl_python.sh" -py-version "${py_version}" + wheelhouse_dir="${repo_root}/wheelhouse" +fi + +# minimal-cu* extra intentionally avoids numba/numba-cuda (which re-enable the +# GIL). pytest-run-parallel drives the concurrent sweep. +CUDA_CCCL_WHEEL_PATH="$(ls "${wheelhouse_dir}"/cuda_cccl-*.whl)" +python -m pip install "${CUDA_CCCL_WHEEL_PATH}[minimal-cu${cuda_major_version}]" +python -m pip install pytest pytest-xdist pytest-run-parallel + +# The instrumented .so links libtsan but keeps it external (auditwheel --exclude), +# so the TSan runtime must be present from process start -- LD_PRELOAD the +# runner's libtsan (same soname/major as the gcc-13 build). Without preload the +# .so fails to load ("cannot allocate memory in static TLS block"). +tsan_runtime="$(gcc -print-file-name=libtsan.so.2)" +if [[ ! -e "${tsan_runtime}" ]]; then + echo "ERROR: libtsan.so.2 not found (gcc -print-file-name returned '${tsan_runtime}')." >&2 + exit 1 +fi + +# setarch -R disables ASLR for the process tree. Required: TSan reserves fixed +# shadow-memory regions and aborts ("unexpected memory mapping") when ASLR drops +# something into them (google/sanitizers#1686). Uses personality(ADDR_NO_RANDOMIZE) +# -- if a runner's seccomp profile blocks it, this call fails and must be allowed +# (e.g. --security-opt seccomp=unconfined on the job's container). +# +# ignore_noninstrumented_modules=1: only c.parallel is instrumented, so ignore +# races inside uninstrumented CPython / CUDA libs (avoids boundary false +# positives). halt_on_error=0: report every race, not just the first. +# exitcode=66: exit non-zero when any (unsuppressed) race is found, failing the +# job even though pytest itself passes. +run_under_tsan() { + setarch -R env \ + LD_PRELOAD="${tsan_runtime}" \ + TSAN_OPTIONS="ignore_noninstrumented_modules=1 halt_on_error=0 history_size=7 exitcode=66" \ + "$@" +} + +# Fail fast if the interpreter is not actually GIL-free (wrong build / +# PYTHON_GIL=1): the sweep would run GIL-serialized and pass vacuously. +run_under_tsan python -c "import sys; assert not sys._is_gil_enabled(), 'GIL is enabled; TSan sweep has no signal'" + +cd "${repo_root}/python/cuda_cccl/tests/" + +# Hand-written free-threaded stress scenarios (they spawn their own worker +# threads) + the serialization node-ids, all sharing one interpreter (-n 0). +run_under_tsan python -m pytest -n 0 -v \ + compute/test_free_threading_stress.py \ + compute/test_multi_cc_serialization.py::test_aot_build_result_load_failure_is_shared_and_retryable \ + compute/test_multi_cc_serialization.py::test_aot_serialization_waits_for_canonical_first_load + +# Broad sweep: run the numba-free functional suite with each test executed +# concurrently across threads (barrier-synchronized), exercising many more +# c.parallel algorithms under contention than the hand-written stress tests. +# -n 0 so the threads share one interpreter; --parallel-threads=2 bounds +# GPU-memory pressure and stays reproducible across runners. +run_under_tsan python -m pytest -n 0 -v --parallel-threads=2 compute/test_no_numba.py diff --git a/ci/util/workflow/get_wheel_artifact_name.sh b/ci/util/workflow/get_wheel_artifact_name.sh index ffd654fe440..b1058c1e75e 100755 --- a/ci/util/workflow/get_wheel_artifact_name.sh +++ b/ci/util/workflow/get_wheel_artifact_name.sh @@ -58,6 +58,11 @@ done suffix="" if [[ "$project" == "python_v2" ]]; then suffix="-v2" +elif [[ "$project" == "python_tsan" ]]; then + # ThreadSanitizer-instrumented wheel (free-threaded TSan nightly lane). Must + # be distinct so its build doesn't clobber the normal wheel and the TSan test + # job doesn't grab an uninstrumented one. + suffix="-tsan" fi echo "wheel-cccl${suffix}-$os-$arch-py$py_version" From fff293af45d8fb209dd51324e5ad3d680d1f454e Mon Sep 17 00:00:00 2001 From: Nader Al Awar Date: Thu, 16 Jul 2026 16:15:59 -0500 Subject: [PATCH 4/4] Fix TSan bug --- c/parallel/src/segmented_sort.cu | 53 ++++++++++++++------------------ 1 file changed, 23 insertions(+), 30 deletions(-) diff --git a/c/parallel/src/segmented_sort.cu b/c/parallel/src/segmented_sort.cu index 1570c6a0cde..e7357294bed 100644 --- a/c/parallel/src/segmented_sort.cu +++ b/c/parallel/src/segmented_sort.cu @@ -177,7 +177,6 @@ struct selector_state_t }; cccl_op_t make_segments_selector_op( - OffsetT offset, cccl_iterator_t begin_offset_iterator, cccl_iterator_t end_offset_iterator, const char* selector_op_name, @@ -188,7 +187,6 @@ cccl_op_t make_segments_selector_op( size_t num_lto_opts) { cccl_op_t selector_op{}; - auto selector_op_state = std::make_unique(); std::string offset_t; check(cccl_type_name_from_nvrtc(&offset_t)); @@ -241,15 +239,8 @@ extern "C" __device__ void {0}(void* state_ptr, const void* arg_ptr, void* resul selector_op.code_type = CCCL_OP_LTOIR; selector_op.size = sizeof(selector_state_t); selector_op.alignment = alignof(selector_state_t); - - selector_op_state->initialize(offset, begin_offset_iterator, end_offset_iterator); - auto* state_copy = static_cast(std::malloc(sizeof(selector_state_t))); - if (!state_copy) - { - throw std::bad_alloc(); - } - std::memcpy(state_copy, selector_op_state.get(), sizeof(selector_state_t)); - selector_op.state = state_copy; + // No state in the build result: it is per-call (see segmented_sort_kernel_source). + selector_op.state = nullptr; return selector_op; } @@ -279,6 +270,8 @@ struct owned_selector_op struct segmented_sort_kernel_source { cccl_device_segmented_sort_build_result_t& build; + selector_state_t* large_state; + selector_state_t* small_state; CUkernel SegmentedSortFallbackKernel() const { @@ -304,17 +297,20 @@ struct segmented_sort_kernel_source indirect_arg_t LargeSegmentsSelector( OffsetT offset, indirect_iterator_t begin_offset_iterator, indirect_iterator_t end_offset_iterator) const { - static_cast(build.large_segments_selector_op.state) - ->initialize(offset, begin_offset_iterator, end_offset_iterator); - return indirect_arg_t(build.large_segments_selector_op); + large_state->initialize(offset, begin_offset_iterator, end_offset_iterator); + // Same compiled op, pointed at this call's state instead of the shared one. + cccl_op_t op = build.large_segments_selector_op; + op.state = large_state; + return indirect_arg_t(op); } indirect_arg_t SmallSegmentsSelector( OffsetT offset, indirect_iterator_t begin_offset_iterator, indirect_iterator_t end_offset_iterator) const { - static_cast(build.small_segments_selector_op.state) - ->initialize(offset, begin_offset_iterator, end_offset_iterator); - return indirect_arg_t(build.small_segments_selector_op); + small_state->initialize(offset, begin_offset_iterator, end_offset_iterator); + cccl_op_t op = build.small_segments_selector_op; + op.state = small_state; + return indirect_arg_t(op); } void SetSegmentOffset(indirect_arg_t& selector, long long base_segment_offset) const @@ -487,7 +483,6 @@ try // DispatchThreeWayPartition eventually. This causes increased compilation // times, which might be avoidable. segmented_sort::owned_selector_op large_selector_op{segmented_sort::make_segments_selector_op( - 0, start_offset_it, end_offset_it, "cccl_large_segments_selector_op", @@ -497,7 +492,6 @@ try lopts, num_lto_args)}; segmented_sort::owned_selector_op small_selector_op{segmented_sort::make_segments_selector_op( - 0, start_offset_it, end_offset_it, "cccl_small_segments_selector_op", @@ -908,6 +902,11 @@ CUresult cccl_device_segmented_sort_impl( cub::DoubleBuffer d_values_double_buffer( *static_cast(&val_arg_in), *static_cast(&val_arg_out)); + // Per-call state, stack-local so concurrent launches stay independent. Safe under + // async: copied by value into kernel params at each launch within this dispatch. + segmented_sort::selector_state_t large_selector_state{}; + segmented_sort::selector_state_t small_selector_state{}; + auto exec_status = cub::detail::segmented_sort::dispatch( d_temp_storage, *temp_storage_bytes, @@ -923,7 +922,8 @@ CUresult cccl_device_segmented_sort_impl( *static_cast(build.runtime_policy), /* partition_policy_selector */ *static_cast(build.partition_runtime_policy), - /* kernel_source */ segmented_sort::segmented_sort_kernel_source{build}, + /* kernel_source */ + segmented_sort::segmented_sort_kernel_source{build, &large_selector_state, &small_selector_state}, /* partition_kernel_source */ segmented_sort::partition_kernel_source{build}, /* launcher_factory */ cub::detail::CudaDriverLauncherFactory{cu_device, build.cc}); @@ -1119,23 +1119,16 @@ inline cccl_op_t deserialize_selector_op(cccl::serialization::buffer_reader& r, r.read_bytes(code_owner.get(), static_cast(code_n)); } - // The selector state is runtime-specific (device pointers, segment offsets) and is - // fully overwritten by selector_state_t::initialize() before each launch, so it is - // not part of the blob. Reconstruct a fresh, zero-initialized state of the correct - // size/alignment here; a malformed blob can no longer dictate the state size. + // No selector state to reconstruct (it is per-call). op.size stays hardcoded so a + // malformed blob cannot dictate it. using segmented_sort::selector_state_t; - std::unique_ptr state_owner(std::calloc(1, sizeof(selector_state_t)), std::free); - if (!state_owner) - { - throw std::bad_alloc{}; - } // commit — no throws past this point op.size = sizeof(selector_state_t); op.alignment = alignof(selector_state_t); op.code = static_cast(code_owner.release()); op.code_size = static_cast(code_n); - op.state = state_owner.release(); + op.state = nullptr; // op.name is not freed by cleanup; safe to point at a string literal. op.name = fixed_name; // Selector ops never carry extra ltoirs in this codepath.