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

257 lines
12 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: Generate PTX assembly code for the kernel modules Run the NVPTX backend over the GPUModule IR and write the resulting assembly code in a string. To work correctly, it is important to invalidate analysis results that still reference the IR in the kernel module. Hence, this change clears all references to dominators, loop info, and scalar evolution. Finally, the NVPTX backend has troubles to generate code for various special floating point types (not surprising), but also for uncommon integer types. This commit does not resolve these issues, but pulls out problematic test cases into separate files to XFAIL them individually and resolve them in future (not immediate) changes one by one. llvm-svn: 276396
2016-07-22 15:11:12 +08:00
; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-kernel-asm \
; RUN: -disable-output < %s | \
; RUN: FileCheck %s -check-prefix=KERNEL-ASM
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:
GPGPU: Dynamically ensure 'sufficient compute' Offloading to a GPU is only beneficial if there is a sufficient amount of compute that can be accelerated. Many kernels just have a very small number of dynamic compute, which means GPU acceleration is not beneficial. We compute at run-time an approximation of how many dynamic instructions will be executed and fall back to CPU code in case this number is not sufficiently large. To keep the run-time checking code simple, we over-approximate the number of instructions executed in each statement by computing the volume of the rectangular hull of its iteration space. llvm-svn: 281848
2016-09-18 14:50:35 +08:00
; IR-NEXT: %0 = call { i64, i1 } @llvm.smul.with.overflow.i64(i64 1, i64 1024)
; IR-NEXT: %.obit = extractvalue { i64, i1 } %0, 1
; IR-NEXT: %polly.overflow.state = or i1 false, %.obit
; IR-NEXT: %.res = extractvalue { i64, i1 } %0, 0
; IR-NEXT: %1 = call { i64, i1 } @llvm.smul.with.overflow.i64(i64 %.res, i64 1024)
; IR-NEXT: %.obit1 = extractvalue { i64, i1 } %1, 1
; IR-NEXT: %polly.overflow.state2 = or i1 %polly.overflow.state, %.obit1
; IR-NEXT: %.res3 = extractvalue { i64, i1 } %1, 0
; IR-NEXT: %2 = call { i64, i1 } @llvm.smul.with.overflow.i64(i64 7, i64 %.res3)
; IR-NEXT: %.obit4 = extractvalue { i64, i1 } %2, 1
; IR-NEXT: %polly.overflow.state5 = or i1 %polly.overflow.state2, %.obit4
; IR-NEXT: %.res6 = extractvalue { i64, i1 } %2, 0
; IR-NEXT: %3 = call { i64, i1 } @llvm.sadd.with.overflow.i64(i64 0, i64 %.res6)
; IR-NEXT: %.obit7 = extractvalue { i64, i1 } %3, 1
; IR-NEXT: %polly.overflow.state8 = or i1 %polly.overflow.state5, %.obit7
; IR-NEXT: %.res9 = extractvalue { i64, i1 } %3, 0
; IR-NEXT: %4 = icmp sge i64 %.res9, 2621440
; IR-NEXT: %5 = and i1 true, %4
; IR-NEXT: %polly.rtc.overflown = xor i1 %polly.overflow.state8, true
; IR-NEXT: %polly.rtc.result = and i1 %5, %polly.rtc.overflown
; IR-NEXT: br i1 %polly.rtc.result, label %polly.start, label %bb2
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.start:
[GPGPU] Use separate basic block for GPU initialization code This increases the readability of the IR and also clarifies that the GPU inititialization is executed _after_ the scalar initialization which needs to before the code of the transformed scop is executed. Besides increased readability, the IR should not change. Specifically, I do not expect any changes in program semantics due to this patch. llvm-svn: 278125
2016-08-09 23:35:03 +08:00
; IR-NEXT: br label %polly.acc.initialize
; IR: polly.acc.initialize:
GPGPU: initialize GPU context and simplify the corresponding GPURuntime interface. There is no need to expose the selected device at the moment. We also pass back pointers as return values, as this simplifies the interface. llvm-svn: 276623
2016-07-25 17:16:01 +08:00
; IR-NEXT: [[GPUContext:%.*]] = call i8* @polly_initContext()
GPGPU: Complete code to allocate and free device arrays At the beginning of each SCoP, we allocate device arrays for all arrays used on the GPU and we free such arrays after the SCoP has been executed. llvm-svn: 276635
2016-07-25 20:47:33 +08:00
; IR-NEXT: %p_dev_array_MemRef_A = call i8* @polly_allocateMemoryForDevice(i64 4194304)
GPGPU: Emit data-transfer code Also factor out getArraySize() to avoid code dupliciation and reorder some function arguments to indicate the direction into which data is transferred. llvm-svn: 276636
2016-07-25 20:47:39 +08:00
; IR-NEXT: [[HostPtr:%.*]] = bitcast [1024 x float]* %A to i8*
; IR-NEXT: call void @polly_copyFromHostToDevice(i8* [[HostPtr]], i8* %p_dev_array_MemRef_A, i64 4194304)
GPGPU: Add basic support for kernel launches llvm-svn: 276863
2016-07-27 21:20:16 +08:00
; IR-NEXT: [[DevPtr:%.*]] = call i8* @polly_getDevicePtr(i8* %p_dev_array_MemRef_A)
GPGPU: Generate kernel parameter allocation with right size Before this change we miscounted the number of function parameters. llvm-svn: 276960
2016-07-28 14:47:50 +08:00
; IR-NEXT: [[ParamSlot:%.*]] = getelementptr [1 x i8*], [1 x i8*]* %polly_launch_0_params, i64 0, i64 0
GPGPU: Detect read-only scalar arrays ... and pass these by value rather than by reference. llvm-svn: 281837
2016-09-18 03:22:18 +08:00
; IR-NEXT: store i8* [[DevPtr]], i8** %polly_launch_0_param_0
GPGPU: Add basic support for kernel launches llvm-svn: 276863
2016-07-27 21:20:16 +08:00
; IR-NEXT: [[ParamTyped:%.*]] = bitcast i8** %polly_launch_0_param_0 to i8*
; IR-NEXT: store i8* [[ParamTyped]], i8** [[ParamSlot]]
GPGPU: Load GPU kernels We embed the PTX code into the host IR as a global variable and compile it at run-time into a GPU kernel. llvm-svn: 276645
2016-07-26 00:31:21 +08:00
; IR-NEXT: call i8* @polly_getKernel
GPGPU: Dynamically ensure 'sufficient compute' Offloading to a GPU is only beneficial if there is a sufficient amount of compute that can be accelerated. Many kernels just have a very small number of dynamic compute, which means GPU acceleration is not beneficial. We compute at run-time an approximation of how many dynamic instructions will be executed and fall back to CPU code in case this number is not sufficiently large. To keep the run-time checking code simple, we over-approximate the number of instructions executed in each statement by computing the volume of the rectangular hull of its iteration space. llvm-svn: 281848
2016-09-18 14:50:35 +08:00
; IR-NEXT: call void @polly_launchKernel(i8* %11, i32 32, i32 32, i32 32, i32 16, i32 1, i8* %polly_launch_0_params_i8ptr)
GPGPU: Load GPU kernels We embed the PTX code into the host IR as a global variable and compile it at run-time into a GPU kernel. llvm-svn: 276645
2016-07-26 00:31:21 +08:00
; IR-NEXT: call void @polly_freeKernel
GPGPU: Emit data-transfer code Also factor out getArraySize() to avoid code dupliciation and reorder some function arguments to indicate the direction into which data is transferred. llvm-svn: 276636
2016-07-25 20:47:39 +08:00
; IR-NEXT: [[HostPtr2:%.*]] = bitcast [1024 x float]* %A to i8*
; IR-NEXT: call void @polly_copyFromDeviceToHost(i8* %p_dev_array_MemRef_A, i8* [[HostPtr2]], i64 4194304)
GPGPU: Complete code to allocate and free device arrays At the beginning of each SCoP, we allocate device arrays for all arrays used on the GPU and we free such arrays after the SCoP has been executed. llvm-svn: 276635
2016-07-25 20:47:33 +08:00
; IR-NEXT: call void @polly_freeDeviceMemory(i8* %p_dev_array_MemRef_A)
GPGPU: initialize GPU context and simplify the corresponding GPURuntime interface. There is no need to expose the selected device at the moment. We also pass back pointers as return values, as this simplifies the interface. llvm-svn: 276623
2016-07-25 17:16:01 +08:00
; IR-NEXT: call void @polly_freeContext(i8* [[GPUContext]])
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-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: Mark kernel functions as polly.skip Otherwise, we would try to re-optimize them with Polly-ACC and possibly even generate kernels that try to offload themselves, which does not work as the GPURuntime is not available on the accelerator and also does not make any sense. llvm-svn: 277589
2016-08-03 20:00:07 +08:00
; KERNEL-IR-LABEL: define ptx_kernel void @kernel_0(i8* %MemRef_A) #0 {
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
[tests] Adjust test output to recent changed SCEV canonocalization [NFC] LLVM recently changed the SCEV canonicalization which changed the output of one of our GPGPU test cases. llvm-svn: 286770
2016-11-14 03:27:17 +08:00
; KERNEL-IR-NEXT: %10 = mul i64 %5, %9
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-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: Mark kernel functions as polly.skip Otherwise, we would try to re-optimize them with Polly-ACC and possibly even generate kernels that try to offload themselves, which does not work as the GPURuntime is not available on the accelerator and also does not make any sense. llvm-svn: 277589
2016-08-03 20:00:07 +08:00
; KERNEL-IR: attributes #0 = { "polly.skip.fn" }
GPGPU: add intrinsic functions to obtain a kernels thread and block ids llvm-svn: 275953
2016-07-19 15:32:44 +08:00
GPGPU: Generate PTX assembly code for the kernel modules Run the NVPTX backend over the GPUModule IR and write the resulting assembly code in a string. To work correctly, it is important to invalidate analysis results that still reference the IR in the kernel module. Hence, this change clears all references to dominators, loop info, and scalar evolution. Finally, the NVPTX backend has troubles to generate code for various special floating point types (not surprising), but also for uncommon integer types. This commit does not resolve these issues, but pulls out problematic test cases into separate files to XFAIL them individually and resolve them in future (not immediate) changes one by one. llvm-svn: 276396
2016-07-22 15:11:12 +08:00
; KERNEL-ASM: .version 3.2
; KERNEL-ASM-NEXT: .target sm_30
; KERNEL-ASM-NEXT: .address_size 64
; KERNEL-ASM: // .globl kernel_0
; KERNEL-ASM: .visible .entry kernel_0(
; KERNEL-ASM-NEXT: .param .u64 kernel_0_param_0
; KERNEL-ASM-NEXT: )
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
}