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

203 lines
9.0 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 host control flow Create LLVM-IR for all host-side control flow of a given GPU AST. We implement this by introducing a new GPUNodeBuilder class derived from IslNodeBuilder. The IslNodeBuilder will take care of generating all general-purpose ast nodes, but we provide our own createUser implementation to handle the different GPU specific user statements. For now, we just skip any user statement and only generate a host-code sceleton, but in subsequent commits we will add handling of normal ScopStmt's performing computations, kernel calls, as well as host-device data transfers. We will also introduce run-time check generation and LICM in subsequent commits. llvm-svn: 275783
2016-07-18 19:56:39 +08:00
; RUN: opt %loadPolly -polly-codegen-ppcg -S < %s | \
; RUN: FileCheck %s -check-prefix=IR
GPGPU: add intrinsic functions to obtain a kernels thread and block ids llvm-svn: 275953
2016-07-19 15:32:44 +08:00
; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-kernel-ir \
; RUN: -disable-output < %s | \
; RUN: FileCheck %s -check-prefix=KERNEL-IR
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:
GPGPU: Model array access information This allows us to derive host-device and device-host data-transfers. llvm-svn: 275535
2016-07-15 15:05:54 +08:00
; SCHED-NEXT: extension: "{ [] -> from_device_MemRef_A[]; [] -> to_device_MemRef_A[] }"
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: sequence:
GPGPU: Model array access information This allows us to derive host-device and device-host data-transfers. llvm-svn: 275535
2016-07-15 15:05:54 +08:00
; SCHED-NEXT: - filter: "{ to_device_MemRef_A[] }"
; SCHED-NEXT: child:
; SCHED-NEXT: set:
; SCHED-NEXT: - filter: "{ to_device_MemRef_A[] }"
; SCHED-NEXT: child:
; SCHED-NEXT: guard: "{ [] }"
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: - 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 ]
GPGPU: Model array access information This allows us to derive host-device and device-host data-transfers. llvm-svn: 275535
2016-07-15 15:05:54 +08:00
; SCHED-NEXT: - filter: "{ from_device_MemRef_A[] }"
; SCHED-NEXT: child:
; SCHED-NEXT: set:
; SCHED-NEXT: - filter: "{ from_device_MemRef_A[] }"
; SCHED-NEXT: child:
; SCHED-NEXT: guard: "{ [] }"
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: {
GPGPU: Model array access information This allows us to derive host-device and device-host data-transfers. llvm-svn: 275535
2016-07-15 15:05:54 +08:00
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * (1024) * sizeof(float), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(16, 32);
; CODE-NEXT: dim3 k0_dimGrid(32, 32);
GPGPU: collect array references Initialize the list of references to a GPU array to ensure that the arrays that need to be passed to kernel calls are computed correctly. Furthermore, the very same information is also necessary to compute synchronization correctly. As the functionality to compute these references is already available, what is left for us to do is only to connect the necessary functionality to compute array reference information. llvm-svn: 275798
2016-07-18 23:44:32 +08:00
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
GPGPU: Model array access information This allows us to derive host-device and device-host data-transfers. llvm-svn: 275535
2016-07-15 15:05:54 +08:00
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * (1024) * sizeof(float), cudaMemcpyDeviceToHost));
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: }
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 host control flow Create LLVM-IR for all host-side control flow of a given GPU AST. We implement this by introducing a new GPUNodeBuilder class derived from IslNodeBuilder. The IslNodeBuilder will take care of generating all general-purpose ast nodes, but we provide our own createUser implementation to handle the different GPU specific user statements. For now, we just skip any user statement and only generate a host-code sceleton, but in subsequent commits we will add handling of normal ScopStmt's performing computations, kernel calls, as well as host-device data transfers. We will also introduce run-time check generation and LICM in subsequent commits. llvm-svn: 275783
2016-07-18 19:56:39 +08:00
; IR: polly.split_new_and_old:
; IR-NEXT: br i1 true, label %polly.start, label %bb2
; IR: polly.start:
; IR-NEXT: br label %polly.exiting
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 host control flow Create LLVM-IR for all host-side control flow of a given GPU AST. We implement this by introducing a new GPUNodeBuilder class derived from IslNodeBuilder. The IslNodeBuilder will take care of generating all general-purpose ast nodes, but we provide our own createUser implementation to handle the different GPU specific user statements. For now, we just skip any user statement and only generate a host-code sceleton, but in subsequent commits we will add handling of normal ScopStmt's performing computations, kernel calls, as well as host-device data transfers. We will also introduce run-time check generation and LICM in subsequent commits. llvm-svn: 275783
2016-07-18 19:56:39 +08:00
; IR: polly.exiting:
; IR-NEXT: br label %polly.merge_new_and_old
GPGPU: Model array access information This allows us to derive host-device and device-host data-transfers. llvm-svn: 275535
2016-07-15 15:05:54 +08:00
GPGPU: generate control flow within the kernel llvm-svn: 275956
2016-07-19 15:33:11 +08:00
; KERNEL-IR-LABEL: define ptx_kernel void @kernel_0(i8* %MemRef_A) {
GPGPU: add intrinsic functions to obtain a kernels thread and block ids llvm-svn: 275953
2016-07-19 15:32:44 +08:00
; KERNEL-IR-NEXT: entry:
; KERNEL-IR-NEXT: %0 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
; KERNEL-IR-NEXT: %b0 = zext i32 %0 to i64
; KERNEL-IR-NEXT: %1 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
; KERNEL-IR-NEXT: %b1 = zext i32 %1 to i64
; KERNEL-IR-NEXT: %2 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
; KERNEL-IR-NEXT: %t0 = zext i32 %2 to i64
; KERNEL-IR-NEXT: %3 = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
; KERNEL-IR-NEXT: %t1 = zext i32 %3 to i64
GPGPU: generate control flow within the kernel llvm-svn: 275956
2016-07-19 15:33:11 +08:00
; KERNEL-IR-NEXT: br label %polly.loop_preheader
GPGPU: generate code for ScopStatements This change introduces the actual compute code in the GPU kernels. To ensure all values referenced from the statements in the GPU kernel are indeed available we scan all ScopStmts in the GPU kernel for references to llvm::Values that are not yet covered by already modeled outer loop iterators, parameters, or array base pointers and also pass these additional llvm::Values to the GPU kernel. For arrays used in the GPU kernel we introduce a new ScopArrayInfo object, which is referenced by the newly generated access functions within the GPU kernel and which is used to help with code generation. llvm-svn: 276270
2016-07-21 21:15:59 +08:00
; KERNEL-IR-LABEL: polly.loop_exit: ; preds = %polly.stmt.bb5
GPGPU: add intrinsic functions to obtain a kernels thread and block ids llvm-svn: 275953
2016-07-19 15:32:44 +08:00
; KERNEL-IR-NEXT: ret void
GPGPU: generate control flow within the kernel llvm-svn: 275956
2016-07-19 15:33:11 +08:00
GPGPU: generate code for ScopStatements This change introduces the actual compute code in the GPU kernels. To ensure all values referenced from the statements in the GPU kernel are indeed available we scan all ScopStmts in the GPU kernel for references to llvm::Values that are not yet covered by already modeled outer loop iterators, parameters, or array base pointers and also pass these additional llvm::Values to the GPU kernel. For arrays used in the GPU kernel we introduce a new ScopArrayInfo object, which is referenced by the newly generated access functions within the GPU kernel and which is used to help with code generation. llvm-svn: 276270
2016-07-21 21:15:59 +08:00
; KERNEL-IR-LABEL: polly.loop_header: ; preds = %polly.stmt.bb5, %polly.loop_preheader
; KERNEL-IR-NEXT: %polly.indvar = phi i64 [ 0, %polly.loop_preheader ], [ %polly.indvar_next, %polly.stmt.bb5 ]
; KERNEL-IR-NEXT: %4 = mul nsw i64 32, %b0
; KERNEL-IR-NEXT: %5 = add nsw i64 %4, %t0
; KERNEL-IR-NEXT: %6 = mul nsw i64 32, %b1
; KERNEL-IR-NEXT: %7 = add nsw i64 %6, %t1
; KERNEL-IR-NEXT: %8 = mul nsw i64 16, %polly.indvar
; KERNEL-IR-NEXT: %9 = add nsw i64 %7, %8
; KERNEL-IR-NEXT: br label %polly.stmt.bb5
; KERNEL-IR-LABEL: polly.stmt.bb5: ; preds = %polly.loop_header
; KERNEL-IR-NEXT: %10 = mul i64 %9, %5
; KERNEL-IR-NEXT: %p_tmp6 = sitofp i64 %10 to float
; KERNEL-IR-NEXT: %polly.access.cast.MemRef_A = bitcast i8* %MemRef_A to float*
; KERNEL-IR-NEXT: %11 = mul nsw i64 32, %b0
; KERNEL-IR-NEXT: %12 = add nsw i64 %11, %t0
; KERNEL-IR-NEXT: %polly.access.mul.MemRef_A = mul nsw i64 %12, 1024
; KERNEL-IR-NEXT: %13 = mul nsw i64 32, %b1
; KERNEL-IR-NEXT: %14 = add nsw i64 %13, %t1
; KERNEL-IR-NEXT: %15 = mul nsw i64 16, %polly.indvar
; KERNEL-IR-NEXT: %16 = add nsw i64 %14, %15
; KERNEL-IR-NEXT: %polly.access.add.MemRef_A = add nsw i64 %polly.access.mul.MemRef_A, %16
; KERNEL-IR-NEXT: %polly.access.MemRef_A = getelementptr float, float* %polly.access.cast.MemRef_A, i64 %polly.access.add.MemRef_A
; KERNEL-IR-NEXT: %tmp8_p_scalar_ = load float, float* %polly.access.MemRef_A, align 4
; KERNEL-IR-NEXT: %p_tmp9 = fadd float %tmp8_p_scalar_, %p_tmp6
; KERNEL-IR-NEXT: %polly.access.cast.MemRef_A1 = bitcast i8* %MemRef_A to float*
; KERNEL-IR-NEXT: %17 = mul nsw i64 32, %b0
; KERNEL-IR-NEXT: %18 = add nsw i64 %17, %t0
; KERNEL-IR-NEXT: %polly.access.mul.MemRef_A2 = mul nsw i64 %18, 1024
; KERNEL-IR-NEXT: %19 = mul nsw i64 32, %b1
; KERNEL-IR-NEXT: %20 = add nsw i64 %19, %t1
; KERNEL-IR-NEXT: %21 = mul nsw i64 16, %polly.indvar
; KERNEL-IR-NEXT: %22 = add nsw i64 %20, %21
; KERNEL-IR-NEXT: %polly.access.add.MemRef_A3 = add nsw i64 %polly.access.mul.MemRef_A2, %22
; KERNEL-IR-NEXT: %polly.access.MemRef_A4 = getelementptr float, float* %polly.access.cast.MemRef_A1, i64 %polly.access.add.MemRef_A3
; KERNEL-IR-NEXT: store float %p_tmp9, float* %polly.access.MemRef_A4, align 4
GPGPU: generate control flow within the kernel llvm-svn: 275956
2016-07-19 15:33:11 +08:00
; KERNEL-IR-NEXT: %polly.indvar_next = add nsw i64 %polly.indvar, 1
; KERNEL-IR-NEXT: %polly.loop_cond = icmp sle i64 %polly.indvar, 0
; KERNEL-IR-NEXT: br i1 %polly.loop_cond, label %polly.loop_header, label %polly.loop_exit
GPGPU: generate code for ScopStatements This change introduces the actual compute code in the GPU kernels. To ensure all values referenced from the statements in the GPU kernel are indeed available we scan all ScopStmts in the GPU kernel for references to llvm::Values that are not yet covered by already modeled outer loop iterators, parameters, or array base pointers and also pass these additional llvm::Values to the GPU kernel. For arrays used in the GPU kernel we introduce a new ScopArrayInfo object, which is referenced by the newly generated access functions within the GPU kernel and which is used to help with code generation. llvm-svn: 276270
2016-07-21 21:15:59 +08:00
; KERNEL-IR-LABEL: polly.loop_preheader: ; preds = %entry
GPGPU: generate control flow within the kernel llvm-svn: 275956
2016-07-19 15:33:11 +08:00
; KERNEL-IR-NEXT: br label %polly.loop_header
GPGPU: generate code for ScopStatements This change introduces the actual compute code in the GPU kernels. To ensure all values referenced from the statements in the GPU kernel are indeed available we scan all ScopStmts in the GPU kernel for references to llvm::Values that are not yet covered by already modeled outer loop iterators, parameters, or array base pointers and also pass these additional llvm::Values to the GPU kernel. For arrays used in the GPU kernel we introduce a new ScopArrayInfo object, which is referenced by the newly generated access functions within the GPU kernel and which is used to help with code generation. llvm-svn: 276270
2016-07-21 21:15:59 +08:00
GPGPU: add intrinsic functions to obtain a kernels thread and block ids llvm-svn: 275953
2016-07-19 15:32:44 +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
}