1 ; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s 2 3 ; Verify that we correctly emit code for i8 ldg/ldu. We do not expose 8-bit 4 ; registers in the backend, so these loads need special handling. 5 6 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64" 7 target triple = "nvptx64-unknown-unknown" 8 9 ; CHECK-LABEL: ex_zext 10 define void @ex_zext(i8* noalias readonly %data, i32* %res) { 11 entry: 12 ; CHECK: ld.global.nc.u8 13 %val = load i8, i8* %data 14 ; CHECK: cvt.u32.u8 15 %valext = zext i8 %val to i32 16 store i32 %valext, i32* %res 17 ret void 18 } 19 20 ; CHECK-LABEL: ex_sext 21 define void @ex_sext(i8* noalias readonly %data, i32* %res) { 22 entry: 23 ; CHECK: ld.global.nc.u8 24 %val = load i8, i8* %data 25 ; CHECK: cvt.s32.s8 26 %valext = sext i8 %val to i32 27 store i32 %valext, i32* %res 28 ret void 29 } 30 31 ; CHECK-LABEL: ex_zext_v2 32 define void @ex_zext_v2(<2 x i8>* noalias readonly %data, <2 x i32>* %res) { 33 entry: 34 ; CHECK: ld.global.nc.v2.u8 35 %val = load <2 x i8>, <2 x i8>* %data 36 ; CHECK: cvt.u32.u16 37 %valext = zext <2 x i8> %val to <2 x i32> 38 store <2 x i32> %valext, <2 x i32>* %res 39 ret void 40 } 41 42 ; CHECK-LABEL: ex_sext_v2 43 define void @ex_sext_v2(<2 x i8>* noalias readonly %data, <2 x i32>* %res) { 44 entry: 45 ; CHECK: ld.global.nc.v2.u8 46 %val = load <2 x i8>, <2 x i8>* %data 47 ; CHECK: cvt.s32.s8 48 %valext = sext <2 x i8> %val to <2 x i32> 49 store <2 x i32> %valext, <2 x i32>* %res 50 ret void 51 } 52 53 !nvvm.annotations = !{!0,!1,!2,!3} 54 !0 = !{void (i8*, i32*)* @ex_zext, !"kernel", i32 1} 55 !1 = !{void (i8*, i32*)* @ex_sext, !"kernel", i32 1} 56 !2 = !{void (<2 x i8>*, <2 x i32>*)* @ex_zext_v2, !"kernel", i32 1} 57 !3 = !{void (<2 x i8>*, <2 x i32>*)* @ex_sext_v2, !"kernel", i32 1} 58
