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