maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/polly/test/GPGPU/double-parallel-loop.ll

114 lines
4.4 KiB
LLVM
Raw Normal View History

GPGPU: create default initialized PPCG scop and gpu program At this stage, we do not yet modify the IR but just generate a default initialized ppcg_scop and gpu_prog and free both immediately. Both will later be filled with data from the polly::Scop and are needed to use PPCG for GPU schedule generation. This commit does not yet perform any GPU code generation, but ensures that the basic infrastructure has been put in place. We also add a simple test case to ensure the new code is run and use this opportunity to verify that GPU_CODEGEN tests are only run if GPU code generation has been enabled in cmake. llvm-svn: 275389
2016-07-14 18:22:19 +08:00
; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
GPGPU: compute new schedule from polly scop To do so we copy the necessary information to compute an initial schedule from polly::Scop to ppcg's scop. Most of the necessary information is directly available and only needs to be passed on to ppcg, with the exception of 'tagged' access relations, access relations that additionally carry information about which memory access an access relation originates from. We could possibly perform the construction of tagged accesses as part of ScopInfo, but as this format is currently specific to ppcg we do not do this yet, but keep this functionality local to our GPU code generation. After the scop has been initialized, we compute data dependences and ask ppcg to compute an initial schedule. Some of this functionality is already available in polly::DependenceInfo and polly::ScheduleOptimizer, but to keep differences to ppcg small we use ppcg's functionality here. We may later investiage if a closer integration of these tools makes sense. llvm-svn: 275390
2016-07-14 18:22:25 +08:00
; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-schedule \
; RUN: -disable-output < %s | \
; RUN: FileCheck -check-prefix=SCHED %s
GPGPU: Generate an AST for the GPU-mapped schedule For this we need to provide an explicit list of statements as they occur in the polly::Scop to ppcg. We also setup basic AST printing facilities to facilitate debugging. To allow code reuse some (minor) changes in ppcg are have been necessary. llvm-svn: 275436
2016-07-14 23:51:37 +08:00
; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \
; RUN: -disable-output < %s | \
; RUN: FileCheck -check-prefix=CODE %s
GPGPU: create default initialized PPCG scop and gpu program At this stage, we do not yet modify the IR but just generate a default initialized ppcg_scop and gpu_prog and free both immediately. Both will later be filled with data from the polly::Scop and are needed to use PPCG for GPU schedule generation. This commit does not yet perform any GPU code generation, but ensures that the basic infrastructure has been put in place. We also add a simple test case to ensure the new code is run and use this opportunity to verify that GPU_CODEGEN tests are only run if GPU code generation has been enabled in cmake. llvm-svn: 275389
2016-07-14 18:22:19 +08:00
; REQUIRES: pollyacc
; CHECK: Stmt_bb5
GPGPU: Use CHECK-NEXT to harden test cases A sequence of CHECK lines allows additional statements to appear in the output of the tested program without any test failures appearing. As we do not want this to happen, switch this test case to use CHECK-NEXT. llvm-svn: 275534
2016-07-15 15:05:49 +08:00
; CHECK-NEXT: Domain :=
; CHECK-NEXT: { Stmt_bb5[i0, i1] : 0 <= i0 <= 1023 and 0 <= i1 <= 1023 };
; CHECK-NEXT: Schedule :=
; CHECK-NEXT: { Stmt_bb5[i0, i1] -> [i0, i1] };
; CHECK-NEXT: ReadAccess := [Reduction Type: NONE] [Scalar: 0]
; CHECK-NEXT: { Stmt_bb5[i0, i1] -> MemRef_A[i0, i1] };
; CHECK-NEXT: MustWriteAccess := [Reduction Type: NONE] [Scalar: 0]
; CHECK-NEXT: { Stmt_bb5[i0, i1] -> MemRef_A[i0, i1] };
GPGPU: compute new schedule from polly scop To do so we copy the necessary information to compute an initial schedule from polly::Scop to ppcg's scop. Most of the necessary information is directly available and only needs to be passed on to ppcg, with the exception of 'tagged' access relations, access relations that additionally carry information about which memory access an access relation originates from. We could possibly perform the construction of tagged accesses as part of ScopInfo, but as this format is currently specific to ppcg we do not do this yet, but keep this functionality local to our GPU code generation. After the scop has been initialized, we compute data dependences and ask ppcg to compute an initial schedule. Some of this functionality is already available in polly::DependenceInfo and polly::ScheduleOptimizer, but to keep differences to ppcg small we use ppcg's functionality here. We may later investiage if a closer integration of these tools makes sense. llvm-svn: 275390
2016-07-14 18:22:25 +08:00
; SCHED: domain: "{ Stmt_bb5[i0, i1] : 0 <= i0 <= 1023 and 0 <= i1 <= 1023 }"
GPGPU: Use CHECK-NEXT to harden test cases A sequence of CHECK lines allows additional statements to appear in the output of the tested program without any test failures appearing. As we do not want this to happen, switch this test case to use CHECK-NEXT. llvm-svn: 275534
2016-07-15 15:05:49 +08:00
; SCHED-NEXT: child:
; SCHED-NEXT: context: "{ [] }"
; SCHED-NEXT: child:
; SCHED-NEXT: extension: "{ }"
; SCHED-NEXT: child:
; SCHED-NEXT: sequence:
; SCHED-NEXT: - filter: "{ }"
; SCHED-NEXT: - filter: "{ Stmt_bb5[i0, i1] }"
; SCHED-NEXT: child:
; SCHED-NEXT: guard: "{ [] }"
; SCHED-NEXT: child:
; SCHED-NEXT: mark: "kernel"
; SCHED-NEXT: child:
; SCHED-NEXT: context: "[b0, b1, t0, t1] -> { [] : 0 <= b0 <= 31 and 0 <= b1 <= 31 and 0 <= t0 <= 31 and 0 <= t1 <= 15 }"
; SCHED-NEXT: child:
; SCHED-NEXT: filter: "[b0, b1] -> { Stmt_bb5[i0, i1] : -31 - 32b0 + i0 <= 8192*floor((i0)/8192) <= -32b0 + i0 and -31 - 32b1 + i1 <= 8192*floor((i1)/8192) <= -32b1 + i1 }"
; SCHED-NEXT: child:
; SCHED-NEXT: schedule: "[{ Stmt_bb5[i0, i1] -> [(floor((i0)/8192))] }, { Stmt_bb5[i0, i1] -> [(floor((i1)/8192))] }]"
; SCHED-NEXT: permutable: 1
; SCHED-NEXT: coincident: [ 1, 1 ]
; SCHED-NEXT: child:
; SCHED-NEXT: filter: "[t0, t1] -> { Stmt_bb5[i0, i1] : 32*floor((-t0 + i0)/32) = -t0 + i0 and 16*floor((-t1 + i1)/16) = -t1 + i1 and 0 <= t0 <= 31 and 0 <= t1 <= 15 }"
; SCHED-NEXT: child:
; SCHED-NEXT: schedule: "[{ Stmt_bb5[i0, i1] -> [(0)] }, { Stmt_bb5[i0, i1] -> [(floor((i1)/16) - 2*floor((i1)/32))] }]"
; SCHED-NEXT: permutable: 1
; SCHED-NEXT: coincident: [ 1, 1 ]
; SCHED-NEXT: - filter: "{ }"
GPGPU: Map initial schedule to GPU schedule This change now applies ppcg's GPU mapping on our initial schedule. For this to work, we need to also initialize the set of all names (isl_ids) used in the scop as well as the program context. llvm-svn: 275396
2016-07-14 18:51:52 +08:00
GPGPU: Generate an AST for the GPU-mapped schedule For this we need to provide an explicit list of statements as they occur in the polly::Scop to ppcg. We also setup basic AST printing facilities to facilitate debugging. To allow code reuse some (minor) changes in ppcg are have been necessary. llvm-svn: 275436
2016-07-14 23:51:37 +08:00
; CODE: Code
GPGPU: Use CHECK-NEXT to harden test cases A sequence of CHECK lines allows additional statements to appear in the output of the tested program without any test failures appearing. As we do not want this to happen, switch this test case to use CHECK-NEXT. llvm-svn: 275534
2016-07-15 15:05:49 +08:00
; CODE-NEXT: ====
; CODE-NEXT: # host
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(16, 32);
; CODE-NEXT: dim3 k0_dimGrid(32, 32);
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> ();
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
GPGPU: Generate an AST for the GPU-mapped schedule For this we need to provide an explicit list of statements as they occur in the polly::Scop to ppcg. We also setup basic AST printing facilities to facilitate debugging. To allow code reuse some (minor) changes in ppcg are have been necessary. llvm-svn: 275436
2016-07-14 23:51:37 +08:00
; CODE: # kernel0
GPGPU: Use CHECK-NEXT to harden test cases A sequence of CHECK lines allows additional statements to appear in the output of the tested program without any test failures appearing. As we do not want this to happen, switch this test case to use CHECK-NEXT. llvm-svn: 275534
2016-07-15 15:05:49 +08:00
; CODE-NEXT: for (int c3 = 0; c3 <= 1; c3 += 1)
; CODE-NEXT: Stmt_bb5(32 * b0 + t0, 32 * b1 + t1 + 16 * c3);
GPGPU: Generate an AST for the GPU-mapped schedule For this we need to provide an explicit list of statements as they occur in the polly::Scop to ppcg. We also setup basic AST printing facilities to facilitate debugging. To allow code reuse some (minor) changes in ppcg are have been necessary. llvm-svn: 275436
2016-07-14 23:51:37 +08:00
GPGPU: create default initialized PPCG scop and gpu program At this stage, we do not yet modify the IR but just generate a default initialized ppcg_scop and gpu_prog and free both immediately. Both will later be filled with data from the polly::Scop and are needed to use PPCG for GPU schedule generation. This commit does not yet perform any GPU code generation, but ensures that the basic infrastructure has been put in place. We also add a simple test case to ensure the new code is run and use this opportunity to verify that GPU_CODEGEN tests are only run if GPU code generation has been enabled in cmake. llvm-svn: 275389
2016-07-14 18:22:19 +08:00
; void double_parallel_loop(float A[][1024]) {
; for (long i = 0; i < 1024; i++)
; for (long j = 0; j < 1024; j++)
; A[i][j] += i * j;
; }
;
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
define void @double_parallel_loop([1024 x float]* %A) {
bb:
br label %bb2
bb2: ; preds = %bb13, %bb
%i.0 = phi i64 [ 0, %bb ], [ %tmp14, %bb13 ]
%exitcond1 = icmp ne i64 %i.0, 1024
br i1 %exitcond1, label %bb3, label %bb15
bb3: ; preds = %bb2
br label %bb4
bb4: ; preds = %bb10, %bb3
%j.0 = phi i64 [ 0, %bb3 ], [ %tmp11, %bb10 ]
%exitcond = icmp ne i64 %j.0, 1024
br i1 %exitcond, label %bb5, label %bb12
bb5: ; preds = %bb4
%tmp = mul nuw nsw i64 %i.0, %j.0
%tmp6 = sitofp i64 %tmp to float
%tmp7 = getelementptr inbounds [1024 x float], [1024 x float]* %A, i64 %i.0, i64 %j.0
%tmp8 = load float, float* %tmp7, align 4
%tmp9 = fadd float %tmp8, %tmp6
store float %tmp9, float* %tmp7, align 4
br label %bb10
bb10: ; preds = %bb5
%tmp11 = add nuw nsw i64 %j.0, 1
br label %bb4
bb12: ; preds = %bb4
br label %bb13
bb13: ; preds = %bb12
%tmp14 = add nuw nsw i64 %i.0, 1
br label %bb2
bb15: ; preds = %bb2
ret void
}