xref: /aosp_15_r20/external/llvm/test/CodeGen/NVPTX/pr13291-i1-store.ll (revision 9880d6810fe72a1726cb53787c6711e909410d58)
1*9880d681SAndroid Build Coastguard Worker; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32
2*9880d681SAndroid Build Coastguard Worker; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
3*9880d681SAndroid Build Coastguard Worker
4*9880d681SAndroid Build Coastguard Workerdefine ptx_kernel void @t1(i1* %a) {
5*9880d681SAndroid Build Coastguard Worker; PTX32:      mov.u16 %rs{{[0-9]+}}, 0;
6*9880d681SAndroid Build Coastguard Worker; PTX32-NEXT: st.global.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}};
7*9880d681SAndroid Build Coastguard Worker; PTX64:      mov.u16 %rs{{[0-9]+}}, 0;
8*9880d681SAndroid Build Coastguard Worker; PTX64-NEXT: st.global.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}};
9*9880d681SAndroid Build Coastguard Worker  store i1 false, i1* %a
10*9880d681SAndroid Build Coastguard Worker  ret void
11*9880d681SAndroid Build Coastguard Worker}
12*9880d681SAndroid Build Coastguard Worker
13*9880d681SAndroid Build Coastguard Worker
14*9880d681SAndroid Build Coastguard Workerdefine ptx_kernel void @t2(i1* %a, i8* %b) {
15*9880d681SAndroid Build Coastguard Worker; PTX32: ld.global.u8 %rs{{[0-9]+}}, [%r{{[0-9]+}}]
16*9880d681SAndroid Build Coastguard Worker; PTX32: and.b16 %rs{{[0-9]+}}, %rs{{[0-9]+}}, 1;
17*9880d681SAndroid Build Coastguard Worker; PTX32: setp.eq.b16 %p{{[0-9]+}}, %rs{{[0-9]+}}, 1;
18*9880d681SAndroid Build Coastguard Worker; PTX64: ld.global.u8 %rs{{[0-9]+}}, [%rd{{[0-9]+}}]
19*9880d681SAndroid Build Coastguard Worker; PTX64: and.b16 %rs{{[0-9]+}}, %rs{{[0-9]+}}, 1;
20*9880d681SAndroid Build Coastguard Worker; PTX64: setp.eq.b16 %p{{[0-9]+}}, %rs{{[0-9]+}}, 1;
21*9880d681SAndroid Build Coastguard Worker
22*9880d681SAndroid Build Coastguard Worker  %t1 = load i1, i1* %a
23*9880d681SAndroid Build Coastguard Worker  %t2 = select i1 %t1, i8 1, i8 2
24*9880d681SAndroid Build Coastguard Worker  store i8 %t2, i8* %b
25*9880d681SAndroid Build Coastguard Worker  ret void
26*9880d681SAndroid Build Coastguard Worker}
27