From 218caa410aa38c29984be31a5229b9fa717560ee Mon Sep 17 00:00:00 2001 From: Daniel Baumann Date: Wed, 17 Apr 2024 14:19:13 +0200 Subject: Merging upstream version 1.68.2+dfsg1. Signed-off-by: Daniel Baumann --- tests/assembly/nvptx-linking-binary.rs | 40 ++++++++++++++++++++++++++++++++++ 1 file changed, 40 insertions(+) create mode 100644 tests/assembly/nvptx-linking-binary.rs (limited to 'tests/assembly/nvptx-linking-binary.rs') diff --git a/tests/assembly/nvptx-linking-binary.rs b/tests/assembly/nvptx-linking-binary.rs new file mode 100644 index 000000000..64b9c2f17 --- /dev/null +++ b/tests/assembly/nvptx-linking-binary.rs @@ -0,0 +1,40 @@ +// assembly-output: ptx-linker +// compile-flags: --crate-type bin +// only-nvptx64 +// ignore-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 -- cgit v1.2.3