[FEA]: Support for finding `libcudadevrt.a` through pathfinder

[brandon-b-miller]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

Area

cuda.bindings

Is your feature request related to a problem? Please describe.

numba-cuda needs to be able to find libcudadevrt.a to launch kernels involving cooperative groups, but needs to find it itself via get_cudalib. Ideally the pathfinder could provide this for us.

Describe the solution you'd like

Something like path_finder._get_nvidia_static_library("cudadevrt")

Describe alternatives you've considered

Numba's existing logic

Additional context

NVIDIA/numba-cuda#302

[leofang]

numba-cuda needs to be able to find libcudadevrt.a to launch kernels involving cooperative groups

@brandon-b-miller could you confirm if this is true? AFAIK cooperative launch does not require libcudadevrt. The only CUDA feature that needs it is CDP1 and CDP2 (CUDA Dynamic Parallelism).

[brandon-b-miller]

In the relevant codepath, there's this note in the numba-cuda source code:

# We need to link against cudadevrt if grid sync is being used.

https://github.com/NVIDIA/numba-cuda/blob/main/numba_cuda/numba/cuda/dispatcher.py#L157-L161

[leofang]

That's why I would like us to test it locally before determining the priority, because we have the same grid.sync() test in the CI and I am pretty sure libcudadevrt is not needed for CG, only for CDP1/2 (which I added to CuPy) 🙂

cuda-python/cuda_core/tests/test_launcher.py

Lines 157 to 199 in 24fde17

	@skipif_need_cuda_headers # cg
	def test_cooperative_launch():
	dev = Device()
	dev.set_current()
	s = dev.create_stream(options={"nonblocking": True})
	# CUDA kernel templated on type T
	code = r"""
	#include <cooperative_groups.h>
	extern "C" __global__ void test_grid_sync() {
	namespace cg = cooperative_groups;
	auto grid = cg::this_grid();
	grid.sync();
	}
	"""
	# Compile and force instantiation for this type
	arch = "".join(f"{i}" for i in dev.compute_capability)
	include_path = str(pathlib.Path(os.environ["CUDA_PATH"]) / pathlib.Path("include"))
	pro_opts = ProgramOptions(std="c++17", arch=f"sm_{arch}", include_path=include_path)
	prog = Program(code, code_type="c++", options=pro_opts)
	ker = prog.compile("cubin").get_kernel("test_grid_sync")
	# # Launch without setting cooperative_launch
	# # Commented out as this seems to be a sticky error...
	# config = LaunchConfig(grid=1, block=1)
	# launch(s, config, ker)
	# from cuda.core.experimental._utils.cuda_utils import CUDAError
	# with pytest.raises(CUDAError) as e:
	# s.sync()
	# assert "CUDA_ERROR_LAUNCH_FAILED" in str(e)
	# Crazy grid sizes would not work
	block = 128
	config = LaunchConfig(grid=dev.properties.max_grid_dim_x // block + 1, block=block, cooperative_launch=True)
	with pytest.raises(ValueError):
	launch(s, config, ker)
	# This works just fine
	config = LaunchConfig(grid=1, block=1, cooperative_launch=True)
	launch(s, config, ker)
	s.sync()

[brandon-b-miller]

Without adding libcudadevrt.a to the link, the cooperative groups tests fail for me locally:

python -m numba.runtests numba.cuda.tests.cudapy.test_cooperative_groups

pynvjitlink.api.NvJitLinkError: NVJITLINK_ERROR_INTERNAL error when calling nvJitLinkComplete
error   : Undefined reference to 'cudaCGGetIntrinsicHandle' in '<cudapy-ptx>'
ERROR 9: finish

[leofang]

It seems this is needed when -rdc=true is passed. Does numba-cuda do that always? I forgot.