tests/assembly/nvptx-linking-binary.rs

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