aboutsummaryrefslogtreecommitdiffstats
path: root/src/intel/compiler
diff options
context:
space:
mode:
Diffstat (limited to 'src/intel/compiler')
-rw-r--r--src/intel/compiler/brw_compiler.c1
-rw-r--r--src/intel/compiler/brw_fs.cpp5
-rw-r--r--src/intel/compiler/brw_nir_lower_cs_intrinsics.c113
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) {
generated by cgit v1.2.3 (git 2.39.1) at 2024-12-13 23:49:43 +0000