mirror of
https://github.com/c64scene-ar/llvm-6502.git
synced 2024-12-27 13:30:05 +00:00
1d56cda023
Summary: Converting outermost zext(a) to sext(a) causes worse code when the computation of zext(a) could be reused. For example, after converting ... = array[zext(a)] ... = array[zext(a) + 1] to ... = array[sext(a)] ... = array[zext(a) + 1], the program computes sext(a), which is actually unnecessary. I added one test in split-gep-and-gvn.ll to illustrate this scenario. Also, with r211281 and r211084, we annotate more "nuw" tags to computation involving CUDA intrinsics such as threadIdx.x. These annotations help with splitting GEP a lot, rendering the benefit we get from this reverted optimization only marginal. Test Plan: make check-all Reviewers: eliben, meheff Reviewed By: meheff Subscribers: jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D4542 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@213209 91177308-0d34-0410-b5e6-96231b3b80d8
197 lines
8.6 KiB
LLVM
197 lines
8.6 KiB
LLVM
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX
|
|
; RUN: opt < %s -S -separate-const-offset-from-gep -gvn -dce | FileCheck %s --check-prefix=IR
|
|
|
|
; Verifies the SeparateConstOffsetFromGEP pass.
|
|
; The following code computes
|
|
; *output = array[x][y] + array[x][y+1] + array[x+1][y] + array[x+1][y+1]
|
|
;
|
|
; We expect SeparateConstOffsetFromGEP to transform it to
|
|
;
|
|
; float *base = &a[x][y];
|
|
; *output = base[0] + base[1] + base[32] + base[33];
|
|
;
|
|
; so the backend can emit PTX that uses fewer virtual registers.
|
|
|
|
target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
|
|
target triple = "nvptx64-unknown-unknown"
|
|
|
|
@array = internal addrspace(3) constant [32 x [32 x float]] zeroinitializer, align 4
|
|
|
|
define void @sum_of_array(i32 %x, i32 %y, float* nocapture %output) {
|
|
.preheader:
|
|
%0 = sext i32 %y to i64
|
|
%1 = sext i32 %x to i64
|
|
%2 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0
|
|
%3 = addrspacecast float addrspace(3)* %2 to float*
|
|
%4 = load float* %3, align 4
|
|
%5 = fadd float %4, 0.000000e+00
|
|
%6 = add i32 %y, 1
|
|
%7 = sext i32 %6 to i64
|
|
%8 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %7
|
|
%9 = addrspacecast float addrspace(3)* %8 to float*
|
|
%10 = load float* %9, align 4
|
|
%11 = fadd float %5, %10
|
|
%12 = add i32 %x, 1
|
|
%13 = sext i32 %12 to i64
|
|
%14 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %0
|
|
%15 = addrspacecast float addrspace(3)* %14 to float*
|
|
%16 = load float* %15, align 4
|
|
%17 = fadd float %11, %16
|
|
%18 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %7
|
|
%19 = addrspacecast float addrspace(3)* %18 to float*
|
|
%20 = load float* %19, align 4
|
|
%21 = fadd float %17, %20
|
|
store float %21, float* %output, align 4
|
|
ret void
|
|
}
|
|
; PTX-LABEL: sum_of_array(
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rd|r)[0-9]+]]{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}}
|
|
|
|
; IR-LABEL: @sum_of_array(
|
|
; IR: [[BASE_PTR:%[a-zA-Z0-9]+]] = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 1
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 32
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 33
|
|
|
|
; @sum_of_array2 is very similar to @sum_of_array. The only difference is in
|
|
; the order of "sext" and "add" when computing the array indices. @sum_of_array
|
|
; computes add before sext, e.g., array[sext(x + 1)][sext(y + 1)], while
|
|
; @sum_of_array2 computes sext before add,
|
|
; e.g., array[sext(x) + 1][sext(y) + 1]. SeparateConstOffsetFromGEP should be
|
|
; able to extract constant offsets from both forms.
|
|
define void @sum_of_array2(i32 %x, i32 %y, float* nocapture %output) {
|
|
.preheader:
|
|
%0 = sext i32 %y to i64
|
|
%1 = sext i32 %x to i64
|
|
%2 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0
|
|
%3 = addrspacecast float addrspace(3)* %2 to float*
|
|
%4 = load float* %3, align 4
|
|
%5 = fadd float %4, 0.000000e+00
|
|
%6 = add i64 %0, 1
|
|
%7 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %6
|
|
%8 = addrspacecast float addrspace(3)* %7 to float*
|
|
%9 = load float* %8, align 4
|
|
%10 = fadd float %5, %9
|
|
%11 = add i64 %1, 1
|
|
%12 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %11, i64 %0
|
|
%13 = addrspacecast float addrspace(3)* %12 to float*
|
|
%14 = load float* %13, align 4
|
|
%15 = fadd float %10, %14
|
|
%16 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %11, i64 %6
|
|
%17 = addrspacecast float addrspace(3)* %16 to float*
|
|
%18 = load float* %17, align 4
|
|
%19 = fadd float %15, %18
|
|
store float %19, float* %output, align 4
|
|
ret void
|
|
}
|
|
; PTX-LABEL: sum_of_array2(
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rd|r)[0-9]+]]{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}}
|
|
|
|
; IR-LABEL: @sum_of_array2(
|
|
; IR: [[BASE_PTR:%[a-zA-Z0-9]+]] = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 1
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 32
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 33
|
|
|
|
|
|
; This function loads
|
|
; array[zext(x)][zext(y)]
|
|
; array[zext(x)][zext(y +nuw 1)]
|
|
; array[zext(x +nuw 1)][zext(y)]
|
|
; array[zext(x +nuw 1)][zext(y +nuw 1)].
|
|
;
|
|
; This function is similar to @sum_of_array, but it
|
|
; 1) extends array indices using zext instead of sext;
|
|
; 2) annotates the addition with "nuw"; otherwise, zext(x + 1) => zext(x) + 1
|
|
; may be invalid.
|
|
define void @sum_of_array3(i32 %x, i32 %y, float* nocapture %output) {
|
|
.preheader:
|
|
%0 = zext i32 %y to i64
|
|
%1 = zext i32 %x to i64
|
|
%2 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0
|
|
%3 = addrspacecast float addrspace(3)* %2 to float*
|
|
%4 = load float* %3, align 4
|
|
%5 = fadd float %4, 0.000000e+00
|
|
%6 = add nuw i32 %y, 1
|
|
%7 = zext i32 %6 to i64
|
|
%8 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %7
|
|
%9 = addrspacecast float addrspace(3)* %8 to float*
|
|
%10 = load float* %9, align 4
|
|
%11 = fadd float %5, %10
|
|
%12 = add nuw i32 %x, 1
|
|
%13 = zext i32 %12 to i64
|
|
%14 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %0
|
|
%15 = addrspacecast float addrspace(3)* %14 to float*
|
|
%16 = load float* %15, align 4
|
|
%17 = fadd float %11, %16
|
|
%18 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %7
|
|
%19 = addrspacecast float addrspace(3)* %18 to float*
|
|
%20 = load float* %19, align 4
|
|
%21 = fadd float %17, %20
|
|
store float %21, float* %output, align 4
|
|
ret void
|
|
}
|
|
; PTX-LABEL: sum_of_array3(
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rd|r)[0-9]+]]{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}}
|
|
|
|
; IR-LABEL: @sum_of_array3(
|
|
; IR: [[BASE_PTR:%[a-zA-Z0-9]+]] = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 1
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 32
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 33
|
|
|
|
|
|
; This function loads
|
|
; array[zext(x)][zext(y)]
|
|
; array[zext(x)][zext(y)]
|
|
; array[zext(x) + 1][zext(y) + 1]
|
|
; array[zext(x) + 1][zext(y) + 1].
|
|
;
|
|
; We expect the generated code to reuse the computation of
|
|
; &array[zext(x)][zext(y)]. See the expected IR and PTX for details.
|
|
define void @sum_of_array4(i32 %x, i32 %y, float* nocapture %output) {
|
|
.preheader:
|
|
%0 = zext i32 %y to i64
|
|
%1 = zext i32 %x to i64
|
|
%2 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0
|
|
%3 = addrspacecast float addrspace(3)* %2 to float*
|
|
%4 = load float* %3, align 4
|
|
%5 = fadd float %4, 0.000000e+00
|
|
%6 = add i64 %0, 1
|
|
%7 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %6
|
|
%8 = addrspacecast float addrspace(3)* %7 to float*
|
|
%9 = load float* %8, align 4
|
|
%10 = fadd float %5, %9
|
|
%11 = add i64 %1, 1
|
|
%12 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %11, i64 %0
|
|
%13 = addrspacecast float addrspace(3)* %12 to float*
|
|
%14 = load float* %13, align 4
|
|
%15 = fadd float %10, %14
|
|
%16 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %11, i64 %6
|
|
%17 = addrspacecast float addrspace(3)* %16 to float*
|
|
%18 = load float* %17, align 4
|
|
%19 = fadd float %15, %18
|
|
store float %19, float* %output, align 4
|
|
ret void
|
|
}
|
|
; PTX-LABEL: sum_of_array4(
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rd|r)[0-9]+]]{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}}
|
|
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}}
|
|
|
|
; IR-LABEL: @sum_of_array4(
|
|
; IR: [[BASE_PTR:%[a-zA-Z0-9]+]] = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 1
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 32
|
|
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 33
|