Merge pull request #162 from frasercrmck/masked-group-ops

[vecz] Don't mask work-group collective operations

modules/compiler/spirv-ll/source/builder_core.cpp

@@ -6291,6 +6291,7 @@ void Builder::generateReduction(const T *op, const std::string &opName,
     // Always inline the function, this means for constant execution scope the
     // optimizer can remove the branches.
     reductionWrapper->addFnAttr(llvm::Attribute::AlwaysInline);
+    reductionWrapper->addFnAttr(llvm::Attribute::Convergent);
     // Restore the original insert point.
     IRBuilder.SetInsertPoint(insertBB, insertPoint);
   }

modules/compiler/vecz/source/transform/control_flow_conversion_pass.cpp

@@ -1292,7 +1292,8 @@ bool ControlFlowConversionState::Impl::applyMaskToCall(CallInst *CI,
   }
 
   // Builtins without side effects do not need to be masked.
-  auto const props = Ctx.builtins().analyzeBuiltin(*callee).properties;
+  auto const builtin = Ctx.builtins().analyzeBuiltin(*callee);
+  auto const props = builtin.properties;
   if (props & compiler::utils::eBuiltinPropertyNoSideEffects) {
     LLVM_DEBUG(dbgs() << "vecz-cf: Called function is an pure builtin\n");
     return true;
@@ -1313,6 +1314,18 @@ bool ControlFlowConversionState::Impl::applyMaskToCall(CallInst *CI,
         dbgs() << "vecz-cf: Called function does not have any side-effects\n");
     return true;
   }
+  // We don't want to mask work-group collective builtins, because they are
+  // barriers (see above). This should actually be a rare situation, as these
+  // builtins are required to be uniform/convergent and so either all
+  // work-items or no work-items should hit them. Most of the time, this
+  // situation relies on the vectorizer failing to trace the branch flow and
+  // failing to realize the conditions are in fact uniform.
+  if (auto info = Ctx.builtins().isMuxGroupCollective(builtin.ID);
+      info && info->isWorkGroupScope()) {
+    LLVM_DEBUG(
+        dbgs() << "vecz-cf: Called function is a work-group collective\n");
+    return true;
+  }
 
   // Create the new function and replace the old one with it
   CallInst *newCI = emitMaskedVersion(CI, mask);

modules/compiler/vecz/test/lit/llvm/masked_group_collective.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 -vecz-passes="cfg-convert" -vecz-simd-width=4 -S < %s | FileCheck %s
+
+target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target triple = "spir64-unknown-unknown"
+
+declare i64 @__mux_get_local_id()
+declare i32 @__mux_work_group_scan_inclusive_smax_i32(i32, i32)
+
+; CHECK-LABEL: define spir_kernel void @__vecz_v4_foo()
+; CHECK-NOT: @__vecz_b_masked___mux_work_group_scan_inclusive_smax_i32
+define spir_kernel void @foo() {
+entry:
+  %0 = call i64 @__mux_get_local_id()
+  br i1 false, label %for.body.i11, label %if.end.i105.i
+
+for.body.i11:
+  %1 = icmp slt i64 %0, 0
+  br i1 %1, label %if.end.i13, label %if.end.i13
+
+if.end.i13:
+  br i1 false, label %exit, label %if.end.i105.i
+
+if.end.i105.i:
+  %2 = call i32 @__mux_work_group_scan_inclusive_smax_i32(i32 0, i32 0)
+  br label %exit
+
+exit:
+  ret void
+}