[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:

I'm surprised SPIR-V verifier doesn't complain that this op cannot have a Workgroup scope.

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