Static link CUDA Runtime

[kmaehashi]

This is a part of #7620.

• Static link CUDA Runtime in each CuPy moduels (Runtime, Jitify, cuRAND, Thrust)
• Lazy import Jitify
• Removes cupy.cuda.profiler.initialize() function, which was removed in CUDA 12 and marked deprecated since CuPy v12 (CUDA 12: Deprecate (or remove) cudaProfilerInitialize #7332). Code search yields no use-case for this function.

[kmaehashi]

/test mini

[kmaehashi]

/test mini

[kmaehashi]

@leofang
This code relies on locally-installed CUDA Toolkit version (specifically cooperative_groups.h). Wondering how we can rewrite this...

cupy/cupyx/jit/cg.py

Lines 102 to 103 in 1ef4036

	if _runtime.runtimeGetVersion() < 11060:
	raise RuntimeError("block_rank() is supported on CUDA 11.6+")

[kmaehashi]

This code relies on locally-installed CUDA Toolkit version (specifically cooperative_groups.h).

Fixed in e89e35a. Could you review this commit @asi1024?

[leofang]

Yes it should be a compile time error. LGTM.

[leofang]

FYI. I think this change is OK. For a more robust treatment for _get_arch(), I suggest to take a look at this helper function that I added recently. We don't have to do it in this PR or even in v13, since right now we almost always generate SASS which is problem free. But it is possible that the nvcc/nvrtc version is newer than the driver version, and the generated PTX would not be loadable. (I hit this when updating the PTX tests in test_raw.py, which prompted me to write that helper 🙂)

[kmaehashi]

/test mini

[takagi]

Would you have a look at cuda-python test failure?

[kmaehashi]

/test mini

[takagi]

Do we also need to include include '_runtime_extern.pxi' when we're using CUDA Python? And, cupy.linux.cuda122's failure is not related.

[kmaehashi]

Do we also need to include include '_runtime_extern.pxi' when we're using CUDA Python? And, cupy.linux.cuda122's failure is not related.

Thanks, that's right! Fixed.

/test mini

[leofang]

I can take a final look later today. Reminder to add a CI test: #7954 (comment)

[takagi]

/test mini

[leofang]

Thanks for the work, LGTM except for one question: I see the version check for JIT was reverted (7a4a4bb), is there any particular reason? I would think checking the header version is safer, instead of using libcuda.so as a proxy. But this is not a big showstopper, we can follow-up later.

See also my above reminder.

[kmaehashi]

I can take a final look later today. Reminder to add a CI test: #7954 (comment)

Thanks for the reminder! Added.

[kmaehashi]

I see the version check for JIT was reverted (7a4a4bb), is there any particular reason? I would think checking the header version is safer, instead of using libcuda.so as a proxy. But this is not a big showstopper, we can follow-up later.

The check has been rewritten using _getLocalRuntimeVersion() (e3f5873). This is because (1) I thought Python-level error checking (as originally implemented) is more user-friendly, and (2) I noticed CUDA_VERSION requires cuda.h which cannot be compiled with NVRTC.

[kmaehashi]

I can take a final look later today. Reminder to add a CI test: #7954 (comment)

Thanks for the reminder! Added.

I noticed that we can't test this PR like other modules, because cupy_backends.cuda.api.runtime uses static link and is not lazy-imported. After everything is done I'll add more sophisticated test, e.g., try installing/importing the generated CuPy binary within a vanilla Python docker image.

[leofang]

(2) I noticed CUDA_VERSION requires cuda.h which cannot be compiled with NVRTC.

This is new to me, thanks for sharing.

[leofang]

LGTM, merge?

[leofang]

I noticed that we can't test this PR like other modules, because cupy_backends.cuda.api.runtime uses static link and is not lazy-imported. After everything is done I'll add more sophisticated test, e.g., try installing/importing the generated CuPy binary within a vanilla Python docker image.

Another way to test it (but only on Linux) is to do LD_DEBUG=libs python -c "import cupy", parse the stdout, and check if libcudart.so is loaded. It's a bit hacky, though.

[takagi]

LGTM!