[Mesa-dev] [PATCH 35/56] i965/fs: Implement basic SPIR-V subgroup intrinsics

[Jason Ekstrand]

Reviewed-by: Samuel Iglesias Gonsálvez <siglesias at igalia.com>
Reviewed-by: Iago Toral Quiroga <itoral at igalia.com>
---
 src/intel/compiler/brw_fs_nir.cpp                |  8 ++++++++
 src/intel/compiler/brw_nir_lower_cs_intrinsics.c | 18 ++++++++++++++++++
 2 files changed, 26 insertions(+)

diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp
index 554d61d..651997b 100644
--- a/src/intel/compiler/brw_fs_nir.cpp
+++ b/src/intel/compiler/brw_fs_nir.cpp
@@ -4501,6 +4501,14 @@ fs_visitor::nir_emit_intrinsic(const fs_builder &bld, nir_intrinsic_instr *instr
       break;
    }
 
+   case nir_intrinsic_first_invocation: {
+      fs_reg tmp = bld.vgrf(BRW_REGISTER_TYPE_UD);
+      bld.exec_all().emit(SHADER_OPCODE_FIND_LIVE_CHANNEL, tmp);
+      bld.MOV(retype(dest, BRW_REGISTER_TYPE_UD),
+              fs_reg(component(tmp, 0)));
+      break;
+   }
+
    default:
       unreachable("unknown intrinsic");
    }
diff --git a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
index 66eef6b..bfbdea0 100644
--- a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
+++ b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
@@ -103,6 +103,24 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
          break;
       }
 
+      case nir_intrinsic_load_subgroup_id:
+         if (state->local_workgroup_size > 8)
+            continue;
+
+         /* For small workgroup sizes, we know subgroup_id will be zero */
+         sysval = nir_imm_int(b, 0);
+         break;
+
+      case nir_intrinsic_load_num_subgroups: {
+         unsigned local_workgroup_size =
+            nir->info.cs.local_size[0] * nir->info.cs.local_size[1] *
+            nir->info.cs.local_size[2];
+         unsigned num_subgroups =
+            DIV_ROUND_UP(local_workgroup_size, state->dispatch_width);
+         sysval = nir_imm_int(b, num_subgroups);
+         break;
+      }
+
       default:
          continue;
       }
-- 
2.5.0.400.gff86faf