276 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_20 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}4 5target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"6target triple = "nvptx64-nvidia-cuda"7 8define <4 x float> @t1(ptr %p1) {9; CHECK-LABEL: t1(10; CHECK: {11; CHECK-NEXT: .reg .b32 %r<41>;12; CHECK-NEXT: .reg .b64 %rd<2>;13; CHECK-EMPTY:14; CHECK-NEXT: // %bb.0:15; CHECK-NEXT: ld.param.b64 %rd1, [t1_param_0];16; CHECK-NEXT: ld.b8 %r1, [%rd1+12];17; CHECK-NEXT: ld.b8 %r2, [%rd1+13];18; CHECK-NEXT: shl.b32 %r3, %r2, 8;19; CHECK-NEXT: or.b32 %r4, %r3, %r1;20; CHECK-NEXT: ld.b8 %r5, [%rd1+14];21; CHECK-NEXT: shl.b32 %r6, %r5, 16;22; CHECK-NEXT: ld.b8 %r7, [%rd1+15];23; CHECK-NEXT: shl.b32 %r8, %r7, 24;24; CHECK-NEXT: or.b32 %r9, %r8, %r6;25; CHECK-NEXT: or.b32 %r10, %r9, %r4;26; CHECK-NEXT: ld.b8 %r11, [%rd1+8];27; CHECK-NEXT: ld.b8 %r12, [%rd1+9];28; CHECK-NEXT: shl.b32 %r13, %r12, 8;29; CHECK-NEXT: or.b32 %r14, %r13, %r11;30; CHECK-NEXT: ld.b8 %r15, [%rd1+10];31; CHECK-NEXT: shl.b32 %r16, %r15, 16;32; CHECK-NEXT: ld.b8 %r17, [%rd1+11];33; CHECK-NEXT: shl.b32 %r18, %r17, 24;34; CHECK-NEXT: or.b32 %r19, %r18, %r16;35; CHECK-NEXT: or.b32 %r20, %r19, %r14;36; CHECK-NEXT: ld.b8 %r21, [%rd1+4];37; CHECK-NEXT: ld.b8 %r22, [%rd1+5];38; CHECK-NEXT: shl.b32 %r23, %r22, 8;39; CHECK-NEXT: or.b32 %r24, %r23, %r21;40; CHECK-NEXT: ld.b8 %r25, [%rd1+6];41; CHECK-NEXT: shl.b32 %r26, %r25, 16;42; CHECK-NEXT: ld.b8 %r27, [%rd1+7];43; CHECK-NEXT: shl.b32 %r28, %r27, 24;44; CHECK-NEXT: or.b32 %r29, %r28, %r26;45; CHECK-NEXT: or.b32 %r30, %r29, %r24;46; CHECK-NEXT: ld.b8 %r31, [%rd1];47; CHECK-NEXT: ld.b8 %r32, [%rd1+1];48; CHECK-NEXT: shl.b32 %r33, %r32, 8;49; CHECK-NEXT: or.b32 %r34, %r33, %r31;50; CHECK-NEXT: ld.b8 %r35, [%rd1+2];51; CHECK-NEXT: shl.b32 %r36, %r35, 16;52; CHECK-NEXT: ld.b8 %r37, [%rd1+3];53; CHECK-NEXT: shl.b32 %r38, %r37, 24;54; CHECK-NEXT: or.b32 %r39, %r38, %r36;55; CHECK-NEXT: or.b32 %r40, %r39, %r34;56; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r40, %r30, %r20, %r10};57; CHECK-NEXT: ret;58 %r = load <4 x float>, ptr %p1, align 159 ret <4 x float> %r60}61 62define <4 x float> @t2(ptr %p1) {63; CHECK-LABEL: t2(64; CHECK: {65; CHECK-NEXT: .reg .b32 %r<5>;66; CHECK-NEXT: .reg .b64 %rd<2>;67; CHECK-EMPTY:68; CHECK-NEXT: // %bb.0:69; CHECK-NEXT: ld.param.b64 %rd1, [t2_param_0];70; CHECK-NEXT: ld.b32 %r1, [%rd1+12];71; CHECK-NEXT: ld.b32 %r2, [%rd1+8];72; CHECK-NEXT: ld.b32 %r3, [%rd1+4];73; CHECK-NEXT: ld.b32 %r4, [%rd1];74; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r4, %r3, %r2, %r1};75; CHECK-NEXT: ret;76 %r = load <4 x float>, ptr %p1, align 477 ret <4 x float> %r78}79 80define <4 x float> @t3(ptr %p1) {81; CHECK-LABEL: t3(82; CHECK: {83; CHECK-NEXT: .reg .b32 %r<5>;84; CHECK-NEXT: .reg .b64 %rd<2>;85; CHECK-EMPTY:86; CHECK-NEXT: // %bb.0:87; CHECK-NEXT: ld.param.b64 %rd1, [t3_param_0];88; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1+8];89; CHECK-NEXT: ld.v2.b32 {%r3, %r4}, [%rd1];90; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r3, %r4, %r1, %r2};91; CHECK-NEXT: ret;92 %r = load <4 x float>, ptr %p1, align 893 ret <4 x float> %r94}95 96define <4 x float> @t4(ptr %p1) {97; CHECK-LABEL: t4(98; CHECK: {99; CHECK-NEXT: .reg .b32 %r<5>;100; CHECK-NEXT: .reg .b64 %rd<2>;101; CHECK-EMPTY:102; CHECK-NEXT: // %bb.0:103; CHECK-NEXT: ld.param.b64 %rd1, [t4_param_0];104; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];105; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r1, %r2, %r3, %r4};106; CHECK-NEXT: ret;107 %r = load <4 x float>, ptr %p1, align 16108 ret <4 x float> %r109}110 111define void @test_v1halfp0a1(ptr noalias readonly %from, ptr %to) {112; CHECK-LABEL: test_v1halfp0a1(113; CHECK: {114; CHECK-NEXT: .reg .b16 %rs<3>;115; CHECK-NEXT: .reg .b64 %rd<3>;116; CHECK-EMPTY:117; CHECK-NEXT: // %bb.0:118; CHECK-NEXT: ld.param.b64 %rd1, [test_v1halfp0a1_param_0];119; CHECK-NEXT: ld.b8 %rs1, [%rd1];120; CHECK-NEXT: ld.b8 %rs2, [%rd1+1];121; CHECK-NEXT: ld.param.b64 %rd2, [test_v1halfp0a1_param_1];122; CHECK-NEXT: st.b8 [%rd2+1], %rs2;123; CHECK-NEXT: st.b8 [%rd2], %rs1;124; CHECK-NEXT: ret;125 %1 = load <1 x half>, ptr %from , align 1126 store <1 x half> %1, ptr %to , align 1127 ret void128}129 130define void @test_v2halfp0a1(ptr noalias readonly %from, ptr %to) {131; CHECK-LABEL: test_v2halfp0a1(132; CHECK: {133; CHECK-NEXT: .reg .b32 %r<5>;134; CHECK-NEXT: .reg .b64 %rd<3>;135; CHECK-EMPTY:136; CHECK-NEXT: // %bb.0:137; CHECK-NEXT: ld.param.b64 %rd1, [test_v2halfp0a1_param_0];138; CHECK-NEXT: ld.b8 %r1, [%rd1+1];139; CHECK-NEXT: ld.b8 %r2, [%rd1];140; CHECK-NEXT: ld.b8 %r3, [%rd1+3];141; CHECK-NEXT: ld.b8 %r4, [%rd1+2];142; CHECK-NEXT: ld.param.b64 %rd2, [test_v2halfp0a1_param_1];143; CHECK-NEXT: st.b8 [%rd2+2], %r4;144; CHECK-NEXT: st.b8 [%rd2+3], %r3;145; CHECK-NEXT: st.b8 [%rd2], %r2;146; CHECK-NEXT: st.b8 [%rd2+1], %r1;147; CHECK-NEXT: ret;148 %1 = load <2 x half>, ptr %from , align 1149 store <2 x half> %1, ptr %to , align 1150 ret void151}152 153define void @test_v4halfp0a1(ptr noalias readonly %from, ptr %to) {154; CHECK-LABEL: test_v4halfp0a1(155; CHECK: {156; CHECK-NEXT: .reg .b32 %r<9>;157; CHECK-NEXT: .reg .b64 %rd<3>;158; CHECK-EMPTY:159; CHECK-NEXT: // %bb.0:160; CHECK-NEXT: ld.param.b64 %rd1, [test_v4halfp0a1_param_0];161; CHECK-NEXT: ld.b8 %r1, [%rd1+1];162; CHECK-NEXT: ld.b8 %r2, [%rd1];163; CHECK-NEXT: ld.b8 %r3, [%rd1+3];164; CHECK-NEXT: ld.b8 %r4, [%rd1+2];165; CHECK-NEXT: ld.b8 %r5, [%rd1+5];166; CHECK-NEXT: ld.b8 %r6, [%rd1+4];167; CHECK-NEXT: ld.b8 %r7, [%rd1+7];168; CHECK-NEXT: ld.b8 %r8, [%rd1+6];169; CHECK-NEXT: ld.param.b64 %rd2, [test_v4halfp0a1_param_1];170; CHECK-NEXT: st.b8 [%rd2+6], %r8;171; CHECK-NEXT: st.b8 [%rd2+7], %r7;172; CHECK-NEXT: st.b8 [%rd2+4], %r6;173; CHECK-NEXT: st.b8 [%rd2+5], %r5;174; CHECK-NEXT: st.b8 [%rd2+2], %r4;175; CHECK-NEXT: st.b8 [%rd2+3], %r3;176; CHECK-NEXT: st.b8 [%rd2], %r2;177; CHECK-NEXT: st.b8 [%rd2+1], %r1;178; CHECK-NEXT: ret;179 %1 = load <4 x half>, ptr %from , align 1180 store <4 x half> %1, ptr %to , align 1181 ret void182}183 184 185define void @s1(ptr %p1, <4 x float> %v) {186; CHECK-LABEL: s1(187; CHECK: {188; CHECK-NEXT: .reg .b32 %r<17>;189; CHECK-NEXT: .reg .b64 %rd<2>;190; CHECK-EMPTY:191; CHECK-NEXT: // %bb.0:192; CHECK-NEXT: ld.param.b64 %rd1, [s1_param_0];193; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [s1_param_1];194; CHECK-NEXT: st.b8 [%rd1+12], %r4;195; CHECK-NEXT: st.b8 [%rd1+8], %r3;196; CHECK-NEXT: st.b8 [%rd1+4], %r2;197; CHECK-NEXT: st.b8 [%rd1], %r1;198; CHECK-NEXT: shr.u32 %r5, %r4, 24;199; CHECK-NEXT: st.b8 [%rd1+15], %r5;200; CHECK-NEXT: shr.u32 %r6, %r4, 16;201; CHECK-NEXT: st.b8 [%rd1+14], %r6;202; CHECK-NEXT: shr.u32 %r7, %r4, 8;203; CHECK-NEXT: st.b8 [%rd1+13], %r7;204; CHECK-NEXT: shr.u32 %r8, %r3, 24;205; CHECK-NEXT: st.b8 [%rd1+11], %r8;206; CHECK-NEXT: shr.u32 %r9, %r3, 16;207; CHECK-NEXT: st.b8 [%rd1+10], %r9;208; CHECK-NEXT: shr.u32 %r10, %r3, 8;209; CHECK-NEXT: st.b8 [%rd1+9], %r10;210; CHECK-NEXT: shr.u32 %r11, %r2, 24;211; CHECK-NEXT: st.b8 [%rd1+7], %r11;212; CHECK-NEXT: shr.u32 %r12, %r2, 16;213; CHECK-NEXT: st.b8 [%rd1+6], %r12;214; CHECK-NEXT: shr.u32 %r13, %r2, 8;215; CHECK-NEXT: st.b8 [%rd1+5], %r13;216; CHECK-NEXT: shr.u32 %r14, %r1, 24;217; CHECK-NEXT: st.b8 [%rd1+3], %r14;218; CHECK-NEXT: shr.u32 %r15, %r1, 16;219; CHECK-NEXT: st.b8 [%rd1+2], %r15;220; CHECK-NEXT: shr.u32 %r16, %r1, 8;221; CHECK-NEXT: st.b8 [%rd1+1], %r16;222; CHECK-NEXT: ret;223 store <4 x float> %v, ptr %p1, align 1224 ret void225}226 227define void @s2(ptr %p1, <4 x float> %v) {228; CHECK-LABEL: s2(229; CHECK: {230; CHECK-NEXT: .reg .b32 %r<5>;231; CHECK-NEXT: .reg .b64 %rd<2>;232; CHECK-EMPTY:233; CHECK-NEXT: // %bb.0:234; CHECK-NEXT: ld.param.b64 %rd1, [s2_param_0];235; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [s2_param_1];236; CHECK-NEXT: st.b32 [%rd1+12], %r4;237; CHECK-NEXT: st.b32 [%rd1+8], %r3;238; CHECK-NEXT: st.b32 [%rd1+4], %r2;239; CHECK-NEXT: st.b32 [%rd1], %r1;240; CHECK-NEXT: ret;241 store <4 x float> %v, ptr %p1, align 4242 ret void243}244 245define void @s3(ptr %p1, <4 x float> %v) {246; CHECK-LABEL: s3(247; CHECK: {248; CHECK-NEXT: .reg .b32 %r<5>;249; CHECK-NEXT: .reg .b64 %rd<2>;250; CHECK-EMPTY:251; CHECK-NEXT: // %bb.0:252; CHECK-NEXT: ld.param.b64 %rd1, [s3_param_0];253; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [s3_param_1];254; CHECK-NEXT: st.v2.b32 [%rd1+8], {%r3, %r4};255; CHECK-NEXT: st.v2.b32 [%rd1], {%r1, %r2};256; CHECK-NEXT: ret;257 store <4 x float> %v, ptr %p1, align 8258 ret void259}260 261define void @s4(ptr %p1, <4 x float> %v) {262; CHECK-LABEL: s4(263; CHECK: {264; CHECK-NEXT: .reg .b32 %r<5>;265; CHECK-NEXT: .reg .b64 %rd<2>;266; CHECK-EMPTY:267; CHECK-NEXT: // %bb.0:268; CHECK-NEXT: ld.param.b64 %rd1, [s4_param_0];269; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [s4_param_1];270; CHECK-NEXT: st.v4.b32 [%rd1], {%r1, %r2, %r3, %r4};271; CHECK-NEXT: ret;272 store <4 x float> %v, ptr %p1, align 16273 ret void274}275 276