Fixup and re-enable many tests for NVPTX

src/doc/rustc-dev-guide/src/tests/best-practices.md

@@ -195,8 +195,30 @@ See [LLVM FileCheck guide][FileCheck] for details.
 >
 > Pending concrete advice.
 
+## Target-specific tests
+
+When adding a test for target-specific behavior, prefer writing the test so that
+it is run by an existing CI test-suite invocation.
+
+In particular, for targets that are not used as hosts in CI, prefer
+cross-building the test to that target instead of making the whole test
+`only-X` for that target. For example, a UI test can use a target-specific
+revision with `compile-flags: --target ...`:
+
+```rust
+//@ revisions: some-target
+//@[some-target] compile-flags: --target SomeTarget
+//@[some-target] needs-llvm-components: some_backend
+```
+
+This keeps the test covered by the regular test-suite invocation while still
+checking the target-specific behavior. If the target requires a specific LLVM
+backend, add the appropriate `needs-llvm-components` directive. See the
+[compiletest directives] for more information about these directives.
+
 [compiletest]: ./compiletest.md
 [compiletest directives]: ./directives.md
 [`run-make`]: ./compiletest.md#run-make-tests
 [FileCheck]: https://llvm.org/docs/CommandGuide/FileCheck.html
 [crash tests]: ./compiletest.md#crash-tests
+

src/doc/rustc-dev-guide/src/tests/directives.md

@@ -45,9 +45,13 @@ Directives can generally be found by browsing the
 
 <!-- date-check: Oct 2024 -->
 
-| Directive         | Explanation                   | Supported test suites | Possible values                        |
-|-------------------|-------------------------------|-----------------------|----------------------------------------|
-| `assembly-output` | Assembly output kind to check | `assembly`            | `emit-asm`, `bpf-linker`, `ptx-linker` |
+| Directive         | Explanation                   | Supported test suites | Possible values                             |
+|-------------------|-------------------------------|-----------------------|---------------------------------------------|
+| `assembly-output` | Assembly output kind to check | `assembly`            | `emit-asm`, `bpf-linker`, `linker-asm`[^la] |
+
+[^la]: For targets where the linker produces the final assembly artifact, either 
+by default or when explicitly instructed to do so. This allows assembly tests to 
+inspect linked code, rather than only rustc's pre-link assembly output.
 
 ### Auxiliary builds
 
@@ -123,6 +127,23 @@ means the test won't be compiled or run.
 
 * `ignore-X` where `X` is a target detail or other criteria on which to ignore the test (see below)
 * `only-X` is like `ignore-X`, but will *only* run the test on that target or stage
