src/test/run-make/nvptx-dylib-crate/kernel.rs

   1 #![no_std]
   2 #![deny(warnings)]
   3 #![feature(abi_ptx, core_intrinsics)]
   4
   5 extern crate dep;
   6
   7 // Verify the default CUDA arch.
   8 // CHECK: .target sm_20
   9 // CHECK: .address_size 64
  10
  11 // Make sure declarations are there.
  12 // CHECK: .func (.param .b32 func_retval0) wrapping_external_fn
  13 // CHECK: .func (.param .b32 func_retval0) panicking_external_fn
  14 // CHECK: .func [[PANIC_HANDLER:_ZN4core9panicking5panic[a-zA-Z0-9]+]]
  15
  16 // CHECK-LABEL: .visible .entry top_kernel(
  17 #[no_mangle]
  18 pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) {
  19     // CHECK:      call.uni (retval0),
  20     // CHECK-NEXT: wrapping_external_fn
  21     // CHECK:      ld.param.b32 %[[LHS:r[0-9]+]], [retval0+0];
  22     let lhs = dep::wrapping_external_fn(*a);
  23
  24     // CHECK:      call.uni (retval0),
  25     // CHECK-NEXT: panicking_external_fn
  26     // CHECK:      ld.param.b32 %[[RHS:r[0-9]+]], [retval0+0];
  27     let rhs = dep::panicking_external_fn(*a);
  28
  29     // CHECK: add.s32 %[[RES:r[0-9]+]], %[[RHS]], %[[LHS]];
  30     // CHECK: st.global.u32 [%{{rd[0-9]+}}], %[[RES]];
  31     *b = lhs + rhs;
  32 }
  33
  34 // Verify that external function bodies are available.
  35 // CHECK-LABEL: .func (.param .b32 func_retval0) wrapping_external_fn
  36 // CHECK: {
  37 // CHECK:   st.param.b32 [func_retval0+0], %{{r[0-9]+}};
  38 // CHECK: }
  39
  40 // Also verify panic behavior.
  41 // CHECK-LABEL: .func (.param .b32 func_retval0) panicking_external_fn
  42 // CHECK: {
  43 // CHECK:   %{{p[0-9]+}} bra [[PANIC_LABEL:[a-zA-Z0-9_]+]];
  44 // CHECK: [[PANIC_LABEL]]:
  45 // CHECK:   call.uni
  46 // CHECK:   [[PANIC_HANDLER]]
  47 // CHECK: }
  48
  49 // Verify whether out dummy panic formatter has a correct body.
  50 // CHECK: .func [[PANIC_FMT:_ZN4core9panicking9panic_fmt[a-zA-Z0-9]+]]()
  51 // CHECK: {
  52 // CHECK:   trap;
  53 // CHECK: }
  54
  55 #[panic_handler]
  56 unsafe fn breakpoint_panic_handler(_: &::core::panic::PanicInfo) -> ! {
  57     core::intrinsics::breakpoint();
  58     core::hint::unreachable_unchecked();
  59 }