1 // assembly-output: ptx-linker
2 // compile-flags: --crate-type cdylib
8 // aux-build: breakpoint-panic-handler.rs
9 extern crate breakpoint_panic_handler;
11 // aux-build: non-inline-dependency.rs
12 extern crate non_inline_dependency as dep;
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
18 // CHECK-LABEL: .visible .entry top_kernel(
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);
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);
31 // CHECK: add.s32 %[[RES:r[0-9]+]], %[[RHS]], %[[LHS]];
32 // CHECK: st.global.u32 [%{{rd[0-9]+}}], %[[RES]];
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