+
+<div class="warning">
+
+Using an `only-X` directive can cause a test to not be run in CI at all.
+
+For example, a UI test containing `//@ only-SomeTarget` will only be run in CI if
+the UI test suite is actually run for *SomeTarget*. Otherwise, the test will be
+skipped in every CI job.
+
+The same can happen when combining multiple `only-X` directives. A test is only
+run if *all* of its `only-X` requirements match the same CI job.
+
+For target-specific tests, consider whether the test should be cross-built to
+the target instead; see the corresponding
+[best practices for target-specific tests](./best-practices.md#target-specific-tests).
+</div>
+
 * `ignore-auxiliary` is intended for files that *participate* in one or more other
   main test files but that `compiletest` should not try to build the file itself.
   Please backlink to which main test is actually using the auxiliary file.

src/tools/compiletest/src/runtest/assembly.rs

@@ -31,6 +31,7 @@ impl TestCx<'_> {
         let emit = match self.props.assembly_output.as_deref() {
             Some("emit-asm") => Emit::Asm,
             Some("bpf-linker") => Emit::LinkArgsAsm,
+            Some("linker-asm") => Emit::None,
             Some(other) => self.fatal(&format!("unknown 'assembly-output' directive: {other}")),
             None => self.fatal("missing 'assembly-output' directive"),
         };

tests/assembly-llvm/auxiliary/non-inline-dependency.rs

@@ -1,14 +1,28 @@
-#![no_std]
-#![deny(warnings)]
+//@ add-minicore
+//@ no-prefer-dynamic
+//@ compile-flags: --target nvptx64-nvidia-cuda
+//@ needs-llvm-components: nvptx
+#![feature(no_core, intrinsics)]
+#![crate_type = "rlib"]
+#![no_core]
+
+extern crate minicore;
+use minicore::*;
+
+#[rustc_intrinsic]
+pub const fn wrapping_mul<T: Copy>(a: T, b: T) -> T;
+
+#[rustc_intrinsic]
+pub const fn mul_with_overflow<T: Copy>(x: T, y: T) -> (T, bool);
 
 #[inline(never)]
 #[no_mangle]
 pub fn wrapping_external_fn(a: u32) -> u32 {
-    a.wrapping_mul(a)
+    wrapping_mul(a, a)
 }
 
 #[inline(never)]
 #[no_mangle]
-pub fn panicking_external_fn(a: u32) -> u32 {
-    a * a
+pub fn overflowing_external_fn(a: u32) -> u32 {
+    mul_with_overflow(a, a).0
 }

tests/assembly-llvm/nvptx-arch-default.rs

@@ -1,11 +1,8 @@
 //@ assembly-output: emit-asm
-//@ compile-flags: --crate-type cdylib
-//@ only-nvptx64
-
-#![no_std]
-
-//@ aux-build: breakpoint-panic-handler.rs
-extern crate breakpoint_panic_handler;
+//@ compile-flags: --target nvptx64-nvidia-cuda --crate-type cdylib
+//@ needs-llvm-components: nvptx
+#![feature(no_core)]
+#![no_core]
 
 // Verify default arch with llvm-bitcode-linker.
 // CHECK: .version 7.0

tests/assembly-llvm/nvptx-arch-emit-asm.rs

@@ -1,8 +1,8 @@
 //@ assembly-output: emit-asm
-//@ compile-flags: --crate-type rlib
-//@ only-nvptx64
-
-#![no_std]
+//@ compile-flags: --target nvptx64-nvidia-cuda --crate-type rlib
+//@ needs-llvm-components: nvptx
+#![feature(no_core)]
+#![no_core]
 
 // Verify default arch without llvm-bitcode-linker involved.
 // CHECK: .version 7.0

tests/assembly-llvm/nvptx-arch-target-cpu.rs

@@ -1,11 +1,8 @@
 //@ assembly-output: emit-asm
-//@ compile-flags: --crate-type cdylib -C target-cpu=sm_87
-//@ only-nvptx64
-
-#![no_std]
-
-//@ aux-build: breakpoint-panic-handler.rs
-extern crate breakpoint_panic_handler;
+//@ compile-flags: --target nvptx64-nvidia-cuda --crate-type cdylib -C target-cpu=sm_87
+//@ needs-llvm-components: nvptx
+#![feature(no_core)]
+#![no_core]
 
 // Verify target arch override via `target-cpu`.
 // CHECK: .target sm_87

tests/assembly-llvm/nvptx-c-abi-arg-v7.rs

@@ -1,6 +1,6 @@
 //@ assembly-output: emit-asm
-//@ compile-flags: --crate-type cdylib -C target-cpu=sm_86
-//@ only-nvptx64
+//@ compile-flags: --target nvptx64-nvidia-cuda --crate-type cdylib -C target-cpu=sm_86
+//@ needs-llvm-components: nvptx
 
 // The PTX ABI stability is tied to major versions of the PTX ISA
 // These tests assume major version 7

tests/assembly-llvm/nvptx-c-abi-ret-v7.rs

@@ -1,6 +1,6 @@
 //@ assembly-output: emit-asm
-//@ compile-flags: --crate-type cdylib -C target-cpu=sm_86
-//@ only-nvptx64
+//@ compile-flags: --target nvptx64-nvidia-cuda --crate-type cdylib -C target-cpu=sm_86
+//@ needs-llvm-components: nvptx
 
 // The PTX ABI stability is tied to major versions of the PTX ISA
 // These tests assume major version 7

tests/assembly-llvm/nvptx-debug.rs

@@ -1,19 +1,17 @@
+//@ add-minicore
 //@ assembly-output: emit-asm
-//@ compile-flags: --crate-type cdylib -C debuginfo=2
-//@ only-nvptx64
+//@ compile-flags: --target=nvptx64-nvidia-cuda --crate-type cdylib -C debuginfo=2
+//@ needs-llvm-components: nvptx
 
 // Tests related to debug symbols for nvptx
 
-#![feature(abi_ptx)]
-#![no_std]
+#![feature(no_core, abi_ptx)]
+#![no_core]
 
-//@ aux-build: breakpoint-panic-handler.rs
-extern crate breakpoint_panic_handler;
+extern crate minicore;
 
 #[no_mangle]
-pub extern "ptx-kernel" fn foo() {
-    panic!("bar");
-}
+pub extern "ptx-kernel" fn foo() {}
 
 // We make sure that all debug sections are available and visit them
 // CHECK: .section .debug_abbrev

tests/assembly-llvm/nvptx-internalizing.rs

@@ -1,15 +1,16 @@
+//@ add-minicore
 //@ assembly-output: emit-asm
-//@ compile-flags: --crate-type cdylib
-//@ only-nvptx64
+//@ compile-flags: --target nvptx64-nvidia-cuda --crate-type cdylib
+//@ needs-llvm-components: nvptx
 
-#![feature(abi_ptx)]
-#![no_std]
+#![feature(abi_ptx, no_core, intrinsics)]
+#![no_core]
 
-//@ aux-build: breakpoint-panic-handler.rs
-extern crate breakpoint_panic_handler;
+extern crate minicore;
+use minicore::*;
 
-//@ aux-build: non-inline-dependency.rs
-extern crate non_inline_dependency as dep;
+#[rustc_intrinsic]
+pub const unsafe fn unchecked_add<T: Copy>(x: T, y: T) -> T;
 
 // Verify that no extra function declarations are present.
 // CHECK-NOT: .func
@@ -18,7 +19,7 @@ extern crate non_inline_dependency as dep;
 #[no_mangle]
 pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) {
     // CHECK: add.s32 %{{r[0-9]+}}, %{{r[0-9]+}}, 5;
-    *b = *a + 5;
+    *b = unchecked_add(*a, 5);
 }
 
 // Verify that no extra function definitions are here.

tests/assembly-llvm/nvptx-kernel-abi/nvptx-kernel-args-abi-v7.rs

@@ -1,6 +1,6 @@
 //@ assembly-output: emit-asm
-//@ compile-flags: --crate-type cdylib -C target-cpu=sm_86
-//@ only-nvptx64
+//@ compile-flags: --target=nvptx64-nvidia-cuda --crate-type cdylib -C target-cpu=sm_86
+//@ needs-llvm-components: nvptx
 
 // The following ABI tests are made with nvcc 11.6 does.
 //

tests/assembly-llvm/nvptx-linking-cdylib.rs

@@ -1,39 +1,43 @@
-//@ assembly-output: emit-asm
-//@ compile-flags: --crate-type cdylib
-//@ only-nvptx64
-//@ ignore-nvptx64
+//@ add-minicore
+//@ assembly-output: linker-asm
+//@ compile-flags: --target nvptx64-nvidia-cuda --crate-type cdylib
+//@ needs-llvm-components: nvptx
 
-#![feature(abi_ptx)]
-#![no_std]
+#![feature(abi_ptx, no_core, intrinsics)]
+#![no_main]
+#![no_core]
 
-//@ aux-build: breakpoint-panic-handler.rs
-extern crate breakpoint_panic_handler;
+extern crate minicore;
+use minicore::*;
+
+#[rustc_intrinsic]
+pub const fn wrapping_add<T: Copy>(a: T, b: T) -> T;
 
 //@ aux-build: non-inline-dependency.rs
 extern crate non_inline_dependency as dep;
 
 // Make sure declarations are there.
-// CHECK: .func (.param .b32 func_retval0) wrapping_external_fn
-// CHECK: .func (.param .b32 func_retval0) panicking_external_fn
+// CHECK: .func (.param .{{[ubs]}}32 func_retval0) wrapping_external_fn
+// CHECK: .func (.param .{{[ubs]}}32 func_retval0) overflowing_external_fn
 
 // CHECK-LABEL: .visible .entry top_kernel(
 #[no_mangle]
 pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) {
     // CHECK:      call.uni (retval0),
-    // CHECK-NEXT: wrapping_external_fn
-    // CHECK:      ld.param.b32 %[[LHS:r[0-9]+]], [retval0+0];
+    // CHECK:      wrapping_external_fn
+    // CHECK:      ld.param.{{[ubs]}}32 %[[LHS:r[0-9]+]], [{{retval0(\+0)?}}];
     let lhs = dep::wrapping_external_fn(*a);
 
     // CHECK:      call.uni (retval0),
-    // CHECK-NEXT: panicking_external_fn
-    // CHECK:      ld.param.b32 %[[RHS:r[0-9]+]], [retval0+0];
-    let rhs = dep::panicking_external_fn(*a);
+    // CHECK:      overflowing_external_fn
+    // CHECK:      ld.param.{{[ubs]}}32 %[[RHS:r[0-9]+]], [{{retval0(\+0)?}}];
+    let rhs = dep::overflowing_external_fn(*a);
 
-    // CHECK: add.s32 %[[RES:r[0-9]+]], %[[RHS]], %[[LHS]];
-    // CHECK: st.global.u32 [%{{rd[0-9]+}}], %[[RES]];
-    *b = lhs + rhs;
+    // CHECK: add.{{[ubs]}}32 %[[RES:r[0-9]+]], %[[RHS]], %[[LHS]];
+    // CHECK: st.global.{{[ubs]}}32 [%{{rd[0-9]+}}], %[[RES]];
+    *b = wrapping_add(lhs, rhs);
 }
 
 // Verify that external function bodies are available.
