[cuda.compute]: add nightly CI job that builds c.parallel with thread sanitizer

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
+      $<$<COMPILE_LANGUAGE:CXX>:-fsanitize=thread;-g>
+      $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-fsanitize=thread;-g>
+  )
+  target_link_options(cccl.c.parallel PRIVATE -fsanitize=thread)
+endif()
+
 cccl_get_cub()
 cccl_get_cudatoolkit()
 cccl_get_thrust()

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<selector_state_t>();
   std::string offset_t;
   check(cccl_type_name_from_nvrtc<OffsetT>(&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<selector_state_t*>(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<selector_state_t*>(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<selector_state_t*>(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<indirect_arg_t> d_values_double_buffer(
       *static_cast<indirect_arg_t**>(&val_arg_in), *static_cast<indirect_arg_t**>(&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<Order, OffsetT>(
       d_temp_storage,
       *temp_storage_bytes,
@@ -923,7 +922,8 @@ CUresult cccl_device_segmented_sort_impl(
       *static_cast<cub::detail::segmented_sort::policy_selector*>(build.runtime_policy),
       /* partition_policy_selector */
       *static_cast<cub::detail::three_way_partition::policy_selector*>(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<size_t>(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<void, decltype(&std::free)> 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<const char*>(code_owner.release());
   op.code_size = static_cast<size_t>(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.

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

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" "$@"

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 .
 

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]

ci/test_cuda_compute_minimal_python.sh

@@ -44,4 +44,28 @@ 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).
+  #
+  # 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
+  # *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

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 <python_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