xref: /llvm-project/llvm/test/CodeGen/NVPTX/i1-load-lower.ll (revision 310e79875752886a7713911e2a1ec14bc75bd4b3)
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --function foo --extra_scrub --default-march nvptx64 --filter-out ".*//.*" --filter-out "[\(\)\{\}]" --version 5
2
3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
4; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
5
6target triple = "nvptx-nvidia-cuda"
7
8@i1g = addrspace(1) global i1 false, align 2
9
10define void @foo() {
11; CHECK-LABEL: foo(
12; CHECK:    .reg .pred %p<2>;
13; CHECK:    .reg .b16 %rs<4>;
14; CHECK-EMPTY:
15; CHECK:    ld.global.u8 %rs1, [i1g];
16; CHECK:    and.b16 %rs2, %rs1, 1;
17; CHECK:    setp.eq.b16 %p1, %rs2, 1;
18; CHECK:    @%p1 bra $L__BB0_2;
19; CHECK:    mov.b16 %rs3, 1;
20; CHECK:    st.global.u8 [i1g], %rs3;
21; CHECK:    ret;
22  %tmp = load i1, ptr addrspace(1) @i1g, align 2
23  br i1 %tmp, label %if.end, label %if.then
24
25if.then:
26  store i1 true, ptr addrspace(1) @i1g, align 2
27  br label %if.end
28
29if.end:
30  ret void
31}
32