[Polly] [PPCGCodeGeneration] Error out if verifyModule fails on GPUModule.

[bollu]

Summary

• We used to silently return an empty kernel string on verifyModule failure
• Change this behaviour to an llvm_unreachable.

[bollu]

• Make VerifyModule failure an llvm_unreachable state
• add command line option to enable bailing out if verifyModule fails

[grosser]

LGTM, but please leave the test case unchanged. It still covers the default behavior.

lib/CodeGen/PPCGCodeGeneration.cpp

1517

Drop debug stuff.

[bollu]

Closed by r306290

lib/CodeGen/PPCGCodeGeneration.cpp

 
 static cl::opt<bool> ManagedMemory("polly-acc-codegen-managed-memory",
                                    cl::desc("Generate Host kernel code assuming"
                                             " that all memory has been"
                                             " declared as managed memory"),
                                    cl::Hidden, cl::init(false), cl::ZeroOrMore,
                                    cl::cat(PollyCategory));
 
+static cl::opt<bool>
+    FailOnVerifyModuleFailure("polly-acc-fail-on-verify-module-failure",
+                              cl::desc("Fail and generate a backtrace if"
+                                       " verifyModule fails on the GPU "
+                                       " kernel module."),
+                              cl::Hidden, cl::init(false), cl::ZeroOrMore,
+                              cl::cat(PollyCategory));
+
 static cl::opt<std::string>
     CudaVersion("polly-acc-cuda-version",
                 cl::desc("The CUDA version to compile for"), cl::Hidden,
                 cl::init("sm_30"), cl::ZeroOrMore, cl::cat(PollyCategory));
 
 static cl::opt<int>
     MinCompute("polly-acc-mincompute",
                cl::desc("Minimal number of compute statements to run on GPU."),
void GPUNodeBuilder::createKernel(__isl_take isl_ast_node *KernelStmt) {
   IDToSAI.clear();
   Annotator.resetAlternativeAliasBases();
   for (auto &BasePtr : LocalArrays)
     S.invalidateScopArrayInfo(BasePtr, MemoryKind::Array);
   LocalArrays.clear();
 
   std::string ASMString = finalizeKernelFunction();
   Builder.SetInsertPoint(&HostInsertPoint);
   Value *Parameters = createLaunchParameters(Kernel, F, SubtreeValues);
 
   std::string Name = "kernel_" + std::to_string(Kernel->id);
   Value *KernelString = Builder.CreateGlobalStringPtr(ASMString, Name);
   Value *NameString = Builder.CreateGlobalStringPtr(Name, Name + "_name");
   Value *GPUKernel = createCallGetKernel(KernelString, NameString);
 
   Value *GridDimX, *GridDimY;
   std::tie(GridDimX, GridDimY) = getGridSizes(Kernel);
std::string GPUNodeBuilder::createKernelASM() {
   }
 
   PM.run(*GPUModule);
 
   return ASMStream.str();
 }
 
 std::string GPUNodeBuilder::finalizeKernelFunction() {
+
   if (verifyModule(*GPUModule)) {
+    DEBUG(dbgs() << "verifyModule failed on module:\n";
+          GPUModule->print(dbgs(), nullptr); dbgs() << "\n";);
+
+    if (FailOnVerifyModuleFailure)
+      llvm_unreachable("VerifyModule failed.");
+
     BuildSuccessful = false;
     return "";
   }
 
   if (DumpKernelIR)
     outs() << *GPUModule << "\n";
 
   // Optimize module.

test/GPGPU/invalid-kernel.ll

-; 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-code -polly-acc-fail-on-verify-module-failure \
+; RUN: -disable-output < %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: -disable-output < %s
 
-; RUN: opt %loadPolly -polly-codegen-ppcg -S < %s | \
-; RUN: FileCheck %s -check-prefix=IR
+; RUN: opt %loadPolly -polly-codegen-ppcg -S < %s
 
 ; REQUIRES: pollyacc
+; XFAIL: *
 ;
 ;    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
+; RUN: opt %loadPolly -polly-codegen-ppcg -S < %s 
 
 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