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