add nvptx_target_feature

[jedbrown]

Tracking issue: #141468 (nvptx), which is part of #44839 (catch-all arches)
The feature gate is #![feature(nvptx_target_feature)]

This exposes the target features sm_20 through sm_120a as defined by LLVM.

Cc: @gonzalobg
@rustbot label +O-NVPTX +A-target-feature

[rustbot]

r? @wesleywiser

rustbot has assigned @wesleywiser.
They will have a look at your PR within the next two weeks and either review your PR or reassign to another reviewer.

Use r? to explicitly pick a reviewer

[jedbrown]

Do I need to create a dedicated tracking issue or can I reuse #44839?

I have an associated patch for rust-lang/stdarch that provides a fix for #137086. Can I submit that while this PR is still in review?

[jedbrown]

It looks like the failing build environment is using LLVM-18 (sm_100 onward were added recently llvm/llvm-project#124155). I could remove those entirely (such hardware isn't available yet) or guard them based on LLVM version (how?).

[gonzalobg]

This exposes the target features sm_20 through sm_120a as defined by LLVM.

IIUC LLVM exposes these as target CPUs (Proc) above, not target features, but I think it is fine to expose these as target features in Rust.

[gonzalobg]

I think it'd be good to have more testing (could be punted into an issue and done later):

1. A basic fail test that tries to compile for sm_xx but uses an op that is only available in sm_yy.
   • A pass version of this test that guards the op correctly using cfg!
2. A test that compiles for sm_yy, uses an op that is supported for sm_xx but not in sm_yy even though yy > xx, and therefore fails to compile. That is guarding the op on sm_yy does not suffice, it needs to be precisely guarded on sm_xx.
3. A test that compiles for sm_xx, uses an op that is supported for sm_xx, but the PTX ISA version used does not support this op.
   • I think the proposed solution is not sufficient to guard these ops. The proposed scheme may need to be extended, e.g., to be able to also guard on PTX ISA versions and treating "ptx82" as a target feature), or maybe we need a different solution (exposing target-features describing functionality and do the SM+PTX math in rustc; or expose ptx ISA version and then build in stdarch, etc.).

I think it'd also be good to have more docs:

• The nvptx target docs should describe these:

  rust/src/doc/rustc/src/platform-support/nvptx64-nvidia-cuda.md

  Lines 25 to 26 in 1370611

  	Are there notable `#[target_feature(...)]` or `-C target-feature=` values that
  	programs may wish to use?

• The stdarch docs mention "CPU" throughout, but then list the nvptx target, I think "CPU" should be removed: https://github.com/rust-lang/stdarch/blob/9426bb56586c6ae4095a2dcbd66c570253e6fb32/crates/core_arch/src/core_arch_docs.md

[taiki-e]

@jedbrown

