From 2e171b52eeac53d4ece15816688b8c259cd1fea1 Mon Sep 17 00:00:00 2001 From: Philip Pfaffe Date: Wed, 23 May 2018 14:56:57 +0000 Subject: [PATCH] [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 --- polly/test/GPGPU/host-statement.ll | 2 +- polly/test/GPGPU/live-range-reordering-with-privatization.ll | 2 +- polly/test/GPGPU/parametric-loop-bound.ll | 2 +- polly/test/GPGPU/scalar-writes-in-scop-requires-abort.ll | 2 +- polly/test/GPGPU/size-cast.ll | 4 ++-- 5 files changed, 6 insertions(+), 6 deletions(-) diff --git a/polly/test/GPGPU/host-statement.ll b/polly/test/GPGPU/host-statement.ll index e54d3689e14c..9fe851ea8905 100644 --- a/polly/test/GPGPU/host-statement.ll +++ b/polly/test/GPGPU/host-statement.ll @@ -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: } diff --git a/polly/test/GPGPU/live-range-reordering-with-privatization.ll b/polly/test/GPGPU/live-range-reordering-with-privatization.ll index 9fcaed3b0f52..d7d9e0c4626b 100644 --- a/polly/test/GPGPU/live-range-reordering-with-privatization.ll +++ b/polly/test/GPGPU/live-range-reordering-with-privatization.ll @@ -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); diff --git a/polly/test/GPGPU/parametric-loop-bound.ll b/polly/test/GPGPU/parametric-loop-bound.ll index 6e8842d653b0..610fc401181d 100644 --- a/polly/test/GPGPU/parametric-loop-bound.ll +++ b/polly/test/GPGPU/parametric-loop-bound.ll @@ -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 <<>> (dev_MemRef_A, n); ; CODE-NEXT: cudaCheckKernel(); ; CODE-NEXT: } diff --git a/polly/test/GPGPU/scalar-writes-in-scop-requires-abort.ll b/polly/test/GPGPU/scalar-writes-in-scop-requires-abort.ll index 4d0b7c2a49a5..ff04efef766a 100644 --- a/polly/test/GPGPU/scalar-writes-in-scop-requires-abort.ll +++ b/polly/test/GPGPU/scalar-writes-in-scop-requires-abort.ll @@ -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: } diff --git a/polly/test/GPGPU/size-cast.ll b/polly/test/GPGPU/size-cast.ll index 0312dc81a85d..d4f928ee9c1c 100644 --- a/polly/test/GPGPU/size-cast.ll +++ b/polly/test/GPGPU/size-cast.ll @@ -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 <<>> (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 -- GitLab