Merging upstream version 1.68.2+dfsg1.

Signed-off-by: Daniel Baumann <daniel.baumann@progress-linux.org>

1 files changed, 40 insertions, 0 deletions

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