-// CHECK: .func (.param .b32 func_retval0) wrapping_external_fn
-// CHECK: .func (.param .b32 func_retval0) panicking_external_fn
+// CHECK: .func (.param .{{[ubs]}}32 func_retval0) wrapping_external_fn
+// CHECK: .func (.param .{{[ubs]}}32 func_retval0) overflowing_external_fn

tests/assembly-llvm/nvptx-safe-naming.rs

@@ -1,14 +1,19 @@
+//@ add-minicore
 //@ assembly-output: emit-asm
-//@ compile-flags: --crate-type cdylib
-//@ only-nvptx64
+//@ compile-flags: --target=nvptx64-nvidia-cuda --crate-type cdylib
+//@ needs-llvm-components: nvptx
 
-#![feature(abi_ptx)]
-#![no_std]
+#![feature(abi_ptx, no_core, intrinsics)]
+#![no_core]
 
-//@ aux-build: breakpoint-panic-handler.rs
-extern crate breakpoint_panic_handler;
+extern crate minicore;
+use minicore::*;
 
-// Verify function name doesn't contain unacceaptable characters.
+#[rustc_intrinsic]
+pub const fn wrapping_mul<T: Copy>(a: T, b: T) -> T;
+
+// We want to make sure neither we nor LLVM introduce characters invalid for
+// NVPTX.
 // CHECK: .func (.param .b32 func_retval0) [[IMPL_FN:[a-zA-Z0-9$_]+square]]
 
 // CHECK-LABEL: .visible .entry top_kernel(
@@ -21,6 +26,7 @@ pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) {
 
 pub mod deep {
     pub mod private {
+        use crate::wrapping_mul;
         pub struct MyStruct<T>(T);
 
         impl MyStruct<u32> {
@@ -30,7 +36,7 @@ pub mod deep {
 
             #[inline(never)]
             pub fn square(&self) -> u32 {
-                self.0.wrapping_mul(self.0)
+                wrapping_mul(self.0, self.0)
             }
         }
     }