[NVIDIA][CUDA][THOR] Device Support is outdated->wrong due to usage of 101/101a

[nWEIdia]

NVIDIA https://www.nvidia.com/en-us/autonomous-machines/embedded-systems/jetson-thor/ is GA now.

However, #124155 seems to have introduced SM101, which later was changed to SM110 (See cuda13 release notes here: https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html#id2)

Therefore, we need a follow-up PR to rename all the 101 101a etc to their new name.

This would impact downstream Triton -> PyTorch projects.
e.g. using

pytorch cuda 13 nightly: pip3 install --pre torch torchvision --index-url https://download.pytorch.org/whl/nightly/cu130
with the following triton wheel
https://download.pytorch.org/whl/nightly/pytorch_triton-3.4.0%2Bgitf7888497-cp310-cp310-linux_aarch64.whl

the following errors would be encountered:

`python test/inductor/test_control_flow.py CondTests.test_cond_mismatched_branch_output_size_device_cuda_dynamic_False
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
'sm_110a' is not a recognized processor for this target (ignoring processor)
LLVM ERROR: Cannot select: intrinsic %llvm.nvvm.shfl.sync.bfly.i32

`

cc @ptrblck @tinglvv @eqy @Aidyn-A @atalman @malfet @davidberard98
cc @durga4github

[Artem-B]

Therefore, we need a follow-up PR to rename all the 101 101a etc to their new name.

That may be a bit more complicated than that.

sm_101 does exist in the older CUDA versions that we still need to support, so we can't just rename it and hope that the universe will just forget about existence of cuda-12.x.

GPU renaming also throws a wrench into how we constrain availability of the instructions.

If some instruction is available on sm_101/sm_110, is it available on sm_103? The answer would depend on whether we're using CUDA-12 or cuda-13 (or rather ptx 9.0 or 8.x). At the very least we'll need to change all those constraints to account for these disjoint availability ranges.

Similar constraint determination issue will affect whatever sm_101/sm_110-specific builtins we may introduce in clang. Unfortunately for us, target builtin constraints are rather inflexible and we may not be able to make them work in all cases.

Fortunately we do not seem to have any of those at the moment, so we may be able to get away with it by introducing sm_100 while keeping sm_101 around. We may need to add extra checks to error out if sm_101 is used with newer PTX along with the similar checks on clang side.

cc: @AlexMaclean

[nWEIdia]

sm_101 does exist in the older CUDA versions that we still need to support

How do you qualify your support? Is making sure the generated code being runnable on THOR device part of the metrics?

[Artem-B]

How do you qualify your support? Is making sure the generated code being runnable on THOR device part of the metrics?

That's a rather open-ended question.

Are you asking about the policy on removing support for something in LLVM? In general the timeline spans bout two major releases -- announce/document deprecation in one release, remove it in the next one.

Or do you have something else in mind? E.g. how we determine which GPUs we do support? The answer is generally "whatever GPUs are supported by CUDA SDK", and until now it was always a continuous set, where the GPUs were removed at the edges of the set only. Renaming sm_101 to sm_110 makes sm_101 conditionally disappear, depending on the PTX version we use, and we have no good mechanism to handle it. Our best bet is to continue to accept it as the input, and convert it, as necessary to sm_110 if we're targeting recent enough PTX. It's all probably doable, I'm just pointing that it's unlikely to be a 2-line "just rename it" patch.

Plan B is to pretend that sm_101 never existed (and happens to work now by accident), and just tell existing users "too bad, use sm_110, PTX-9.0 and cuda-13". I have no clue how many existing users of sm_101/sm_110 GPU are out there in the wild. If there are none, this may be an option, but lacking visibility into that, the default is to assume that they do exist and that removing the already released sm_101 support abruptly may be a showstopper for them. Not everyone can "just switch" to cuda-13.

Is making sure the generated code being runnable on THOR device part of the metrics?

That would be a question for NVIDIA, I guess. For LLVM all GPU variants are about the same. Some end users care more about particular GPU variants, but that's about it.

Consider outcomes of the following scenarios:

• assume that someone uses just released LLVM 21, and does wand to JIT-compile something for sm_101.
• they currently build their app with cuda-12.9, which reports the GPU as having CC 10.1, they map it to sm_101 target when they JIT their code, they pass it to ptxas from cuda_12.9 which accepts it. So far so good.

Now imagine that they want to update LLVM past the patch which renames sm_101->sm_110

• cuda runtime still reports CC 10.1, existing code still maps to sm_101, byt LLVM now complains that it does not now of such GPU.
• now the user has to add a special case, saying that if CUDA runtime reports sm_101, they need to tell LLVM to target sm_110.
• now LLVM can compile stuff to PTX, but then ptxas complains because in cuda-12 it has no idea about sm_110.
• Ugh. The user now has to intercept generated PTX, and replace sm_110-> back to sm_101, before feeding it to ptxas.
• Fine things kind of work now. Problem solved.

Except the next thing to happen is that they do switch to cuda-13, and all of a sudden things break again.

• now cuda runtime reports CC 11.0, user has to add the mapping to sm_110, but it's a normal cosmetic change.
• LLVM still happily compiles the PTX. So far so good.
• But now ptxas fails because we've patched the sm name back to sm_101. So we have to change the code, again, and apply patching only to the builds that use cuda-12.
• Problem fixed. Again.

It's all fixable, but the problem is that it has to be fixable by each and every end user, and that's a very bad choice logistically.

I wish NVIDIA would just let sm_101 be and avoided this mess.
However, given that the sm_101 disappearance is also public now, we do have to deal with it.

The best thing to do here is to see if we can insulate the end users from this without adding too much mess to LLVM.
Ideally I want the two scenarios above to continue working without forcing users to add conditional special cases for certain combinations of PTX versions and GPU variants.

[nWEIdia]

do you have something else in mind?

Yes, I am thinking if the user encounters failure due to LLVM renaming 101 -> 110, they can blame NVIDIA
But if a user encounters failure due to LLVM not changing 101 -> 110, they will blame LLVM.

[nWEIdia]

it's unlikely to be a 2-line "just rename it" patch.

Partially agree. It is not 2-line.
"Just rename it" is the PR that we are planning to send for review.

[Aidyn-A]

Wouldn't it be easier to keep both sm_101 and sm_110 at the same time, but gate keep the version based on PTX version? Basically if PTX version is 8.x -- use sm_101, and if PTX version 9.x use sm_110?

[Artem-B]

Wouldn't it be easier to keep both sm_101 and sm_110 at the same time, but gate keep the version based on PTX version?

We will have to keep sm_101 alive, because LLVM does support PTX 8.6-8.9 and therefore must accept sm_101 if those PTX versions are targeted. We'll need similar approach in clang driver -- keep sm_101 around, but mark it as unsupported by cuda-13+