AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

panfrost: Lower sysvals in GL

Browse Source 
Drop the backend compiler sysval handling in favour of the pass in the GL
driver, bringing us into compliance with Ekstrand's rule.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/20906>
This commit is contained in:
Alyssa Rosenzweig
2023-02-10 21:22:49 -05:00
committed by Marge Bot
parent 455a76f10a
commit ffb9919c2f
13 changed files with 10 additions and 387 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/gallium/drivers/panfrost/pan_shader.c
+4 -1
View File
@@ -89,7 +89,6 @@ panfrost_shader_compile(struct panfrost_screen *screen, const nir_shader *ir,
struct panfrost_compile_inputs inputs = {
.debug = dbg,
.gpu_id = dev->gpu_id,
.fixed_sysval_ubo = -1,
};
/* Lower this early so the backends don't have to worry about it */
@@ -134,7 +133,11 @@ panfrost_shader_compile(struct panfrost_screen *screen, const nir_shader *ir,
dev->gpu_id < 0x700);
}
struct panfrost_sysvals sysvals = {0};
NIR_PASS_V(s, panfrost_nir_lower_sysvals, &sysvals);
screen->vtbl.compile_shader(s, &inputs, &out->binary, &out->info);
out->info.sysvals = sysvals;
assert(req_local_mem >= out->info.wls_size);
out->info.wls_size = req_local_mem;
src/panfrost/compiler/bifrost_compile.c
+5 -85
View File
@@ -684,30 +684,6 @@ bi_load_ubo_to(bi_builder *b, unsigned bitsize, bi_index dest0, bi_index src0,
return I;
}
static bi_instr *
bi_load_sysval_to(bi_builder *b, bi_index dest, int sysval,
unsigned nr_components, unsigned offset)
{
unsigned sysval_ubo = b->shader->inputs->fixed_sysval_ubo >= 0
? b->shader->inputs->fixed_sysval_ubo
: b->shader->nir->info.num_ubos;
unsigned uniform = pan_lookup_sysval(b->shader->sysval_to_id,
b->shader->info.sysvals, sysval);
unsigned idx = (uniform * 16) + offset;
return bi_load_ubo_to(b, nr_components * 32, dest, bi_imm_u32(idx),
bi_imm_u32(sysval_ubo));
}
static void
bi_load_sysval_nir(bi_builder *b, nir_intrinsic_instr *intr,
unsigned nr_components, unsigned offset)
{
bi_load_sysval_to(b, bi_dest_index(&intr->dest),
panfrost_sysval_for_instr(&intr->instr, NULL),
nr_components, offset);
}
static void
bi_load_sample_id_to(bi_builder *b, bi_index dst)
{
@@ -1758,44 +1734,6 @@ bi_emit_intrinsic(bi_builder *b, nir_intrinsic_instr *instr)
bi_discard_f32(b, bi_zero(), bi_zero(), BI_CMPF_EQ);
break;
case nir_intrinsic_load_ssbo_address:
case nir_intrinsic_load_xfb_address:
bi_load_sysval_nir(b, instr, 2, 0);
break;
case nir_intrinsic_load_work_dim:
case nir_intrinsic_load_num_vertices:
case nir_intrinsic_load_first_vertex:
case nir_intrinsic_load_draw_id:
case nir_intrinsic_load_multisampled_pan:
case nir_intrinsic_load_rt_conversion_pan:
bi_load_sysval_nir(b, instr, 1, 0);
break;
case nir_intrinsic_load_base_vertex:
bi_load_sysval_nir(b, instr, 1, 4);
break;
case nir_intrinsic_load_base_instance:
case nir_intrinsic_get_ssbo_size:
bi_load_sysval_nir(b, instr, 1, 8);
break;
case nir_intrinsic_load_viewport_scale:
case nir_intrinsic_load_viewport_offset:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_workgroup_size:
bi_load_sysval_nir(b, instr, 3, 0);
break;
case nir_intrinsic_image_size:
bi_load_sysval_nir(b, instr, nir_dest_num_components(instr->dest), 0);
break;
case nir_intrinsic_load_blend_const_color_rgba:
bi_load_sysval_nir(b, instr, nir_dest_num_components(instr->dest), 0);
break;
case nir_intrinsic_load_sample_positions_pan:
bi_collect_v2i32_to(b, dst, bi_fau(BIR_FAU_SAMPLE_POS_ARRAY, false),
bi_fau(BIR_FAU_SAMPLE_POS_ARRAY, true));
@@ -3841,13 +3779,6 @@ bi_is_simple_tex(nir_tex_instr *instr)
static void
bi_emit_tex(bi_builder *b, nir_tex_instr *instr)
{
if (instr->op == nir_texop_txs) {
bi_load_sysval_to(b, bi_dest_index(&instr->dest),
panfrost_sysval_for_instr(&instr->instr, NULL),
nir_dest_num_components(instr->dest), 0);
return;
}
if (b->shader->arch >= 9)
bi_emit_tex_valhall(b, instr);
else if (bi_is_simple_tex(instr))
@@ -4846,7 +4777,6 @@ static bi_context *
bi_compile_variant_nir(nir_shader *nir,
const struct panfrost_compile_inputs *inputs,
struct util_dynarray *binary,
struct hash_table_u64 *sysval_to_id,
struct bi_shader_info info, enum bi_idvs_mode idvs)
{
bi_context *ctx = rzalloc(NULL, bi_context);
@@ -4854,7 +4784,6 @@ bi_compile_variant_nir(nir_shader *nir,
/* There may be another program in the dynarray, start at the end */
unsigned offset = binary->size;
ctx->sysval_to_id = sysval_to_id;
ctx->inputs = inputs;
ctx->nir = nir;
ctx->stage = nir->info.stage;
@@ -5113,14 +5042,12 @@ static void
bi_compile_variant(nir_shader *nir,
const struct panfrost_compile_inputs *inputs,
struct util_dynarray *binary,
struct hash_table_u64 *sysval_to_id,
struct pan_shader_info *info, enum bi_idvs_mode idvs)
{
struct bi_shader_info local_info = {
.push = &info->push,
.bifrost = &info->bifrost,
.tls_size = info->tls_size,
.sysvals = &info->sysvals,
.push_offset = info->push.count,
};
@@ -5138,8 +5065,8 @@ bi_compile_variant(nir_shader *nir,
* offset, to keep the ABI simple. */
assert((offset == 0) ^ (idvs == BI_IDVS_VARYING));
bi_context *ctx = bi_compile_variant_nir(nir, inputs, binary, sysval_to_id,
local_info, idvs);
bi_context *ctx =
bi_compile_variant_nir(nir, inputs, binary, local_info, idvs);
/* A register is preloaded <==> it is live before the first block */
bi_block *first_block = list_first_entry(&ctx->blocks, bi_block, link);
@@ -5233,21 +5160,16 @@ bifrost_compile_shader_nir(nir_shader *nir,
bi_optimize_nir(nir, inputs->gpu_id, inputs->is_blend);
struct hash_table_u64 *sysval_to_id =
panfrost_init_sysvals(&info->sysvals, inputs->fixed_sysval_layout, NULL);
info->tls_size = nir->scratch_size;
info->vs.idvs = bi_should_idvs(nir, inputs);
pan_nir_collect_varyings(nir, info);
if (info->vs.idvs) {
bi_compile_variant(nir, inputs, binary, sysval_to_id, info,
BI_IDVS_POSITION);
bi_compile_variant(nir, inputs, binary, sysval_to_id, info,
BI_IDVS_VARYING);
bi_compile_variant(nir, inputs, binary, info, BI_IDVS_POSITION);
bi_compile_variant(nir, inputs, binary, info, BI_IDVS_VARYING);
} else {
bi_compile_variant(nir, inputs, binary, sysval_to_id, info, BI_IDVS_NONE);
bi_compile_variant(nir, inputs, binary, info, BI_IDVS_NONE);
}
if (gl_shader_stage_is_compute(nir->info.stage)) {
@@ -5262,6 +5184,4 @@ bifrost_compile_shader_nir(nir_shader *nir,
}
info->ubo_mask &= (1 << nir->info.num_ubos) - 1;
_mesa_hash_table_u64_destroy(sysval_to_id);
}
src/panfrost/compiler/compiler.h
-2
View File
@@ -803,7 +803,6 @@ bi_block_add_successor(bi_block *block, bi_block *successor)
struct bi_shader_info {
struct panfrost_ubo_push *push;
struct bifrost_shader_info *bifrost;
struct panfrost_sysvals *sysvals;
unsigned tls_size;
unsigned work_reg_count;
unsigned push_offset;
@@ -827,7 +826,6 @@ typedef struct {
struct bi_shader_info info;
gl_shader_stage stage;
struct list_head blocks; /* list of bi_block */
struct hash_table_u64 *sysval_to_id;
uint32_t quirks;
unsigned arch;
enum bi_idvs_mode idvs;
src/panfrost/lib/pan_blend.c
-1
View File
@@ -868,7 +868,6 @@ GENX(pan_blend_get_shader_locked)(const struct panfrost_device *dev,
.gpu_id = dev->gpu_id,
.is_blend = true,
.blend.nr_samples = key.nr_samples,
.fixed_sysval_ubo = -1,
};
enum pipe_format rt_formats[8] = {0};
src/panfrost/lib/pan_blitter.c
-1
View File
@@ -618,7 +618,6 @@ pan_blitter_get_blit_shader(struct panfrost_device *dev,
.gpu_id = dev->gpu_id,
.is_blit = true,
.no_idvs = true,
.fixed_sysval_ubo = -1,
};
struct util_dynarray binary;
src/panfrost/lib/pan_indirect_dispatch.c
-1
View File
@@ -123,7 +123,6 @@ pan_indirect_dispatch_init(struct panfrost_device *dev)
struct panfrost_compile_inputs inputs = {
.gpu_id = dev->gpu_id,
.fixed_sysval_ubo = -1,
.no_ubo_to_push = true,
};
struct pan_shader_info shader_info;
src/panfrost/lib/pan_shader.c
+1 -6
View File
@@ -206,14 +206,9 @@ GENX(pan_shader_compile)(nir_shader *s, struct panfrost_compile_inputs *inputs,
}
info->outputs_written = s->info.outputs_written;
/* Sysvals have dedicated UBO */
info->ubo_count = s->info.num_ubos;
if (info->sysvals.sysval_count && inputs->fixed_sysval_ubo < 0)
info->ubo_count++;
info->attribute_count += BITSET_LAST_BIT(s->info.images_used);
info->writes_global = s->info.writes_memory;
info->ubo_count = s->info.num_ubos;
info->sampler_count = info->texture_count =
BITSET_LAST_BIT(s->info.textures_used);
src/panfrost/midgard/compiler.h
-2
View File
@@ -306,8 +306,6 @@ typedef struct compiler_context {
midgard_instruction
*writeout_branch[MIDGARD_NUM_RTS][MIDGARD_MAX_SAMPLE_ITER];
struct hash_table_u64 *sysval_to_id;
/* Mask of UBOs that need to be uploaded */
uint32_t ubo_mask;
} compiler_context;
src/panfrost/midgard/midgard_compile.c
-72
View File
@@ -1559,29 +1559,6 @@ emit_attr_read(compiler_context *ctx, unsigned dest, unsigned offset,
emit_mir_instruction(ctx, ins);
}
static void
emit_sysval_read(compiler_context *ctx, nir_instr *instr,
unsigned nr_components, unsigned offset)
{
nir_dest nir_dest;
/* Figure out which uniform this is */
unsigned sysval_ubo = ctx->inputs->fixed_sysval_ubo >= 0
? ctx->inputs->fixed_sysval_ubo
: ctx->nir->info.num_ubos;
int sysval = panfrost_sysval_for_instr(instr, &nir_dest);
unsigned dest = nir_dest_index(&nir_dest);
unsigned uniform =
pan_lookup_sysval(ctx->sysval_to_id, &ctx->info->sysvals, sysval);
/* Emit the read itself -- this is never indirect */
midgard_instruction *ins =
emit_ubo_read(ctx, instr, dest, (uniform * 16) + offset, NULL, 0,
sysval_ubo, nr_components);
ins->mask = mask_of(nr_components);
}
static unsigned
compute_builtin_arg(nir_intrinsic_op op)
{
@@ -1770,12 +1747,6 @@ emit_intrinsic(compiler_context *ctx, nir_intrinsic_instr *instr)
emit_image_op(ctx, instr, false);
break;
case nir_intrinsic_image_size: {
unsigned nr_comp = nir_intrinsic_dest_components(instr);
emit_sysval_read(ctx, &instr->instr, nr_comp, 0);
break;
}
case nir_intrinsic_load_ubo:
case nir_intrinsic_load_global:
case nir_intrinsic_load_global_constant:
@@ -2074,42 +2045,6 @@ emit_intrinsic(compiler_context *ctx, nir_intrinsic_instr *instr)
emit_global(ctx, &instr->instr, false, reg, &instr->src[1], seg);
break;
case nir_intrinsic_load_ssbo_address:
case nir_intrinsic_load_xfb_address:
emit_sysval_read(ctx, &instr->instr, 2, 0);
break;
case nir_intrinsic_load_first_vertex:
case nir_intrinsic_load_work_dim:
case nir_intrinsic_load_num_vertices:
emit_sysval_read(ctx, &instr->instr, 1, 0);
break;
case nir_intrinsic_load_base_vertex:
emit_sysval_read(ctx, &instr->instr, 1, 4);
break;
case nir_intrinsic_load_base_instance:
case nir_intrinsic_get_ssbo_size:
emit_sysval_read(ctx, &instr->instr, 1, 8);
break;
case nir_intrinsic_load_sample_positions_pan:
emit_sysval_read(ctx, &instr->instr, 2, 0);
break;
case nir_intrinsic_load_viewport_scale:
case nir_intrinsic_load_viewport_offset:
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_sampler_lod_parameters_pan:
case nir_intrinsic_load_workgroup_size:
emit_sysval_read(ctx, &instr->instr, 3, 0);
break;
case nir_intrinsic_load_blend_const_color_rgba:
emit_sysval_read(ctx, &instr->instr, 4, 0);
break;
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_global_invocation_id:
@@ -2502,9 +2437,6 @@ emit_tex(compiler_context *ctx, nir_tex_instr *instr)
case nir_texop_txf_ms:
emit_texop_native(ctx, instr, midgard_tex_op_fetch);
break;
case nir_texop_txs:
emit_sysval_read(ctx, &instr->instr, 4, 0);
break;
default: {
fprintf(stderr, "Unhandled texture op: %d\n", instr->op);
assert(0);
@@ -3135,8 +3067,6 @@ midgard_compile_shader_nir(nir_shader *nir,
/* TODO: Bound against what? */
compiler_context *ctx = rzalloc(NULL, compiler_context);
ctx->sysval_to_id =
panfrost_init_sysvals(&info->sysvals, inputs->fixed_sysval_layout, ctx);
ctx->inputs = inputs;
ctx->nir = nir;
@@ -3350,7 +3280,5 @@ midgard_compile_shader_nir(nir_shader *nir,
}
_mesa_hash_table_u64_destroy(ctx->ssa_constants);
_mesa_hash_table_u64_destroy(ctx->sysval_to_id);
ralloc_free(ctx);
}
src/panfrost/util/meson.build
-1
View File
@@ -34,7 +34,6 @@ libpanfrost_util_files = files(
'pan_lower_writeout.c',
'pan_lower_xfb.c',
'pan_lower_64bit_intrin.c',
'pan_sysval.c',
)
libpanfrost_util = static_library(
src/panfrost/util/pan_ir.h
-11
View File
@@ -160,15 +160,6 @@ struct panfrost_ubo_push {
unsigned pan_lookup_pushed_ubo(struct panfrost_ubo_push *push, unsigned ubo,
unsigned offs);
struct hash_table_u64 *
panfrost_init_sysvals(struct panfrost_sysvals *sysvals,
struct panfrost_sysvals *fixed_sysvals, void *memctx);
unsigned pan_lookup_sysval(struct hash_table_u64 *sysval_to_id,
struct panfrost_sysvals *sysvals, int sysval);
int panfrost_sysval_for_instr(nir_instr *instr, nir_dest *dest);
struct panfrost_compile_inputs {
struct util_debug_callback *debug;
@@ -178,8 +169,6 @@ struct panfrost_compile_inputs {
unsigned nr_samples;
uint64_t bifrost_blend_desc;
} blend;
int fixed_sysval_ubo;
struct panfrost_sysvals *fixed_sysval_layout;
bool no_idvs;
bool no_ubo_to_push;
src/panfrost/util/pan_sysval.c
-203
View File
@@ -1,203 +0,0 @@
/*
* Copyright (C) 2020 Collabora Ltd.
*
* Permission is hereby granted, free of charge, to any person obtaining a
* copy of this software and associated documentation files (the "Software"),
* to deal in the Software without restriction, including without limitation
* the rights to use, copy, modify, merge, publish, distribute, sublicense,
* and/or sell copies of the Software, and to permit persons to whom the
* Software is furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice (including the next
* paragraph) shall be included in all copies or substantial portions of the
* Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
* THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*
* Authors (Collabora):
* Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
*/
#include "compiler/nir/nir_builder.h"
#include "pan_ir.h"
/* TODO: ssbo_size */
static int
panfrost_sysval_for_ssbo(nir_intrinsic_instr *instr)
{
nir_src index = instr->src[0];
assert(nir_src_is_const(index));
uint32_t uindex = nir_src_as_uint(index);
return PAN_SYSVAL(SSBO, uindex);
}
static int
panfrost_sysval_for_sampler(nir_intrinsic_instr *instr)
{
/* TODO: indirect samplers !!! */
nir_src index = instr->src[0];
assert(nir_src_is_const(index));
uint32_t uindex = nir_src_as_uint(index);
return PAN_SYSVAL(SAMPLER, uindex);
}
static int
panfrost_sysval_for_image_size(nir_intrinsic_instr *instr)
{
nir_src index = instr->src[0];
assert(nir_src_is_const(index));
bool is_array = nir_intrinsic_image_array(instr);
uint32_t uindex = nir_src_as_uint(index);
unsigned dim = nir_intrinsic_dest_components(instr) - is_array;
return PAN_SYSVAL(IMAGE_SIZE, PAN_TXS_SYSVAL_ID(uindex, dim, is_array));
}
static int
panfrost_sysval_for_rt_conversion(nir_intrinsic_instr *instr)
{
unsigned size = nir_alu_type_get_type_size(nir_intrinsic_src_type(instr));
unsigned rt = nir_intrinsic_base(instr);
return PAN_SYSVAL(RT_CONVERSION, rt | (size << 4));
}
static unsigned
panfrost_nir_sysval_for_intrinsic(nir_intrinsic_instr *instr)
{
switch (instr->intrinsic) {
case nir_intrinsic_load_viewport_scale:
return PAN_SYSVAL_VIEWPORT_SCALE;
case nir_intrinsic_load_viewport_offset:
return PAN_SYSVAL_VIEWPORT_OFFSET;
case nir_intrinsic_load_num_workgroups:
return PAN_SYSVAL_NUM_WORK_GROUPS;
case nir_intrinsic_load_workgroup_size:
return PAN_SYSVAL_LOCAL_GROUP_SIZE;
case nir_intrinsic_load_work_dim:
return PAN_SYSVAL_WORK_DIM;
case nir_intrinsic_load_sample_positions_pan:
return PAN_SYSVAL_SAMPLE_POSITIONS;
case nir_intrinsic_load_multisampled_pan:
return PAN_SYSVAL_MULTISAMPLED;
case nir_intrinsic_load_first_vertex:
case nir_intrinsic_load_base_vertex:
case nir_intrinsic_load_base_instance:
return PAN_SYSVAL_VERTEX_INSTANCE_OFFSETS;
case nir_intrinsic_load_draw_id:
return PAN_SYSVAL_DRAWID;
case nir_intrinsic_load_ssbo_address:
case nir_intrinsic_get_ssbo_size:
return panfrost_sysval_for_ssbo(instr);
case nir_intrinsic_load_xfb_address:
return PAN_SYSVAL(XFB, nir_intrinsic_base(instr));
case nir_intrinsic_load_num_vertices:
return PAN_SYSVAL_NUM_VERTICES;
case nir_intrinsic_load_sampler_lod_parameters_pan:
return panfrost_sysval_for_sampler(instr);
case nir_intrinsic_image_size:
return panfrost_sysval_for_image_size(instr);
case nir_intrinsic_load_blend_const_color_rgba:
return PAN_SYSVAL_BLEND_CONSTANTS;
case nir_intrinsic_load_rt_conversion_pan:
return panfrost_sysval_for_rt_conversion(instr);
default:
return ~0;
}
}
int
panfrost_sysval_for_instr(nir_instr *instr, nir_dest *dest)
{
nir_intrinsic_instr *intr;
nir_dest *dst = NULL;
nir_tex_instr *tex;
unsigned sysval = ~0;
switch (instr->type) {
case nir_instr_type_intrinsic:
intr = nir_instr_as_intrinsic(instr);
sysval = panfrost_nir_sysval_for_intrinsic(intr);
dst = &intr->dest;
break;
case nir_instr_type_tex:
tex = nir_instr_as_tex(instr);
if (tex->op != nir_texop_txs)
break;
sysval = PAN_SYSVAL(TEXTURE_SIZE,
PAN_TXS_SYSVAL_ID(tex->texture_index,
nir_tex_instr_dest_size(tex) -
(tex->is_array ? 1 : 0),
tex->is_array));
dst = &tex->dest;
break;
default:
break;
}
if (dest && dst)
*dest = *dst;
return sysval;
}
static unsigned
pan_add_sysval(struct hash_table_u64 *sysval_to_id,
struct panfrost_sysvals *sysvals, int sysval, unsigned id)
{
assert(id < MAX_SYSVAL_COUNT);
_mesa_hash_table_u64_insert(sysval_to_id, sysval,
(void *)((uintptr_t)id + 1));
sysvals->sysvals[id] = sysval;
return id;
}
unsigned
pan_lookup_sysval(struct hash_table_u64 *sysval_to_id,
struct panfrost_sysvals *sysvals, int sysval)
{
/* Try to lookup */
void *cached = _mesa_hash_table_u64_search(sysval_to_id, sysval);
if (cached) {
unsigned id = ((uintptr_t)cached) - 1;
assert(id < MAX_SYSVAL_COUNT);
assert(sysvals->sysvals[id] == sysval);
return id;
}
/* Else assign */
return pan_add_sysval(sysval_to_id, sysvals, sysval,
sysvals->sysval_count++);
}
struct hash_table_u64 *
panfrost_init_sysvals(struct panfrost_sysvals *sysvals,
struct panfrost_sysvals *fixed_sysvals, void *memctx)
{
memset(sysvals, 0, sizeof(*sysvals));
struct hash_table_u64 *sysval_to_id = _mesa_hash_table_u64_create(memctx);
if (fixed_sysvals) {
for (unsigned i = 0; i < fixed_sysvals->sysval_count; i++) {
if (!fixed_sysvals->sysvals[i])
continue;
pan_add_sysval(sysval_to_id, sysvals, fixed_sysvals->sysvals[i], i);
}
sysvals->sysval_count = fixed_sysvals->sysval_count;
}
return sysval_to_id;
}
src/panfrost/vulkan/panvk_vX_shader.c
-1
View File
@@ -273,7 +273,6 @@ panvk_per_arch(shader_create)(struct panvk_device *dev,
.gpu_id = pdev->gpu_id,
.no_ubo_to_push = true,
.no_idvs = true, /* TODO */
.fixed_sysval_ubo = -1,
};
NIR_PASS_V(nir, nir_lower_indirect_derefs,