i965/fs: Implement basic SPIR-V subgroup intrinsics

Reviewed-by: Samuel Iglesias Gonsálvez <[email protected]>
Reviewed-by: Iago Toral Quiroga <[email protected]>

2 files changed, 26 insertions, 0 deletions

diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp
index 554d61d71af..651997bb6ff 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 66eef6be0a6..bfbdea0e8fa 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;
       }