103 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}4 5; Verify that we correctly emit code for i8 ldg/ldu. We do not expose 8-bit6; registers in the backend, so these loads need special handling.7 8target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"9target triple = "nvptx64-unknown-unknown"10 11define ptx_kernel void @ex_zext(ptr noalias readonly %data, ptr %res) {12; CHECK-LABEL: ex_zext(13; CHECK: {14; CHECK-NEXT: .reg .b32 %r<2>;15; CHECK-NEXT: .reg .b64 %rd<5>;16; CHECK-EMPTY:17; CHECK-NEXT: // %bb.0: // %entry18; CHECK-NEXT: ld.param.b64 %rd1, [ex_zext_param_0];19; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;20; CHECK-NEXT: ld.param.b64 %rd3, [ex_zext_param_1];21; CHECK-NEXT: cvta.to.global.u64 %rd4, %rd3;22; CHECK-NEXT: ld.global.nc.b8 %r1, [%rd2];23; CHECK-NEXT: st.global.b32 [%rd4], %r1;24; CHECK-NEXT: ret;25entry:26 %val = load i8, ptr %data27 %valext = zext i8 %val to i3228 store i32 %valext, ptr %res29 ret void30}31 32define ptx_kernel void @ex_sext(ptr noalias readonly %data, ptr %res) {33; CHECK-LABEL: ex_sext(34; CHECK: {35; CHECK-NEXT: .reg .b32 %r<2>;36; CHECK-NEXT: .reg .b64 %rd<5>;37; CHECK-EMPTY:38; CHECK-NEXT: // %bb.0: // %entry39; CHECK-NEXT: ld.param.b64 %rd1, [ex_sext_param_0];40; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;41; CHECK-NEXT: ld.param.b64 %rd3, [ex_sext_param_1];42; CHECK-NEXT: cvta.to.global.u64 %rd4, %rd3;43; CHECK-NEXT: ld.global.nc.s8 %r1, [%rd2];44; CHECK-NEXT: st.global.b32 [%rd4], %r1;45; CHECK-NEXT: ret;46entry:47 %val = load i8, ptr %data48 %valext = sext i8 %val to i3249 store i32 %valext, ptr %res50 ret void51}52 53define ptx_kernel void @ex_zext_v2(ptr noalias readonly %data, ptr %res) {54; CHECK-LABEL: ex_zext_v2(55; CHECK: {56; CHECK-NEXT: .reg .b16 %rs<3>;57; CHECK-NEXT: .reg .b32 %r<3>;58; CHECK-NEXT: .reg .b64 %rd<5>;59; CHECK-EMPTY:60; CHECK-NEXT: // %bb.0: // %entry61; CHECK-NEXT: ld.param.b64 %rd1, [ex_zext_v2_param_0];62; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;63; CHECK-NEXT: ld.param.b64 %rd3, [ex_zext_v2_param_1];64; CHECK-NEXT: cvta.to.global.u64 %rd4, %rd3;65; CHECK-NEXT: ld.global.nc.v2.b8 {%rs1, %rs2}, [%rd2];66; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;67; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;68; CHECK-NEXT: st.global.v2.b32 [%rd4], {%r2, %r1};69; CHECK-NEXT: ret;70entry:71 %val = load <2 x i8>, ptr %data72 %valext = zext <2 x i8> %val to <2 x i32>73 store <2 x i32> %valext, ptr %res74 ret void75}76 77define ptx_kernel void @ex_sext_v2(ptr noalias readonly %data, ptr %res) {78; CHECK-LABEL: ex_sext_v2(79; CHECK: {80; CHECK-NEXT: .reg .b16 %rs<3>;81; CHECK-NEXT: .reg .b32 %r<5>;82; CHECK-NEXT: .reg .b64 %rd<5>;83; CHECK-EMPTY:84; CHECK-NEXT: // %bb.0: // %entry85; CHECK-NEXT: ld.param.b64 %rd1, [ex_sext_v2_param_0];86; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;87; CHECK-NEXT: ld.param.b64 %rd3, [ex_sext_v2_param_1];88; CHECK-NEXT: cvta.to.global.u64 %rd4, %rd3;89; CHECK-NEXT: ld.global.nc.v2.b8 {%rs1, %rs2}, [%rd2];90; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;91; CHECK-NEXT: cvt.s32.s8 %r2, %r1;92; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;93; CHECK-NEXT: cvt.s32.s8 %r4, %r3;94; CHECK-NEXT: st.global.v2.b32 [%rd4], {%r4, %r2};95; CHECK-NEXT: ret;96entry:97 %val = load <2 x i8>, ptr %data98 %valext = sext <2 x i8> %val to <2 x i32>99 store <2 x i32> %valext, ptr %res100 ret void101}102 103