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 <ian.d.romanick@intel.com> (v3) Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
This commit is contained in:
@@ -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, \
|
||||
|
@@ -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);
|
||||
}
|
||||
|
||||
|
@@ -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 *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);
|
||||
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));
|
||||
}
|
||||
}
|
||||
|
||||
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) {
|
||||
|
Reference in New Issue
Block a user