Skip to content

Commit 82e2f3e

Browse files
borsbors
committed
Auto merge of #58791 - denzp:asm-compile-tests, r=alexcrichton
Introduce assembly tests suite The change introduces a new test suite - **Assembly** tests. The motivation behind this is an ability to perform end-to-end codegen testing with LLVM backend. Turned out, NVPTX backend sometimes missing common Rust features (`i128` and libcalls in the past, and still full atomics support) due to different reasons. Prior to this change, basic NVPTX assembly tests were implemented within `run-make` suite. Now, it's easier to write additional and maintain existing tests for the target. cc @gnzlbg @peterhj cc @eddyb I adjusted mangling scheme expectation, so there is no need to change the tests for #57967
2 parents 9c499cc + 60f1644 commit 82e2f3e

23 files changed

+327
-140
lines changed

‎src/bootstrap/builder.rs

+1
Original file line numberDiff line numberDiff line change
@@ -374,6 +374,7 @@ impl<'a> Builder<'a> {
374374
test::MirOpt,
375375
test::Codegen,
376376
test::CodegenUnits,
377+
test::Assembly,
377378
test::Incremental,
378379
test::Debuginfo,
379380
test::UiFullDeps,

‎src/bootstrap/test.rs

+6
Original file line numberDiff line numberDiff line change
@@ -936,6 +936,12 @@ host_test!(RunMakeFullDeps {
936936
suite: "run-make-fulldeps"
937937
});
938938

939+
default_test!(Assembly {
940+
path: "src/test/assembly",
941+
mode: "assembly",
942+
suite: "assembly"
943+
});
944+
939945
#[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)]
940946
struct Compiletest {
941947
compiler: Compiler,

‎src/ci/docker/test-various/Dockerfile

+2-1
Original file line numberDiff line numberDiff line change
@@ -53,7 +53,8 @@ ENV WASM_SCRIPT python2.7 /checkout/x.py test --target $WASM_TARGETS \
5353

5454
ENV NVPTX_TARGETS=nvptx64-nvidia-cuda
5555
ENV NVPTX_SCRIPT python2.7 /checkout/x.py test --target $NVPTX_TARGETS \
56-
src/test/run-make
56+
src/test/run-make \
57+
src/test/assembly
5758

5859
ENV MUSL_TARGETS=x86_64-unknown-linux-musl \
5960
CC_x86_64_unknown_linux_musl=x86_64-linux-musl-gcc \

‎src/test/assembly/auxiliary/breakpoint-panic-handler.rs

+8
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,8 @@
1+
#![feature(core_intrinsics)]
2+
#![no_std]
3+
4+
#[panic_handler]
5+
unsafe fn breakpoint_panic_handler(_: &::core::panic::PanicInfo) -> ! {
6+
core::intrinsics::breakpoint();
7+
core::hint::unreachable_unchecked();
8+
}

‎src/test/run-make/nvptx-dylib-crate/dep.rs renamed to ‎src/test/assembly/auxiliary/non-inline-dependency.rs

File renamed without changes.

‎src/test/assembly/nvptx-arch-default.rs

+12
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,12 @@
1+
// assembly-output: ptx-linker
2+
// compile-flags: --crate-type cdylib
3+
// only-nvptx64
4+
5+
#![no_std]
6+
7+
// aux-build: breakpoint-panic-handler.rs
8+
extern crate breakpoint_panic_handler;
9+
10+
// Verify default target arch with ptx-linker.
11+
// CHECK: .target sm_30
12+
// CHECK: .address_size 64

‎src/test/assembly/nvptx-arch-emit-asm.rs

+9
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,9 @@
1+
// assembly-output: emit-asm
2+
// compile-flags: --crate-type rlib
3+
// only-nvptx64
4+
5+
#![no_std]
6+
7+
// Verify default arch without ptx-linker involved.
8+
// CHECK: .target sm_30
9+
// CHECK: .address_size 64

‎src/test/assembly/nvptx-arch-link-arg.rs

+12
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,12 @@
1+
// assembly-output: ptx-linker
2+
// compile-flags: --crate-type cdylib -C link-arg=--arch=sm_60
3+
// only-nvptx64
4+
5+
#![no_std]
6+
7+
// aux-build: breakpoint-panic-handler.rs
8+
extern crate breakpoint_panic_handler;
9+
10+
// Verify target arch override via `link-arg`.
11+
// CHECK: .target sm_60
12+
// CHECK: .address_size 64

‎src/test/assembly/nvptx-arch-target-cpu.rs

+12
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,12 @@
1+
// assembly-output: ptx-linker
2+
// compile-flags: --crate-type cdylib -C target-cpu=sm_50
3+
// only-nvptx64
4+
5+
#![no_std]
6+
7+
// aux-build: breakpoint-panic-handler.rs
8+
extern crate breakpoint_panic_handler;
9+
10+
// Verify target arch override via `target-cpu`.
11+
// CHECK: .target sm_50
12+
// CHECK: .address_size 64

‎src/test/assembly/nvptx-atomics.rs

+85
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,85 @@
1+
// assembly-output: ptx-linker
2+
// compile-flags: --crate-type cdylib
3+
// only-nvptx64
4+
5+
#![feature(abi_ptx, core_intrinsics)]
6+
#![no_std]
7+
8+
use core::intrinsics::*;
9+
10+
// aux-build: breakpoint-panic-handler.rs
11+
extern crate breakpoint_panic_handler;
12+
13+
// Currently, LLVM NVPTX backend can only emit atomic instructions with
14+
// `relaxed` (PTX default) ordering. But it's also useful to make sure
15+
// the backend won't fail with other orders. Apparently, the backend
16+
// doesn't support fences as well. As a workaround `llvm.nvvm.membar.*`
17+
// could work, and perhaps on the long run, all the atomic operations
18+
// should rather be provided by `core::arch::nvptx`.
19+
20+
// Also, PTX ISA doesn't have atomic `load`, `store` and `nand`.
21+
22+
// FIXME(denzp): add tests for `core::sync::atomic::*`.
23+
24+
#[no_mangle]
25+
pub unsafe extern "ptx-kernel" fn atomics_kernel(a: *mut u32) {
26+
// CHECK: atom.global.and.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
27+
// CHECK: atom.global.and.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
28+
atomic_and(a, 1);
29+
atomic_and_relaxed(a, 1);
30+
31+
// CHECK: atom.global.cas.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1, 2;
32+
// CHECK: atom.global.cas.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1, 2;
33+
atomic_cxchg(a, 1, 2);
34+
atomic_cxchg_relaxed(a, 1, 2);
35+
36+
// CHECK: atom.global.max.s32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
37+
// CHECK: atom.global.max.s32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
38+
atomic_max(a, 1);
39+
atomic_max_relaxed(a, 1);
40+
41+
// CHECK: atom.global.min.s32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
42+
// CHECK: atom.global.min.s32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
43+
atomic_min(a, 1);
44+
atomic_min_relaxed(a, 1);
45+
46+
// CHECK: atom.global.or.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
47+
// CHECK: atom.global.or.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
48+
atomic_or(a, 1);
49+
atomic_or_relaxed(a, 1);
50+
51+
// CHECK: atom.global.max.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
52+
// CHECK: atom.global.max.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
53+
atomic_umax(a, 1);
54+
atomic_umax_relaxed(a, 1);
55+
56+
// CHECK: atom.global.min.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
57+
// CHECK: atom.global.min.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
58+
atomic_umin(a, 1);
59+
atomic_umin_relaxed(a, 1);
60+
61+
// CHECK: atom.global.add.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
62+
// CHECK: atom.global.add.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
63+
atomic_xadd(a, 1);
64+
atomic_xadd_relaxed(a, 1);
65+
66+
// CHECK: atom.global.exch.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
67+
// CHECK: atom.global.exch.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
68+
atomic_xchg(a, 1);
69+
atomic_xchg_relaxed(a, 1);
70+
71+
// CHECK: atom.global.xor.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
72+
// CHECK: atom.global.xor.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1;
73+
atomic_xor(a, 1);
74+
atomic_xor_relaxed(a, 1);
75+
76+
// CHECK: mov.u32 %[[sub_0_arg:r[0-9]+]], 100;
77+
// CHECK: neg.s32 temp, %[[sub_0_arg]];
78+
// CHECK: atom.global.add.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], temp;
79+
atomic_xsub(a, 100);
80+
81+
// CHECK: mov.u32 %[[sub_1_arg:r[0-9]+]], 200;
82+
// CHECK: neg.s32 temp, %[[sub_1_arg]];
83+
// CHECK: atom.global.add.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], temp;
84+
atomic_xsub_relaxed(a, 200);
85+
}

‎src/test/assembly/nvptx-internalizing.rs

+27
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
// assembly-output: ptx-linker
2+
// compile-flags: --crate-type cdylib
3+
// only-nvptx64
4+
5+
#![feature(abi_ptx)]
6+
#![no_std]
7+
8+
// aux-build: breakpoint-panic-handler.rs
9+
extern crate breakpoint_panic_handler;
10+
11+
// aux-build: non-inline-dependency.rs
12+
extern crate non_inline_dependency as dep;
13+
14+
// Verify that no extra function declarations are present.
15+
// CHECK-NOT: .func
16+
17+
// CHECK: .visible .entry top_kernel(
18+
#[no_mangle]
19+
pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) {
20+
// CHECK: add.s32 %{{r[0-9]+}}, %{{r[0-9]+}}, 5;
21+
*b = *a + 5;
22+
}
23+
24+
// Verify that no extra function definitions are here.
25+
// CHECK-NOT: .func
26+
// CHECK-NOT: .entry
27+

‎src/test/assembly/nvptx-linking-binary.rs

+39
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,39 @@
1+
// assembly-output: ptx-linker
2+
// compile-flags: --crate-type bin
3+
// only-nvptx64
4+
5+
#![feature(abi_ptx)]
6+
#![no_main]
7+
#![no_std]
8+
9+
// aux-build: breakpoint-panic-handler.rs
10+
extern crate breakpoint_panic_handler;
11+
12+
// aux-build: non-inline-dependency.rs
13+
extern crate non_inline_dependency as dep;
14+
15+
// Make sure declarations are there.
16+
// CHECK: .func (.param .b32 func_retval0) wrapping_external_fn
17+
// CHECK: .func (.param .b32 func_retval0) panicking_external_fn
18+
19+
// CHECK-LABEL: .visible .entry top_kernel(
20+
#[no_mangle]
21+
pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) {
22+
// CHECK: call.uni (retval0),
23+
// CHECK-NEXT: wrapping_external_fn
24+
// CHECK: ld.param.b32 %[[LHS:r[0-9]+]], [retval0+0];
25+
let lhs = dep::wrapping_external_fn(*a);
26+
27+
// CHECK: call.uni (retval0),
28+
// CHECK-NEXT: panicking_external_fn
29+
// CHECK: ld.param.b32 %[[RHS:r[0-9]+]], [retval0+0];
30+
let rhs = dep::panicking_external_fn(*a);
31+
32+
// CHECK: add.s32 %[[RES:r[0-9]+]], %[[RHS]], %[[LHS]];
33+
// CHECK: st.global.u32 [%{{rd[0-9]+}}], %[[RES]];
34+
*b = lhs + rhs;
35+
}
36+
37+
// Verify that external function bodies are available.
38+
// CHECK: .func (.param .b32 func_retval0) wrapping_external_fn
39+
// CHECK: .func (.param .b32 func_retval0) panicking_external_fn

