Merge branch 'main' into vk/dispatcher

Author: VijayKandiah

README.md

@@ -37,12 +37,19 @@ they are compiled for the appropriate compute capability.
 ```
 cd testing
 # Optionally, build test binaries and point to their location for the test suite
-make
+make -j $(nproc)
 export NUMBA_CUDA_TEST_BIN_DIR=`pwd`
 # Execute tests
 pytest -n auto -v
 ```
 
+Alternatively, you can use [pixi](https://pixi.sh/latest/installation/) to wrap all of that up for you:
+
+```
+# run tests against CUDA 13
+pixi run -e cu13 test -n auto -v
+```
+
 
 Testing should discover the `numba.cuda` module from the `numba_cuda` package. You
 can check where `numba.cuda` files are being located by running

ci/coverage_report.sh

@@ -20,7 +20,7 @@ python -m pip install \
 rapids-logger "Build test binaries"
 export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing
 pushd $NUMBA_CUDA_TEST_BIN_DIR
-make
+make -j $(nproc)
 
 rapids-logger "Check GPU usage"
 nvidia-smi

ci/test_conda.sh

@@ -66,7 +66,7 @@ nvidia-smi
 rapids-logger "Build test binaries"
 export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing
 pushd $NUMBA_CUDA_TEST_BIN_DIR
-make
+make -j $(nproc)
 
 rapids-logger "Show Numba system info"
 python -m numba --sysinfo

ci/test_conda_ctypes_binding.sh

@@ -52,7 +52,7 @@ nvidia-smi
 rapids-logger "Build test binaries"
 export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing
 pushd $NUMBA_CUDA_TEST_BIN_DIR
-make
+make -j $(nproc)
 
 rapids-logger "Show Numba system info"
 python -m numba --sysinfo

ci/test_wheel.sh

@@ -28,7 +28,7 @@ python -m pip install "${DEPENDENCIES[@]}"
 rapids-logger "Build tests"
 export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing
 pushd $NUMBA_CUDA_TEST_BIN_DIR
-make
+make -j $(nproc)
 
 rapids-logger "Test importing numba.cuda"
 python -c "from numba import cuda"

ci/test_wheel_ctypes_binding.sh

@@ -19,7 +19,7 @@ python -m pip install \
 rapids-logger "Copy and cd into test binaries dir"
 export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing
 pushd $NUMBA_CUDA_TEST_BIN_DIR
-# make
+# make -j $(nproc)
 
 # Prevent the testsuite trying to use the test binaries
 unset NUMBA_CUDA_TEST_BIN_DIR

ci/test_wheel_deps_wheels.sh

@@ -17,7 +17,7 @@ rapids-logger "Build test binaries"
 
 export NUMBA_CUDA_TEST_BIN_DIR=`pwd`/testing
 pushd $NUMBA_CUDA_TEST_BIN_DIR
-make
+make -j $(nproc)
 
 rapids-logger "Check GPU usage"
 nvidia-smi

testing/Makefile

@@ -1,5 +1,8 @@
 # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: BSD-2-Clause
+.PHONY: all nrt_extern test_device_functions clean
+
+.DEFAULT_GOAL := all
 
 # Generates the input files used by the nvjitlink tests
 
@@ -45,28 +48,79 @@ OUTPUT_DIR := ./
 
 NRT_INCLUDE_DIR := $(shell python -c "from numba.cuda.memory_management.nrt import get_include; print(get_include())")
 
-all:
-	@echo "GPU CC: $(GPU_CC)"
-	@echo "Alternative CC: $(ALT_CC)"
-	# Compile all test objects
-	nvcc $(NVCC_FLAGS) $(CUBIN_FLAGS) -o $(OUTPUT_DIR)/undefined_extern.cubin undefined_extern.cu
-	nvcc $(NVCC_FLAGS) $(CUBIN_FLAGS) -o $(OUTPUT_DIR)/test_device_functions.cubin test_device_functions.cu
-	nvcc $(NVCC_FLAGS) $(FATBIN_FLAGS) -o $(OUTPUT_DIR)/test_device_functions.fatbin test_device_functions.cu
-	nvcc $(NVCC_FLAGS) $(MULTI_FATBIN_FLAGS) -o $(OUTPUT_DIR)/test_device_functions_multi.fatbin test_device_functions.cu
-	nvcc $(NVCC_FLAGS) $(PTX_FLAGS) -o $(OUTPUT_DIR)/test_device_functions.ptx test_device_functions.cu
-	nvcc $(NVCC_FLAGS) $(OBJECT_FLAGS) -o $(OUTPUT_DIR)/test_device_functions.o test_device_functions.cu
-	nvcc $(NVCC_FLAGS) $(LIBRARY_FLAGS) -o $(OUTPUT_DIR)/test_device_functions.a test_device_functions.cu
-
-	nvcc $(NVCC_FLAGS) $(CUBIN_FLAGS) -o $(OUTPUT_DIR)/nrt_extern.cubin nrt_extern.cu -I$(NRT_INCLUDE_DIR)
-	nvcc $(NVCC_FLAGS) $(FATBIN_FLAGS) -o $(OUTPUT_DIR)/nrt_extern.fatbin nrt_extern.cu -I$(NRT_INCLUDE_DIR)
-	nvcc $(NVCC_FLAGS) $(MULTI_FATBIN_FLAGS) -o $(OUTPUT_DIR)/nrt_extern_multi.fatbin nrt_extern.cu -I$(NRT_INCLUDE_DIR)
-	nvcc $(NVCC_FLAGS) $(PTX_FLAGS) -o $(OUTPUT_DIR)/nrt_extern.ptx nrt_extern.cu -I$(NRT_INCLUDE_DIR)
-	nvcc $(NVCC_FLAGS) $(OBJECT_FLAGS) -o $(OUTPUT_DIR)/nrt_extern.o nrt_extern.cu -I$(NRT_INCLUDE_DIR)
-	nvcc $(NVCC_FLAGS) $(LIBRARY_FLAGS) -o $(OUTPUT_DIR)/nrt_extern.a nrt_extern.cu -I$(NRT_INCLUDE_DIR)
-
-	# Generate LTO-IR wrapped in a fatbin
-	nvcc $(NVCC_FLAGS) $(LTOIR_FLAGS) -o $(OUTPUT_DIR)/test_device_functions.ltoir.o test_device_functions.cu
-	nvcc $(NVCC_FLAGS) $(LTOIR_FLAGS) -o $(OUTPUT_DIR)/nrt_extern.ltoir.o nrt_extern.cu -I$(NRT_INCLUDE_DIR)
+$(OUTPUT_DIR)/undefined_extern.cubin: undefined_extern.cu
+	nvcc $(NVCC_FLAGS) $(CUBIN_FLAGS) -o $@ $<
+
+$(OUTPUT_DIR)/test_device_functions.cubin: test_device_functions.cu
+	nvcc $(NVCC_FLAGS) $(CUBIN_FLAGS) -o $@ $<
+
+$(OUTPUT_DIR)/test_device_functions.fatbin: test_device_functions.cu
+	nvcc $(NVCC_FLAGS) $(FATBIN_FLAGS) -o $@ $<
+
+$(OUTPUT_DIR)/test_device_functions_multi.fatbin: test_device_functions.cu
+	nvcc $(NVCC_FLAGS) $(MULTI_FATBIN_FLAGS) -o $@ $<
+
+$(OUTPUT_DIR)/test_device_functions.ptx: test_device_functions.cu
+	nvcc $(NVCC_FLAGS) $(PTX_FLAGS) -o $@ $<
+
+$(OUTPUT_DIR)/test_device_functions.o: test_device_functions.cu
+	nvcc $(NVCC_FLAGS) $(OBJECT_FLAGS) -o $@ $<
+
+$(OUTPUT_DIR)/test_device_functions.a: test_device_functions.cu
+	nvcc $(NVCC_FLAGS) $(LIBRARY_FLAGS) -o $@ $<
+
+$(OUTPUT_DIR)/test_device_functions.ltoir.o: test_device_functions.cu
+	nvcc $(NVCC_FLAGS) $(LTOIR_FLAGS) -o $@ $<
+
+$(OUTPUT_DIR)/test_device_functions.ltoir: test_device_functions.cu
 	# Generate LTO-IR in a "raw" LTO-IR container
-	python generate_raw_ltoir.py --arch sm_$(GPU_CC) -o $(OUTPUT_DIR)/test_device_functions.ltoir test_device_functions.cu
-	python generate_raw_ltoir.py --arch sm_$(GPU_CC) -o $(OUTPUT_DIR)/nrt_extern.ltoir nrt_extern.cu --nrt
+	python generate_raw_ltoir.py --arch sm_$(GPU_CC) -o $(OUTPUT_DIR)/test_device_functions.ltoir $<
+
+test_device_functions: $(OUTPUT_DIR)/test_device_functions.cubin \
+    $(OUTPUT_DIR)/test_device_functions.fatbin \
+    $(OUTPUT_DIR)/test_device_functions_multi.fatbin \
+    $(OUTPUT_DIR)/test_device_functions.ptx \
+    $(OUTPUT_DIR)/test_device_functions.o \
+    $(OUTPUT_DIR)/test_device_functions.a \
+    $(OUTPUT_DIR)/test_device_functions.ltoir.o \
+    $(OUTPUT_DIR)/test_device_functions.ltoir
+
+$(OUTPUT_DIR)/nrt_extern.cubin: nrt_extern.cu
+	nvcc $(NVCC_FLAGS) $(CUBIN_FLAGS) -o $@ $< -I$(NRT_INCLUDE_DIR)
+
+$(OUTPUT_DIR)/nrt_extern.fatbin: nrt_extern.cu
+	nvcc $(NVCC_FLAGS) $(FATBIN_FLAGS) -o $@ $< -I$(NRT_INCLUDE_DIR)
+
+$(OUTPUT_DIR)/nrt_extern_multi.fatbin: nrt_extern.cu
+	nvcc $(NVCC_FLAGS) $(MULTI_FATBIN_FLAGS) -o $@ $< -I$(NRT_INCLUDE_DIR)
+
+$(OUTPUT_DIR)/nrt_extern.ptx: nrt_extern.cu
+	nvcc $(NVCC_FLAGS) $(PTX_FLAGS) -o $@ $< -I$(NRT_INCLUDE_DIR)
+
+$(OUTPUT_DIR)/nrt_extern.o: nrt_extern.cu
+	nvcc $(NVCC_FLAGS) $(OBJECT_FLAGS) -o $@ $< -I$(NRT_INCLUDE_DIR)
+
+$(OUTPUT_DIR)/nrt_extern.a: nrt_extern.cu
+	nvcc $(NVCC_FLAGS) $(LIBRARY_FLAGS) -o $@ $< -I$(NRT_INCLUDE_DIR)
+
+$(OUTPUT_DIR)/nrt_extern.ltoir.o: nrt_extern.cu
+	nvcc $(NVCC_FLAGS) $(LTOIR_FLAGS) -o $@ $< -I$(NRT_INCLUDE_DIR)
+
+$(OUTPUT_DIR)/nrt_extern.ltoir: nrt_extern.cu
+	python generate_raw_ltoir.py --arch sm_$(GPU_CC) -o $@ $< --nrt
+
+nrt_extern: $(OUTPUT_DIR)/nrt_extern.cubin \
+    $(OUTPUT_DIR)/nrt_extern.fatbin \
+    $(OUTPUT_DIR)/nrt_extern_multi.fatbin \
+    $(OUTPUT_DIR)/nrt_extern.ptx \
+    $(OUTPUT_DIR)/nrt_extern.o \
+    $(OUTPUT_DIR)/nrt_extern.a \
+    $(OUTPUT_DIR)/nrt_extern.ltoir.o \
+    $(OUTPUT_DIR)/nrt_extern.ltoir
+
+all: test_device_functions nrt_extern $(OUTPUT_DIR)/undefined_extern.cubin
+	@>&2 echo "GPU CC: $(GPU_CC)"
+	@>&2 echo "Alternative CC: $(ALT_CC)"
+
+clean:
+	rm -f *.cubin *.fatbin *.ptx *.o *.a *.ltoir