diff --git a/polly/test/GPGPU/host-statement.ll b/polly/test/GPGPU/host-statement.ll index e54d3689e14cb85e3005beb23f5258c7a858bd80..9fe851ea8905ff36fd3b932252c9ba98b7fc2187 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 9fcaed3b0f5286eed9dfed9be1b8a521d91f78b1..d7d9e0c4626b1656f9b96793f19daf54520dce2a 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 6e8842d653b0e863b00980934b1bf65a4de88b6f..610fc401181d24eea6c1998dc3e38c75f5212ba2 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 4d0b7c2a49a58087c0738f4a48423d07f904860e..ff04efef766accac45b487abfeb16a55ed2e2abb 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 0312dc81a85ddc05f4200df3cc6ac8d68882ad4f..d4f928ee9c1c98e1e4cf1e8b3a8f4d146e10104f 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