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 6target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64" 7target triple = "nvptx64-unknown-unknown" 8 9; CHECK-LABEL: ex_zext 10define void @ex_zext(i8* noalias readonly %data, i32* %res) { 11entry: 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 21define void @ex_sext(i8* noalias readonly %data, i32* %res) { 22entry: 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 32define void @ex_zext_v2(<2 x i8>* noalias readonly %data, <2 x i32>* %res) { 33entry: 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 43define void @ex_sext_v2(<2 x i8>* noalias readonly %data, <2 x i32>* %res) { 44entry: 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