From 974daec495eae05b3c3179cd6c131a65ff2efcc7 Mon Sep 17 00:00:00 2001 From: Jason Ekstrand Date: Mon, 21 Aug 2017 22:17:37 -0700 Subject: i965/fs: Implement basic SPIR-V subgroup intrinsics MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Reviewed-by: Samuel Iglesias Gonsálvez Reviewed-by: Iago Toral Quiroga --- 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 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; } -- cgit v1.2.3