From 74be809237df22d48a835f354cdba7b135382af2 Mon Sep 17 00:00:00 2001 From: Caio Oliveira Date: Fri, 30 Aug 2024 08:44:23 -0700 Subject: [PATCH] compiler: Allow derivative_group to be used for all stages in shader_info MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit These will now also be used by stages that have workgroups. Reviewed-by: Daniel Schürmann Part-of: --- src/amd/llvm/ac_nir_to_llvm.c | 4 ++-- src/compiler/glsl/linker.cpp | 12 ++++++------ src/compiler/nir/nir.c | 2 +- src/compiler/nir/nir_builder.h | 2 +- src/compiler/nir/nir_lower_system_values.c | 6 +++--- src/compiler/nir/nir_opt_algebraic.py | 2 +- src/compiler/nir/nir_print.c | 4 ++-- src/compiler/shader_info.h | 12 ++++++------ src/compiler/spirv/spirv_to_nir.c | 4 ++-- src/gallium/drivers/radeonsi/si_compute.c | 8 ++++---- src/gallium/drivers/radeonsi/si_shader_nir.c | 6 +++--- .../drivers/zink/nir_to_spirv/nir_to_spirv.c | 6 +++--- src/intel/compiler/brw_nir_lower_cs_intrinsics.c | 15 ++++++--------- .../compiler/elk/elk_nir_lower_cs_intrinsics.c | 6 +++--- src/mesa/main/compute.c | 4 ++-- src/nouveau/compiler/nak_nir.c | 4 ++-- 16 files changed, 47 insertions(+), 50 deletions(-) diff --git a/src/amd/llvm/ac_nir_to_llvm.c b/src/amd/llvm/ac_nir_to_llvm.c index fed3bd80054..29d0042507f 100644 --- a/src/amd/llvm/ac_nir_to_llvm.c +++ b/src/amd/llvm/ac_nir_to_llvm.c @@ -1479,7 +1479,7 @@ static LLVMValueRef build_tex_intrinsic(struct ac_nir_context *ctx, const nir_te case nir_texop_tex: if (ctx->stage != MESA_SHADER_FRAGMENT && (!gl_shader_stage_is_compute(ctx->stage) || - ctx->info->cs.derivative_group == DERIVATIVE_GROUP_NONE)) { + ctx->info->derivative_group == DERIVATIVE_GROUP_NONE)) { assert(!args->lod); args->level_zero = true; } @@ -1514,7 +1514,7 @@ static LLVMValueRef build_tex_intrinsic(struct ac_nir_context *ctx, const nir_te args->attributes = AC_ATTR_INVARIANT_LOAD; bool cs_derivs = - gl_shader_stage_is_compute(ctx->stage) && ctx->info->cs.derivative_group != DERIVATIVE_GROUP_NONE; + gl_shader_stage_is_compute(ctx->stage) && ctx->info->derivative_group != DERIVATIVE_GROUP_NONE; if (ctx->stage == MESA_SHADER_FRAGMENT || cs_derivs) { /* Prevent texture instructions with implicit derivatives from being * sinked into branches. */ diff --git a/src/compiler/glsl/linker.cpp b/src/compiler/glsl/linker.cpp index 6993fe0a6c4..9a8fa9fe864 100644 --- a/src/compiler/glsl/linker.cpp +++ b/src/compiler/glsl/linker.cpp @@ -1394,7 +1394,7 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog, gl_prog->info.workgroup_size_variable = false; - gl_prog->info.cs.derivative_group = DERIVATIVE_GROUP_NONE; + gl_prog->info.derivative_group = DERIVATIVE_GROUP_NONE; /* From the ARB_compute_shader spec, in the section describing local size * declarations: @@ -1442,13 +1442,13 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog, enum gl_derivative_group group = shader->info.Comp.DerivativeGroup; if (group != DERIVATIVE_GROUP_NONE) { - if (gl_prog->info.cs.derivative_group != DERIVATIVE_GROUP_NONE && - gl_prog->info.cs.derivative_group != group) { + if (gl_prog->info.derivative_group != DERIVATIVE_GROUP_NONE && + gl_prog->info.derivative_group != group) { linker_error(prog, "compute shader defined with conflicting " "derivative groups\n"); return; } - gl_prog->info.cs.derivative_group = group; + gl_prog->info.derivative_group = group; } } @@ -1463,7 +1463,7 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog, return; } - if (gl_prog->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) { + if (gl_prog->info.derivative_group == DERIVATIVE_GROUP_QUADS) { if (gl_prog->info.workgroup_size[0] % 2 != 0) { linker_error(prog, "derivative_group_quadsNV must be used with a " "local group size whose first dimension " @@ -1476,7 +1476,7 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog, "is a multiple of 2\n"); return; } - } else if (gl_prog->info.cs.derivative_group == DERIVATIVE_GROUP_LINEAR) { + } else if (gl_prog->info.derivative_group == DERIVATIVE_GROUP_LINEAR) { if ((gl_prog->info.workgroup_size[0] * gl_prog->info.workgroup_size[1] * gl_prog->info.workgroup_size[2]) % 4 != 0) { diff --git a/src/compiler/nir/nir.c b/src/compiler/nir/nir.c index 3f52baee24b..d7755958927 100644 --- a/src/compiler/nir/nir.c +++ b/src/compiler/nir/nir.c @@ -2199,7 +2199,7 @@ nir_shader_supports_implicit_lod(nir_shader *shader) { return (shader->info.stage == MESA_SHADER_FRAGMENT || (shader->info.stage == MESA_SHADER_COMPUTE && - shader->info.cs.derivative_group != DERIVATIVE_GROUP_NONE)); + shader->info.derivative_group != DERIVATIVE_GROUP_NONE)); } nir_intrinsic_op diff --git a/src/compiler/nir/nir_builder.h b/src/compiler/nir/nir_builder.h index 90213f9dbaf..f3ad4586f74 100644 --- a/src/compiler/nir/nir_builder.h +++ b/src/compiler/nir/nir_builder.h @@ -2012,7 +2012,7 @@ nir_build_deriv(nir_builder *b, nir_def *x, nir_op alu, nir_intrinsic_op intrin) * move this to glsl-to-nir. */ if (b->shader->info.stage == MESA_SHADER_COMPUTE && - b->shader->info.cs.derivative_group == DERIVATIVE_GROUP_NONE) { + b->shader->info.derivative_group == DERIVATIVE_GROUP_NONE) { return nir_imm_zero(b, x->num_components, x->bit_size); } diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c index a3833b233f0..9293406ad13 100644 --- a/src/compiler/nir/nir_lower_system_values.c +++ b/src/compiler/nir/nir_lower_system_values.c @@ -539,7 +539,7 @@ lower_compute_system_value_instr(nir_builder *b, return lower_id_to_index(b, local_index, local_size, bit_size); } if (options && options->shuffle_local_ids_for_quad_derivatives && - b->shader->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS && + b->shader->info.derivative_group == DERIVATIVE_GROUP_QUADS && _mesa_set_search(state->lower_once_list, instr) == NULL) { nir_def *ids = nir_load_local_invocation_id(b); _mesa_set_add(state->lower_once_list, ids->parent_instr); @@ -803,8 +803,8 @@ nir_lower_compute_system_values(nir_shader *shader, /* Update this so as not to lower it again. */ if (options && options->shuffle_local_ids_for_quad_derivatives && - shader->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) - shader->info.cs.derivative_group = DERIVATIVE_GROUP_LINEAR; + shader->info.derivative_group == DERIVATIVE_GROUP_QUADS) + shader->info.derivative_group = DERIVATIVE_GROUP_LINEAR; return progress; } diff --git a/src/compiler/nir/nir_opt_algebraic.py b/src/compiler/nir/nir_opt_algebraic.py index 4e4f5f2309f..2a8f01dc7e2 100644 --- a/src/compiler/nir/nir_opt_algebraic.py +++ b/src/compiler/nir/nir_opt_algebraic.py @@ -2887,7 +2887,7 @@ for op in ['fadd', 'fmul', 'fmulz', 'iadd', 'imul']: for op in ['fddx', 'fddx_fine', 'fddx_coarse', 'fddy', 'fddy_fine', 'fddy_coarse']: optimizations += [ - ((op, 'a'), 0.0, 'info->stage == MESA_SHADER_COMPUTE && info->cs.derivative_group == DERIVATIVE_GROUP_NONE') + ((op, 'a'), 0.0, 'info->stage == MESA_SHADER_COMPUTE && info->derivative_group == DERIVATIVE_GROUP_NONE') ] # Some optimizations for ir3-specific instructions. diff --git a/src/compiler/nir/nir_print.c b/src/compiler/nir/nir_print.c index 5320b4e0b9e..a5a0e30acbc 100644 --- a/src/compiler/nir/nir_print.c +++ b/src/compiler/nir/nir_print.c @@ -1648,7 +1648,7 @@ print_intrinsic_instr(nir_intrinsic_instr *instr, print_state *state) nir_foreach_variable_with_modes(var, state->shader, var_mode) { if (!var->name) continue; - + bool match; if (instr->intrinsic == nir_intrinsic_load_uniform) { match = var->data.driver_location == nir_intrinsic_base(instr); @@ -2570,6 +2570,7 @@ print_shader_info(const struct shader_info *info, FILE *fp) print_nz_bool(fp, "flrp_lowered", info->flrp_lowered); print_nz_bool(fp, "io_lowered", info->io_lowered); print_nz_bool(fp, "writes_memory", info->writes_memory); + print_nz_unsigned(fp, "derivative_group", info->derivative_group); switch (info->stage) { case MESA_SHADER_VERTEX: @@ -2649,7 +2650,6 @@ print_shader_info(const struct shader_info *info, FILE *fp) info->cs.workgroup_size_hint[1], info->cs.workgroup_size_hint[2]); print_nz_unsigned(fp, "user_data_components_amd", info->cs.user_data_components_amd); - print_nz_unsigned(fp, "derivative_group", info->cs.derivative_group); fprintf(fp, "ptr_size: %u\n", info->cs.ptr_size); break; diff --git a/src/compiler/shader_info.h b/src/compiler/shader_info.h index 62206d8616d..8de52485cf6 100644 --- a/src/compiler/shader_info.h +++ b/src/compiler/shader_info.h @@ -285,6 +285,12 @@ typedef struct shader_info { */ bool use_legacy_math_rules; + /* + * Arrangement of invocations used to calculate derivatives in + * compute/task/mesh shaders. From KHR_compute_shader_derivatives. + */ + enum gl_derivative_group derivative_group:2; + union { struct { /* Which inputs are doubles */ @@ -441,12 +447,6 @@ typedef struct shader_info { uint8_t user_data_components_amd:4; - /* - * Arrangement of invocations used to calculate derivatives in a compute - * shader. From NV_compute_shader_derivatives. - */ - enum gl_derivative_group derivative_group:2; - /* * If the shader might run with shared mem on top of `shared_size`. */ diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c index 62406018ab3..ba709f9310f 100644 --- a/src/compiler/spirv/spirv_to_nir.c +++ b/src/compiler/spirv/spirv_to_nir.c @@ -5274,12 +5274,12 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point, case SpvExecutionModeDerivativeGroupQuadsNV: vtn_assert(b->shader->info.stage == MESA_SHADER_COMPUTE); - b->shader->info.cs.derivative_group = DERIVATIVE_GROUP_QUADS; + b->shader->info.derivative_group = DERIVATIVE_GROUP_QUADS; break; case SpvExecutionModeDerivativeGroupLinearNV: vtn_assert(b->shader->info.stage == MESA_SHADER_COMPUTE); - b->shader->info.cs.derivative_group = DERIVATIVE_GROUP_LINEAR; + b->shader->info.derivative_group = DERIVATIVE_GROUP_LINEAR; break; case SpvExecutionModePixelInterlockOrderedEXT: diff --git a/src/gallium/drivers/radeonsi/si_compute.c b/src/gallium/drivers/radeonsi/si_compute.c index ee249d94228..22a64784fe4 100644 --- a/src/gallium/drivers/radeonsi/si_compute.c +++ b/src/gallium/drivers/radeonsi/si_compute.c @@ -999,7 +999,7 @@ static void si_emit_dispatch_packets(struct si_context *sctx, const struct pipe_ } /* Thread tiling within a workgroup. */ - switch (sctx->cs_shader_state.program->shader.selector->info.base.cs.derivative_group) { + switch (sctx->cs_shader_state.program->shader.selector->info.base.derivative_group) { case DERIVATIVE_GROUP_LINEAR: break; case DERIVATIVE_GROUP_QUADS: @@ -1233,10 +1233,10 @@ static void si_launch_grid(struct pipe_context *ctx, const struct pipe_grid_info NULL); } } - + if (u_trace_perfetto_active(&sctx->ds.trace_context)) trace_si_begin_compute(&sctx->trace); - + if (sctx->bo_list_add_all_compute_resources) si_compute_resources_add_all_to_bo_list(sctx); @@ -1310,7 +1310,7 @@ static void si_launch_grid(struct pipe_context *ctx, const struct pipe_grid_info if (u_trace_perfetto_active(&sctx->ds.trace_context)) trace_si_end_compute(&sctx->trace, info->grid[0], info->grid[1], info->grid[2]); - + if (cs_regalloc_hang) { sctx->flags |= SI_CONTEXT_CS_PARTIAL_FLUSH; si_mark_atom_dirty(sctx, &sctx->atoms.s.cache_flush); diff --git a/src/gallium/drivers/radeonsi/si_shader_nir.c b/src/gallium/drivers/radeonsi/si_shader_nir.c index 3c234e59693..7834eff1039 100644 --- a/src/gallium/drivers/radeonsi/si_shader_nir.c +++ b/src/gallium/drivers/radeonsi/si_shader_nir.c @@ -341,16 +341,16 @@ static void si_lower_nir(struct si_screen *sscreen, struct nir_shader *nir) * divisible by 2. */ options.lower_local_invocation_index = - nir->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS || + nir->info.derivative_group == DERIVATIVE_GROUP_QUADS || (sscreen->info.gfx_level >= GFX12 && - nir->info.cs.derivative_group == DERIVATIVE_GROUP_NONE && + nir->info.derivative_group == DERIVATIVE_GROUP_NONE && (nir->info.workgroup_size_variable || (nir->info.workgroup_size[0] % 2 == 0 && nir->info.workgroup_size[1] % 2 == 0))); NIR_PASS_V(nir, nir_lower_compute_system_values, &options); /* Gfx12 supports this in hw. */ if (sscreen->info.gfx_level < GFX12 && - nir->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) { + nir->info.derivative_group == DERIVATIVE_GROUP_QUADS) { nir_opt_cse(nir); /* CSE load_local_invocation_id */ memset(&options, 0, sizeof(options)); options.shuffle_local_ids_for_quad_derivatives = true; diff --git a/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c b/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c index f1148cbb36f..e5e4d9c04f1 100644 --- a/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c +++ b/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c @@ -4763,12 +4763,12 @@ nir_to_spirv(struct nir_shader *s, const struct zink_shader_info *sinfo, const s spirv_builder_emit_specid(&ctx.builder, ctx.shared_mem_size, ZINK_VARIABLE_SHARED_MEM); spirv_builder_emit_name(&ctx.builder, ctx.shared_mem_size, "variable_shared_mem"); } - if (s->info.cs.derivative_group) { + if (s->info.derivative_group) { SpvCapability caps[] = { 0, SpvCapabilityComputeDerivativeGroupQuadsNV, SpvCapabilityComputeDerivativeGroupLinearNV }; SpvExecutionMode modes[] = { 0, SpvExecutionModeDerivativeGroupQuadsNV, SpvExecutionModeDerivativeGroupLinearNV }; spirv_builder_emit_extension(&ctx.builder, "SPV_NV_compute_shader_derivatives"); - spirv_builder_emit_cap(&ctx.builder, caps[s->info.cs.derivative_group]); - spirv_builder_emit_exec_mode(&ctx.builder, entry_point, modes[s->info.cs.derivative_group]); + spirv_builder_emit_cap(&ctx.builder, caps[s->info.derivative_group]); + spirv_builder_emit_exec_mode(&ctx.builder, entry_point, modes[s->info.derivative_group]); ctx.explicit_lod = false; } break; diff --git a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c index 8722a89cca0..c507b0f8d10 100644 --- a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c +++ b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c @@ -27,7 +27,6 @@ struct lower_intrinsics_state { nir_shader *nir; nir_function_impl *impl; - enum gl_derivative_group derivative_group; bool progress; bool hw_generated_local_id; nir_builder builder; @@ -62,7 +61,7 @@ compute_local_index_id(struct lower_intrinsics_state *state, nir_intrinsic_instr } if (state->hw_generated_local_id) { - assert(state->derivative_group != DERIVATIVE_GROUP_QUADS); + assert(nir->info.derivative_group != DERIVATIVE_GROUP_QUADS); nir_def *local_id_vec = nir_load_local_invocation_id(b); nir_def *local_id[3] = { nir_channel(b, local_id_vec, 0), @@ -132,7 +131,7 @@ compute_local_index_id(struct lower_intrinsics_state *state, nir_intrinsic_instr */ nir_def *id_x, *id_y, *id_z; - switch (state->derivative_group) { + switch (nir->info.derivative_group) { case DERIVATIVE_GROUP_NONE: if (nir->info.num_images == 0 && nir->info.num_textures == 0) { @@ -333,16 +332,14 @@ brw_nir_lower_cs_intrinsics(nir_shader *nir, struct lower_intrinsics_state state = { .nir = nir, .hw_generated_local_id = false, - .derivative_group = gl_shader_stage_is_compute(nir->info.stage) ? - nir->info.cs.derivative_group : DERIVATIVE_GROUP_NONE, }; /* Constraints from NV_compute_shader_derivatives. */ if (!nir->info.workgroup_size_variable) { - if (state.derivative_group == DERIVATIVE_GROUP_QUADS) { + if (nir->info.derivative_group == DERIVATIVE_GROUP_QUADS) { assert(nir->info.workgroup_size[0] % 2 == 0); assert(nir->info.workgroup_size[1] % 2 == 0); - } else if (state.derivative_group == DERIVATIVE_GROUP_LINEAR) { + } else if (nir->info.derivative_group == DERIVATIVE_GROUP_LINEAR) { ASSERTED unsigned workgroup_size = nir->info.workgroup_size[0] * nir->info.workgroup_size[1] * @@ -353,7 +350,7 @@ brw_nir_lower_cs_intrinsics(nir_shader *nir, if (devinfo->verx10 >= 125 && prog_data && nir->info.stage == MESA_SHADER_COMPUTE && - state.derivative_group != DERIVATIVE_GROUP_QUADS && + nir->info.derivative_group != DERIVATIVE_GROUP_QUADS && !nir->info.workgroup_size_variable && util_is_power_of_two_nonzero(nir->info.workgroup_size[0]) && util_is_power_of_two_nonzero(nir->info.workgroup_size[1])) { @@ -362,7 +359,7 @@ brw_nir_lower_cs_intrinsics(nir_shader *nir, /* TODO: more heuristics about 1D/SLM access vs. 2D access */ bool linear = - state.derivative_group == DERIVATIVE_GROUP_LINEAR || + nir->info.derivative_group == DERIVATIVE_GROUP_LINEAR || BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_LOCAL_INVOCATION_INDEX) || (nir->info.workgroup_size[1] == 1 && diff --git a/src/intel/compiler/elk/elk_nir_lower_cs_intrinsics.c b/src/intel/compiler/elk/elk_nir_lower_cs_intrinsics.c index 4c8d0165f72..3d9cd99cd24 100644 --- a/src/intel/compiler/elk/elk_nir_lower_cs_intrinsics.c +++ b/src/intel/compiler/elk/elk_nir_lower_cs_intrinsics.c @@ -75,7 +75,7 @@ compute_local_index_id(nir_builder *b, */ nir_def *id_x, *id_y, *id_z; - switch (nir->info.cs.derivative_group) { + switch (nir->info.derivative_group) { case DERIVATIVE_GROUP_NONE: if (nir->info.num_images == 0 && nir->info.num_textures == 0) { @@ -308,10 +308,10 @@ elk_nir_lower_cs_intrinsics(nir_shader *nir, /* Constraints from NV_compute_shader_derivatives. */ if (gl_shader_stage_is_compute(nir->info.stage) && !nir->info.workgroup_size_variable) { - if (nir->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) { + if (nir->info.derivative_group == DERIVATIVE_GROUP_QUADS) { assert(nir->info.workgroup_size[0] % 2 == 0); assert(nir->info.workgroup_size[1] % 2 == 0); - } else if (nir->info.cs.derivative_group == DERIVATIVE_GROUP_LINEAR) { + } else if (nir->info.derivative_group == DERIVATIVE_GROUP_LINEAR) { ASSERTED unsigned workgroup_size = nir->info.workgroup_size[0] * nir->info.workgroup_size[1] * diff --git a/src/mesa/main/compute.c b/src/mesa/main/compute.c index e601f19d533..299ba097f65 100644 --- a/src/mesa/main/compute.c +++ b/src/mesa/main/compute.c @@ -197,7 +197,7 @@ validate_DispatchComputeGroupSizeARB(struct gl_context *ctx, * of , , and is not a multiple * of four." */ - if (prog->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS && + if (prog->info.derivative_group == DERIVATIVE_GROUP_QUADS && ((info->block[0] & 1) || (info->block[1] & 1))) { _mesa_error(ctx, GL_INVALID_VALUE, "glDispatchComputeGroupSizeARB(derivative_group_quadsNV " @@ -206,7 +206,7 @@ validate_DispatchComputeGroupSizeARB(struct gl_context *ctx, return GL_FALSE; } - if (prog->info.cs.derivative_group == DERIVATIVE_GROUP_LINEAR && + if (prog->info.derivative_group == DERIVATIVE_GROUP_LINEAR && total_invocations & 3) { _mesa_error(ctx, GL_INVALID_VALUE, "glDispatchComputeGroupSizeARB(derivative_group_linearNV " diff --git a/src/nouveau/compiler/nak_nir.c b/src/nouveau/compiler/nak_nir.c index 4e5f827b282..641ae34c006 100644 --- a/src/nouveau/compiler/nak_nir.c +++ b/src/nouveau/compiler/nak_nir.c @@ -306,7 +306,7 @@ nak_nir_lower_subgroup_id_intrin(nir_builder *b, nir_intrinsic_instr *intrin, if (nak_nir_workgroup_has_one_subgroup(b->shader)) { num_subgroups = nir_imm_int(b, 1); } else { - assert(b->shader->info.cs.derivative_group == DERIVATIVE_GROUP_NONE); + assert(b->shader->info.derivative_group == DERIVATIVE_GROUP_NONE); nir_def *workgroup_size = nir_load_workgroup_size(b); workgroup_size = @@ -327,7 +327,7 @@ nak_nir_lower_subgroup_id_intrin(nir_builder *b, nir_intrinsic_instr *intrin, if (nak_nir_workgroup_has_one_subgroup(b->shader)) { subgroup_id = nir_imm_int(b, 0); } else { - assert(b->shader->info.cs.derivative_group == DERIVATIVE_GROUP_NONE); + assert(b->shader->info.derivative_group == DERIVATIVE_GROUP_NONE); nir_def *invocation_index = nir_load_local_invocation_index(b); nir_def *subgroup_size = nir_load_subgroup_size(b);