diff options
| -rw-r--r-- | polly/lib/CodeGen/PPCGCodeGeneration.cpp | 17 | ||||
| -rw-r--r-- | polly/test/GPGPU/invalid-kernel.ll | 78 |
2 files changed, 91 insertions, 4 deletions
diff --git a/polly/lib/CodeGen/PPCGCodeGeneration.cpp b/polly/lib/CodeGen/PPCGCodeGeneration.cpp index 66ffa730acc..783d256a32c 100644 --- a/polly/lib/CodeGen/PPCGCodeGeneration.cpp +++ b/polly/lib/CodeGen/PPCGCodeGeneration.cpp @@ -152,6 +152,12 @@ public: /// Finalize the generated scop. virtual void finalize(); + /// Track if the full build process was successful. + /// + /// This value is set to false, if throughout the build process an error + /// occurred which prevents us from generating valid GPU code. + bool BuildSuccessful = true; + private: /// A vector of array base pointers for which a new ScopArrayInfo was created. /// @@ -1409,10 +1415,10 @@ std::string GPUNodeBuilder::createKernelASM() { } std::string GPUNodeBuilder::finalizeKernelFunction() { - // Verify module. - llvm::legacy::PassManager Passes; - Passes.add(createVerifierPass()); - Passes.run(*GPUModule); + if (verifyModule(*GPUModule)) { + BuildSuccessful = false; + return ""; + } if (DumpKernelIR) outs() << *GPUModule << "\n"; @@ -2139,6 +2145,9 @@ public: NodeBuilder.initializeAfterRTH(); NodeBuilder.create(Root); NodeBuilder.finalize(); + + if (!NodeBuilder.BuildSuccessful) + SplitBlock->getTerminator()->setOperand(0, Builder.getFalse()); } bool runOnScop(Scop &CurrentScop) override { diff --git a/polly/test/GPGPU/invalid-kernel.ll b/polly/test/GPGPU/invalid-kernel.ll new file mode 100644 index 00000000000..298f6a589a6 --- /dev/null +++ b/polly/test/GPGPU/invalid-kernel.ll @@ -0,0 +1,78 @@ +; 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 -polly-acc-dump-kernel-ir \ +; RUN: -disable-output < %s | \ +; RUN: not FileCheck %s -check-prefix=KERNEL-IR + +; RUN: opt %loadPolly -polly-codegen-ppcg -S < %s | \ +; RUN: FileCheck %s -check-prefix=IR + +; REQUIRES: pollyacc +; +; void foo(long A[1024], long B[1024]) { +; for (long i = 0; i < 1024; i++) +; A[i] += (B[i] + (long)&B[i]); +; } + +; This kernel loads/stores a pointer address we model. This is a rare case, +; were we still lack proper code-generation support. We check here that we +; detect the invalid IR and bail out gracefully. + +; CODE: Code +; CODE-NEXT: ==== +; CODE-NEXT: # host +; CODE-NEXT: { +; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_B, MemRef_B, (1024) * sizeof(i64), cudaMemcpyHostToDevice)); +; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * sizeof(i64), cudaMemcpyHostToDevice)); +; CODE-NEXT: { +; CODE-NEXT: dim3 k0_dimBlock(32); +; CODE-NEXT: dim3 k0_dimGrid(32); +; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_B, dev_MemRef_A); +; CODE-NEXT: cudaCheckKernel(); +; CODE-NEXT: } + +; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * sizeof(i64), cudaMemcpyDeviceToHost)); +; CODE-NEXT: } + +; CODE: # kernel0 +; CODE-NEXT: Stmt_bb2(32 * b0 + t0); + +; RUN: opt %loadPolly -polly-codegen-ppcg -S < %s | \ +; RUN: FileCheck %s -check-prefix=IR + +; KERNEL-IR: kernel + +; IR: br i1 false, label %polly.start, label %bb1 + +target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" + +define void @foo(i64* %A, i64* %B) { +bb: + br label %bb1 + +bb1: ; preds = %bb10, %bb + %i.0 = phi i64 [ 0, %bb ], [ %tmp11, %bb10 ] + %exitcond = icmp ne i64 %i.0, 1024 + br i1 %exitcond, label %bb2, label %bb12 + +bb2: ; preds = %bb1 + %tmp = getelementptr inbounds i64, i64* %B, i64 %i.0 + %tmp3 = load i64, i64* %tmp, align 8 + %tmp4 = getelementptr inbounds i64, i64* %B, i64 %i.0 + %tmp5 = ptrtoint i64* %tmp4 to i64 + %tmp6 = add nsw i64 %tmp3, %tmp5 + %tmp7 = getelementptr inbounds i64, i64* %A, i64 %i.0 + %tmp8 = load i64, i64* %tmp7, align 8 + %tmp9 = add nsw i64 %tmp8, %tmp6 + store i64 %tmp9, i64* %tmp7, align 8 + br label %bb10 + +bb10: ; preds = %bb2 + %tmp11 = add nuw nsw i64 %i.0, 1 + br label %bb1 + +bb12: ; preds = %bb1 + ret void +} |
