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/size-cast.ll

64 lines
2.4 KiB
LLVM
Raw Normal View History

GPGPU: Extend types when array sizes have smaller types This prevents a compiler crash. llvm-svn: 281303
2016-09-13 16:02:14 +08:00
; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \
; RUN: -disable-output < %s | \
; RUN: FileCheck -check-prefix=CODE %s
; RUN: opt %loadPolly -polly-codegen-ppcg -S < %s | \
; RUN: FileCheck %s -check-prefix=IR
; REQUIRES: pollyacc
; This test case ensures that we properly sign-extend the types we are using.
[PPCGCodeGen] [3/3] Update PPCGCodeGen + tests to latest ppcg. This commit *WILL COMPILE*. 1. `PPCG` now uses `isl_multi_pw_aff` instead of an array of `pw_aff`. This needs us to adjust how we index array bounds and how we construct array bounds. 2. `PPCG` introduces two new kinds of nodes: `init_device` and `clear_device`. We should investigate what the correct way to handle these are. 3. `PPCG` has gotten smarter with its use of live range reordering, so some of the tests have a qualitative improvement. 4. `PPCG` changed its output style, so many test cases need to be updated to fit the new style for `polly-acc-dump-code` checks. Differential Revision: https://reviews.llvm.org/D35677 llvm-svn: 308625
2017-07-20 23:48:36 +08:00
; CODE: if (arg >= 1 && arg1 == 0) {
; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_arg3, MemRef_arg3, (arg) * sizeof(double), cudaMemcpyHostToDevice));
GPGPU: Extend types when array sizes have smaller types This prevents a compiler crash. llvm-svn: 281303
2016-09-13 16:02:14 +08:00
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
[Acc] Update testcases for minor changes in the PPCG mapper and statement naming - A recent ppcg/isl update caused the grid/block size upper bounds to deviate by one from the oracle. This is not an effect that's visible at runtime. - Statement naming changed in polly. Update the testcases. llvm-svn: 333090
2018-05-23 22:56:57 +08:00
; CODE-NEXT: dim3 k0_dimGrid(arg >= 1048545 ? 32768 : (arg + 31) / 32);
GPGPU: Extend types when array sizes have smaller types This prevents a compiler crash. llvm-svn: 281303
2016-09-13 16:02:14 +08:00
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_arg3, dev_MemRef_arg2, arg, arg1);
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_arg2, dev_MemRef_arg2, (arg) * sizeof(double), cudaMemcpyDeviceToHost));
[PPCGCodeGen] [3/3] Update PPCGCodeGen + tests to latest ppcg. This commit *WILL COMPILE*. 1. `PPCG` now uses `isl_multi_pw_aff` instead of an array of `pw_aff`. This needs us to adjust how we index array bounds and how we construct array bounds. 2. `PPCG` introduces two new kinds of nodes: `init_device` and `clear_device`. We should investigate what the correct way to handle these are. 3. `PPCG` has gotten smarter with its use of live range reordering, so some of the tests have a qualitative improvement. 4. `PPCG` changed its output style, so many test cases need to be updated to fit the new style for `polly-acc-dump-code` checks. Differential Revision: https://reviews.llvm.org/D35677 llvm-svn: 308625
2017-07-20 23:48:36 +08:00
; CODE-NEXT cudaCheckReturn(cudaFree(dev_MemRef_arg3));
; CODE-NEXT cudaCheckReturn(cudaFree(dev_MemRef_arg2));
GPGPU: Extend types when array sizes have smaller types This prevents a compiler crash. llvm-svn: 281303
2016-09-13 16:02:14 +08:00
; CODE: # kernel0
; CODE-NEXT: for (int c0 = 0; c0 <= (arg - 32 * b0 - 1) / 1048576; c0 += 1)
; CODE-NEXT: if (arg >= 32 * b0 + t0 + 1048576 * c0 + 1)
; CODE-NEXT: Stmt_bb6(0, 32 * b0 + t0 + 1048576 * c0);
[NFC] In GPGPU testcases, replace numeric registers in CHECK directives. Using numeric registers is flaky, since as soon as one additional instruction is generated by us, all the tests need to be adapted. llvm-svn: 319544
2017-12-01 22:16:39 +08:00
; IR-LABEL: call i8* @polly_initContextCUDA()
; IR: sext i32 %arg to i64
GPGPU: Extend types when array sizes have smaller types This prevents a compiler crash. llvm-svn: 281303
2016-09-13 16:02:14 +08:00
; IR-NEXT: mul i64
; IR-NEXT: @polly_allocateMemoryForDevice
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"
define void @hoge(i32 %arg, i32 %arg1, [1000 x double]* %arg2, double* %arg3) {
bb:
br label %bb4
bb4: ; preds = %bb13, %bb
br label %bb6
bb5: ; preds = %bb13
ret void
bb6: ; preds = %bb6, %bb4
%tmp = phi i64 [ 0, %bb4 ], [ %tmp10, %bb6 ]
%tmp7 = getelementptr inbounds double, double* %arg3, i64 %tmp
%tmp8 = load double, double* %tmp7, align 8
%tmp9 = getelementptr inbounds [1000 x double], [1000 x double]* %arg2, i64 0, i64 %tmp
[Acc] Update testcases for minor changes in the PPCG mapper and statement naming - A recent ppcg/isl update caused the grid/block size upper bounds to deviate by one from the oracle. This is not an effect that's visible at runtime. - Statement naming changed in polly. Update the testcases. llvm-svn: 333090
2018-05-23 22:56:57 +08:00
store double %tmp8, double* %tmp9, align 8
GPGPU: Extend types when array sizes have smaller types This prevents a compiler crash. llvm-svn: 281303
2016-09-13 16:02:14 +08:00
%tmp10 = add nuw nsw i64 %tmp, 1
%tmp11 = zext i32 %arg to i64
%tmp12 = icmp ne i64 %tmp10, %tmp11
br i1 %tmp12, label %bb6, label %bb13
bb13: ; preds = %bb6
%tmp14 = zext i32 %arg1 to i64
%tmp15 = icmp ne i64 0, %tmp14
br i1 %tmp15, label %bb4, label %bb5
}