From 3ee3024804f9817dfa4f9ee4fa3d6b963a84c9cb Mon Sep 17 00:00:00 2001 From: Caio Marcelo de Oliveira Filho Date: Wed, 27 Mar 2019 15:07:59 -0700 Subject: [PATCH] intel/fs: Add support for CS to group invocations in quads When using quads, instead of mapping the elements to the next 4 local invocation indices, we map the two next in the "current" row and two next in the "next row". A side effect is that a thread will execute the indices in a different order. We now perform the lowering of both local invocation ID and index together -- and don't rely anymore on lowering done by nir_lower_system_values. That is convenient when doing the math for quads, because we need X and Y to get the right invocation index. When the pass progresses, fold the constants and clean up to reduce the noise from the indexing math. This implements the derivative_group_quadsNV semantics from NV_compute_shader_derivatives. v2: Take subgroup_id into account, otherwise only values in the first subgroup would be used. (Jason) v3: Calculate invocation index and ID together, to avoid duplicating some math in the quads case when both index and ID are used. (Jason) v4: Don't call cleanup passes as part of the lowering, let that to the call site. (Jason) Change calculation to use less instructions. (Jason) Reviewed-by: Ian Romanick (v3) Reviewed-by: Jason Ekstrand --- src/intel/compiler/brw_compiler.c | 1 - src/intel/compiler/brw_fs.cpp | 5 + .../compiler/brw_nir_lower_cs_intrinsics.c | 113 +++++++++++++++--- 3 files changed, 103 insertions(+), 16 deletions(-) diff --git a/src/intel/compiler/brw_compiler.c b/src/intel/compiler/brw_compiler.c index a3a0a393fad..d3f8c7ef1e0 100644 --- a/src/intel/compiler/brw_compiler.c +++ b/src/intel/compiler/brw_compiler.c @@ -45,7 +45,6 @@ .lower_flrp64 = true, \ .lower_isign = true, \ .lower_ldexp = true, \ - .lower_cs_local_id_from_index = true, \ .lower_device_index_to_zero = true, \ .native_integers = true, \ .use_interpolated_input_intrinsics = true, \ diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp index 0c2439d9daf..a637ee3422f 100644 --- a/src/intel/compiler/brw_fs.cpp +++ b/src/intel/compiler/brw_fs.cpp @@ -8017,6 +8017,11 @@ compile_cs_to_nir(const struct brw_compiler *compiler, nir_shader *shader = nir_shader_clone(mem_ctx, src_shader); shader = brw_nir_apply_sampler_key(shader, compiler, &key->tex, true); brw_nir_lower_cs_intrinsics(shader, dispatch_width); + + /* Clean up after the local index and ID calculations. */ + nir_opt_constant_folding(shader); + nir_opt_dce(shader); + return brw_postprocess_nir(shader, compiler, true); } diff --git a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c index fab5edc893f..6180e6a706b 100644 --- a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c +++ b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c @@ -41,6 +41,10 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state, nir_builder *b = &state->builder; nir_shader *nir = state->nir; + /* Reuse calculated values inside the block. */ + nir_ssa_def *local_index = NULL; + nir_ssa_def *local_id = NULL; + nir_foreach_instr_safe(instr, block) { if (instr->type != nir_instr_type_intrinsic) continue; @@ -51,22 +55,91 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state, nir_ssa_def *sysval; switch (intrinsic->intrinsic) { - case nir_intrinsic_load_local_invocation_index: { - /* We construct the local invocation index from: - * - * gl_LocalInvocationIndex = - * cs_thread_local_id + subgroup_invocation; - */ - nir_ssa_def *subgroup_id; - if (state->local_workgroup_size <= state->dispatch_width) - subgroup_id = nir_imm_int(b, 0); - else - subgroup_id = nir_load_subgroup_id(b); + case nir_intrinsic_load_local_invocation_index: + case nir_intrinsic_load_local_invocation_id: { + /* First time we are using those, so let's calculate them. */ + if (!local_index) { + assert(!local_id); + + nir_ssa_def *subgroup_id; + if (state->local_workgroup_size <= state->dispatch_width) + subgroup_id = nir_imm_int(b, 0); + else + subgroup_id = nir_load_subgroup_id(b); + + nir_ssa_def *thread_local_id = + nir_imul_imm(b, subgroup_id, state->dispatch_width); + nir_ssa_def *channel = nir_load_subgroup_invocation(b); + nir_ssa_def *linear = nir_iadd(b, channel, thread_local_id); + + nir_ssa_def *size_x = nir_imm_int(b, nir->info.cs.local_size[0]); + nir_ssa_def *size_y = nir_imm_int(b, nir->info.cs.local_size[1]); + + /* The local invocation index and ID must respect the following + * + * gl_LocalInvocationID.x = + * gl_LocalInvocationIndex % gl_WorkGroupSize.x; + * gl_LocalInvocationID.y = + * (gl_LocalInvocationIndex / gl_WorkGroupSize.x) % + * gl_WorkGroupSize.y; + * gl_LocalInvocationID.z = + * (gl_LocalInvocationIndex / + * (gl_WorkGroupSize.x * gl_WorkGroupSize.y)) % + * gl_WorkGroupSize.z; + * + * However, the final % gl_WorkGroupSize.z does nothing unless we + * accidentally end up with a gl_LocalInvocationIndex that is too + * large so it can safely be omitted. + */ + + if (state->nir->info.cs.derivative_group != DERIVATIVE_GROUP_QUADS) { + /* If we are not grouping in quads, just set the local invocatio + * index linearly, and calculate local invocation ID from that. + */ + local_index = linear; + + nir_ssa_def *id_x, *id_y, *id_z; + id_x = nir_umod(b, local_index, size_x); + id_y = nir_umod(b, nir_udiv(b, local_index, size_x), size_y); + id_z = nir_udiv(b, local_index, nir_imul(b, size_x, size_y)); + local_id = nir_vec3(b, id_x, id_y, id_z); + } else { + /* For quads, first we figure out the 2x2 grid the invocation + * belongs to -- treating extra Z layers as just more rows. + * Then map that into local invocation ID (trivial) and local + * invocation index. Skipping Z simplify index calculation. + */ + + nir_ssa_def *one = nir_imm_int(b, 1); + nir_ssa_def *double_size_x = nir_ishl(b, size_x, one); + + /* ID within a pair of rows, where each group of 4 is 2x2 quad. */ + nir_ssa_def *row_pair_id = nir_umod(b, linear, double_size_x); + nir_ssa_def *y_row_pairs = nir_udiv(b, linear, double_size_x); + + nir_ssa_def *x = + nir_ior(b, + nir_iand(b, row_pair_id, one), + nir_iand(b, nir_ishr(b, row_pair_id, one), + nir_imm_int(b, 0xfffffffe))); + nir_ssa_def *y = + nir_ior(b, + nir_ishl(b, y_row_pairs, one), + nir_iand(b, nir_ishr(b, row_pair_id, one), one)); + + local_id = nir_vec3(b, x, + nir_umod(b, y, size_y), + nir_udiv(b, y, size_y)); + local_index = nir_iadd(b, x, nir_imul(b, y, size_x)); + } + } - nir_ssa_def *thread_local_id = - nir_imul(b, subgroup_id, nir_imm_int(b, state->dispatch_width)); - nir_ssa_def *channel = nir_load_subgroup_invocation(b); - sysval = nir_iadd(b, channel, thread_local_id); + assert(local_id); + assert(local_index); + if (intrinsic->intrinsic == nir_intrinsic_load_local_invocation_id) + sysval = local_id; + else + sysval = local_index; break; } @@ -125,10 +198,20 @@ brw_nir_lower_cs_intrinsics(nir_shader *nir, memset(&state, 0, sizeof(state)); state.nir = nir; state.dispatch_width = dispatch_width; + + assert(!nir->info.cs.local_size_variable); state.local_workgroup_size = nir->info.cs.local_size[0] * nir->info.cs.local_size[1] * nir->info.cs.local_size[2]; + /* Constraints from NV_compute_shader_derivatives. */ + if (nir->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) { + assert(nir->info.cs.local_size[0] % 2 == 0); + assert(nir->info.cs.local_size[1] % 2 == 0); + } else if (nir->info.cs.derivative_group == DERIVATIVE_GROUP_LINEAR) { + assert(state.local_workgroup_size % 4 == 0); + } + do { state.progress = false; nir_foreach_function(function, nir) { -- 2.30.2