48 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -verify-machineinstrs | FileCheck %s3; RUN: %if ptxas %{ llc < %s | %ptxas-verify %}4 5target triple = "nvptx64-nvidia-cuda"6 7declare ptr @bar(i64)8declare i64 @baz()9 10define ptr @foo(i1 %cond) {11; CHECK-LABEL: foo(12; CHECK: {13; CHECK-NEXT: .reg .pred %p<2>;14; CHECK-NEXT: .reg .b16 %rs<3>;15; CHECK-NEXT: .reg .b64 %rd<2>;16; CHECK-EMPTY:17; CHECK-NEXT: // %bb.0: // %entry18; CHECK-NEXT: ld.param.b8 %rs1, [foo_param_0];19; CHECK-NEXT: and.b16 %rs2, %rs1, 1;20; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;21; CHECK-NEXT: { // callseq 0, 022; CHECK-NEXT: .param .b64 retval0;23; CHECK-NEXT: call.uni (retval0), baz, ();24; CHECK-NEXT: ld.param.b64 %rd1, [retval0];25; CHECK-NEXT: } // callseq 026; CHECK-NEXT: @%p1 bra $L__BB0_2;27; CHECK-NEXT: // %bb.1: // %bb28; CHECK-NEXT: { // callseq 1, 029; CHECK-NEXT: .param .b64 param0;30; CHECK-NEXT: .param .b64 retval0;31; CHECK-NEXT: st.param.b64 [param0], %rd1;32; CHECK-NEXT: call.uni (retval0), bar, (param0);33; CHECK-NEXT: } // callseq 134; CHECK-NEXT: $L__BB0_2: // %common.ret35; CHECK-NEXT: st.param.b64 [func_retval0], 0;36; CHECK-NEXT: ret;37entry:38 %call = call i64 @baz()39 br i1 %cond, label %common.ret, label %bb40 41bb:42 %tmp = call ptr @bar(i64 %call)43 br label %common.ret44 45common.ret:46 ret ptr null47}48