diff options
| author | Guillaume Gomez <guillaume1.gomez@gmail.com> | 2025-07-21 14:22:51 +0200 |
|---|---|---|
| committer | Guillaume Gomez <guillaume1.gomez@gmail.com> | 2025-07-22 14:27:48 +0200 |
| commit | ed93c1783b404d728d4809973a0550eb33cd293f (patch) | |
| tree | 2b62aab8b2482fcde1f9128ff53be8cd66bde847 /tests/assembly-llvm/nvptx-linking-cdylib.rs | |
| parent | c0b282f0ccdab7523cdb8dfa41b23bed5573da76 (diff) | |
| download | rust-ed93c1783b404d728d4809973a0550eb33cd293f.tar.gz rust-ed93c1783b404d728d4809973a0550eb33cd293f.zip | |
Rename `tests/assembly` into `tests/assembly-llvm`
Diffstat (limited to 'tests/assembly-llvm/nvptx-linking-cdylib.rs')
| -rw-r--r-- | tests/assembly-llvm/nvptx-linking-cdylib.rs | 39 |
1 files changed, 39 insertions, 0 deletions
diff --git a/tests/assembly-llvm/nvptx-linking-cdylib.rs b/tests/assembly-llvm/nvptx-linking-cdylib.rs new file mode 100644 index 00000000000..9742e26fb31 --- /dev/null +++ b/tests/assembly-llvm/nvptx-linking-cdylib.rs @@ -0,0 +1,39 @@ +//@ assembly-output: ptx-linker +//@ compile-flags: --crate-type cdylib +//@ only-nvptx64 +//@ ignore-nvptx64 + +#![feature(abi_ptx)] +#![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 |