Merge pull request intel#254 from frasercrmck/better-cfg-convert-diags

[vecz] Provide more context when CFG conversion fails

Author: frasercrmck

llvm/lib/SYCLNativeCPUUtils/compiler_passes/vecz/source/debugging.cpp

@@ -26,31 +26,39 @@ namespace vecz {
 ///
 /// @param[in] V The value (can be `nullptr`) to be included in the remark
 /// @param[in] Msg The main remark message
+/// @param[in] Note An optional additional note to provide more context/info.
 /// @return The remark message as it is to be printed
-static std::string createRemarkMessage(const Value *V, StringRef Msg) {
+static std::string createRemarkMessage(const Value *V, StringRef Msg,
+                                       StringRef Note = "") {
   std::string helper_str("Vecz: ");
   raw_string_ostream helper_stream(helper_str);
   helper_stream << Msg;
   if (V) {
     if (isa<Instruction>(V)) {
       // Instructions are already prefixed by two spaces when printed
-      V->print(helper_stream, true);
+      V->print(helper_stream, /*IsForDebug=*/true);
     } else if (const Function *F = dyn_cast<Function>(V)) {
-      // Printing a functions leads to it's whole body being printed
+      // Printing a functions leads to its whole body being printed
       helper_stream << " function \"" << F->getName() << "\"";
     } else {
       helper_stream << " ";
-      V->print(helper_stream, true);
+      V->print(helper_stream, /*IsForDebug=*/true);
     }
   }
   helper_stream << '\n';
 
+  // Provide extra context, if supplied
+  if (!Note.empty()) {
+    helper_stream << "  note: " << Note << '\n';
+  }
+
   return helper_stream.str();
 }
 
-void emitVeczRemarkMissed(const Function *F, const Value *V, StringRef Msg) {
+void emitVeczRemarkMissed(const Function *F, const Value *V, StringRef Msg,
+                          StringRef Note) {
   const Instruction *I = V ? dyn_cast<Instruction>(V) : nullptr;
-  auto RemarkMsg = createRemarkMessage(V, Msg);
+  auto RemarkMsg = createRemarkMessage(V, Msg, Note);
   OptimizationRemarkEmitter ORE(F);
   if (I) {
     ORE.emit(OptimizationRemarkMissed("vecz", "vecz", I) << RemarkMsg);
@@ -61,8 +69,8 @@ void emitVeczRemarkMissed(const Function *F, const Value *V, StringRef Msg) {
   }
 }
 
-void emitVeczRemarkMissed(const Function *F, StringRef Msg) {
-  emitVeczRemarkMissed(F, nullptr, Msg);
+void emitVeczRemarkMissed(const Function *F, StringRef Msg, StringRef Note) {
+  emitVeczRemarkMissed(F, nullptr, Msg, Note);
 }
 
 void emitVeczRemark(const Function *F, const Value *V, StringRef Msg) {

llvm/lib/SYCLNativeCPUUtils/compiler_passes/vecz/source/include/debugging.h

@@ -28,6 +28,7 @@
 #include <llvm/IR/Module.h>
 #include <llvm/IR/PassManager.h>
 #include <llvm/IR/Value.h>
+#include <llvm/Support/Error.h>
 #include <llvm/Support/raw_ostream.h>
 
 #include <cstdlib>
@@ -64,6 +65,12 @@ struct VeczFailResult {
   operator std::optional<T>() const {
     return std::nullopt;
   }
+
+  /// @brief For functions that return an llvm::Error
+  operator llvm::Error() const {
+    return llvm::make_error<llvm::StringError>("Unknown VeczFailResult",
+                                               llvm::inconvertibleErrorCode());
+  }
 };
 
 struct AnalysisFailResult : public internal::VeczFailResult {
@@ -175,13 +182,16 @@ struct AnalysisFailResult : public internal::VeczFailResult {
 /// @param[in] F The function in which we are currently working
 /// @param[in] V The value (can be `nullptr`) to be included in the message
 /// @param[in] Msg The main remark message text
+/// @param[in] Note An optional additional note to provide more context/info.
 void emitVeczRemarkMissed(const llvm::Function *F, const llvm::Value *V,
-                          llvm::StringRef Msg);
+                          llvm::StringRef Msg, llvm::StringRef Note = "");
 /// @brief Emit a RemarkMissed message
 ///
 /// @param[in] F The function in which we are currently working
 /// @param[in] Msg The main remark message text
-void emitVeczRemarkMissed(const llvm::Function *F, llvm::StringRef Msg);
+/// @param[in] Note An optional additional note to provide more context/info.
+void emitVeczRemarkMissed(const llvm::Function *F, llvm::StringRef Msg,
+                          llvm::StringRef Note = "");
 /// @brief Emit a Remark message
 ///
 /// @param[in] F The function in which we are currently working

llvm/lib/SYCLNativeCPUUtils/compiler_passes/vecz/source/transform/control_flow_conversion_pass.cpp

@@ -30,6 +30,7 @@
 #include <llvm/IR/Dominators.h>
 #include <llvm/IR/IRBuilder.h>
 #include <llvm/Support/Debug.h>
+#include <llvm/Support/Error.h>
 #include <llvm/Support/raw_ostream.h>
 
 #include <queue>
@@ -145,17 +146,19 @@ class ControlFlowConversionState::Impl : public ControlFlowConversionState {
   /// @brief Apply masks to basic blocks in the function, to prevent
   /// side-effects for inactive instances.
   ///
-  /// @return true if masks were applied successfully, false otherwise.
-  bool applyMasks();
+  /// @return llvm::Error::success if masks were applied successfully, an error
+  /// message explaining the failure otherwise.
+  Error applyMasks();
 
   /// @brief Apply a mask to the given basic block, to prevent side-effects for
   /// inactive instances.
   ///
   /// @param[in] BB Basic block to apply masks to.
   /// @param[in] mask Mask to apply.
   ///
-  /// @return true if masks were applied successfully, false otherwise.
-  bool applyMask(BasicBlock &BB, Value *mask);
+  /// @return llvm::Error::success if masks were applied successfully, an error
+  /// message explaining the failure otherwise.
+  Error applyMask(BasicBlock &BB, Value *mask);
 
   /// @brief Emit a call instructions to the masked version of the called
   /// function.
@@ -378,6 +381,14 @@ Instruction *copyExitMask(Value *mask, StringRef base, BasicBlock &BB) {
   VECZ_ERROR_IF(!mask, "Trying to copy exit mask with invalid arguments");
   return copyMask(mask, base + ".exit_mask", BB.getTerminator());
 }
+
+/// Wrap a string into an llvm::StringError, pointing to an instruction.
+static inline Error makeStringError(const Twine &message, Instruction &I) {
+  std::string helper_str = message.str();
+  raw_string_ostream helper_stream(helper_str);
+  helper_stream << " " << I;
+  return make_error<StringError>(helper_stream.str(), inconvertibleErrorCode());
+}
 }  // namespace
 
 ////////////////////////////////////////////////////////////////////////////////
@@ -538,8 +549,9 @@ bool ControlFlowConversionState::Impl::convertToDataFlow() {
                          "Could not generate masks for");
     return false;
   }
-  if (!applyMasks()) {
-    emitVeczRemarkMissed(&F, VU.scalarFunction(), "Could not apply masks for");
+  if (auto err = applyMasks()) {
+    emitVeczRemarkMissed(&F, VU.scalarFunction(), "Could not apply masks for",
+                         llvm::toString(std::move(err)));
     return false;
   }
 
@@ -1075,19 +1087,21 @@ bool ControlFlowConversionState::Impl::createCombinedLoopExitMask(
   return true;
 }
 
-bool ControlFlowConversionState::Impl::applyMasks() {
+Error ControlFlowConversionState::Impl::applyMasks() {
   for (auto &BB : F) {
     // Use masks with instructions that have side-effects.
     if (!DR->isUniform(BB) && !DR->isByAll(BB)) {
       auto *const entryMask = MaskInfos[&BB].entryMask;
       VECZ_ERROR_IF(!entryMask, "BasicBlock should have an entry mask");
-      VECZ_FAIL_IF(!applyMask(BB, entryMask));
+      if (auto err = applyMask(BB, entryMask)) {
+        return err;
+      }
     }
   }
-  return true;
+  return Error::success();
 }
 
-bool ControlFlowConversionState::Impl::applyMask(BasicBlock &BB, Value *mask) {
+Error ControlFlowConversionState::Impl::applyMask(BasicBlock &BB, Value *mask) {
   // Packetization hasn't happened yet so this better be a scalar 1 bit int.
   assert(mask->getType()->isIntegerTy(1) && "CFG mask type should be int1");
   // Map the unmasked instruction with the masked one.
@@ -1102,17 +1116,17 @@ bool ControlFlowConversionState::Impl::applyMask(BasicBlock &BB, Value *mask) {
     // Turn loads and stores into masked loads and stores.
     if (memOp && (memOp->isLoad() || memOp->isStore())) {
       if (!tryApplyMaskToMemOp(*memOp, mask, toDelete)) {
-        return false;
+        return makeStringError("Could not apply mask to MemOp", I);
       }
     } else if (auto *CI = dyn_cast<CallInst>(&I)) {
       // Turn calls into masked calls if possible.
       if (!applyMaskToCall(CI, mask, toDelete)) {
-        return false;
+        return makeStringError("Could not apply mask to call instruction", I);
       }
     } else if (I.isAtomic() && !isa<FenceInst>(&I)) {
       // We need to apply masks to atomic functions, but it is currently not
       // implemented. See CA-3294.
-      return false;
+      return makeStringError("Could not apply mask to atomic instruction", I);
     } else if (auto *branch = dyn_cast<BranchInst>(&I)) {
       // We have to be careful with infinite loops, because if they exist on a
       // divergent code path, they will always be entered and will hang the
@@ -1138,7 +1152,8 @@ bool ControlFlowConversionState::Impl::applyMask(BasicBlock &BB, Value *mask) {
     updateMaps(unmasked, masked);
     IRCleanup::deleteInstructionNow(unmasked);
   }
-  return true;
+
+  return Error::success();
 }
 
 CallInst *ControlFlowConversionState::Impl::emitMaskedVersion(CallInst *CI,

llvm/lib/SYCLNativeCPUUtils/compiler_passes/vecz/test/lit/llvm/diverging_atomic.ll

@@ -0,0 +1,45 @@
+; Copyright (C) Codeplay Software Limited
+;
+; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
+; Exceptions; you may not use this file except in compliance with the License.
+; You may obtain a copy of the License at
+;
+;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
+;
+; Unless required by applicable law or agreed to in writing, software
+; distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
+; WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
+; License for the specific language governing permissions and limitations
+; under the License.
+;
+; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+; RUN: veczc -w 4 -vecz-passes=cfg-convert,verify -S \
+; RUN:   --pass-remarks-missed=vecz < %s 2>&1 | FileCheck %s
+
+target triple = "spir64-unknown-unknown"
+target datalayout = "e-p:64:64:64-m:e-i64:64-f80:128-n8:16:32:64-S128"
+
+; CHECK: Vecz: Could not apply masks for function "kernel"
+; CHECK-NEXT: note: Could not apply mask to atomic instruction
+; CHECK-SAME:  %atomic = atomicrmw add ptr %arrayidx.in, i32 2 monotonic, align 4
+
+define spir_kernel void @kernel(ptr %in, ptr %out) {
+entry:
+  %gid = tail call i64 @__mux_get_global_id(i32 0)
+  %cmp = icmp eq i64 %gid, 0
+  br i1 %cmp, label %if.then, label %end
+
+if.then:
+  %arrayidx.in = getelementptr inbounds i32, ptr %in, i64 %gid
+  %atomic = atomicrmw add ptr %arrayidx.in, i32 2 monotonic, align 4
+  br label %end
+
+end:
+  %merge = phi i32 [ 0, %entry ], [ %atomic, %if.then ]
+  %arrayidx.out = getelementptr inbounds i32, ptr %out, i64 %gid
+  store i32 %merge, ptr %arrayidx.out, align 4
+  ret void
+}
+
+declare i64 @__mux_get_global_id(i32)