diff options
| -rw-r--r-- | src/intel/compiler/brw_fs_nir.cpp | 8 | ||||
| -rw-r--r-- | src/intel/compiler/brw_nir_lower_cs_intrinsics.c | 18 |
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; } |