It looks like the failing build environment is using LLVM-18 (sm_100 onward were added recently llvm/llvm-project#124155). I could remove those entirely (such hardware isn't available yet) or guard them based on LLVM version (how?).

You can address this issue by modifying to_llvm_features function in rustc_codegen_llvm/src/llvm_util.rs like some riscv features do:

rust/compiler/rustc_codegen_llvm/src/llvm_util.rs

Lines 280 to 285 in c4b38a5

	// Filter out features that are not supported by the current LLVM version
	("riscv32" | "riscv64", "zaamo") if get_version().0 < 19 => None,
	("riscv32" | "riscv64", "zabha") if get_version().0 < 19 => None,
	("riscv32" | "riscv64", "zalrsc") if get_version().0 < 19 => None,
	("riscv32" | "riscv64", "zama16b") if get_version().0 < 19 => None,
	("riscv32" | "riscv64", "zacas") if get_version().0 < 20 => None,

---

@gonzalobg

This exposes the target features sm_20 through sm_120a as defined by LLVM.

IIUC LLVM exposes these as target CPUs (Proc) above, not target features, but I think it is fine to expose these as target features in Rust.

LLVM actually supports sm_* on both the features and CPUs. (However, their behavior differs as from the sm_* CPUs both sm_* and ptx* features are enabled, but from the sm_* features only sm_* features are enabled.)

Outputs of both --print target-cpus and --print target-features also show them

$ rustc --target nvptx64-nvidia-cuda --print target-cpus
Available CPUs for this target:
    sm_100
    sm_100a
    sm_101
    sm_101a
    sm_120
    sm_120a
    sm_20
    sm_21
    sm_30   - This is the default target CPU for the current build target (currently nvptx64-nvidia-cuda).
    sm_32
    sm_35
    sm_37
    sm_50
    sm_52
    sm_53
    sm_60
    sm_61
    sm_62
    sm_70
    sm_72
    sm_75
    sm_80
    sm_86
    sm_87
    sm_89
    sm_90
    sm_90a

$ rustc --target nvptx64-nvidia-cuda --print target-features
Features supported by rustc for this target:
    crt-static - Enables C Run-time Libraries to be statically linked.

Code-generation features supported by LLVM for this target:
    ptx32      - Use PTX version 32.
    ptx40      - Use PTX version 40.
    ptx41      - Use PTX version 41.
    ptx42      - Use PTX version 42.
    ptx43      - Use PTX version 43.
    ptx50      - Use PTX version 50.
    ptx60      - Use PTX version 60.
    ptx61      - Use PTX version 61.
    ptx62      - Use PTX version 62.
    ptx63      - Use PTX version 63.
    ptx64      - Use PTX version 64.
    ptx65      - Use PTX version 65.
    ptx70      - Use PTX version 70.
    ptx71      - Use PTX version 71.
    ptx72      - Use PTX version 72.
    ptx73      - Use PTX version 73.
    ptx74      - Use PTX version 74.
    ptx75      - Use PTX version 75.
    ptx76      - Use PTX version 76.
    ptx77      - Use PTX version 77.
    ptx78      - Use PTX version 78.
    ptx80      - Use PTX version 80.
    ptx81      - Use PTX version 81.
    ptx82      - Use PTX version 82.
    ptx83      - Use PTX version 83.
    ptx84      - Use PTX version 84.
    ptx85      - Use PTX version 85.
    ptx86      - Use PTX version 86.
    ptx87      - Use PTX version 87.
    sm_100     - Target SM 100.
    sm_100a    - Target SM 100a.
    sm_101     - Target SM 101.
    sm_101a    - Target SM 101a.
    sm_120     - Target SM 120.
    sm_120a    - Target SM 120a.
    sm_20      - Target SM 20.
    sm_21      - Target SM 21.
    sm_30      - Target SM 30.
    sm_32      - Target SM 32.
    sm_35      - Target SM 35.
    sm_37      - Target SM 37.
    sm_50      - Target SM 50.
    sm_52      - Target SM 52.
    sm_53      - Target SM 53.
    sm_60      - Target SM 60.
    sm_61      - Target SM 61.
    sm_62      - Target SM 62.
    sm_70      - Target SM 70.
    sm_72      - Target SM 72.
    sm_75      - Target SM 75.
    sm_80      - Target SM 80.
    sm_86      - Target SM 86.
    sm_87      - Target SM 87.
    sm_89      - Target SM 89.
    sm_90      - Target SM 90.
    sm_90a     - Target SM 90a.

Use +feature to enable a feature, or -feature to disable it.
For example, rustc -C target-cpu=mycpu -C target-feature=+feature1,-feature2

Code-generation features cannot be used in cfg or #[target_feature],
and may be renamed or removed in a future version of LLVM or rustc.

[taiki-e]

Btw, is it intentional that the ptx* target features are not added? My understanding is that we need to check both sm_* and ptx* to use some of the functionality, like LLVM does here.

[jedbrown]

Thanks @taiki-e. I'll update accordingly. I'm seeing PTX 7.8 being written in an otherwise-default configuration with target sm_89. Do you know where such defaults are set?

[jedbrown]

I updated to add the ptx* features and to fix dependencies within the sm_* chain (and deal with the LLVM 20 requirement).

One remaining issue (for the tests @gonzalobg requested) is that #[target_feature(enable = "sm_70")] does not prevent compilation when that target feature isn't present. You can still call the intrinsic unsafely. With respect to PTX requirements, this leads to confusing linker errors. Is there a proper way to enforce this?

[bors]

☔ The latest upstream changes (presumably #139229) made this pull request unmergeable. Please resolve the merge conflicts.

[rustbot]

Some changes occurred in src/doc/rustc/src/platform-support

cc @Noratrieb

[jedbrown]

@wesleywiser 👋 Checks have passed and this is ready for your review.

[bors]

☔ The latest upstream changes (presumably #140887) made this pull request unmergeable. Please resolve the merge conflicts.

[jedbrown]

I rebased and have questions about implied LLVM target features for clang and rustc.

$ echo '__device__ void f(){}' | ~/src/rust-enzyme/build/host/llvm/bin/clang++ --cuda-gpu-arch=sm_89 --cuda-device-only -S -emit-llvm -x cu - -o - | rg target-features
clang++: warning: CUDA version 12.8 is only partially supported [-Wunknown-cuda-version]
attributes #0 = { convergent mustprogress noinline nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="sm_89" "target-features"="+ptx87,+sm_89" }

$ rustc +stage1 --print cfg --target=nvptx64-nvidia-cuda -Ctarget-cpu=sm_89 | rg target_feature
target_feature="ptx78"
target_feature="sm_89"

1. clang selects the most recent PTX ISA known to both clang and the found CUDA toolkit, regardless of the specified arch.
2. rustc selects the oldest PTX ISA that supports the specified arch (sm_89) in this case.
3. Neither activate the implied features, which is different from CPU arches.

Is this all as intended? I think choosing the oldest ISA and activating the implied features would be preferable, but that seems to be an LLVM choice.

[gonzalobg]

Neither activate the implied features

Which ones are those?

[jedbrown]

I talked offline with Gonzalo, but the issue is that, for example, -Ctarget-feature=+avx512f enables all the implicated features (like avx2), but -Ctarget-feature=+sm_89 does not activate any of the implied features. Intrinsics like _barrier_sync (unaligned) require at least ptx60 and I'd like to be able to write #[cfg(target_feature = "ptx60")] (the ISA is a compile-time property so we can't merely enable the target feature).

In the current model, either the user would need to manually enable -Ctarget-feature=+ptx78,+ptx77,...,+ptx60 or intrinsics would need to use #[cfg(any(target_feature = "ptx60", target_feature = "ptx61", ..., target_feature = "ptx78"))]. It seems to me that activating all the implied features would be more analogous to how features are handled on CPUs, and avoid the need for unsustainable things like this.

An alternative would be to teach cfg to handle >=, as in #[cfg(target_feature >= "ptx60")]. If that's preferred, where would I start to implement that?

[kulst]

I was able to find the logic Clang uses to lift the ptx* feature to fit the local CUDA version. So this is plausible.

LLVM defines the logic for NVPTX target features here.
sm_* and ptx* are defined by using FeatureSM and FeaturePTX classes. These classes use the SubtargetFeature class which is defined here. For implied features the fifth argument of the SubtargetFeature is used. However, FeatureSM and FeaturePTX leave this argument empty.
avx512f on the other hand mentions three implied features where avx2 is one of them.

In the long run, I would prefer ptx* implying all previous ptx* features (same for sm*), provided that this can be changed in LLVM without any problems.

[kulst]

I talked offline with Gonzalo, but the issue is that, for example, -Ctarget-feature=+avx512f enables all the implicated features (like avx2), but -Ctarget-feature=+sm_89 does not activate any of the implied features. Intrinsics like _barrier_sync (unaligned) require at least ptx60 and I'd like to be able to write #[cfg(target_feature = "ptx60")] (the ISA is a compile-time property so we can't merely enable the target feature).

I was wondering why this does not already work using your branch.
I think if you explicitly set a target feature with -Ctarget-feature the target features implied by rustc should be added to the rustc session cfg. However, if you only implicitely set target features by specifying -Ctarget-cpu the implied rustc features are not added. Can you confirm this?

The key function is probably target_config which is used to set the session's config:

As far as I can tell it does the following:

1. First it creates an informational target machine without handing -Ctarget-feature to LLVM. The target machine's features therefore contains the features implied by LLVM but not those implied by rustc.
2. It then takes all possible target features which are specified in rustc and keeps only features that exist in the target machine.
3. This list is extended by all activated features which are specified in -Ctarget-features as well as their features implied by rustc.

--> The third step adds target features implied by rustc but only for those target features that are specified with -Ctarget-feature. I think we could add a step 2.5 to add target features implied by rustc to solve this. Therefore no LLVM change would be necessary.

[jedbrown]

I think if you explicitly set a target feature with -Ctarget-feature the target features implied by rustc should be added to the rustc session cfg. However, if you only implicitely set target features by specifying -Ctarget-cpu the implied rustc features are not added. Can you confirm this?

Yes

$ rustc +stage1 --print cfg --target=nvptx64-nvidia-cuda -Ctarget-cpu=sm_89 | rg target_feature
target_feature="ptx78"
target_feature="sm_89"
$ rustc +stage1 --print cfg --target=nvptx64-nvidia-cuda -Ctarget-cpu=sm_89 -Ctarget-feature=+sm_80,+ptx78 | rg target_feature
target_feature="ptx32"
target_feature="ptx40"
target_feature="ptx41"
target_feature="ptx42"
target_feature="ptx43"
target_feature="ptx50"
target_feature="ptx60"
target_feature="ptx61"
target_feature="ptx62"
target_feature="ptx63"
target_feature="ptx64"
target_feature="ptx65"
target_feature="ptx70"
target_feature="ptx71"
target_feature="ptx72"
target_feature="ptx73"
target_feature="ptx74"
target_feature="ptx75"
target_feature="ptx76"
target_feature="ptx77"
target_feature="ptx78"
target_feature="sm_20"
target_feature="sm_21"
target_feature="sm_30"
target_feature="sm_32"
target_feature="sm_35"
target_feature="sm_37"
target_feature="sm_50"
target_feature="sm_52"
target_feature="sm_53"
target_feature="sm_60"
target_feature="sm_61"
target_feature="sm_62"
target_feature="sm_70"
target_feature="sm_72"
target_feature="sm_75"
target_feature="sm_80"
target_feature="sm_89"

This works with #[cfg(target_feature = "ptx70")] in user crates, but seemingly not in stdarch when building using -Zbuild-std=core (an issue I'm still trying to understand as it doesn't seem related to activation of implicit features). And thanks for finding where this relevant logic lives. Based on success running code with -Ctarget-feature=+ptx86,+sm_89, I think eagerly activating those implied features won't cause a problem.

[jedbrown]

I implemented the suggested logic to activate implied features, and a test.

[gonzalobg]

This looks good to me now.
There is more work to be done to continue improving this in the future, but this is all still very unstable.
@jedbrown can you open a tracking issue just for this nvptx_target_feature to start capturing that?

@wesleywiser I can't approve from this github account.
It may make sense to assign a second pair of eyes to take a look.

[jedbrown]

Thanks. I made the tracking issue #141468 and have updated the documentation to warn against use of #[target_feature(enable = "ptx80")], as well as clarifying how to build. Any further discussion of intrinsics in the documentation can happen after rust-lang/stdarch#1805.

[wesleywiser]

cc target maintainers @RDambrosio016 and @kjetilkjeka for review as well

[bors]

☔ The latest upstream changes (presumably #141984) made this pull request unmergeable. Please resolve the merge conflicts.

[kjetilkjeka]

Sorry for taking so long before getting to this.

The changes look good, useful and is definitely a step in the right direction 👍

A small annoyance is that actually using target-feature=+ptx80 together with the implied futures leads to redundant warnings (one per implied feature). I'm not sure it's easy to do something with that as both the actual warning and the implied features are useful on their own and it will anyway go away when they are made stable so I don't think it necessarily makes sense to insist on tackling this here and now.

---

This is how the warning looked before:

And this is what happens with the implied features

[jedbrown]

Thanks, I added a note about the redundant warnings to the tracking issue. I think many users will be content with -Ctarget-cpu=sm_80 (i.e., without explicitly upgrading the PTX version) and thus not see the warnings.