Make RTC compatible with CUDA enhanced compatibility

[ptrendx]

Description

Starting with CUDA 11.1 it is possible to run programs compiled with newer CUDA toolkit with older driver (as long as the major version is the same, e.g. CUDA 11.1 works with CUDA 11.0 driver) without the compat library. This requires a few changes to API used by nvRTC however, which are addressed by this PR.

Checklist

Essentials

• Changes are complete (i.e. I finished coding on this PR)
• Code is well-documented

Comments

• Change was tested in internal CI using CUDA 11.1 toolkit and Titan RTX with 450.80.02 driver.

[mxnet-bot]

Hey @ptrendx , Thanks for submitting the PR
All tests are already queued to run once. If tests fail, you can trigger one or more tests again with the following commands:

• To trigger all jobs: @mxnet-bot run ci [all]
• To trigger specific jobs: @mxnet-bot run ci [job1, job2]

---

CI supported jobs: [sanity, windows-cpu, centos-gpu, miscellaneous, website, unix-gpu, edge, windows-gpu, centos-cpu, clang, unix-cpu]

---

Note:
Only following 3 categories can trigger CI :PR Author, MXNet Committer, Jenkins Admin.
All CI tests must pass before the PR can be merged.

[ptrendx]

Note: this does not touch the legacy RTC part (https://mxnet.apache.org/versions/1.6/api/python/docs/api/mxnet/rtc/index.html) - what is the plan for it @szha?

[szha]

I think we need to continue to support mx.rtc

[ptrendx]

Are there any people using it? The interface is not great, since the CudaModule from there is not even an operator so can't be used in a model. We could reuse the recent RTC stuff to make it much better experience (and actually potentially pretty useful).

That said, this PR does not touch that functionality (because the compilation options there are set by the user). I could make it so if you specify the proper option (--gpu-architecture=sm_XX insteaf of compute_XX yourself) it gets the cubin instead of ptx so it works with enhanced compatibility.

[szha]

Agreed. In 2.0 we can change the interface.

[ptrendx]

@mxnet-bot run ci [centos-cpu, centos-gpu, edge, miscellaneous]

[mxnet-bot]

Jenkins CI successfully triggered : [edge, centos-cpu, miscellaneous, centos-gpu]

[ptrendx]

@mxnet-bot run ci [centos-cpu, unix-gpu, edge, website]

[mxnet-bot]

Jenkins CI successfully triggered : [edge, website, unix-gpu, centos-cpu]

[DickJC123]

For others interested in understanding better the motivation behind this PR, I suggest https://docs.nvidia.com/deploy/pdf/CUDA_Compatibility.pdf . One paragraph worth repeating from that doc is:

To use other CUDA APIs introduced in a minor release (that require a new
driver), one would have to implement fallbacks or fail gracefully. This situation
is not different from what is available today where developers use macros to
compile out features based on CUDA versions. Users should refer to the CUDA
headers and documentation for new CUDA APIs introduced in a release.

Thus, it's fair to use an 11.1 feature that is supported by both 11.1 and 11.0 kernel-mode drivers. Before using an 11.1 feature that requires an 11.1 kernel-mode driver, one should check dynamically for that feature's presence at runtime, as suggested in the document section 3.2 "Handling New CUDA Features." This is particularly important to pay attention to while the upstream CI testing has no enhanced-compatibility build.

[DickJC123]

FWIW, while nvrtcGetCUBINSize() and nvrtcGetCUBIN() are not yet in the nvrtc docs, their use is described in https://docs.nvidia.com/deploy/cuda-compatibility/

[ptrendx]

@mxnet-bot run ci [unix-cpu]

[mxnet-bot]

Jenkins CI successfully triggered : [unix-cpu]