[vecz] Provide more context when CFG conversion fails

This should help debugging (using `--pass-remarks-missed=vecz`) why
control-flow conversion failed to apply masks to the CFG.  The previous
diagnostics would only print the name of the function that couldn't be
converted, but not any more specific information.

This commit adds an extra level of information via a 'note', which is
optionally printed on the line after the main diagnostic.

This is not intended to be a full solution to better vecz diagnostics,
but a good first step.

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)