diff options
Diffstat (limited to 'src/test')
16 files changed, 250 insertions, 130 deletions
diff --git a/src/test/assembly/auxiliary/breakpoint-panic-handler.rs b/src/test/assembly/auxiliary/breakpoint-panic-handler.rs new file mode 100644 index 00000000000..d54c1181e1a --- /dev/null +++ b/src/test/assembly/auxiliary/breakpoint-panic-handler.rs @@ -0,0 +1,8 @@ +#![feature(core_intrinsics)] +#![no_std] + +#[panic_handler] +unsafe fn breakpoint_panic_handler(_: &::core::panic::PanicInfo) -> ! { + core::intrinsics::breakpoint(); + core::hint::unreachable_unchecked(); +} diff --git a/src/test/run-make/nvptx-dylib-crate/dep.rs b/src/test/assembly/auxiliary/non-inline-dependency.rs index 57f3ee87cdb..57f3ee87cdb 100644 --- a/src/test/run-make/nvptx-dylib-crate/dep.rs +++ b/src/test/assembly/auxiliary/non-inline-dependency.rs diff --git a/src/test/assembly/nvptx-arch-default.rs b/src/test/assembly/nvptx-arch-default.rs new file mode 100644 index 00000000000..7fe71c33521 --- /dev/null +++ b/src/test/assembly/nvptx-arch-default.rs @@ -0,0 +1,12 @@ +// assembly-output: ptx-linker +// compile-flags: --crate-type cdylib +// only-nvptx64 + +#![no_std] + +// aux-build: breakpoint-panic-handler.rs +extern crate breakpoint_panic_handler; + +// Verify default target arch with ptx-linker. +// CHECK: .target sm_30 +// CHECK: .address_size 64 diff --git a/src/test/assembly/nvptx-arch-emit-asm.rs b/src/test/assembly/nvptx-arch-emit-asm.rs new file mode 100644 index 00000000000..0ca17729c02 --- /dev/null +++ b/src/test/assembly/nvptx-arch-emit-asm.rs @@ -0,0 +1,9 @@ +// assembly-output: emit-asm +// compile-flags: --crate-type rlib +// only-nvptx64 + +#![no_std] + +// Verify default arch without ptx-linker involved. +// CHECK: .target sm_30 +// CHECK: .address_size 64 diff --git a/src/test/assembly/nvptx-arch-link-arg.rs b/src/test/assembly/nvptx-arch-link-arg.rs new file mode 100644 index 00000000000..f6b6e8ccaa1 --- /dev/null +++ b/src/test/assembly/nvptx-arch-link-arg.rs @@ -0,0 +1,12 @@ +// assembly-output: ptx-linker +// compile-flags: --crate-type cdylib -C link-arg=--arch=sm_60 +// only-nvptx64 + +#![no_std] + +// aux-build: breakpoint-panic-handler.rs +extern crate breakpoint_panic_handler; + +// Verify target arch override via `link-arg`. +// CHECK: .target sm_60 +// CHECK: .address_size 64 diff --git a/src/test/assembly/nvptx-arch-target-cpu.rs b/src/test/assembly/nvptx-arch-target-cpu.rs new file mode 100644 index 00000000000..08a7a193bbd --- /dev/null +++ b/src/test/assembly/nvptx-arch-target-cpu.rs @@ -0,0 +1,12 @@ +// assembly-output: ptx-linker +// compile-flags: --crate-type cdylib -C target-cpu=sm_50 +// only-nvptx64 + +#![no_std] + +// aux-build: breakpoint-panic-handler.rs +extern crate breakpoint_panic_handler; + +// Verify target arch override via `target-cpu`. +// CHECK: .target sm_50 +// CHECK: .address_size 64 diff --git a/src/test/assembly/nvptx-atomics.rs b/src/test/assembly/nvptx-atomics.rs new file mode 100644 index 00000000000..3bbd7b3d12d --- /dev/null +++ b/src/test/assembly/nvptx-atomics.rs @@ -0,0 +1,85 @@ +// assembly-output: ptx-linker +// compile-flags: --crate-type cdylib +// only-nvptx64 + +#![feature(abi_ptx, core_intrinsics)] +#![no_std] + +use core::intrinsics::*; + +// aux-build: breakpoint-panic-handler.rs +extern crate breakpoint_panic_handler; + +// Currently, LLVM NVPTX backend can only emit atomic instructions with +// `relaxed` (PTX default) ordering. But it's also useful to make sure +// the backend won't fail with other orders. Apparently, the backend +// doesn't support fences as well. As a workaround `llvm.nvvm.membar.*` +// could work, and perhaps on the long run, all the atomic operations +// should rather be provided by `core::arch::nvptx`. + +// Also, PTX ISA doesn't have atomic `load`, `store` and `nand`. + +// FIXME(denzp): add tests for `core::sync::atomic::*`. + +#[no_mangle] +pub unsafe extern "ptx-kernel" fn atomics_kernel(a: *mut u32) { + // CHECK: atom.global.and.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + // CHECK: atom.global.and.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + atomic_and(a, 1); + atomic_and_relaxed(a, 1); + + // CHECK: atom.global.cas.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1, 2; + // CHECK: atom.global.cas.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1, 2; + atomic_cxchg(a, 1, 2); + atomic_cxchg_relaxed(a, 1, 2); + + // CHECK: atom.global.max.s32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + // CHECK: atom.global.max.s32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + atomic_max(a, 1); + atomic_max_relaxed(a, 1); + + // CHECK: atom.global.min.s32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + // CHECK: atom.global.min.s32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + atomic_min(a, 1); + atomic_min_relaxed(a, 1); + + // CHECK: atom.global.or.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + // CHECK: atom.global.or.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + atomic_or(a, 1); + atomic_or_relaxed(a, 1); + + // CHECK: atom.global.max.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + // CHECK: atom.global.max.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + atomic_umax(a, 1); + atomic_umax_relaxed(a, 1); + + // CHECK: atom.global.min.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + // CHECK: atom.global.min.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + atomic_umin(a, 1); + atomic_umin_relaxed(a, 1); + + // CHECK: atom.global.add.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + // CHECK: atom.global.add.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + atomic_xadd(a, 1); + atomic_xadd_relaxed(a, 1); + + // CHECK: atom.global.exch.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + // CHECK: atom.global.exch.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + atomic_xchg(a, 1); + atomic_xchg_relaxed(a, 1); + + // CHECK: atom.global.xor.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + // CHECK: atom.global.xor.b32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], 1; + atomic_xor(a, 1); + atomic_xor_relaxed(a, 1); + + // CHECK: mov.u32 %[[sub_0_arg:r[0-9]+]], 100; + // CHECK: neg.s32 temp, %[[sub_0_arg]]; + // CHECK: atom.global.add.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], temp; + atomic_xsub(a, 100); + + // CHECK: mov.u32 %[[sub_1_arg:r[0-9]+]], 200; + // CHECK: neg.s32 temp, %[[sub_1_arg]]; + // CHECK: atom.global.add.u32 %{{r[0-9]+}}, [%{{rd[0-9]+}}], temp; + atomic_xsub_relaxed(a, 200); +} diff --git a/src/test/assembly/nvptx-internalizing.rs b/src/test/assembly/nvptx-internalizing.rs new file mode 100644 index 00000000000..db822642632 --- /dev/null +++ b/src/test/assembly/nvptx-internalizing.rs @@ -0,0 +1,27 @@ +// assembly-output: ptx-linker +// compile-flags: --crate-type cdylib +// only-nvptx64 + +#![feature(abi_ptx)] +#![no_std] + +// aux-build: breakpoint-panic-handler.rs +extern crate breakpoint_panic_handler; + +// aux-build: non-inline-dependency.rs +extern crate non_inline_dependency as dep; + +// Verify that no extra function declarations are present. +// CHECK-NOT: .func + +// CHECK: .visible .entry top_kernel( +#[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; +} + +// Verify that no extra function definitions are here. +// CHECK-NOT: .func +// CHECK-NOT: .entry + diff --git a/src/test/assembly/nvptx-linking-binary.rs b/src/test/assembly/nvptx-linking-binary.rs new file mode 100644 index 00000000000..d88ed9139ca --- /dev/null +++ b/src/test/assembly/nvptx-linking-binary.rs @@ -0,0 +1,39 @@ +// assembly-output: ptx-linker +// compile-flags: --crate-type bin +// only-nvptx64 + +#![feature(abi_ptx)] +#![no_main] +#![no_std] + +// aux-build: breakpoint-panic-handler.rs +extern crate breakpoint_panic_handler; + +// 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-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]; + 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: add.s32 %[[RES:r[0-9]+]], %[[RHS]], %[[LHS]]; + // CHECK: st.global.u32 [%{{rd[0-9]+}}], %[[RES]]; + *b = 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 diff --git a/src/test/assembly/nvptx-linking-cdylib.rs b/src/test/assembly/nvptx-linking-cdylib.rs new file mode 100644 index 00000000000..1145f567d8c --- /dev/null +++ b/src/test/assembly/nvptx-linking-cdylib.rs @@ -0,0 +1,38 @@ +// assembly-output: ptx-linker +// compile-flags: --crate-type cdylib +// only-nvptx64 + +#![feature(abi_ptx)] +#![no_std] + +// aux-build: breakpoint-panic-handler.rs +extern crate breakpoint_panic_handler; + +// 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-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]; + 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: add.s32 %[[RES:r[0-9]+]], %[[RHS]], %[[LHS]]; + // CHECK: st.global.u32 [%{{rd[0-9]+}}], %[[RES]]; + *b = 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 diff --git a/src/test/run-make/nvptx-emit-asm/kernel.rs b/src/test/assembly/nvptx-safe-naming.rs index b71e18d9103..ab6f91423aa 100644 --- a/src/test/run-make/nvptx-emit-asm/kernel.rs +++ b/src/test/assembly/nvptx-safe-naming.rs @@ -1,13 +1,15 @@ -#![no_std] -#![deny(warnings)] +// assembly-output: ptx-linker +// compile-flags: --crate-type cdylib +// only-nvptx64 + #![feature(abi_ptx)] +#![no_std] -// Verify the default CUDA arch. -// CHECK: .target sm_30 -// CHECK: .address_size 64 +// aux-build: breakpoint-panic-handler.rs +extern crate breakpoint_panic_handler; // Verify function name doesn't contain unacceaptable characters. -// CHECK: .func (.param .b32 func_retval0) [[IMPL_FN:_ZN[a-zA-Z0-9$_]+square[a-zA-Z0-9$_]+]] +// CHECK: .func (.param .b32 func_retval0) [[IMPL_FN:[a-zA-Z0-9$_]+square[a-zA-Z0-9$_]+]]( // CHECK-LABEL: .visible .entry top_kernel( #[no_mangle] @@ -33,9 +35,3 @@ pub mod deep { } } } - -// Verify that external function bodies are available. -// CHECK: .func (.param .b32 func_retval0) [[IMPL_FN]] -// CHECK: { -// CHECK: mul.lo.s32 %{{r[0-9]+}}, %{{r[0-9]+}}, %{{r[0-9]+}} -// CHECK: } diff --git a/src/test/run-make/nvptx-binary-crate/Makefile b/src/test/run-make/nvptx-binary-crate/Makefile deleted file mode 100644 index 2c211b5c785..00000000000 --- a/src/test/run-make/nvptx-binary-crate/Makefile +++ /dev/null @@ -1,12 +0,0 @@ --include ../../run-make-fulldeps/tools.mk - -ifeq ($(TARGET),nvptx64-nvidia-cuda) -all: - $(RUSTC) main.rs --crate-type="bin" --target $(TARGET) -O -C link-arg=--arch=sm_60 -o $(TMPDIR)/main.link_arg.ptx - $(RUSTC) main.rs --crate-type="bin" --target $(TARGET) -O -C target-cpu=sm_60 -o $(TMPDIR)/main.target_cpu.ptx - - FileCheck main.rs --input-file $(TMPDIR)/main.link_arg.ptx - FileCheck main.rs --input-file $(TMPDIR)/main.target_cpu.ptx -else -all: -endif diff --git a/src/test/run-make/nvptx-binary-crate/main.rs b/src/test/run-make/nvptx-binary-crate/main.rs deleted file mode 100644 index 826bc3a47bb..00000000000 --- a/src/test/run-make/nvptx-binary-crate/main.rs +++ /dev/null @@ -1,28 +0,0 @@ -#![no_std] -#![no_main] -#![deny(warnings)] -#![feature(abi_ptx, core_intrinsics)] - -// Check the overriden CUDA arch. -// CHECK: .target sm_60 -// CHECK: .address_size 64 - -// Verify that no extra function declarations are present. -// CHECK-NOT: .func - -// CHECK-LABEL: .visible .entry top_kernel( -#[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; -} - -// Verify that no extra function definitions are there. -// CHECK-NOT: .func -// CHECK-NOT: .entry - -#[panic_handler] -unsafe fn breakpoint_panic_handler(_: &::core::panic::PanicInfo) -> ! { - core::intrinsics::breakpoint(); - core::hint::unreachable_unchecked(); -} diff --git a/src/test/run-make/nvptx-dylib-crate/Makefile b/src/test/run-make/nvptx-dylib-crate/Makefile deleted file mode 100644 index 7284e9d1a7c..00000000000 --- a/src/test/run-make/nvptx-dylib-crate/Makefile +++ /dev/null @@ -1,10 +0,0 @@ --include ../../run-make-fulldeps/tools.mk - -ifeq ($(TARGET),nvptx64-nvidia-cuda) -all: - $(RUSTC) dep.rs --crate-type="rlib" --target $(TARGET) - $(RUSTC) kernel.rs --crate-type="cdylib" -O --target $(TARGET) - FileCheck kernel.rs --input-file $(TMPDIR)/kernel.ptx -else -all: -endif diff --git a/src/test/run-make/nvptx-dylib-crate/kernel.rs b/src/test/run-make/nvptx-dylib-crate/kernel.rs deleted file mode 100644 index 63fd6b063dd..00000000000 --- a/src/test/run-make/nvptx-dylib-crate/kernel.rs +++ /dev/null @@ -1,59 +0,0 @@ -#![no_std] -#![deny(warnings)] -#![feature(abi_ptx, core_intrinsics)] - -extern crate dep; - -// Verify the default CUDA arch. -// CHECK: .target sm_30 -// CHECK: .address_size 64 - -// 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 [[PANIC_HANDLER:_ZN4core9panicking5panic[a-zA-Z0-9]+]] - -// 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]; - 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: add.s32 %[[RES:r[0-9]+]], %[[RHS]], %[[LHS]]; - // CHECK: st.global.u32 [%{{rd[0-9]+}}], %[[RES]]; - *b = lhs + rhs; -} - -// Verify that external function bodies are available. -// CHECK-LABEL: .func (.param .b32 func_retval0) wrapping_external_fn -// CHECK: { -// CHECK: st.param.b32 [func_retval0+0], %{{r[0-9]+}}; -// CHECK: } - -// Also verify panic behavior. -// CHECK-LABEL: .func (.param .b32 func_retval0) panicking_external_fn -// CHECK: { -// CHECK: %{{p[0-9]+}} bra [[PANIC_LABEL:[a-zA-Z0-9_]+]]; -// CHECK: [[PANIC_LABEL]]: -// CHECK: call.uni -// CHECK: [[PANIC_HANDLER]] -// CHECK: } - -// Verify whether out dummy panic formatter has a correct body. -// CHECK: .func [[PANIC_FMT:_ZN4core9panicking9panic_fmt[a-zA-Z0-9]+]]() -// CHECK: { -// CHECK: trap; -// CHECK: } - -#[panic_handler] -unsafe fn breakpoint_panic_handler(_: &::core::panic::PanicInfo) -> ! { - core::intrinsics::breakpoint(); - core::hint::unreachable_unchecked(); -} diff --git a/src/test/run-make/nvptx-emit-asm/Makefile b/src/test/run-make/nvptx-emit-asm/Makefile deleted file mode 100644 index e03601878bd..00000000000 --- a/src/test/run-make/nvptx-emit-asm/Makefile +++ /dev/null @@ -1,9 +0,0 @@ --include ../../run-make-fulldeps/tools.mk - -ifeq ($(TARGET),nvptx64-nvidia-cuda) -all: - $(RUSTC) kernel.rs --crate-type="rlib" --emit asm,llvm-ir -O --target $(TARGET) - FileCheck kernel.rs --input-file $(TMPDIR)/kernel.s -else -all: -endif |