204 lines · plain
1; RUN: llc -O0 < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck %s -check-prefixes=ALL,CLS322; RUN: llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s -check-prefixes=ALL,NOPTRCONV,CLS643; RUN: llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s -check-prefixes=ALL,PTRCONV,CLS644; RUN: %if ptxas-ptr32 %{ llc -O0 < %s -mtriple=nvptx -mcpu=sm_20 | %ptxas-verify %}5; RUN: %if ptxas %{ llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}6; RUN: %if ptxas %{ llc -O0 < %s -mtriple=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | %ptxas-verify %}7 8; ALL-LABEL: conv19define i32 @conv1(ptr addrspace(1) %ptr) {10; CLS32: cvta.global.u3211; ALL-NOT: cvt.u64.u3212; CLS64: cvta.global.u6413; ALL: ld.b3214 %genptr = addrspacecast ptr addrspace(1) %ptr to ptr15 %val = load i32, ptr %genptr16 ret i32 %val17}18 19; ALL-LABEL: conv220define i32 @conv2(ptr addrspace(3) %ptr) {21; CLS32: cvta.shared.u3222; PTRCONV: cvt.u64.u3223; NOPTRCONV-NOT: cvt.u64.u3224; CLS64: cvta.shared.u6425; ALL: ld.b3226 %genptr = addrspacecast ptr addrspace(3) %ptr to ptr27 %val = load i32, ptr %genptr28 ret i32 %val29}30 31; ALL-LABEL: conv332define i32 @conv3(ptr addrspace(4) %ptr) {33; CLS32: cvta.const.u3234; PTRCONV: cvt.u64.u3235; NOPTRCONV-NOT: cvt.u64.u3236; CLS64: cvta.const.u6437; ALL: ld.b3238 %genptr = addrspacecast ptr addrspace(4) %ptr to ptr39 %val = load i32, ptr %genptr40 ret i32 %val41}42 43; ALL-LABEL: conv444define i32 @conv4(ptr addrspace(5) %ptr) {45; CLS32: cvta.local.u3246; PTRCONV: cvt.u64.u3247; NOPTRCONV-NOT: cvt.u64.u3248; CLS64: cvta.local.u6449; ALL: ld.b3250 %genptr = addrspacecast ptr addrspace(5) %ptr to ptr51 %val = load i32, ptr %genptr52 ret i32 %val53}54 55; ALL-LABEL: conv556define i32 @conv5(ptr %ptr) {57; CLS32: cvta.to.global.u3258; ALL-NOT: cvt.u64.u3259; CLS64: cvta.to.global.u6460; ALL: ld.global.b3261 %specptr = addrspacecast ptr %ptr to ptr addrspace(1)62 %val = load i32, ptr addrspace(1) %specptr63 ret i32 %val64}65 66; ALL-LABEL: conv667define i32 @conv6(ptr %ptr) {68; CLS32: cvta.to.shared.u3269; CLS64: cvta.to.shared.u6470; PTRCONV: cvt.u32.u6471; NOPTRCONV-NOT: cvt.u32.u6472; ALL: ld.shared.b3273 %specptr = addrspacecast ptr %ptr to ptr addrspace(3)74 %val = load i32, ptr addrspace(3) %specptr75 ret i32 %val76}77 78; ALL-LABEL: conv779define i32 @conv7(ptr %ptr) {80; CLS32: cvta.to.const.u3281; CLS64: cvta.to.const.u6482; PTRCONV: cvt.u32.u6483; NOPTRCONV-NOT: cvt.u32.u6484; ALL: ld.const.b3285 %specptr = addrspacecast ptr %ptr to ptr addrspace(4)86 %val = load i32, ptr addrspace(4) %specptr87 ret i32 %val88}89 90; ALL-LABEL: conv891define i32 @conv8(ptr %ptr) {92; CLS32: cvta.to.local.u3293; CLS64: cvta.to.local.u6494; PTRCONV: cvt.u32.u6495; NOPTRCONV-NOT: cvt.u32.u6496; ALL: ld.local.b3297 %specptr = addrspacecast ptr %ptr to ptr addrspace(5)98 %val = load i32, ptr addrspace(5) %specptr99 ret i32 %val100}101 102; ALL-LABEL: conv9103define i32 @conv9(ptr addrspace(1) %ptr) {104; CLS32: // implicit-def: %[[ADDR:r[0-9]+]]105; PTRCONV: // implicit-def: %[[ADDR:r[0-9]+]]106; NOPTRCONV: // implicit-def: %[[ADDR:rd[0-9]+]]107; ALL: ld.shared.b32 %r{{[0-9]+}}, [%[[ADDR]]]108 %specptr = addrspacecast ptr addrspace(1) %ptr to ptr addrspace(3)109 %val = load i32, ptr addrspace(3) %specptr110 ret i32 %val111}112 113; Check that we support addrspacecast when splitting the vector114; result (<2 x ptr> => 2 x <1 x ptr>).115; This also checks that scalarization works for addrspacecast116; (when going from <1 x ptr> to ptr.)117; ALL-LABEL: split1To0118define void @split1To0(ptr nocapture noundef readonly %xs) {119; CLS32: cvta.global.u32120; CLS32: cvta.global.u32121; CLS64: cvta.global.u64122; CLS64: cvta.global.u64123; ALL: st.b32124; ALL: st.b32125 %vec_addr = load <2 x ptr addrspace(1)>, ptr %xs, align 16126 %addrspacecast = addrspacecast <2 x ptr addrspace(1)> %vec_addr to <2 x ptr>127 %extractelement0 = extractelement <2 x ptr> %addrspacecast, i64 0128 store float 0.5, ptr %extractelement0, align 4129 %extractelement1 = extractelement <2 x ptr> %addrspacecast, i64 1130 store float 1.0, ptr %extractelement1, align 4131 ret void132}133 134; Same as split1To0 but from 0 to 1, to make sure the addrspacecast preserve135; the source and destination addrspaces properly.136; ALL-LABEL: split0To1137define void @split0To1(ptr nocapture noundef readonly %xs) {138; CLS32: cvta.to.global.u32139; CLS32: cvta.to.global.u32140; CLS64: cvta.to.global.u64141; CLS64: cvta.to.global.u64142; ALL: st.global.b32143; ALL: st.global.b32144 %vec_addr = load <2 x ptr>, ptr %xs, align 16145 %addrspacecast = addrspacecast <2 x ptr> %vec_addr to <2 x ptr addrspace(1)>146 %extractelement0 = extractelement <2 x ptr addrspace(1)> %addrspacecast, i64 0147 store float 0.5, ptr addrspace(1) %extractelement0, align 4148 %extractelement1 = extractelement <2 x ptr addrspace(1)> %addrspacecast, i64 1149 store float 1.0, ptr addrspace(1) %extractelement1, align 4150 ret void151}152 153; Check that we support addrspacecast when a widening is required154; (3 x ptr => 4 x ptr).155; ALL-LABEL: widen1To0156define void @widen1To0(ptr nocapture noundef readonly %xs) {157; CLS32: cvta.global.u32158; CLS32: cvta.global.u32159; CLS32: cvta.global.u32160 161; CLS64: cvta.global.u64162; CLS64: cvta.global.u64163; CLS64: cvta.global.u64164 165; ALL: st.b32166; ALL: st.b32167; ALL: st.b32168 %vec_addr = load <3 x ptr addrspace(1)>, ptr %xs, align 16169 %addrspacecast = addrspacecast <3 x ptr addrspace(1)> %vec_addr to <3 x ptr>170 %extractelement0 = extractelement <3 x ptr> %addrspacecast, i64 0171 store float 0.5, ptr %extractelement0, align 4172 %extractelement1 = extractelement <3 x ptr> %addrspacecast, i64 1173 store float 1.0, ptr %extractelement1, align 4174 %extractelement2 = extractelement <3 x ptr> %addrspacecast, i64 2175 store float 1.5, ptr %extractelement2, align 4176 ret void177}178 179; Same as widen1To0 but from 0 to 1, to make sure the addrspacecast preserve180; the source and destination addrspaces properly.181; ALL-LABEL: widen0To1182define void @widen0To1(ptr nocapture noundef readonly %xs) {183; CLS32: cvta.to.global.u32184; CLS32: cvta.to.global.u32185; CLS32: cvta.to.global.u32186 187; CLS64: cvta.to.global.u64188; CLS64: cvta.to.global.u64189; CLS64: cvta.to.global.u64190 191; ALL: st.global.b32192; ALL: st.global.b32193; ALL: st.global.b32194 %vec_addr = load <3 x ptr>, ptr %xs, align 16195 %addrspacecast = addrspacecast <3 x ptr> %vec_addr to <3 x ptr addrspace(1)>196 %extractelement0 = extractelement <3 x ptr addrspace(1)> %addrspacecast, i64 0197 store float 0.5, ptr addrspace(1) %extractelement0, align 4198 %extractelement1 = extractelement <3 x ptr addrspace(1)> %addrspacecast, i64 1199 store float 1.0, ptr addrspace(1) %extractelement1, align 4200 %extractelement2 = extractelement <3 x ptr addrspace(1)> %addrspacecast, i64 2201 store float 1.5, ptr addrspace(1) %extractelement2, align 4202 ret void203}204