‎src/test/assembly/nvptx-linking-cdylib.rs

+38
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
// assembly-output: ptx-linker
2+
// compile-flags: --crate-type cdylib
3+
// only-nvptx64
4+
5+
#![feature(abi_ptx)]
6+
#![no_std]
7+
8+
// aux-build: breakpoint-panic-handler.rs
9+
extern crate breakpoint_panic_handler;
10+
11+
// aux-build: non-inline-dependency.rs
12+
extern crate non_inline_dependency as dep;
13+
14+
// Make sure declarations are there.
15+
// CHECK: .func (.param .b32 func_retval0) wrapping_external_fn
16+
// CHECK: .func (.param .b32 func_retval0) panicking_external_fn
17+
18+
// CHECK-LABEL: .visible .entry top_kernel(
19+
#[no_mangle]
20+
pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) {
21+
// CHECK: call.uni (retval0),
22+
// CHECK-NEXT: wrapping_external_fn
23+
// CHECK: ld.param.b32 %[[LHS:r[0-9]+]], [retval0+0];
24+
let lhs = dep::wrapping_external_fn(*a);
25+
26+
// CHECK: call.uni (retval0),
27+
// CHECK-NEXT: panicking_external_fn
28+
// CHECK: ld.param.b32 %[[RHS:r[0-9]+]], [retval0+0];
29+
let rhs = dep::panicking_external_fn(*a);
30+
31+
// CHECK: add.s32 %[[RES:r[0-9]+]], %[[RHS]], %[[LHS]];
32+
// CHECK: st.global.u32 [%{{rd[0-9]+}}], %[[RES]];
33+
*b = lhs + rhs;
34+
}
35+
36+
// Verify that external function bodies are available.
37+
// CHECK: .func (.param .b32 func_retval0) wrapping_external_fn
38+
// CHECK: .func (.param .b32 func_retval0) panicking_external_fn

