maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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
This commit is contained in:
Philip Pfaffe 2018-05-23 14:56:57 +00:00
parent c06a6380a0
commit 2e171b52ee
5 changed files with 6 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
polly/test/GPGPU/host-statement.ll
View File

@ -47,7 +47,7 @@ declare void @llvm.lifetime.start(i64, i8* nocapture) #0
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (512) * (512) * sizeof(double), cudaMemcpyDeviceToHost));
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(MemRef_R, dev_MemRef_R, (p_0 + 1) * (512) * sizeof(double), cudaMemcpyDeviceToHost));
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(MemRef_Q, dev_MemRef_Q, (512) * (512) * sizeof(double), cudaMemcpyDeviceToHost));
; CODE-NEXT: Stmt_for_cond33_preheader();
; CODE-NEXT: Stmt_for_cond33_preheader_last();
; CODE: }

2
polly/test/GPGPU/live-range-reordering-with-privatization.ll
View File

@ -26,7 +26,7 @@
; CODE: # kernel0
; CODE-NEXT: for (int c0 = 0; c0 <= (tmp - 32 * b0 - 1) / 1048576; c0 += 1)
; CODE-NEXT: if (tmp >= 32 * b0 + t0 + 1048576 * c0 + 1) {
; CODE-NEXT: Stmt_for_body(32 * b0 + t0 + 1048576 * c0);
; CODE-NEXT: Stmt_for_body_last(32 * b0 + t0 + 1048576 * c0);
; CODE-NEXT: if (tmp1 >= 4)
; CODE-NEXT: Stmt_if_then(32 * b0 + t0 + 1048576 * c0);
; CODE-NEXT: Stmt_if_end(32 * b0 + t0 + 1048576 * c0);

2
polly/test/GPGPU/parametric-loop-bound.ll
View File

@ -17,7 +17,7 @@
; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (n) * sizeof(i64), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(n >= 1048546 ? 32768 : (n + 31) / 32);
; CODE-NEXT: dim3 k0_dimGrid(n >= 1048545 ? 32768 : (n + 31) / 32);
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A, n);
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }

2
polly/test/GPGPU/scalar-writes-in-scop-requires-abort.ll
View File

@ -13,7 +13,7 @@
; SCOP: Invariant Accesses: {
; SCOP-NEXT: ReadAccess := [Reduction Type: NONE] [Scalar: 0]
; SCOP-NEXT: { Stmt_loop[i0] -> MemRef_p[0] };
; SCOP-NEXT: { Stmt_loop_a[i0] -> MemRef_p[0] };
; SCOP-NEXT: Execution Context: { : }
; SCOP-NEXT: }

4
polly/test/GPGPU/size-cast.ll
View File

@ -13,7 +13,7 @@
; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_arg3, MemRef_arg3, (arg) * sizeof(double), cudaMemcpyHostToDevice));
; CODE-NEXT: {
; CODE-NEXT: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(arg >= 1048546 ? 32768 : (arg + 31) / 32);
; CODE-NEXT: dim3 k0_dimGrid(arg >= 1048545 ? 32768 : (arg + 31) / 32);
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_arg3, dev_MemRef_arg2, arg, arg1);
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
@ -50,7 +50,7 @@ bb6: ; preds = %bb6, %bb4
%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
store double undef, double* %tmp9, align 8
store double %tmp8, double* %tmp9, align 8
%tmp10 = add nuw nsw i64 %tmp, 1
%tmp11 = zext i32 %arg to i64
%tmp12 = icmp ne i64 %tmp10, %tmp11