Support Python 3.14 free-threaded CUDA wheels

.github/workflows/build-wheel.yml

@@ -33,7 +33,7 @@ jobs:
           - "3.12"
           - "3.13"
           - "3.14"
-          # - "3.14t"
+          - "3.14t"
     name: py${{ matrix.python-version }}
     runs-on: ${{ (inputs.host-platform == 'linux-64' && 'linux-amd64-cpu8') ||
                  (inputs.host-platform == 'linux-aarch64' && 'linux-arm64-cpu8') ||

.github/workflows/ci.yaml

@@ -232,7 +232,17 @@ jobs:
       build_type: pull-request
       script: "ci/test_wheel_deps_wheels.sh"
       matrix: ${{ needs.compute-matrix.outputs.TEST_MATRIX }}
-      matrix_filter: map(select(.ARCH == "amd64" and (.CUDA_VER | split(".") | .[0] | tonumber >= 12))) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
+      matrix_filter: >-
+        map(select(
+          .ARCH == "amd64" and
+          (.CUDA_VER | split(".") | .[0] | tonumber >= 12)
+        ))
+        | group_by(.CUDA_VER | split(".") | map(tonumber) | .[0])
+        | map(max_by([
+          (.PY_VER | sub("t$"; "") | split(".") | map(tonumber)),
+          (.PY_VER | endswith("t") | not),
+          (.CUDA_VER | split(".") | map(tonumber))
+        ]))
 
   test-thirdparty-cudf:
     needs:
@@ -244,7 +254,17 @@ jobs:
       script: "ci/test_thirdparty_cudf.sh"
       matrix: ${{ needs.compute-matrix.outputs.TEST_MATRIX }}
       # TODO: Enable for CUDA 13 when a supporting version of cuDF is available
-      matrix_filter: map(select(.ARCH == "amd64" and (.CUDA_VER | split(".") | .[0] | tonumber == 12))) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
+      matrix_filter: >-
+        map(select(
+          .ARCH == "amd64" and
+          (.CUDA_VER | split(".") | .[0] | tonumber == 12)
+        ))
+        | group_by(.CUDA_VER | split(".") | map(tonumber) | .[0])
+        | map(max_by([
+          (.PY_VER | sub("t$"; "") | split(".") | map(tonumber)),
+          (.PY_VER | endswith("t") | not),
+          (.CUDA_VER | split(".") | map(tonumber))
+        ]))
 
   test-thirdparty-nvmath:
     needs:
@@ -256,7 +276,17 @@ jobs:
       script: "ci/test_thirdparty_nvmath.sh"
       matrix: ${{ needs.compute-matrix.outputs.TEST_MATRIX }}
       # TODO: Enable for CUDA 13 when a supporting version of nvmath-python is available
-      matrix_filter: map(select(.ARCH == "amd64" and (.CUDA_VER | split(".") | .[0] | tonumber == 12))) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
+      matrix_filter: >-
+        map(select(
+          .ARCH == "amd64" and
+          (.CUDA_VER | split(".") | .[0] | tonumber == 12)
+        ))
+        | group_by(.CUDA_VER | split(".") | map(tonumber) | .[0])
+        | map(max_by([
+          (.PY_VER | sub("t$"; "") | split(".") | map(tonumber)),
+          (.PY_VER | endswith("t") | not),
+          (.CUDA_VER | split(".") | map(tonumber))
+        ]))
 
   test-thirdparty-awkward:
     needs:
@@ -268,7 +298,17 @@ jobs:
       script: "ci/test_thirdparty_awkward.sh"
       matrix: ${{ needs.compute-matrix.outputs.TEST_MATRIX }}
       # TODO: Enable for CUDA 13 in future
-      matrix_filter: map(select(.ARCH == "amd64" and (.CUDA_VER | split(".") | .[0] | tonumber == 12))) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
+      matrix_filter: >-
+        map(select(
+          .ARCH == "amd64" and
+          (.CUDA_VER | split(".") | .[0] | tonumber == 12)
+        ))
+        | group_by(.CUDA_VER | split(".") | map(tonumber) | .[0])
+        | map(max_by([
+          (.PY_VER | sub("t$"; "") | split(".") | map(tonumber)),
+          (.PY_VER | endswith("t") | not),
+          (.CUDA_VER | split(".") | map(tonumber))
+        ]))
 
   build-docs:
     needs:
@@ -285,7 +325,13 @@ jobs:
       build_type: pull-request
       script: "ci/coverage_report.sh"
       matrix: ${{ needs.compute-matrix.outputs.TEST_MATRIX }}
-      matrix_filter: 'map(select(.ARCH == "amd64" and .CUDA_VER == "12.9.1" and .PY_VER == "3.12")) | .[0:1]'
+      matrix_filter: >-
+        map(select(
+          .ARCH == "amd64" and
+          .CUDA_VER == "12.9.1" and
+          .PY_VER == "3.12"
+        ))
+        | .[0:1]
 
   # =============================================================
 

.gitignore

@@ -6,6 +6,7 @@ __pycache__
 build
 .*.swp
 *.so
+*.pyd
 numba_cuda/numba/cuda/tests/cudadrv/test_device_functions.*
 numba_cuda/numba/cuda/tests/cudadrv/undefined_extern.*
 testing/*.a

ci/test-matrix.yml

@@ -34,9 +34,7 @@ linux:
     - { ARCH: 'amd64', PY_VER: '3.14',  CUDA_VER: '12.9.1', LOCAL_CTK: '0', GPU: 't4',         GPU_COUNT: '1', DRIVER: 'latest' }
     - { ARCH: 'amd64', PY_VER: '3.14',  CUDA_VER: '13.0.2', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest' }
     - { ARCH: 'amd64', PY_VER: '3.14',  CUDA_VER: '13.2.1', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest' }
-    # - { ARCH: 'amd64', PY_VER: '3.14t', CUDA_VER: '12.9.1', LOCAL_CTK: '1', GPU: 't4',         GPU_COUNT: '1', DRIVER: 'latest' }
-    # - { ARCH: 'amd64', PY_VER: '3.14t', CUDA_VER: '13.0.2', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest' }
-    # - { ARCH: 'amd64', PY_VER: '3.14t', CUDA_VER: '13.2.1', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest' }
+    - { ARCH: 'amd64', PY_VER: '3.14t', CUDA_VER: '13.2.1', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest' }
     # linux-aarch64
     - { ARCH: 'arm64', PY_VER: '3.10',  CUDA_VER: '12.9.1', LOCAL_CTK: '1', GPU: 'a100',       GPU_COUNT: '1', DRIVER: 'latest' }
     - { ARCH: 'arm64', PY_VER: '3.10',  CUDA_VER: '13.0.2', LOCAL_CTK: '0', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest' }
@@ -53,9 +51,7 @@ linux:
     - { ARCH: 'arm64', PY_VER: '3.14',  CUDA_VER: '12.9.1', LOCAL_CTK: '0', GPU: 'a100',       GPU_COUNT: '1', DRIVER: 'latest' }
     - { ARCH: 'arm64', PY_VER: '3.14',  CUDA_VER: '13.0.2', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest' }
     - { ARCH: 'arm64', PY_VER: '3.14',  CUDA_VER: '13.2.1', LOCAL_CTK: '1', GPU: 'a100',       GPU_COUNT: '1', DRIVER: 'latest' }
-    # - { ARCH: 'arm64', PY_VER: '3.14t', CUDA_VER: '12.9.1', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest' }
-    # - { ARCH: 'arm64', PY_VER: '3.14t', CUDA_VER: '13.0.2', LOCAL_CTK: '0', GPU: 'a100',       GPU_COUNT: '1', DRIVER: 'latest' }
-    # - { ARCH: 'arm64', PY_VER: '3.14t', CUDA_VER: '13.2.1', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest' }
+    - { ARCH: 'arm64', PY_VER: '3.14t', CUDA_VER: '13.2.1', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest' }
     # special runners
     - { ARCH: 'amd64', PY_VER: '3.13',  CUDA_VER: '13.0.2', LOCAL_CTK: '1', GPU: 'h100',       GPU_COUNT: '1', DRIVER: 'latest' }
     - { ARCH: 'amd64', PY_VER: '3.13',  CUDA_VER: '13.2.1', LOCAL_CTK: '1', GPU: 'h100',       GPU_COUNT: '1', DRIVER: 'latest' }
@@ -83,7 +79,5 @@ windows:
     - { ARCH: 'amd64', PY_VER: '3.14',  CUDA_VER: '12.9.1', LOCAL_CTK: '0', GPU: 'v100',       GPU_COUNT: '1', DRIVER: 'latest', DRIVER_MODE: 'TCC' }
     - { ARCH: 'amd64', PY_VER: '3.14',  CUDA_VER: '13.0.2', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest', DRIVER_MODE: 'MCDM' }
     - { ARCH: 'amd64', PY_VER: '3.14',  CUDA_VER: '13.2.1', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest', DRIVER_MODE: 'MCDM' }
-    # - { ARCH: 'amd64', PY_VER: '3.14t', CUDA_VER: '12.9.1', LOCAL_CTK: '1', GPU: 'l4',         GPU_COUNT: '1', DRIVER: 'latest', DRIVER_MODE: 'TCC' }
-    # - { ARCH: 'amd64', PY_VER: '3.14t', CUDA_VER: '13.0.2', LOCAL_CTK: '0', GPU: 'a100',       GPU_COUNT: '1', DRIVER: 'latest', DRIVER_MODE: 'MCDM' }
-    # - { ARCH: 'amd64', PY_VER: '3.14t', CUDA_VER: '13.2.1', LOCAL_CTK: '0', GPU: 'a100',       GPU_COUNT: '1', DRIVER: 'latest', DRIVER_MODE: 'MCDM' }
+    - { ARCH: 'amd64', PY_VER: '3.14t', CUDA_VER: '13.2.1', LOCAL_CTK: '0', GPU: 'a100',       GPU_COUNT: '1', DRIVER: 'latest', DRIVER_MODE: 'MCDM' }
   nightly: []

docs/source/developer/free_threading.rst

@@ -0,0 +1,77 @@
+..
+   SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+   SPDX-License-Identifier: BSD-2-Clause
+
+Free-threading
+==============
+
+Free-threaded CPython test coverage requires a free-threaded Python build,
+for example Python 3.14t with the ``cp314t`` ABI tag. The regular test suite
+contains free-threaded smoke tests that run automatically in such an
+environment. The heavier stress tests are opt-in because they create many
+threads and subprocesses and are intended for local or dedicated CI runs.
+
+Run the free-threading stress tests with::
+
+   $ PYTHON_GIL=0 NUMBA_CUDA_FT_STRESS=1 \
+       python -m pytest -q --pyargs numba.cuda.tests.stress
+
+``PYTHON_GIL=0`` keeps the GIL disabled for free-threaded CPython builds. The
+stress tests also check the build metadata and skip when Python is not a
+free-threaded build.
+
+Stress test environment variables
+---------------------------------
+
+.. envvar:: NUMBA_CUDA_FT_STRESS
+
+   Enables the opt-in free-threading stress tests when set to ``1``,
+   ``true``, ``yes``, or ``on``. Without this variable, tests in
+   ``numba.cuda.tests.stress.test_free_threading`` are skipped.
+
+.. envvar:: NUMBA_CUDA_FT_STRESS_SECONDS
+
+   Controls the duration, in seconds, of timed CUDA stress tests. The default
+   is ``30`` seconds.
+
+.. envvar:: NUMBA_CUDA_FT_STRESS_WORKERS
+
+   Controls the thread count used by thread-pool stress tests. Defaults vary
+   by test and are capped at the detected CPU count unless this variable is
+   set explicitly.
+
+.. envvar:: NUMBA_CUDA_FT_STRESS_PROCESSES
+
+   Controls the subprocess count used by cache-concurrency stress tests.
+   Defaults vary by test and are capped at the detected CPU count unless this
+   variable is set explicitly.
+
+.. envvar:: NUMBA_CUDA_FT_STRESS_ITERS
+
+   Overrides iteration counts for loop-based stress tests. Defaults vary by
+   test, for example fingerprinting, memoryview buffer helpers, and type
+   conversion stress cases.
+
+Suggested stress profiles
+-------------------------
+
+Use the defaults for a quick local run. On many-core systems, explicitly set
+the worker and process counts to exercise concurrent dispatcher, cache,
+driver, and helper-extension paths more aggressively, for example::
+
+   $ PYTHON_GIL=0 NUMBA_CUDA_FT_STRESS=1 \
+       NUMBA_CUDA_FT_STRESS_SECONDS=120 \
+       NUMBA_CUDA_FT_STRESS_WORKERS=96 \
+       NUMBA_CUDA_FT_STRESS_PROCESSES=24 \
+       NUMBA_CUDA_FT_STRESS_ITERS=10000 \
+       python -m pytest -q --pyargs numba.cuda.tests.stress
+
+If a stress failure is hard to reproduce, increase
+``NUMBA_CUDA_FT_STRESS_SECONDS`` for CUDA tests or
+``NUMBA_CUDA_FT_STRESS_ITERS`` for loop-based no-CUDA tests.
+
+The CUDA dispatch stress tests deliberately avoid a concurrent tight loop of
+full ``gc.collect()`` calls. On many-core CPython 3.14t free-threaded builds,
+that pattern timed out in CPython's ``Python/gc_free_threading.c`` and turned
+the dispatch stress into a CPython full-GC progress test. The dispatch stress
+still performs a full collection after worker shutdown.

docs/source/developer/index.rst

@@ -0,0 +1,10 @@
+..
+   SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+   SPDX-License-Identifier: BSD-2-Clause
+
+Developer documentation
+=======================
+
+.. toctree::
+
+   free_threading.rst

docs/source/index.rst

@@ -28,3 +28,4 @@ Contents
 
    user/index.rst
    reference/index.rst
+   developer/index.rst

docs/source/user/installation.rst

@@ -49,6 +49,21 @@ CUDA 13 dependencies can be installed via ``pip`` with::
 
     $ pip install numba-cuda[cu13]
 
+Free-threaded Python
+--------------------
+
+Free-threaded CPython wheels use the ``cp314t`` ABI tag and are supported for
+Python 3.14t. A conda-forge environment for local smoke testing can be created
+with::
+
+    $ mamba create -n numba-cuda-py314t -c conda-forge \
+        "python=3.14.*=*_cp314t" numba llvmlite numpy \
+        cuda-bindings cuda-core cuda-pathfinder packaging pip
+
+The ``python=3.14.*=*_cp314t`` constraint selects a free-threaded Python build.
+In a supported environment, importing ``numba.cuda`` with ``PYTHON_GIL=0`` must
+leave ``sys._is_gil_enabled()`` false.
+
 If you are not using Conda/pip or if you want to use a different version of CUDA
 toolkit, :ref:`cudatoolkit-lookup` describes how Numba searches for a CUDA toolkit.