252 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -disable-post-ra -frame-pointer=all \3; RUN: -verify-machineinstrs -O0 | FileCheck %s --check-prefixes=O0,COMMON4; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -disable-post-ra -frame-pointer=all \5; RUN: -verify-machineinstrs | FileCheck %s --check-prefixes=O3,COMMON6; RUN: %if ptxas %{ \7; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -disable-post-ra -frame-pointer=all \8; RUN: -verify-machineinstrs -O0 \9; RUN: | %ptxas-verify -arch=sm_90 \10; RUN: %}11; RUN: %if ptxas %{ \12; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -disable-post-ra -frame-pointer=all \13; RUN: -verify-machineinstrs \14; RUN: | %ptxas-verify -arch=sm_90 \15; RUN: %}16 17target triple = "nvptx64-nvidia-cuda"18target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"19 20define i16 @test_bitcast_2xi8_i16(<2 x i8> %a) {21; O0-LABEL: test_bitcast_2xi8_i16(22; O0: {23; O0-NEXT: .reg .b16 %rs<5>;24; O0-NEXT: .reg .b32 %r<3>;25; O0-EMPTY:26; O0-NEXT: // %bb.0:27; O0-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [test_bitcast_2xi8_i16_param_0];28; O0-NEXT: mov.b32 %r1, {%rs1, %rs2};29; O0-NEXT: shl.b16 %rs3, %rs2, 8;30; O0-NEXT: or.b16 %rs4, %rs1, %rs3;31; O0-NEXT: cvt.u32.u16 %r2, %rs4;32; O0-NEXT: st.param.b32 [func_retval0], %r2;33; O0-NEXT: ret;34;35; O3-LABEL: test_bitcast_2xi8_i16(36; O3: {37; O3-NEXT: .reg .b32 %r<2>;38; O3-EMPTY:39; O3-NEXT: // %bb.0:40; O3-NEXT: ld.param.b16 %r1, [test_bitcast_2xi8_i16_param_0];41; O3-NEXT: st.param.b32 [func_retval0], %r1;42; O3-NEXT: ret;43 %res = bitcast <2 x i8> %a to i1644 ret i16 %res45}46 47define <2 x i8> @test_bitcast_i16_2xi8(i16 %a) {48; O0-LABEL: test_bitcast_i16_2xi8(49; O0: {50; O0-NEXT: .reg .b16 %rs<2>;51; O0-EMPTY:52; O0-NEXT: // %bb.0:53; O0-NEXT: ld.param.b16 %rs1, [test_bitcast_i16_2xi8_param_0];54; O0-NEXT: st.param.b16 [func_retval0], %rs1;55; O0-NEXT: ret;56;57; O3-LABEL: test_bitcast_i16_2xi8(58; O3: {59; O3-NEXT: .reg .b16 %rs<2>;60; O3-EMPTY:61; O3-NEXT: // %bb.0:62; O3-NEXT: ld.param.b16 %rs1, [test_bitcast_i16_2xi8_param_0];63; O3-NEXT: st.param.b16 [func_retval0], %rs1;64; O3-NEXT: ret;65 %res = bitcast i16 %a to <2 x i8>66 ret <2 x i8> %res67}68 69define <2 x i8> @test_call_2xi8(<2 x i8> %a) {70; O0-LABEL: test_call_2xi8(71; O0: {72; O0-NEXT: .reg .b16 %rs<5>;73; O0-NEXT: .reg .b32 %r<2>;74; O0-EMPTY:75; O0-NEXT: // %bb.0:76; O0-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [test_call_2xi8_param_0];77; O0-NEXT: mov.b32 %r1, {%rs1, %rs2};78; O0-NEXT: { // callseq 0, 079; O0-NEXT: .param .align 2 .b8 param0[2];80; O0-NEXT: .param .align 2 .b8 retval0[2];81; O0-NEXT: st.param.v2.b8 [param0], {%rs1, %rs2};82; O0-NEXT: call.uni (retval0), test_call_2xi8, (param0);83; O0-NEXT: ld.param.v2.b8 {%rs3, %rs4}, [retval0];84; O0-NEXT: } // callseq 085; O0-NEXT: st.param.v2.b8 [func_retval0], {%rs3, %rs4};86; O0-NEXT: ret;87;88; O3-LABEL: test_call_2xi8(89; O3: {90; O3-NEXT: .reg .b16 %rs<5>;91; O3-EMPTY:92; O3-NEXT: // %bb.0:93; O3-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [test_call_2xi8_param_0];94; O3-NEXT: { // callseq 0, 095; O3-NEXT: .param .align 2 .b8 param0[2];96; O3-NEXT: .param .align 2 .b8 retval0[2];97; O3-NEXT: st.param.v2.b8 [param0], {%rs1, %rs2};98; O3-NEXT: call.uni (retval0), test_call_2xi8, (param0);99; O3-NEXT: ld.param.v2.b8 {%rs3, %rs4}, [retval0];100; O3-NEXT: } // callseq 0101; O3-NEXT: st.param.v2.b8 [func_retval0], {%rs3, %rs4};102; O3-NEXT: ret;103 %res = call <2 x i8> @test_call_2xi8(<2 x i8> %a)104 ret <2 x i8> %res105}106 107define <2 x float> @test_uitofp_2xi8(<2 x i8> %a) {108; O0-LABEL: test_uitofp_2xi8(109; O0: {110; O0-NEXT: .reg .b16 %rs<3>;111; O0-NEXT: .reg .b32 %r<4>;112; O0-EMPTY:113; O0-NEXT: // %bb.0:114; O0-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [test_uitofp_2xi8_param_0];115; O0-NEXT: mov.b32 %r1, {%rs1, %rs2};116; O0-NEXT: cvt.rn.f32.u16 %r2, %rs2;117; O0-NEXT: cvt.rn.f32.u16 %r3, %rs1;118; O0-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};119; O0-NEXT: ret;120;121; O3-LABEL: test_uitofp_2xi8(122; O3: {123; O3-NEXT: .reg .b16 %rs<3>;124; O3-NEXT: .reg .b32 %r<3>;125; O3-EMPTY:126; O3-NEXT: // %bb.0:127; O3-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [test_uitofp_2xi8_param_0];128; O3-NEXT: cvt.rn.f32.u16 %r1, %rs2;129; O3-NEXT: cvt.rn.f32.u16 %r2, %rs1;130; O3-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};131; O3-NEXT: ret;132 %1 = uitofp <2 x i8> %a to <2 x float>133 ret <2 x float> %1134}135 136define void @test_store_i8x2_unaligned(ptr %ptr, <2 x i8> %a) {137; O0-LABEL: test_store_i8x2_unaligned(138; O0: {139; O0-NEXT: .reg .b16 %rs<3>;140; O0-NEXT: .reg .b32 %r<2>;141; O0-NEXT: .reg .b64 %rd<2>;142; O0-EMPTY:143; O0-NEXT: // %bb.0:144; O0-NEXT: ld.param.b64 %rd1, [test_store_i8x2_unaligned_param_0];145; O0-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [test_store_i8x2_unaligned_param_1];146; O0-NEXT: mov.b32 %r1, {%rs1, %rs2};147; O0-NEXT: st.b8 [%rd1+1], %rs2;148; O0-NEXT: st.b8 [%rd1], %rs1;149; O0-NEXT: ret;150;151; O3-LABEL: test_store_i8x2_unaligned(152; O3: {153; O3-NEXT: .reg .b16 %rs<3>;154; O3-NEXT: .reg .b64 %rd<2>;155; O3-EMPTY:156; O3-NEXT: // %bb.0:157; O3-NEXT: ld.param.b64 %rd1, [test_store_i8x2_unaligned_param_0];158; O3-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [test_store_i8x2_unaligned_param_1];159; O3-NEXT: st.b8 [%rd1+1], %rs2;160; O3-NEXT: st.b8 [%rd1], %rs1;161; O3-NEXT: ret;162 store <2 x i8> %a, ptr %ptr, align 1163 ret void164}165 166define void @test_store_i8x2_unaligned_immediate(ptr %ptr) {167; O0-LABEL: test_store_i8x2_unaligned_immediate(168; O0: {169; O0-NEXT: .reg .b64 %rd<2>;170; O0-EMPTY:171; O0-NEXT: // %bb.0:172; O0-NEXT: ld.param.b64 %rd1, [test_store_i8x2_unaligned_immediate_param_0];173; O0-NEXT: st.b8 [%rd1+1], 2;174; O0-NEXT: st.b8 [%rd1], 1;175; O0-NEXT: ret;176;177; O3-LABEL: test_store_i8x2_unaligned_immediate(178; O3: {179; O3-NEXT: .reg .b64 %rd<2>;180; O3-EMPTY:181; O3-NEXT: // %bb.0:182; O3-NEXT: ld.param.b64 %rd1, [test_store_i8x2_unaligned_immediate_param_0];183; O3-NEXT: st.b8 [%rd1+1], 2;184; O3-NEXT: st.b8 [%rd1], 1;185; O3-NEXT: ret;186 store <2 x i8> <i8 1, i8 2>, ptr %ptr, align 1187 ret void188}189 190define i32 @test_zext_load_i8x2_unaligned(ptr %ptr) {191; O0-LABEL: test_zext_load_i8x2_unaligned(192; O0: {193; O0-NEXT: .reg .b16 %rs<3>;194; O0-NEXT: .reg .b64 %rd<2>;195; O0-EMPTY:196; O0-NEXT: // %bb.0:197; O0-NEXT: ld.param.b64 %rd1, [test_zext_load_i8x2_unaligned_param_0];198; O0-NEXT: ld.b8 %rs1, [%rd1+1];199; O0-NEXT: ld.b8 %rs2, [%rd1];200; O0-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};201; O0-NEXT: ret;202;203; O3-LABEL: test_zext_load_i8x2_unaligned(204; O3: {205; O3-NEXT: .reg .b16 %rs<3>;206; O3-NEXT: .reg .b64 %rd<2>;207; O3-EMPTY:208; O3-NEXT: // %bb.0:209; O3-NEXT: ld.param.b64 %rd1, [test_zext_load_i8x2_unaligned_param_0];210; O3-NEXT: ld.b8 %rs1, [%rd1+1];211; O3-NEXT: ld.b8 %rs2, [%rd1];212; O3-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};213; O3-NEXT: ret;214 %a = load <2 x i8>, ptr %ptr, align 1215 %b = zext <2 x i8> %a to <2 x i16>216 %c = bitcast <2 x i16> %b to i32217 ret i32 %c218}219 220define i32 @test_sext_load_i8x2_unaligned(ptr %ptr) {221; O0-LABEL: test_sext_load_i8x2_unaligned(222; O0: {223; O0-NEXT: .reg .b16 %rs<3>;224; O0-NEXT: .reg .b64 %rd<2>;225; O0-EMPTY:226; O0-NEXT: // %bb.0:227; O0-NEXT: ld.param.b64 %rd1, [test_sext_load_i8x2_unaligned_param_0];228; O0-NEXT: ld.s8 %rs1, [%rd1+1];229; O0-NEXT: ld.s8 %rs2, [%rd1];230; O0-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};231; O0-NEXT: ret;232;233; O3-LABEL: test_sext_load_i8x2_unaligned(234; O3: {235; O3-NEXT: .reg .b16 %rs<3>;236; O3-NEXT: .reg .b64 %rd<2>;237; O3-EMPTY:238; O3-NEXT: // %bb.0:239; O3-NEXT: ld.param.b64 %rd1, [test_sext_load_i8x2_unaligned_param_0];240; O3-NEXT: ld.s8 %rs1, [%rd1+1];241; O3-NEXT: ld.s8 %rs2, [%rd1];242; O3-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};243; O3-NEXT: ret;244 %a = load <2 x i8>, ptr %ptr, align 1245 %b = sext <2 x i8> %a to <2 x i16>246 %c = bitcast <2 x i16> %b to i32247 ret i32 %c248}249 250;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:251; COMMON: {{.*}}252