730 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 -verify-machineinstrs | FileCheck -check-prefix=SM20 %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck -check-prefix=SM35 %s4; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}6 7target 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"8target triple = "nvptx64-unknown-unknown"9 10define ptx_kernel void @foo1(ptr noalias readonly %from, ptr %to) {11; SM20-LABEL: foo1(12; SM20: {13; SM20-NEXT: .reg .b32 %r<2>;14; SM20-NEXT: .reg .b64 %rd<5>;15; SM20-EMPTY:16; SM20-NEXT: // %bb.0:17; SM20-NEXT: ld.param.b64 %rd1, [foo1_param_0];18; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;19; SM20-NEXT: ld.param.b64 %rd3, [foo1_param_1];20; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;21; SM20-NEXT: ld.global.b32 %r1, [%rd2];22; SM20-NEXT: st.global.b32 [%rd4], %r1;23; SM20-NEXT: ret;24;25; SM35-LABEL: foo1(26; SM35: {27; SM35-NEXT: .reg .b32 %r<2>;28; SM35-NEXT: .reg .b64 %rd<5>;29; SM35-EMPTY:30; SM35-NEXT: // %bb.0:31; SM35-NEXT: ld.param.b64 %rd1, [foo1_param_0];32; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;33; SM35-NEXT: ld.param.b64 %rd3, [foo1_param_1];34; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;35; SM35-NEXT: ld.global.nc.b32 %r1, [%rd2];36; SM35-NEXT: st.global.b32 [%rd4], %r1;37; SM35-NEXT: ret;38 %1 = load float, ptr %from39 store float %1, ptr %to40 ret void41}42 43define ptx_kernel void @foo2(ptr noalias readonly %from, ptr %to) {44; SM20-LABEL: foo2(45; SM20: {46; SM20-NEXT: .reg .b64 %rd<6>;47; SM20-EMPTY:48; SM20-NEXT: // %bb.0:49; SM20-NEXT: ld.param.b64 %rd1, [foo2_param_0];50; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;51; SM20-NEXT: ld.param.b64 %rd3, [foo2_param_1];52; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;53; SM20-NEXT: ld.global.b64 %rd5, [%rd2];54; SM20-NEXT: st.global.b64 [%rd4], %rd5;55; SM20-NEXT: ret;56;57; SM35-LABEL: foo2(58; SM35: {59; SM35-NEXT: .reg .b64 %rd<6>;60; SM35-EMPTY:61; SM35-NEXT: // %bb.0:62; SM35-NEXT: ld.param.b64 %rd1, [foo2_param_0];63; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;64; SM35-NEXT: ld.param.b64 %rd3, [foo2_param_1];65; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;66; SM35-NEXT: ld.global.nc.b64 %rd5, [%rd2];67; SM35-NEXT: st.global.b64 [%rd4], %rd5;68; SM35-NEXT: ret;69 %1 = load double, ptr %from70 store double %1, ptr %to71 ret void72}73 74define ptx_kernel void @foo3(ptr noalias readonly %from, ptr %to) {75; SM20-LABEL: foo3(76; SM20: {77; SM20-NEXT: .reg .b16 %rs<2>;78; SM20-NEXT: .reg .b64 %rd<5>;79; SM20-EMPTY:80; SM20-NEXT: // %bb.0:81; SM20-NEXT: ld.param.b64 %rd1, [foo3_param_0];82; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;83; SM20-NEXT: ld.param.b64 %rd3, [foo3_param_1];84; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;85; SM20-NEXT: ld.global.b16 %rs1, [%rd2];86; SM20-NEXT: st.global.b16 [%rd4], %rs1;87; SM20-NEXT: ret;88;89; SM35-LABEL: foo3(90; SM35: {91; SM35-NEXT: .reg .b16 %rs<2>;92; SM35-NEXT: .reg .b64 %rd<5>;93; SM35-EMPTY:94; SM35-NEXT: // %bb.0:95; SM35-NEXT: ld.param.b64 %rd1, [foo3_param_0];96; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;97; SM35-NEXT: ld.param.b64 %rd3, [foo3_param_1];98; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;99; SM35-NEXT: ld.global.nc.b16 %rs1, [%rd2];100; SM35-NEXT: st.global.b16 [%rd4], %rs1;101; SM35-NEXT: ret;102 %1 = load i16, ptr %from103 store i16 %1, ptr %to104 ret void105}106 107define ptx_kernel void @foo4(ptr noalias readonly %from, ptr %to) {108; SM20-LABEL: foo4(109; SM20: {110; SM20-NEXT: .reg .b32 %r<2>;111; SM20-NEXT: .reg .b64 %rd<5>;112; SM20-EMPTY:113; SM20-NEXT: // %bb.0:114; SM20-NEXT: ld.param.b64 %rd1, [foo4_param_0];115; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;116; SM20-NEXT: ld.param.b64 %rd3, [foo4_param_1];117; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;118; SM20-NEXT: ld.global.b32 %r1, [%rd2];119; SM20-NEXT: st.global.b32 [%rd4], %r1;120; SM20-NEXT: ret;121;122; SM35-LABEL: foo4(123; SM35: {124; SM35-NEXT: .reg .b32 %r<2>;125; SM35-NEXT: .reg .b64 %rd<5>;126; SM35-EMPTY:127; SM35-NEXT: // %bb.0:128; SM35-NEXT: ld.param.b64 %rd1, [foo4_param_0];129; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;130; SM35-NEXT: ld.param.b64 %rd3, [foo4_param_1];131; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;132; SM35-NEXT: ld.global.nc.b32 %r1, [%rd2];133; SM35-NEXT: st.global.b32 [%rd4], %r1;134; SM35-NEXT: ret;135 %1 = load i32, ptr %from136 store i32 %1, ptr %to137 ret void138}139 140define ptx_kernel void @foo5(ptr noalias readonly %from, ptr %to) {141; SM20-LABEL: foo5(142; SM20: {143; SM20-NEXT: .reg .b64 %rd<6>;144; SM20-EMPTY:145; SM20-NEXT: // %bb.0:146; SM20-NEXT: ld.param.b64 %rd1, [foo5_param_0];147; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;148; SM20-NEXT: ld.param.b64 %rd3, [foo5_param_1];149; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;150; SM20-NEXT: ld.global.b64 %rd5, [%rd2];151; SM20-NEXT: st.global.b64 [%rd4], %rd5;152; SM20-NEXT: ret;153;154; SM35-LABEL: foo5(155; SM35: {156; SM35-NEXT: .reg .b64 %rd<6>;157; SM35-EMPTY:158; SM35-NEXT: // %bb.0:159; SM35-NEXT: ld.param.b64 %rd1, [foo5_param_0];160; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;161; SM35-NEXT: ld.param.b64 %rd3, [foo5_param_1];162; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;163; SM35-NEXT: ld.global.nc.b64 %rd5, [%rd2];164; SM35-NEXT: st.global.b64 [%rd4], %rd5;165; SM35-NEXT: ret;166 %1 = load i64, ptr %from167 store i64 %1, ptr %to168 ret void169}170 171; i128 is non standard integer in nvptx64172define ptx_kernel void @foo6(ptr noalias readonly %from, ptr %to) {173; SM20-LABEL: foo6(174; SM20: {175; SM20-NEXT: .reg .b64 %rd<7>;176; SM20-EMPTY:177; SM20-NEXT: // %bb.0:178; SM20-NEXT: ld.param.b64 %rd1, [foo6_param_0];179; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;180; SM20-NEXT: ld.param.b64 %rd3, [foo6_param_1];181; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;182; SM20-NEXT: ld.global.v2.b64 {%rd5, %rd6}, [%rd2];183; SM20-NEXT: st.global.v2.b64 [%rd4], {%rd5, %rd6};184; SM20-NEXT: ret;185;186; SM35-LABEL: foo6(187; SM35: {188; SM35-NEXT: .reg .b64 %rd<7>;189; SM35-EMPTY:190; SM35-NEXT: // %bb.0:191; SM35-NEXT: ld.param.b64 %rd1, [foo6_param_0];192; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;193; SM35-NEXT: ld.param.b64 %rd3, [foo6_param_1];194; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;195; SM35-NEXT: ld.global.nc.v2.b64 {%rd5, %rd6}, [%rd2];196; SM35-NEXT: st.global.v2.b64 [%rd4], {%rd5, %rd6};197; SM35-NEXT: ret;198 %1 = load i128, ptr %from199 store i128 %1, ptr %to200 ret void201}202 203define ptx_kernel void @foo7(ptr noalias readonly %from, ptr %to) {204; SM20-LABEL: foo7(205; SM20: {206; SM20-NEXT: .reg .b16 %rs<3>;207; SM20-NEXT: .reg .b64 %rd<5>;208; SM20-EMPTY:209; SM20-NEXT: // %bb.0:210; SM20-NEXT: ld.param.b64 %rd1, [foo7_param_0];211; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;212; SM20-NEXT: ld.param.b64 %rd3, [foo7_param_1];213; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;214; SM20-NEXT: ld.global.v2.b8 {%rs1, %rs2}, [%rd2];215; SM20-NEXT: st.global.v2.b8 [%rd4], {%rs1, %rs2};216; SM20-NEXT: ret;217;218; SM35-LABEL: foo7(219; SM35: {220; SM35-NEXT: .reg .b16 %rs<3>;221; SM35-NEXT: .reg .b64 %rd<5>;222; SM35-EMPTY:223; SM35-NEXT: // %bb.0:224; SM35-NEXT: ld.param.b64 %rd1, [foo7_param_0];225; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;226; SM35-NEXT: ld.param.b64 %rd3, [foo7_param_1];227; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;228; SM35-NEXT: ld.global.nc.v2.b8 {%rs1, %rs2}, [%rd2];229; SM35-NEXT: st.global.v2.b8 [%rd4], {%rs1, %rs2};230; SM35-NEXT: ret;231 %1 = load <2 x i8>, ptr %from232 store <2 x i8> %1, ptr %to233 ret void234}235 236define ptx_kernel void @foo8(ptr noalias readonly %from, ptr %to) {237; SM20-LABEL: foo8(238; SM20: {239; SM20-NEXT: .reg .b32 %r<2>;240; SM20-NEXT: .reg .b64 %rd<5>;241; SM20-EMPTY:242; SM20-NEXT: // %bb.0:243; SM20-NEXT: ld.param.b64 %rd1, [foo8_param_0];244; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;245; SM20-NEXT: ld.param.b64 %rd3, [foo8_param_1];246; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;247; SM20-NEXT: ld.global.b32 %r1, [%rd2];248; SM20-NEXT: st.global.b32 [%rd4], %r1;249; SM20-NEXT: ret;250;251; SM35-LABEL: foo8(252; SM35: {253; SM35-NEXT: .reg .b32 %r<2>;254; SM35-NEXT: .reg .b64 %rd<5>;255; SM35-EMPTY:256; SM35-NEXT: // %bb.0:257; SM35-NEXT: ld.param.b64 %rd1, [foo8_param_0];258; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;259; SM35-NEXT: ld.param.b64 %rd3, [foo8_param_1];260; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;261; SM35-NEXT: ld.global.nc.b32 %r1, [%rd2];262; SM35-NEXT: st.global.b32 [%rd4], %r1;263; SM35-NEXT: ret;264 %1 = load <2 x i16>, ptr %from265 store <2 x i16> %1, ptr %to266 ret void267}268 269define ptx_kernel void @foo9(ptr noalias readonly %from, ptr %to) {270; SM20-LABEL: foo9(271; SM20: {272; SM20-NEXT: .reg .b32 %r<3>;273; SM20-NEXT: .reg .b64 %rd<5>;274; SM20-EMPTY:275; SM20-NEXT: // %bb.0:276; SM20-NEXT: ld.param.b64 %rd1, [foo9_param_0];277; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;278; SM20-NEXT: ld.param.b64 %rd3, [foo9_param_1];279; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;280; SM20-NEXT: ld.global.v2.b32 {%r1, %r2}, [%rd2];281; SM20-NEXT: st.global.v2.b32 [%rd4], {%r1, %r2};282; SM20-NEXT: ret;283;284; SM35-LABEL: foo9(285; SM35: {286; SM35-NEXT: .reg .b32 %r<3>;287; SM35-NEXT: .reg .b64 %rd<5>;288; SM35-EMPTY:289; SM35-NEXT: // %bb.0:290; SM35-NEXT: ld.param.b64 %rd1, [foo9_param_0];291; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;292; SM35-NEXT: ld.param.b64 %rd3, [foo9_param_1];293; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;294; SM35-NEXT: ld.global.nc.v2.b32 {%r1, %r2}, [%rd2];295; SM35-NEXT: st.global.v2.b32 [%rd4], {%r1, %r2};296; SM35-NEXT: ret;297 %1 = load <2 x i32>, ptr %from298 store <2 x i32> %1, ptr %to299 ret void300}301 302define ptx_kernel void @foo10(ptr noalias readonly %from, ptr %to) {303; SM20-LABEL: foo10(304; SM20: {305; SM20-NEXT: .reg .b64 %rd<7>;306; SM20-EMPTY:307; SM20-NEXT: // %bb.0:308; SM20-NEXT: ld.param.b64 %rd1, [foo10_param_0];309; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;310; SM20-NEXT: ld.param.b64 %rd3, [foo10_param_1];311; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;312; SM20-NEXT: ld.global.v2.b64 {%rd5, %rd6}, [%rd2];313; SM20-NEXT: st.global.v2.b64 [%rd4], {%rd5, %rd6};314; SM20-NEXT: ret;315;316; SM35-LABEL: foo10(317; SM35: {318; SM35-NEXT: .reg .b64 %rd<7>;319; SM35-EMPTY:320; SM35-NEXT: // %bb.0:321; SM35-NEXT: ld.param.b64 %rd1, [foo10_param_0];322; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;323; SM35-NEXT: ld.param.b64 %rd3, [foo10_param_1];324; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;325; SM35-NEXT: ld.global.nc.v2.b64 {%rd5, %rd6}, [%rd2];326; SM35-NEXT: st.global.v2.b64 [%rd4], {%rd5, %rd6};327; SM35-NEXT: ret;328 %1 = load <2 x i64>, ptr %from329 store <2 x i64> %1, ptr %to330 ret void331}332 333define ptx_kernel void @foo11(ptr noalias readonly %from, ptr %to) {334; SM20-LABEL: foo11(335; SM20: {336; SM20-NEXT: .reg .b32 %r<3>;337; SM20-NEXT: .reg .b64 %rd<5>;338; SM20-EMPTY:339; SM20-NEXT: // %bb.0:340; SM20-NEXT: ld.param.b64 %rd1, [foo11_param_0];341; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;342; SM20-NEXT: ld.param.b64 %rd3, [foo11_param_1];343; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;344; SM20-NEXT: ld.global.v2.b32 {%r1, %r2}, [%rd2];345; SM20-NEXT: st.global.v2.b32 [%rd4], {%r1, %r2};346; SM20-NEXT: ret;347;348; SM35-LABEL: foo11(349; SM35: {350; SM35-NEXT: .reg .b32 %r<3>;351; SM35-NEXT: .reg .b64 %rd<5>;352; SM35-EMPTY:353; SM35-NEXT: // %bb.0:354; SM35-NEXT: ld.param.b64 %rd1, [foo11_param_0];355; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;356; SM35-NEXT: ld.param.b64 %rd3, [foo11_param_1];357; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;358; SM35-NEXT: ld.global.nc.v2.b32 {%r1, %r2}, [%rd2];359; SM35-NEXT: st.global.v2.b32 [%rd4], {%r1, %r2};360; SM35-NEXT: ret;361 %1 = load <2 x float>, ptr %from362 store <2 x float> %1, ptr %to363 ret void364}365 366define ptx_kernel void @foo12(ptr noalias readonly %from, ptr %to) {367; SM20-LABEL: foo12(368; SM20: {369; SM20-NEXT: .reg .b64 %rd<7>;370; SM20-EMPTY:371; SM20-NEXT: // %bb.0:372; SM20-NEXT: ld.param.b64 %rd1, [foo12_param_0];373; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;374; SM20-NEXT: ld.param.b64 %rd3, [foo12_param_1];375; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;376; SM20-NEXT: ld.global.v2.b64 {%rd5, %rd6}, [%rd2];377; SM20-NEXT: st.global.v2.b64 [%rd4], {%rd5, %rd6};378; SM20-NEXT: ret;379;380; SM35-LABEL: foo12(381; SM35: {382; SM35-NEXT: .reg .b64 %rd<7>;383; SM35-EMPTY:384; SM35-NEXT: // %bb.0:385; SM35-NEXT: ld.param.b64 %rd1, [foo12_param_0];386; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;387; SM35-NEXT: ld.param.b64 %rd3, [foo12_param_1];388; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;389; SM35-NEXT: ld.global.nc.v2.b64 {%rd5, %rd6}, [%rd2];390; SM35-NEXT: st.global.v2.b64 [%rd4], {%rd5, %rd6};391; SM35-NEXT: ret;392 %1 = load <2 x double>, ptr %from393 store <2 x double> %1, ptr %to394 ret void395}396 397define ptx_kernel void @foo13(ptr noalias readonly %from, ptr %to) {398; SM20-LABEL: foo13(399; SM20: {400; SM20-NEXT: .reg .b32 %r<2>;401; SM20-NEXT: .reg .b64 %rd<5>;402; SM20-EMPTY:403; SM20-NEXT: // %bb.0:404; SM20-NEXT: ld.param.b64 %rd1, [foo13_param_0];405; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;406; SM20-NEXT: ld.param.b64 %rd3, [foo13_param_1];407; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;408; SM20-NEXT: ld.global.b32 %r1, [%rd2];409; SM20-NEXT: st.global.b32 [%rd4], %r1;410; SM20-NEXT: ret;411;412; SM35-LABEL: foo13(413; SM35: {414; SM35-NEXT: .reg .b32 %r<2>;415; SM35-NEXT: .reg .b64 %rd<5>;416; SM35-EMPTY:417; SM35-NEXT: // %bb.0:418; SM35-NEXT: ld.param.b64 %rd1, [foo13_param_0];419; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;420; SM35-NEXT: ld.param.b64 %rd3, [foo13_param_1];421; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;422; SM35-NEXT: ld.global.nc.b32 %r1, [%rd2];423; SM35-NEXT: st.global.b32 [%rd4], %r1;424; SM35-NEXT: ret;425 %1 = load <4 x i8>, ptr %from426 store <4 x i8> %1, ptr %to427 ret void428}429 430define ptx_kernel void @foo14(ptr noalias readonly %from, ptr %to) {431; SM20-LABEL: foo14(432; SM20: {433; SM20-NEXT: .reg .b32 %r<3>;434; SM20-NEXT: .reg .b64 %rd<5>;435; SM20-EMPTY:436; SM20-NEXT: // %bb.0:437; SM20-NEXT: ld.param.b64 %rd1, [foo14_param_0];438; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;439; SM20-NEXT: ld.param.b64 %rd3, [foo14_param_1];440; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;441; SM20-NEXT: ld.global.v2.b32 {%r1, %r2}, [%rd2];442; SM20-NEXT: st.global.v2.b32 [%rd4], {%r1, %r2};443; SM20-NEXT: ret;444;445; SM35-LABEL: foo14(446; SM35: {447; SM35-NEXT: .reg .b32 %r<3>;448; SM35-NEXT: .reg .b64 %rd<5>;449; SM35-EMPTY:450; SM35-NEXT: // %bb.0:451; SM35-NEXT: ld.param.b64 %rd1, [foo14_param_0];452; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;453; SM35-NEXT: ld.param.b64 %rd3, [foo14_param_1];454; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;455; SM35-NEXT: ld.global.nc.v2.b32 {%r1, %r2}, [%rd2];456; SM35-NEXT: st.global.v2.b32 [%rd4], {%r1, %r2};457; SM35-NEXT: ret;458 %1 = load <4 x i16>, ptr %from459 store <4 x i16> %1, ptr %to460 ret void461}462 463define ptx_kernel void @foo15(ptr noalias readonly %from, ptr %to) {464; SM20-LABEL: foo15(465; SM20: {466; SM20-NEXT: .reg .b32 %r<5>;467; SM20-NEXT: .reg .b64 %rd<5>;468; SM20-EMPTY:469; SM20-NEXT: // %bb.0:470; SM20-NEXT: ld.param.b64 %rd1, [foo15_param_0];471; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;472; SM20-NEXT: ld.param.b64 %rd3, [foo15_param_1];473; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;474; SM20-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd2];475; SM20-NEXT: st.global.v4.b32 [%rd4], {%r1, %r2, %r3, %r4};476; SM20-NEXT: ret;477;478; SM35-LABEL: foo15(479; SM35: {480; SM35-NEXT: .reg .b32 %r<5>;481; SM35-NEXT: .reg .b64 %rd<5>;482; SM35-EMPTY:483; SM35-NEXT: // %bb.0:484; SM35-NEXT: ld.param.b64 %rd1, [foo15_param_0];485; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;486; SM35-NEXT: ld.param.b64 %rd3, [foo15_param_1];487; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;488; SM35-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd2];489; SM35-NEXT: st.global.v4.b32 [%rd4], {%r1, %r2, %r3, %r4};490; SM35-NEXT: ret;491 %1 = load <4 x i32>, ptr %from492 store <4 x i32> %1, ptr %to493 ret void494}495 496define ptx_kernel void @foo16(ptr noalias readonly %from, ptr %to) {497; SM20-LABEL: foo16(498; SM20: {499; SM20-NEXT: .reg .b32 %r<5>;500; SM20-NEXT: .reg .b64 %rd<5>;501; SM20-EMPTY:502; SM20-NEXT: // %bb.0:503; SM20-NEXT: ld.param.b64 %rd1, [foo16_param_0];504; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;505; SM20-NEXT: ld.param.b64 %rd3, [foo16_param_1];506; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;507; SM20-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd2];508; SM20-NEXT: st.global.v4.b32 [%rd4], {%r1, %r2, %r3, %r4};509; SM20-NEXT: ret;510;511; SM35-LABEL: foo16(512; SM35: {513; SM35-NEXT: .reg .b32 %r<5>;514; SM35-NEXT: .reg .b64 %rd<5>;515; SM35-EMPTY:516; SM35-NEXT: // %bb.0:517; SM35-NEXT: ld.param.b64 %rd1, [foo16_param_0];518; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;519; SM35-NEXT: ld.param.b64 %rd3, [foo16_param_1];520; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;521; SM35-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd2];522; SM35-NEXT: st.global.v4.b32 [%rd4], {%r1, %r2, %r3, %r4};523; SM35-NEXT: ret;524 %1 = load <4 x float>, ptr %from525 store <4 x float> %1, ptr %to526 ret void527}528 529define ptx_kernel void @foo17(ptr noalias readonly %from, ptr %to) {530; SM20-LABEL: foo17(531; SM20: {532; SM20-NEXT: .reg .b64 %rd<9>;533; SM20-EMPTY:534; SM20-NEXT: // %bb.0:535; SM20-NEXT: ld.param.b64 %rd1, [foo17_param_0];536; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;537; SM20-NEXT: ld.param.b64 %rd3, [foo17_param_1];538; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;539; SM20-NEXT: ld.global.v2.b64 {%rd5, %rd6}, [%rd2];540; SM20-NEXT: ld.global.v2.b64 {%rd7, %rd8}, [%rd2+16];541; SM20-NEXT: st.global.v2.b64 [%rd4+16], {%rd7, %rd8};542; SM20-NEXT: st.global.v2.b64 [%rd4], {%rd5, %rd6};543; SM20-NEXT: ret;544;545; SM35-LABEL: foo17(546; SM35: {547; SM35-NEXT: .reg .b64 %rd<9>;548; SM35-EMPTY:549; SM35-NEXT: // %bb.0:550; SM35-NEXT: ld.param.b64 %rd1, [foo17_param_0];551; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;552; SM35-NEXT: ld.param.b64 %rd3, [foo17_param_1];553; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;554; SM35-NEXT: ld.global.nc.v2.b64 {%rd5, %rd6}, [%rd2];555; SM35-NEXT: ld.global.nc.v2.b64 {%rd7, %rd8}, [%rd2+16];556; SM35-NEXT: st.global.v2.b64 [%rd4+16], {%rd7, %rd8};557; SM35-NEXT: st.global.v2.b64 [%rd4], {%rd5, %rd6};558; SM35-NEXT: ret;559 %1 = load <4 x double>, ptr %from560 store <4 x double> %1, ptr %to561 ret void562}563 564define ptx_kernel void @foo18(ptr noalias readonly %from, ptr %to) {565; SM20-LABEL: foo18(566; SM20: {567; SM20-NEXT: .reg .b64 %rd<6>;568; SM20-EMPTY:569; SM20-NEXT: // %bb.0:570; SM20-NEXT: ld.param.b64 %rd1, [foo18_param_0];571; SM20-NEXT: cvta.to.global.u64 %rd2, %rd1;572; SM20-NEXT: ld.param.b64 %rd3, [foo18_param_1];573; SM20-NEXT: cvta.to.global.u64 %rd4, %rd3;574; SM20-NEXT: ld.global.b64 %rd5, [%rd2];575; SM20-NEXT: st.global.b64 [%rd4], %rd5;576; SM20-NEXT: ret;577;578; SM35-LABEL: foo18(579; SM35: {580; SM35-NEXT: .reg .b64 %rd<6>;581; SM35-EMPTY:582; SM35-NEXT: // %bb.0:583; SM35-NEXT: ld.param.b64 %rd1, [foo18_param_0];584; SM35-NEXT: cvta.to.global.u64 %rd2, %rd1;585; SM35-NEXT: ld.param.b64 %rd3, [foo18_param_1];586; SM35-NEXT: cvta.to.global.u64 %rd4, %rd3;587; SM35-NEXT: ld.global.nc.b64 %rd5, [%rd2];588; SM35-NEXT: st.global.b64 [%rd4], %rd5;589; SM35-NEXT: ret;590 %1 = load ptr, ptr %from591 store ptr %1, ptr %to592 ret void593}594 595; Test that we can infer a cached load for a pointer induction variable.596define ptx_kernel void @foo19(ptr noalias readonly %from, ptr %to, i32 %n) {597; SM20-LABEL: foo19(598; SM20: {599; SM20-NEXT: .reg .pred %p<2>;600; SM20-NEXT: .reg .b32 %r<4>;601; SM20-NEXT: .reg .b64 %rd<5>;602; SM20-EMPTY:603; SM20-NEXT: // %bb.0: // %entry604; SM20-NEXT: ld.param.b32 %r2, [foo19_param_2];605; SM20-NEXT: ld.param.b64 %rd2, [foo19_param_0];606; SM20-NEXT: cvta.to.global.u64 %rd4, %rd2;607; SM20-NEXT: ld.param.b64 %rd3, [foo19_param_1];608; SM20-NEXT: cvta.to.global.u64 %rd1, %rd3;609; SM20-NEXT: mov.b32 %r3, 0f00000000;610; SM20-NEXT: $L__BB18_1: // %loop611; SM20-NEXT: // =>This Inner Loop Header: Depth=1612; SM20-NEXT: ld.global.b32 %r1, [%rd4];613; SM20-NEXT: add.rn.f32 %r3, %r1, %r3;614; SM20-NEXT: add.s64 %rd4, %rd4, 4;615; SM20-NEXT: add.s32 %r2, %r2, -1;616; SM20-NEXT: setp.ne.b32 %p1, %r2, 0;617; SM20-NEXT: @%p1 bra $L__BB18_1;618; SM20-NEXT: // %bb.2: // %exit619; SM20-NEXT: st.global.b32 [%rd1], %r3;620; SM20-NEXT: ret;621;622; SM35-LABEL: foo19(623; SM35: {624; SM35-NEXT: .reg .pred %p<2>;625; SM35-NEXT: .reg .b32 %r<4>;626; SM35-NEXT: .reg .b64 %rd<5>;627; SM35-EMPTY:628; SM35-NEXT: // %bb.0: // %entry629; SM35-NEXT: ld.param.b32 %r2, [foo19_param_2];630; SM35-NEXT: ld.param.b64 %rd2, [foo19_param_0];631; SM35-NEXT: cvta.to.global.u64 %rd4, %rd2;632; SM35-NEXT: ld.param.b64 %rd3, [foo19_param_1];633; SM35-NEXT: cvta.to.global.u64 %rd1, %rd3;634; SM35-NEXT: mov.b32 %r3, 0f00000000;635; SM35-NEXT: $L__BB18_1: // %loop636; SM35-NEXT: // =>This Inner Loop Header: Depth=1637; SM35-NEXT: ld.global.nc.b32 %r1, [%rd4];638; SM35-NEXT: add.rn.f32 %r3, %r1, %r3;639; SM35-NEXT: add.s64 %rd4, %rd4, 4;640; SM35-NEXT: add.s32 %r2, %r2, -1;641; SM35-NEXT: setp.ne.b32 %p1, %r2, 0;642; SM35-NEXT: @%p1 bra $L__BB18_1;643; SM35-NEXT: // %bb.2: // %exit644; SM35-NEXT: st.global.b32 [%rd1], %r3;645; SM35-NEXT: ret;646entry:647 br label %loop648 649loop:650 %i = phi i32 [ 0, %entry ], [ %nexti, %loop ]651 %sum = phi float [ 0.0, %entry ], [ %nextsum, %loop ]652 %ptr = getelementptr inbounds float, ptr %from, i32 %i653 %value = load float, ptr %ptr, align 4654 %nextsum = fadd float %value, %sum655 %nexti = add nsw i32 %i, 1656 %exitcond = icmp eq i32 %nexti, %n657 br i1 %exitcond, label %exit, label %loop658 659exit:660 store float %nextsum, ptr %to661 ret void662}663 664; This test captures the case of a non-kernel function. In a665; non-kernel function, without interprocedural analysis, we do not666; know that the parameter is global. We also do not know that the667; pointed-to memory is never written to (for the duration of the668; kernel). For both reasons, we cannot use a cached load here.669define void @notkernel(ptr noalias readonly %from, ptr %to) {670; SM20-LABEL: notkernel(671; SM20: {672; SM20-NEXT: .reg .b32 %r<2>;673; SM20-NEXT: .reg .b64 %rd<3>;674; SM20-EMPTY:675; SM20-NEXT: // %bb.0:676; SM20-NEXT: ld.param.b64 %rd1, [notkernel_param_0];677; SM20-NEXT: ld.b32 %r1, [%rd1];678; SM20-NEXT: ld.param.b64 %rd2, [notkernel_param_1];679; SM20-NEXT: st.b32 [%rd2], %r1;680; SM20-NEXT: ret;681;682; SM35-LABEL: notkernel(683; SM35: {684; SM35-NEXT: .reg .b32 %r<2>;685; SM35-NEXT: .reg .b64 %rd<3>;686; SM35-EMPTY:687; SM35-NEXT: // %bb.0:688; SM35-NEXT: ld.param.b64 %rd1, [notkernel_param_0];689; SM35-NEXT: ld.b32 %r1, [%rd1];690; SM35-NEXT: ld.param.b64 %rd2, [notkernel_param_1];691; SM35-NEXT: st.b32 [%rd2], %r1;692; SM35-NEXT: ret;693 %1 = load float, ptr %from694 store float %1, ptr %to695 ret void696}697 698; As @notkernel, but with the parameter explicitly marked as global. We still699; do not know that the parameter is never written to (for the duration of the700; kernel). This case does not currently come up normally since we do not infer701; that pointers are global interprocedurally as of 2015-08-05.702define void @notkernel2(ptr addrspace(1) noalias readonly %from, ptr %to) {703; SM20-LABEL: notkernel2(704; SM20: {705; SM20-NEXT: .reg .b32 %r<2>;706; SM20-NEXT: .reg .b64 %rd<3>;707; SM20-EMPTY:708; SM20-NEXT: // %bb.0:709; SM20-NEXT: ld.param.b64 %rd1, [notkernel2_param_0];710; SM20-NEXT: ld.global.b32 %r1, [%rd1];711; SM20-NEXT: ld.param.b64 %rd2, [notkernel2_param_1];712; SM20-NEXT: st.b32 [%rd2], %r1;713; SM20-NEXT: ret;714;715; SM35-LABEL: notkernel2(716; SM35: {717; SM35-NEXT: .reg .b32 %r<2>;718; SM35-NEXT: .reg .b64 %rd<3>;719; SM35-EMPTY:720; SM35-NEXT: // %bb.0:721; SM35-NEXT: ld.param.b64 %rd1, [notkernel2_param_0];722; SM35-NEXT: ld.global.b32 %r1, [%rd1];723; SM35-NEXT: ld.param.b64 %rd2, [notkernel2_param_1];724; SM35-NEXT: st.b32 [%rd2], %r1;725; SM35-NEXT: ret;726 %1 = load float, ptr addrspace(1) %from727 store float %1, ptr %to728 ret void729}730