‎src/test/run-make/nvptx-emit-asm/kernel.rs renamed to ‎src/test/assembly/nvptx-safe-naming.rs

+8-12
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,15 @@
1-
#![no_std]
2-
#![deny(warnings)]
1+
// assembly-output: ptx-linker
2+
// compile-flags: --crate-type cdylib
3+
// only-nvptx64
4+
35
#![feature(abi_ptx)]
6+
#![no_std]
47

5-
// Verify the default CUDA arch.
6-
// CHECK: .target sm_30
7-
// CHECK: .address_size 64
8+
// aux-build: breakpoint-panic-handler.rs
9+
extern crate breakpoint_panic_handler;
810

911
// Verify function name doesn't contain unacceaptable characters.
10-
// CHECK: .func (.param .b32 func_retval0) [[IMPL_FN:_ZN[a-zA-Z0-9$_]+square[a-zA-Z0-9$_]+]]
12+
// CHECK: .func (.param .b32 func_retval0) [[IMPL_FN:[a-zA-Z0-9$_]+square[a-zA-Z0-9$_]+]](
1113

1214
// CHECK-LABEL: .visible .entry top_kernel(
1315
#[no_mangle]
@@ -33,9 +35,3 @@ pub mod deep {
3335
}
3436
}
3537
}
36-
37-
// Verify that external function bodies are available.
38-
// CHECK: .func (.param .b32 func_retval0) [[IMPL_FN]]
39-
// CHECK: {
40-
// CHECK: mul.lo.s32 %{{r[0-9]+}}, %{{r[0-9]+}}, %{{r[0-9]+}}
41-
// CHECK: }

‎src/test/run-make/nvptx-binary-crate/Makefile

-12
This file was deleted.

‎src/test/run-make/nvptx-binary-crate/main.rs

-28
This file was deleted.

‎src/test/run-make/nvptx-dylib-crate/Makefile

-10
This file was deleted.

0 commit comments

Comments
 (0)