[Mlir-commits] [mlir] [mlir][SPIR-V] Support bitwise and/or/xor in gpu reduction lowering (PR #202351)

[Igor Wodiany]

================
@@ -675,6 +675,73 @@ gpu.module @kernels {
 
 // -----
 
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
+} {
+
+gpu.module @kernels {
+  // CHECK-LABEL:  spirv.func @test
+  //  CHECK-SAME: (%[[ARG:.*]]: i32)
+  gpu.func @test(%arg : i32) kernel
+    attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
+    // CHECK: %{{.*}} = spirv.GroupNonUniformBitwiseAnd <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
----------------
IgWod wrote:

That's a good catch! What I actually meant was: why doesn't MLIR verifier fire here as well? Now, looking at the op definition it looks like it lists both `Workgroup` and `Subgroup`:

https://github.com/llvm/llvm-project/blob/b443896c13aded8b40d7bae4a5a9adbc96fd0d31/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td#L1087.

I wonder if that is something that was allowed at some point or the op definition has just been wrong. We should probably correct that - both description and verification.

https://github.com/llvm/llvm-project/pull/202351