[polly] r333090 - [Acc] Update testcases for minor changes in the PPCG mapper and

[Philip Pfaffe via llvm-commits]

Author: pfaffe
Date: Wed May 23 07:56:57 2018
New Revision: 333090

URL: http://llvm.org/viewvc/llvm-project?rev=333090&view=rev
Log:
[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.

Modified:
    polly/trunk/test/GPGPU/host-statement.ll
    polly/trunk/test/GPGPU/live-range-reordering-with-privatization.ll
    polly/trunk/test/GPGPU/parametric-loop-bound.ll
    polly/trunk/test/GPGPU/scalar-writes-in-scop-requires-abort.ll
    polly/trunk/test/GPGPU/size-cast.ll

Modified: polly/trunk/test/GPGPU/host-statement.ll
URL: http://llvm.org/viewvc/llvm-project/polly/trunk/test/GPGPU/host-statement.ll?rev=333090&r1=333089&r2=333090&view=diff
==============================================================================
--- polly/trunk/test/GPGPU/host-statement.ll (original)
+++ polly/trunk/test/GPGPU/host-statement.ll Wed May 23 07:56:57 2018
@@ -47,7 +47,7 @@ declare void @llvm.lifetime.start(i64, i
 ; 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: }
 

Modified: polly/trunk/test/GPGPU/live-range-reordering-with-privatization.ll
URL: http://llvm.org/viewvc/llvm-project/polly/trunk/test/GPGPU/live-range-reordering-with-privatization.ll?rev=333090&r1=333089&r2=333090&view=diff
==============================================================================
--- polly/trunk/test/GPGPU/live-range-reordering-with-privatization.ll (original)
+++ polly/trunk/test/GPGPU/live-range-reordering-with-privatization.ll Wed May 23 07:56:57 2018
@@ -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);

Modified: polly/trunk/test/GPGPU/parametric-loop-bound.ll
URL: http://llvm.org/viewvc/llvm-project/polly/trunk/test/GPGPU/parametric-loop-bound.ll?rev=333090&r1=333089&r2=333090&view=diff
==============================================================================
--- polly/trunk/test/GPGPU/parametric-loop-bound.ll (original)
+++ polly/trunk/test/GPGPU/parametric-loop-bound.ll Wed May 23 07:56:57 2018
@@ -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:   }

Modified: polly/trunk/test/GPGPU/scalar-writes-in-scop-requires-abort.ll
URL: http://llvm.org/viewvc/llvm-project/polly/trunk/test/GPGPU/scalar-writes-in-scop-requires-abort.ll?rev=333090&r1=333089&r2=333090&view=diff
==============================================================================
--- polly/trunk/test/GPGPU/scalar-writes-in-scop-requires-abort.ll (original)
+++ polly/trunk/test/GPGPU/scalar-writes-in-scop-requires-abort.ll Wed May 23 07:56:57 2018
@@ -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: }
 

Modified: polly/trunk/test/GPGPU/size-cast.ll
URL: http://llvm.org/viewvc/llvm-project/polly/trunk/test/GPGPU/size-cast.ll?rev=333090&r1=333089&r2=333090&view=diff
==============================================================================
--- polly/trunk/test/GPGPU/size-cast.ll (original)
+++ polly/trunk/test/GPGPU/size-cast.ll Wed May 23 07:56:57 2018
@@ -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:
   %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