39 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_70 | FileCheck %s3; RUN: %if ptxas-sm_70 %{ llc < %s -mcpu=sm_70 | %ptxas-verify -arch=sm_70 %}4 5target triple = "nvptx64-nvidia-cuda"6 7%struct.double2 = type { double, double }8 9declare %struct.double2 @add(ptr align(16) byval(%struct.double2), ptr align(16) byval(%struct.double2))10 11define void @call_byval(ptr %out, ptr %in1, ptr %in2) {12; CHECK-LABEL: call_byval(13; CHECK: {14; CHECK-NEXT: .reg .b64 %rd<10>;15; CHECK-EMPTY:16; CHECK-NEXT: // %bb.0:17; CHECK-NEXT: ld.param.b64 %rd1, [call_byval_param_0];18; CHECK-NEXT: { // callseq 0, 019; CHECK-NEXT: .param .align 16 .b8 param0[16];20; CHECK-NEXT: .param .align 16 .b8 param1[16];21; CHECK-NEXT: .param .align 8 .b8 retval0[16];22; CHECK-NEXT: ld.param.b64 %rd2, [call_byval_param_2];23; CHECK-NEXT: ld.v2.b64 {%rd3, %rd4}, [%rd2];24; CHECK-NEXT: st.param.v2.b64 [param1], {%rd3, %rd4};25; CHECK-NEXT: ld.param.b64 %rd5, [call_byval_param_1];26; CHECK-NEXT: ld.v2.b64 {%rd6, %rd7}, [%rd5];27; CHECK-NEXT: st.param.v2.b64 [param0], {%rd6, %rd7};28; CHECK-NEXT: call.uni (retval0), add, (param0, param1);29; CHECK-NEXT: ld.param.b64 %rd8, [retval0+8];30; CHECK-NEXT: ld.param.b64 %rd9, [retval0];31; CHECK-NEXT: } // callseq 032; CHECK-NEXT: st.b64 [%rd1+8], %rd8;33; CHECK-NEXT: st.b64 [%rd1], %rd9;34; CHECK-NEXT: ret;35 %call = call %struct.double2 @add(ptr align(16) byval(%struct.double2) %in1, ptr align(16) byval(%struct.double2) %in2)36 store %struct.double2 %call, ptr %out, align 1637 ret void38}39