| ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 |
| ; RUN: llc < %s -mcpu=sm_70 | FileCheck %s |
| ; RUN: %if ptxas %{ llc < %s -mcpu=sm_70 | %ptxas-verify -arch=sm_70 %} |
| |
| target triple = "nvptx64-nvidia-cuda" |
| |
| %struct = type { [2 x i64] } |
| @G = external constant %struct |
| |
| define void @foo() { |
| ; CHECK-LABEL: foo( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b64 %rd<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.global.b64 %rd1, [G]; |
| ; CHECK-NEXT: ld.global.b64 %rd2, [G+8]; |
| ; CHECK-NEXT: { // callseq 0, 0 |
| ; CHECK-NEXT: .param .align 8 .b8 param0[16]; |
| ; CHECK-NEXT: st.param.b64 [param0], %rd1; |
| ; CHECK-NEXT: st.param.b64 [param0+8], %rd2; |
| ; CHECK-NEXT: call.uni |
| ; CHECK-NEXT: bar, |
| ; CHECK-NEXT: ( |
| ; CHECK-NEXT: param0 |
| ; CHECK-NEXT: ); |
| ; CHECK-NEXT: } // callseq 0 |
| ; CHECK-NEXT: ret; |
| call void @bar(ptr byval(%struct) @G) |
| ret void |
| } |
| |
| declare void @bar(ptr byval(%struct)) |
