test/CodeGen/NVPTX/ld-param-sink.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s | %ptxas-verify %}

target triple = "nvptx64-nvidia-cuda"

declare ptr @bar(i64)
declare i64 @baz()

define ptr @foo(i1 %cond) {
; CHECK-LABEL: foo(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b8 %rs1, [foo_param_0];
; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
; CHECK-NEXT:    setp.ne.b16 %p1, %rs2, 0;
; CHECK-NEXT:    { // callseq 0, 0
; CHECK-NEXT:    .param .b64 retval0;
; CHECK-NEXT:    call.uni (retval0), baz, ();
; CHECK-NEXT:    ld.param.b64 %rd1, [retval0];
; CHECK-NEXT:    } // callseq 0
; CHECK-NEXT:    @%p1 bra $L__BB0_2;
; CHECK-NEXT:  // %bb.1: // %bb
; CHECK-NEXT:    { // callseq 1, 0
; CHECK-NEXT:    .param .b64 param0;
; CHECK-NEXT:    .param .b64 retval0;
; CHECK-NEXT:    st.param.b64 [param0], %rd1;
; CHECK-NEXT:    call.uni (retval0), bar, (param0);
; CHECK-NEXT:    } // callseq 1
; CHECK-NEXT:  $L__BB0_2: // %common.ret
; CHECK-NEXT:    st.param.b64 [func_retval0], 0;
; CHECK-NEXT:    ret;
entry:
  %call = call i64 @baz()
  br i1 %cond, label %common.ret, label %bb

bb:
  %tmp = call ptr @bar(i64 %call)
  br label %common.ret

common.ret:
  ret ptr null
}