325 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 42; RUN: opt < %s -mtriple=nvptx64-nvidia-cuda -S -passes=separate-const-offset-from-gep,gvn \3; RUN: -reassociate-geps-verify-no-dead-code \4; RUN: | FileCheck %s --check-prefix=IR5; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_20 \6; RUN: | FileCheck %s --check-prefix=PTX7 8; Verifies the SeparateConstOffsetFromGEP pass.9; The following code computes10; *output = array[x][y] + array[x][y+1] + array[x+1][y] + array[x+1][y+1]11;12; We expect SeparateConstOffsetFromGEP to transform it to13;14; ptr base = &a[x][y];15; *output = base[0] + base[1] + base[32] + base[33];16;17; so the backend can emit PTX that uses fewer virtual registers.18 19@array = internal addrspace(3) global [32 x [32 x float]] zeroinitializer, align 420 21define void @sum_of_array(i32 %x, i32 %y, ptr nocapture %output) {22; IR-LABEL: define void @sum_of_array(23; IR-SAME: i32 [[X:%.*]], i32 [[Y:%.*]], ptr captures(none) [[OUTPUT:%.*]]) {24; IR-NEXT: .preheader:25; IR-NEXT: [[TMP0:%.*]] = sext i32 [[Y]] to i6426; IR-NEXT: [[TMP1:%.*]] = sext i32 [[X]] to i6427; IR-NEXT: [[TMP2:%.*]] = getelementptr [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 [[TMP1]], i64 [[TMP0]]28; IR-NEXT: [[TMP3:%.*]] = addrspacecast ptr addrspace(3) [[TMP2]] to ptr29; IR-NEXT: [[TMP4:%.*]] = load float, ptr [[TMP3]], align 430; IR-NEXT: [[TMP5:%.*]] = fadd float [[TMP4]], 0.000000e+0031; IR-NEXT: [[TMP6:%.*]] = getelementptr i8, ptr addrspace(3) [[TMP2]], i64 432; IR-NEXT: [[TMP7:%.*]] = addrspacecast ptr addrspace(3) [[TMP6]] to ptr33; IR-NEXT: [[TMP8:%.*]] = load float, ptr [[TMP7]], align 434; IR-NEXT: [[TMP9:%.*]] = fadd float [[TMP5]], [[TMP8]]35; IR-NEXT: [[TMP10:%.*]] = getelementptr i8, ptr addrspace(3) [[TMP2]], i64 12836; IR-NEXT: [[TMP11:%.*]] = addrspacecast ptr addrspace(3) [[TMP10]] to ptr37; IR-NEXT: [[TMP12:%.*]] = load float, ptr [[TMP11]], align 438; IR-NEXT: [[TMP13:%.*]] = fadd float [[TMP9]], [[TMP12]]39; IR-NEXT: [[TMP14:%.*]] = getelementptr i8, ptr addrspace(3) [[TMP2]], i64 13240; IR-NEXT: [[TMP15:%.*]] = addrspacecast ptr addrspace(3) [[TMP14]] to ptr41; IR-NEXT: [[TMP16:%.*]] = load float, ptr [[TMP15]], align 442; IR-NEXT: [[TMP17:%.*]] = fadd float [[TMP13]], [[TMP16]]43; IR-NEXT: store float [[TMP17]], ptr [[OUTPUT]], align 444; IR-NEXT: ret void45;46.preheader:47 %0 = sext i32 %y to i6448 %1 = sext i32 %x to i6449 %2 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %050 %3 = addrspacecast ptr addrspace(3) %2 to ptr51 %4 = load float, ptr %3, align 452 %5 = fadd float %4, 0.000000e+0053 %6 = add i32 %y, 154 %7 = sext i32 %6 to i6455 %8 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %756 %9 = addrspacecast ptr addrspace(3) %8 to ptr57 %10 = load float, ptr %9, align 458 %11 = fadd float %5, %1059 %12 = add i32 %x, 160 %13 = sext i32 %12 to i6461 %14 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %13, i64 %062 %15 = addrspacecast ptr addrspace(3) %14 to ptr63 %16 = load float, ptr %15, align 464 %17 = fadd float %11, %1665 %18 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %13, i64 %766 %19 = addrspacecast ptr addrspace(3) %18 to ptr67 %20 = load float, ptr %19, align 468 %21 = fadd float %17, %2069 store float %21, ptr %output, align 470 ret void71}72; PTX-LABEL: sum_of_array(73; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]74; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+4]75; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+128]76; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+132]77 78; TODO: GVN is unable to preserve the "inbounds" keyword on the first GEP. Need79; some infrastructure changes to enable such optimizations.80 81; @sum_of_array2 is very similar to @sum_of_array. The only difference is in82; the order of "sext" and "add" when computing the array indices. @sum_of_array83; computes add before sext, e.g., array[sext(x + 1)][sext(y + 1)], while84; @sum_of_array2 computes sext before add,85; e.g., array[sext(x) + 1][sext(y) + 1]. SeparateConstOffsetFromGEP should be86; able to extract constant offsets from both forms.87define void @sum_of_array2(i32 %x, i32 %y, ptr nocapture %output) {88; IR-LABEL: define void @sum_of_array2(89; IR-SAME: i32 [[X:%.*]], i32 [[Y:%.*]], ptr captures(none) [[OUTPUT:%.*]]) {90; IR-NEXT: .preheader:91; IR-NEXT: [[TMP0:%.*]] = sext i32 [[Y]] to i6492; IR-NEXT: [[TMP1:%.*]] = sext i32 [[X]] to i6493; IR-NEXT: [[TMP2:%.*]] = getelementptr [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 [[TMP1]], i64 [[TMP0]]94; IR-NEXT: [[TMP3:%.*]] = addrspacecast ptr addrspace(3) [[TMP2]] to ptr95; IR-NEXT: [[TMP4:%.*]] = load float, ptr [[TMP3]], align 496; IR-NEXT: [[TMP5:%.*]] = fadd float [[TMP4]], 0.000000e+0097; IR-NEXT: [[TMP6:%.*]] = getelementptr i8, ptr addrspace(3) [[TMP2]], i64 498; IR-NEXT: [[TMP7:%.*]] = addrspacecast ptr addrspace(3) [[TMP6]] to ptr99; IR-NEXT: [[TMP8:%.*]] = load float, ptr [[TMP7]], align 4100; IR-NEXT: [[TMP9:%.*]] = fadd float [[TMP5]], [[TMP8]]101; IR-NEXT: [[TMP10:%.*]] = getelementptr i8, ptr addrspace(3) [[TMP2]], i64 128102; IR-NEXT: [[TMP11:%.*]] = addrspacecast ptr addrspace(3) [[TMP10]] to ptr103; IR-NEXT: [[TMP12:%.*]] = load float, ptr [[TMP11]], align 4104; IR-NEXT: [[TMP13:%.*]] = fadd float [[TMP9]], [[TMP12]]105; IR-NEXT: [[TMP14:%.*]] = getelementptr i8, ptr addrspace(3) [[TMP2]], i64 132106; IR-NEXT: [[TMP15:%.*]] = addrspacecast ptr addrspace(3) [[TMP14]] to ptr107; IR-NEXT: [[TMP16:%.*]] = load float, ptr [[TMP15]], align 4108; IR-NEXT: [[TMP17:%.*]] = fadd float [[TMP13]], [[TMP16]]109; IR-NEXT: store float [[TMP17]], ptr [[OUTPUT]], align 4110; IR-NEXT: ret void111;112.preheader:113 %0 = sext i32 %y to i64114 %1 = sext i32 %x to i64115 %2 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %0116 %3 = addrspacecast ptr addrspace(3) %2 to ptr117 %4 = load float, ptr %3, align 4118 %5 = fadd float %4, 0.000000e+00119 %6 = add i64 %0, 1120 %7 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %6121 %8 = addrspacecast ptr addrspace(3) %7 to ptr122 %9 = load float, ptr %8, align 4123 %10 = fadd float %5, %9124 %11 = add i64 %1, 1125 %12 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %11, i64 %0126 %13 = addrspacecast ptr addrspace(3) %12 to ptr127 %14 = load float, ptr %13, align 4128 %15 = fadd float %10, %14129 %16 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %11, i64 %6130 %17 = addrspacecast ptr addrspace(3) %16 to ptr131 %18 = load float, ptr %17, align 4132 %19 = fadd float %15, %18133 store float %19, ptr %output, align 4134 ret void135}136; PTX-LABEL: sum_of_array2(137; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]138; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+4]139; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+128]140; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+132]141 142 143 144; This function loads145; array[zext(x)][zext(y)]146; array[zext(x)][zext(y +nuw 1)]147; array[zext(x +nuw 1)][zext(y)]148; array[zext(x +nuw 1)][zext(y +nuw 1)].149;150; This function is similar to @sum_of_array, but it151; 1) extends array indices using zext instead of sext;152; 2) annotates the addition with "nuw"; otherwise, zext(x + 1) => zext(x) + 1153; may be invalid.154define void @sum_of_array3(i32 %x, i32 %y, ptr nocapture %output) {155; IR-LABEL: define void @sum_of_array3(156; IR-SAME: i32 [[X:%.*]], i32 [[Y:%.*]], ptr captures(none) [[OUTPUT:%.*]]) {157; IR-NEXT: .preheader:158; IR-NEXT: [[TMP0:%.*]] = zext i32 [[Y]] to i64159; IR-NEXT: [[TMP1:%.*]] = zext i32 [[X]] to i64160; IR-NEXT: [[TMP2:%.*]] = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 [[TMP1]], i64 [[TMP0]]161; IR-NEXT: [[TMP3:%.*]] = addrspacecast ptr addrspace(3) [[TMP2]] to ptr162; IR-NEXT: [[TMP4:%.*]] = load float, ptr [[TMP3]], align 4163; IR-NEXT: [[TMP5:%.*]] = fadd float [[TMP4]], 0.000000e+00164; IR-NEXT: [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 4165; IR-NEXT: [[TMP7:%.*]] = addrspacecast ptr addrspace(3) [[TMP6]] to ptr166; IR-NEXT: [[TMP8:%.*]] = load float, ptr [[TMP7]], align 4167; IR-NEXT: [[TMP9:%.*]] = fadd float [[TMP5]], [[TMP8]]168; IR-NEXT: [[TMP10:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 128169; IR-NEXT: [[TMP11:%.*]] = addrspacecast ptr addrspace(3) [[TMP10]] to ptr170; IR-NEXT: [[TMP12:%.*]] = load float, ptr [[TMP11]], align 4171; IR-NEXT: [[TMP13:%.*]] = fadd float [[TMP9]], [[TMP12]]172; IR-NEXT: [[TMP14:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 132173; IR-NEXT: [[TMP15:%.*]] = addrspacecast ptr addrspace(3) [[TMP14]] to ptr174; IR-NEXT: [[TMP16:%.*]] = load float, ptr [[TMP15]], align 4175; IR-NEXT: [[TMP17:%.*]] = fadd float [[TMP13]], [[TMP16]]176; IR-NEXT: store float [[TMP17]], ptr [[OUTPUT]], align 4177; IR-NEXT: ret void178;179.preheader:180 %0 = zext i32 %y to i64181 %1 = zext i32 %x to i64182 %2 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %0183 %3 = addrspacecast ptr addrspace(3) %2 to ptr184 %4 = load float, ptr %3, align 4185 %5 = fadd float %4, 0.000000e+00186 %6 = add nuw i32 %y, 1187 %7 = zext i32 %6 to i64188 %8 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %7189 %9 = addrspacecast ptr addrspace(3) %8 to ptr190 %10 = load float, ptr %9, align 4191 %11 = fadd float %5, %10192 %12 = add nuw i32 %x, 1193 %13 = zext i32 %12 to i64194 %14 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %13, i64 %0195 %15 = addrspacecast ptr addrspace(3) %14 to ptr196 %16 = load float, ptr %15, align 4197 %17 = fadd float %11, %16198 %18 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %13, i64 %7199 %19 = addrspacecast ptr addrspace(3) %18 to ptr200 %20 = load float, ptr %19, align 4201 %21 = fadd float %17, %20202 store float %21, ptr %output, align 4203 ret void204}205; PTX-LABEL: sum_of_array3(206; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]207; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+4]208; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+128]209; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+132]210 211 212 213; This function loads214; array[zext(x)][zext(y)]215; array[zext(x)][zext(y)]216; array[zext(x) + 1][zext(y) + 1]217; array[zext(x) + 1][zext(y) + 1].218;219; We expect the generated code to reuse the computation of220; &array[zext(x)][zext(y)]. See the expected IR and PTX for details.221define void @sum_of_array4(i32 %x, i32 %y, ptr nocapture %output) {222; IR-LABEL: define void @sum_of_array4(223; IR-SAME: i32 [[X:%.*]], i32 [[Y:%.*]], ptr captures(none) [[OUTPUT:%.*]]) {224; IR-NEXT: .preheader:225; IR-NEXT: [[TMP0:%.*]] = zext i32 [[Y]] to i64226; IR-NEXT: [[TMP1:%.*]] = zext i32 [[X]] to i64227; IR-NEXT: [[TMP2:%.*]] = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 [[TMP1]], i64 [[TMP0]]228; IR-NEXT: [[TMP3:%.*]] = addrspacecast ptr addrspace(3) [[TMP2]] to ptr229; IR-NEXT: [[TMP4:%.*]] = load float, ptr [[TMP3]], align 4230; IR-NEXT: [[TMP5:%.*]] = fadd float [[TMP4]], 0.000000e+00231; IR-NEXT: [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 4232; IR-NEXT: [[TMP7:%.*]] = addrspacecast ptr addrspace(3) [[TMP6]] to ptr233; IR-NEXT: [[TMP8:%.*]] = load float, ptr [[TMP7]], align 4234; IR-NEXT: [[TMP9:%.*]] = fadd float [[TMP5]], [[TMP8]]235; IR-NEXT: [[TMP10:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 128236; IR-NEXT: [[TMP11:%.*]] = addrspacecast ptr addrspace(3) [[TMP10]] to ptr237; IR-NEXT: [[TMP12:%.*]] = load float, ptr [[TMP11]], align 4238; IR-NEXT: [[TMP13:%.*]] = fadd float [[TMP9]], [[TMP12]]239; IR-NEXT: [[TMP14:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 132240; IR-NEXT: [[TMP15:%.*]] = addrspacecast ptr addrspace(3) [[TMP14]] to ptr241; IR-NEXT: [[TMP16:%.*]] = load float, ptr [[TMP15]], align 4242; IR-NEXT: [[TMP17:%.*]] = fadd float [[TMP13]], [[TMP16]]243; IR-NEXT: store float [[TMP17]], ptr [[OUTPUT]], align 4244; IR-NEXT: ret void245;246.preheader:247 %0 = zext i32 %y to i64248 %1 = zext i32 %x to i64249 %2 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %0250 %3 = addrspacecast ptr addrspace(3) %2 to ptr251 %4 = load float, ptr %3, align 4252 %5 = fadd float %4, 0.000000e+00253 %6 = add i64 %0, 1254 %7 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %6255 %8 = addrspacecast ptr addrspace(3) %7 to ptr256 %9 = load float, ptr %8, align 4257 %10 = fadd float %5, %9258 %11 = add i64 %1, 1259 %12 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %11, i64 %0260 %13 = addrspacecast ptr addrspace(3) %12 to ptr261 %14 = load float, ptr %13, align 4262 %15 = fadd float %10, %14263 %16 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %11, i64 %6264 %17 = addrspacecast ptr addrspace(3) %16 to ptr265 %18 = load float, ptr %17, align 4266 %19 = fadd float %15, %18267 store float %19, ptr %output, align 4268 ret void269}270; PTX-LABEL: sum_of_array4(271; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]272; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+4]273; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+128]274; PTX-DAG: ld.shared.b32 {{%r[0-9]+}}, [[[BASE_REG]]+132]275 276 277 278; The source code is:279; p0 = &input[sext(x + y)];280; p1 = &input[sext(x + (y + 5))];281;282; Without reuniting extensions, SeparateConstOffsetFromGEP would emit283; p0 = &input[sext(x + y)];284; t1 = &input[sext(x) + sext(y)];285; p1 = &t1[5];286;287; With reuniting extensions, it merges p0 and t1 and thus emits288; p0 = &input[sext(x + y)];289; p1 = &p0[5];290define void @reunion(i32 %x, i32 %y, ptr %input) {291; IR-LABEL: define void @reunion(292; IR-SAME: i32 [[X:%.*]], i32 [[Y:%.*]], ptr [[INPUT:%.*]]) {293; IR-NEXT: entry:294; IR-NEXT: [[XY:%.*]] = add nsw i32 [[X]], [[Y]]295; IR-NEXT: [[TMP0:%.*]] = sext i32 [[XY]] to i64296; IR-NEXT: [[P0:%.*]] = getelementptr float, ptr [[INPUT]], i64 [[TMP0]]297; IR-NEXT: [[V0:%.*]] = load float, ptr [[P0]], align 4298; IR-NEXT: call void @use(float [[V0]])299; IR-NEXT: [[P13:%.*]] = getelementptr i8, ptr [[P0]], i64 20300; IR-NEXT: [[V1:%.*]] = load float, ptr [[P13]], align 4301; IR-NEXT: call void @use(float [[V1]])302; IR-NEXT: ret void303;304; PTX-LABEL: reunion(305entry:306 %xy = add nsw i32 %x, %y307 %0 = sext i32 %xy to i64308 %p0 = getelementptr inbounds float, ptr %input, i64 %0309 %v0 = load float, ptr %p0, align 4310; PTX: ld.b32 %r{{[0-9]+}}, [[[p0:%rd[0-9]+]]]311 call void @use(float %v0)312 313 %y5 = add nsw i32 %y, 5314 %xy5 = add nsw i32 %x, %y5315 %1 = sext i32 %xy5 to i64316 %p1 = getelementptr inbounds float, ptr %input, i64 %1317 %v1 = load float, ptr %p1, align 4318; PTX: ld.b32 %r{{[0-9]+}}, [[[p0]]+20]319 call void @use(float %v1)320 321 ret void322}323 324declare void @use(float)325