llvm/test/CodeGen/NVPTX/vaargs.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -O0 -mtriple=nvptx-- -mattr=+ptx60 -mcpu=sm_30 | FileCheck %s --check-prefix=CHECK32
; RUN: llc < %s -O0 -mtriple=nvptx64-- -mattr=+ptx60 -mcpu=sm_30 | FileCheck %s --check-prefix=CHECK64
; RUN: %if ptxas-isa-6.0 && ptxas-ptr32 %{ llc < %s -O0 -mtriple=nvptx -mattr=+ptx60 -mcpu=sm_30 | %ptxas-verify %}
; RUN: %if ptxas-isa-6.0 %{ llc < %s -O0 -mtriple=nvptx64 -mattr=+ptx60 -mcpu=sm_30 | %ptxas-verify %}

; CHECK: .address_size [[BITS:32|64]]

%struct.__va_list_tag = type { ptr, ptr, i32, i32 }

@foo_ptr = internal addrspace(1) global ptr @foo, align 8

define i32 @foo(i32 %a, ...) {
; CHECK32-LABEL: foo(
; CHECK32:       {
; CHECK32-NEXT:    .local .align 8 .b8 __local_depot0[32];
; CHECK32-NEXT:    .reg .b32 %SP;
; CHECK32-NEXT:    .reg .b32 %SPL;
; CHECK32-NEXT:    .reg .b16 %rs<5>;
; CHECK32-NEXT:    .reg .b32 %r<33>;
; CHECK32-NEXT:    .reg .b64 %rd<3>;
; CHECK32-EMPTY:
; CHECK32-NEXT:  // %bb.0: // %entry
; CHECK32-NEXT:    mov.b32 %SPL, __local_depot0;
; CHECK32-NEXT:    ld.param.b32 %r1, [foo_param_0];
; CHECK32-NEXT:    ld.param.b32 %r2, [foo_param_1];
; CHECK32-NEXT:    add.u32 %r3, %SPL, 0;
; CHECK32-NEXT:    cvta.local.u32 %r4, %r3;
; CHECK32-NEXT:    add.u32 %r5, %SPL, 16;
; CHECK32-NEXT:    cvta.local.u32 %r6, %r5;
; CHECK32-NEXT:    cvta.local.u32 %r7, %r2;
; CHECK32-NEXT:    st.b32 [%r4], %r7;
; CHECK32-NEXT:    ld.b8 %rs1, [%r4+3];
; CHECK32-NEXT:    st.b8 [%r6+3], %rs1;
; CHECK32-NEXT:    ld.b8 %rs2, [%r4+2];
; CHECK32-NEXT:    st.b8 [%r6+2], %rs2;
; CHECK32-NEXT:    ld.b8 %rs3, [%r4+1];
; CHECK32-NEXT:    st.b8 [%r6+1], %rs3;
; CHECK32-NEXT:    ld.b8 %rs4, [%r4];
; CHECK32-NEXT:    st.b8 [%r6], %rs4;
; CHECK32-NEXT:    ld.b32 %r8, [%r4];
; CHECK32-NEXT:    add.s32 %r9, %r8, 3;
; CHECK32-NEXT:    and.b32 %r10, %r9, -4;
; CHECK32-NEXT:    add.s32 %r11, %r10, 4;
; CHECK32-NEXT:    st.b32 [%r4], %r11;
; CHECK32-NEXT:    ld.b32 %r12, [%r10];
; CHECK32-NEXT:    ld.b32 %r13, [%r4];
; CHECK32-NEXT:    add.s32 %r14, %r13, 7;
; CHECK32-NEXT:    and.b32 %r15, %r14, -8;
; CHECK32-NEXT:    add.s32 %r16, %r15, 8;
; CHECK32-NEXT:    st.b32 [%r4], %r16;
; CHECK32-NEXT:    ld.b64 %rd1, [%r15];
; CHECK32-NEXT:    ld.b32 %r17, [%r4];
; CHECK32-NEXT:    add.s32 %r18, %r17, 7;
; CHECK32-NEXT:    and.b32 %r19, %r18, -8;
; CHECK32-NEXT:    add.s32 %r20, %r19, 8;
; CHECK32-NEXT:    st.b32 [%r4], %r20;
; CHECK32-NEXT:    ld.b64 %rd2, [%r19];
; CHECK32-NEXT:    ld.b32 %r21, [%r4];
; CHECK32-NEXT:    add.s32 %r22, %r21, 3;
; CHECK32-NEXT:    and.b32 %r23, %r22, -4;
; CHECK32-NEXT:    add.s32 %r24, %r23, 4;
; CHECK32-NEXT:    st.b32 [%r4], %r24;
; CHECK32-NEXT:    ld.b32 %r25, [%r23];
; CHECK32-NEXT:    { // callseq 0, 0
; CHECK32-NEXT:    .param .b32 param0;
; CHECK32-NEXT:    .param .b32 param1;
; CHECK32-NEXT:    .param .b64 param2;
; CHECK32-NEXT:    .param .b64 param3;
; CHECK32-NEXT:    .param .b32 param4;
; CHECK32-NEXT:    .param .b32 retval0;
; CHECK32-NEXT:    st.param.b32 [param4], %r25;
; CHECK32-NEXT:    st.param.b64 [param3], %rd2;
; CHECK32-NEXT:    st.param.b64 [param2], %rd1;
; CHECK32-NEXT:    st.param.b32 [param1], %r12;
; CHECK32-NEXT:    st.param.b32 [param0], %r1;
; CHECK32-NEXT:    call.uni (retval0), bar, (param0, param1, param2, param3, param4);
; CHECK32-NEXT:    ld.param.b32 %r26, [retval0];
; CHECK32-NEXT:    } // callseq 0
; CHECK32-NEXT:    ld.b32 %r27, [%r6];
; CHECK32-NEXT:    add.s32 %r28, %r27, 3;
; CHECK32-NEXT:    and.b32 %r29, %r28, -4;
; CHECK32-NEXT:    add.s32 %r30, %r29, 4;
; CHECK32-NEXT:    st.b32 [%r6], %r30;
; CHECK32-NEXT:    ld.b32 %r31, [%r29];
; CHECK32-NEXT:    add.s32 %r32, %r26, %r31;
; CHECK32-NEXT:    st.param.b32 [func_retval0], %r32;
; CHECK32-NEXT:    ret;
;
; CHECK64-LABEL: foo(
; CHECK64:       {
; CHECK64-NEXT:    .local .align 8 .b8 __local_depot0[48];
; CHECK64-NEXT:    .reg .b64 %SP;
; CHECK64-NEXT:    .reg .b64 %SPL;
; CHECK64-NEXT:    .reg .b16 %rs<9>;
; CHECK64-NEXT:    .reg .b32 %r<6>;
; CHECK64-NEXT:    .reg .b64 %rd<30>;
; CHECK64-EMPTY:
; CHECK64-NEXT:  // %bb.0: // %entry
; CHECK64-NEXT:    mov.b64 %SPL, __local_depot0;
; CHECK64-NEXT:    ld.param.b32 %r1, [foo_param_0];
; CHECK64-NEXT:    ld.param.b64 %rd1, [foo_param_1];
; CHECK64-NEXT:    add.u64 %rd2, %SPL, 0;
; CHECK64-NEXT:    cvta.local.u64 %rd3, %rd2;
; CHECK64-NEXT:    add.u64 %rd4, %SPL, 24;
; CHECK64-NEXT:    cvta.local.u64 %rd5, %rd4;
; CHECK64-NEXT:    cvta.local.u64 %rd6, %rd1;
; CHECK64-NEXT:    st.b64 [%rd3], %rd6;
; CHECK64-NEXT:    ld.b8 %rs1, [%rd3+7];
; CHECK64-NEXT:    st.b8 [%rd5+7], %rs1;
; CHECK64-NEXT:    ld.b8 %rs2, [%rd3+6];
; CHECK64-NEXT:    st.b8 [%rd5+6], %rs2;
; CHECK64-NEXT:    ld.b8 %rs3, [%rd3+5];
; CHECK64-NEXT:    st.b8 [%rd5+5], %rs3;
; CHECK64-NEXT:    ld.b8 %rs4, [%rd3+4];
; CHECK64-NEXT:    st.b8 [%rd5+4], %rs4;
; CHECK64-NEXT:    ld.b8 %rs5, [%rd3+3];
; CHECK64-NEXT:    st.b8 [%rd5+3], %rs5;
; CHECK64-NEXT:    ld.b8 %rs6, [%rd3+2];
; CHECK64-NEXT:    st.b8 [%rd5+2], %rs6;
; CHECK64-NEXT:    ld.b8 %rs7, [%rd3+1];
; CHECK64-NEXT:    st.b8 [%rd5+1], %rs7;
; CHECK64-NEXT:    ld.b8 %rs8, [%rd3];
; CHECK64-NEXT:    st.b8 [%rd5], %rs8;
; CHECK64-NEXT:    ld.b64 %rd7, [%rd3];
; CHECK64-NEXT:    add.s64 %rd8, %rd7, 3;
; CHECK64-NEXT:    and.b64 %rd9, %rd8, -4;
; CHECK64-NEXT:    add.s64 %rd10, %rd9, 4;
; CHECK64-NEXT:    st.b64 [%rd3], %rd10;
; CHECK64-NEXT:    ld.b32 %r2, [%rd9];
; CHECK64-NEXT:    ld.b64 %rd11, [%rd3];
; CHECK64-NEXT:    add.s64 %rd12, %rd11, 7;
; CHECK64-NEXT:    and.b64 %rd13, %rd12, -8;
; CHECK64-NEXT:    add.s64 %rd14, %rd13, 8;
; CHECK64-NEXT:    st.b64 [%rd3], %rd14;
; CHECK64-NEXT:    ld.b64 %rd15, [%rd13];
; CHECK64-NEXT:    ld.b64 %rd16, [%rd3];
; CHECK64-NEXT:    add.s64 %rd17, %rd16, 7;
; CHECK64-NEXT:    and.b64 %rd18, %rd17, -8;
; CHECK64-NEXT:    add.s64 %rd19, %rd18, 8;
; CHECK64-NEXT:    st.b64 [%rd3], %rd19;
; CHECK64-NEXT:    ld.b64 %rd20, [%rd18];
; CHECK64-NEXT:    ld.b64 %rd21, [%rd3];
; CHECK64-NEXT:    add.s64 %rd22, %rd21, 7;
; CHECK64-NEXT:    and.b64 %rd23, %rd22, -8;
; CHECK64-NEXT:    add.s64 %rd24, %rd23, 8;
; CHECK64-NEXT:    st.b64 [%rd3], %rd24;
; CHECK64-NEXT:    ld.b64 %rd25, [%rd23];
; CHECK64-NEXT:    { // callseq 0, 0
; CHECK64-NEXT:    .param .b32 param0;
; CHECK64-NEXT:    .param .b32 param1;
; CHECK64-NEXT:    .param .b64 param2;
; CHECK64-NEXT:    .param .b64 param3;
; CHECK64-NEXT:    .param .b64 param4;
; CHECK64-NEXT:    .param .b32 retval0;
; CHECK64-NEXT:    st.param.b64 [param4], %rd25;
; CHECK64-NEXT:    st.param.b64 [param3], %rd20;
; CHECK64-NEXT:    st.param.b64 [param2], %rd15;
; CHECK64-NEXT:    st.param.b32 [param1], %r2;
; CHECK64-NEXT:    st.param.b32 [param0], %r1;
; CHECK64-NEXT:    call.uni (retval0), bar, (param0, param1, param2, param3, param4);
; CHECK64-NEXT:    ld.param.b32 %r3, [retval0];
; CHECK64-NEXT:    } // callseq 0
; CHECK64-NEXT:    ld.b64 %rd26, [%rd5];
; CHECK64-NEXT:    add.s64 %rd27, %rd26, 3;
; CHECK64-NEXT:    and.b64 %rd28, %rd27, -4;
; CHECK64-NEXT:    add.s64 %rd29, %rd28, 4;
; CHECK64-NEXT:    st.b64 [%rd5], %rd29;
; CHECK64-NEXT:    ld.b32 %r4, [%rd28];
; CHECK64-NEXT:    add.s32 %r5, %r3, %r4;
; CHECK64-NEXT:    st.param.b32 [func_retval0], %r5;
; CHECK64-NEXT:    ret;
entry:
  %al = alloca [1 x %struct.__va_list_tag], align 8
  %al2 = alloca [1 x %struct.__va_list_tag], align 8

; Test va_start

  call void @llvm.va_start(ptr %al)

; Test va_copy()

  call void @llvm.va_copy(ptr %al2, ptr %al)

; Test va_arg(ap, int32_t)

  %0 = va_arg ptr %al, i32

; Test va_arg(ap, int64_t)

  %1 = va_arg ptr %al, i64

; Test va_arg(ap, double)

  %2 = va_arg ptr %al, double

; Test va_arg(ap, ptr)

  %3 = va_arg ptr %al, ptr
  %call = call i32 @bar(i32 %a, i32 %0, i64 %1, double %2, ptr %3)

  call void @llvm.va_end(ptr %al)
  %4 =  va_arg ptr %al2, i32
  call void @llvm.va_end(ptr %al2)
  %5 = add i32 %call, %4
  ret i32 %5
}

define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; Test indirect variadic function call.
; Load arguments to temporary variables
; CHECK32-LABEL: test_foo(
; CHECK32:       {
; CHECK32-NEXT:    .local .align 8 .b8 __local_depot1[32];
; CHECK32-NEXT:    .reg .b32 %SP;
; CHECK32-NEXT:    .reg .b32 %SPL;
; CHECK32-NEXT:    .reg .b32 %r<9>;
; CHECK32-NEXT:    .reg .b64 %rd<3>;
; CHECK32-EMPTY:
; CHECK32-NEXT:  $L__prototype_0:
; CHECK32-NEXT:  .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
; CHECK32-NEXT:  // %bb.0: // %entry
; CHECK32-NEXT:    mov.b32 %SPL, __local_depot1;
; CHECK32-NEXT:    ld.param.b32 %r1, [test_foo_param_0];
; CHECK32-NEXT:    ld.param.b64 %rd1, [test_foo_param_1];
; CHECK32-NEXT:    add.u32 %r2, %SPL, 0;
; CHECK32-NEXT:    cvta.local.u32 %r3, %r2;
; CHECK32-NEXT:    ld.param.b64 %rd2, [test_foo_param_2];
; CHECK32-NEXT:    mov.b32 %r4, foo_ptr;
; CHECK32-NEXT:    cvta.global.u32 %r5, %r4;
; CHECK32-NEXT:    ld.b32 %r6, [%r5];
; CHECK32-NEXT:    ld.param.b32 %r7, [test_foo_param_3];
; CHECK32-NEXT:    st.b32 [%r3], %r1;
; CHECK32-NEXT:    st.b64 [%r3+8], %rd1;
; CHECK32-NEXT:    st.b64 [%r3+16], %rd2;
; CHECK32-NEXT:    st.b32 [%r3+24], %r7;
; CHECK32-NEXT:    { // callseq 1, 0
; CHECK32-NEXT:    .param .b32 param0;
; CHECK32-NEXT:    .param .b32 param1;
; CHECK32-NEXT:    .param .b32 retval0;
; CHECK32-NEXT:    st.param.b32 [param1], %r2;
; CHECK32-NEXT:    st.param.b32 [param0], 4;
; CHECK32-NEXT:    call (retval0), %r6, (param0, param1), $L__prototype_0;
; CHECK32-NEXT:    ld.param.b32 %r8, [retval0];
; CHECK32-NEXT:    } // callseq 1
; CHECK32-NEXT:    st.param.b32 [func_retval0], %r8;
; CHECK32-NEXT:    ret;
;
; CHECK64-LABEL: test_foo(
; CHECK64:       {
; CHECK64-NEXT:    .local .align 8 .b8 __local_depot1[32];
; CHECK64-NEXT:    .reg .b64 %SP;
; CHECK64-NEXT:    .reg .b64 %SPL;
; CHECK64-NEXT:    .reg .b32 %r<3>;
; CHECK64-NEXT:    .reg .b64 %rd<9>;
; CHECK64-EMPTY:
; CHECK64-NEXT:  $L__prototype_0:
; CHECK64-NEXT:  .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
; CHECK64-NEXT:  // %bb.0: // %entry
; CHECK64-NEXT:    mov.b64 %SPL, __local_depot1;
; CHECK64-NEXT:    ld.param.b32 %r1, [test_foo_param_0];
; CHECK64-NEXT:    ld.param.b64 %rd1, [test_foo_param_1];
; CHECK64-NEXT:    add.u64 %rd2, %SPL, 0;
; CHECK64-NEXT:    cvta.local.u64 %rd3, %rd2;
; CHECK64-NEXT:    ld.param.b64 %rd4, [test_foo_param_2];
; CHECK64-NEXT:    mov.b64 %rd5, foo_ptr;
; CHECK64-NEXT:    cvta.global.u64 %rd6, %rd5;
; CHECK64-NEXT:    ld.b64 %rd7, [%rd6];
; CHECK64-NEXT:    ld.param.b64 %rd8, [test_foo_param_3];
; CHECK64-NEXT:    st.b32 [%rd3], %r1;
; CHECK64-NEXT:    st.b64 [%rd3+8], %rd1;
; CHECK64-NEXT:    st.b64 [%rd3+16], %rd4;
; CHECK64-NEXT:    st.b64 [%rd3+24], %rd8;
; CHECK64-NEXT:    { // callseq 1, 0
; CHECK64-NEXT:    .param .b32 param0;
; CHECK64-NEXT:    .param .b64 param1;
; CHECK64-NEXT:    .param .b32 retval0;
; CHECK64-NEXT:    st.param.b64 [param1], %rd2;
; CHECK64-NEXT:    st.param.b32 [param0], 4;
; CHECK64-NEXT:    call (retval0), %rd7, (param0, param1), $L__prototype_0;
; CHECK64-NEXT:    ld.param.b32 %r2, [retval0];
; CHECK64-NEXT:    } // callseq 1
; CHECK64-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK64-NEXT:    ret;

; Store arguments to an array

entry:
  %ptr = load ptr, ptr addrspacecast (ptr addrspace(1) @foo_ptr to ptr), align 8
  %call = call i32 (i32, ...) %ptr(i32 4, i32 %i, i64 %l, double %d, ptr %p)
  ret i32 %call
}

declare void @llvm.va_start(ptr)
declare void @llvm.va_end(ptr)
declare void @llvm.va_copy(ptr, ptr)
declare i32 @bar(i32, i32, i64, double, ptr)