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

asahi: split param structs for GS internal kernel

Browse Source 
this simplifies state management consdierably

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29607>
This commit is contained in:
Alyssa Rosenzweig
2024-05-26 14:06:19 -04:00
committed by Marge Bot
parent d3291ad001
commit 5f72234745
7 changed files with 363 additions and 225 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/asahi/lib/agx_nir_lower_gs.c
+4 -10
View File
@@ -379,10 +379,7 @@ lower_id(nir_builder *b, nir_intrinsic_instr *intr, void *data)
id = load_geometry_param(b, flat_outputs);
else if (intr->intrinsic == nir_intrinsic_load_input_topology_agx)
id = load_geometry_param(b, input_topology);
else if (intr->intrinsic == nir_intrinsic_load_provoking_last) {
id = nir_b2b32(
b, libagx_is_provoking_last(b, nir_load_input_assembly_buffer_agx(b)));
} else
else
return false;
b->cursor = nir_instr_remove(&intr->instr);
@@ -1452,10 +1449,7 @@ agx_nir_gs_setup_indirect(nir_builder *b, const void *data)
{
const struct agx_gs_setup_indirect_key *key = data;
libagx_gs_setup_indirect(b, nir_load_geometry_param_buffer_agx(b),
nir_load_input_assembly_buffer_agx(b),
nir_load_vs_output_buffer_ptr_agx(b),
nir_load_vs_outputs_agx(b),
libagx_gs_setup_indirect(b, nir_load_preamble(b, 1, 64, .base = 0),
nir_imm_int(b, key->prim),
nir_channel(b, nir_load_local_invocation_id(b), 0));
}
@@ -1463,10 +1457,10 @@ agx_nir_gs_setup_indirect(nir_builder *b, const void *data)
void
agx_nir_unroll_restart(nir_builder *b, const void *data)
{
const struct agx_unroll_restart_key *key = data;
b->shader->info.workgroup_size[0] = 1024;
const struct agx_unroll_restart_key *key = data;
nir_def *ia = nir_load_input_assembly_buffer_agx(b);
nir_def *ia = nir_load_preamble(b, 1, 64, .base = 0);
nir_def *draw = nir_channel(b, nir_load_workgroup_id(b), 0);
nir_def *lane = nir_channel(b, nir_load_local_invocation_id(b), 0);
nir_def *mode = nir_imm_int(b, key->prim);
src/asahi/lib/shaders/geometry.cl
+32 -32
View File
@@ -249,8 +249,9 @@ first_true_thread_in_workgroup(bool cond, local uint *scratch)
* sets up most of the new draw descriptor.
*/
static global void *
setup_unroll_for_draw(global struct agx_ia_state *ia, constant uint *in_draw,
uint draw, enum mesa_prim mode, uint index_size_B)
setup_unroll_for_draw(global struct agx_restart_unroll_params *p,
constant uint *in_draw, uint draw, enum mesa_prim mode,
uint index_size_B)
{
/* Determine an upper bound on the memory required for the index buffer.
* Restarts only decrease the unrolled index buffer size, so the maximum size
@@ -263,12 +264,12 @@ setup_unroll_for_draw(global struct agx_ia_state *ia, constant uint *in_draw,
/* Allocate memory from the heap for the unrolled index buffer. Use an atomic
* since multiple threads may be running to handle multidraw in parallel.
*/
global struct agx_geometry_state *heap = ia->heap;
global struct agx_geometry_state *heap = p->heap;
uint old_heap_bottom_B = atomic_fetch_add(
(volatile atomic_uint *)(&heap->heap_bottom), align(alloc_size, 4));
/* Regardless of the input stride, we use tightly packed output draws */
global uint *out = &ia->out_draws[5 * draw];
global uint *out = &p->out_draws[5 * draw];
/* Setup most of the descriptor. Count will be determined after unroll. */
out[1] = in_draw[1]; /* instance count */
@@ -281,28 +282,28 @@ setup_unroll_for_draw(global struct agx_ia_state *ia, constant uint *in_draw,
}
#define UNROLL(INDEX, suffix) \
kernel void libagx_unroll_restart_##suffix(global struct agx_ia_state *ia, \
enum mesa_prim mode, uint draw, \
uint tid) \
kernel void libagx_unroll_restart_##suffix( \
global struct agx_restart_unroll_params *p, enum mesa_prim mode, \
uint draw, uint tid) \
{ \
/* For an indirect multidraw, we are dispatched maxDraws times and \
* terminate trailing invocations. \
*/ \
if (ia->count && draw >= *(ia->count)) \
if (p->count && draw >= *(p->count)) \
return; \
\
constant uint *in_draw = \
(constant uint *)(ia->draws + (draw * ia->draw_stride)); \
(constant uint *)(p->draws + (draw * p->draw_stride)); \
\
uint count = in_draw[0]; \
\
local uintptr_t out_ptr, in_ptr; \
if (tid == 0) { \
out_ptr = (uintptr_t)setup_unroll_for_draw(ia, in_draw, draw, mode, \
out_ptr = (uintptr_t)setup_unroll_for_draw(p, in_draw, draw, mode, \
sizeof(INDEX)); \
\
/* Accessed thru local mem because NIR deref is too aggressive */ \
in_ptr = (uintptr_t)(ia->index_buffer + sizeof(INDEX) * in_draw[2]); \
in_ptr = (uintptr_t)(p->index_buffer + sizeof(INDEX) * in_draw[2]); \
} \
\
barrier(CLK_LOCAL_MEM_FENCE); \
@@ -312,9 +313,9 @@ setup_unroll_for_draw(global struct agx_ia_state *ia, constant uint *in_draw,
local uint scratch[32]; \
\
uint out_prims = 0; \
INDEX restart_idx = ia->restart_index; \
bool flatshade_first = ia->flatshade_first; \
uint in_size_el = ia->index_buffer_size_B / sizeof(INDEX); \
INDEX restart_idx = p->restart_index; \
bool flatshade_first = p->flatshade_first; \
uint in_size_el = p->index_buffer_size_B / sizeof(INDEX); \
\
uint needle = 0; \
uint per_prim = mesa_vertices_per_prim(mode); \
@@ -324,6 +325,7 @@ setup_unroll_for_draw(global struct agx_ia_state *ia, constant uint *in_draw,
for (;;) { \
/* Relies on shortcircuiting */ \
uint idx = next_restart + tid; \
/* XXX: robustness here */ \
bool restart = idx >= count || in[idx] == restart_idx; \
\
uint next_offs = first_true_thread_in_workgroup(restart, scratch); \
@@ -353,7 +355,7 @@ setup_unroll_for_draw(global struct agx_ia_state *ia, constant uint *in_draw,
} \
\
if (tid == 0) \
ia->out_draws[(5 * draw) + 0] = out_prims * per_prim; \
p->out_draws[(5 * draw) + 0] = out_prims * per_prim; \
}
UNROLL(uchar, u8)
@@ -447,16 +449,15 @@ libagx_build_gs_draw(global struct agx_geometry_params *p, uint vertices,
}
void
libagx_gs_setup_indirect(global struct agx_geometry_params *p,
global struct agx_ia_state *ia,
global uintptr_t *vertex_buffer, uint64_t vs_outputs,
libagx_gs_setup_indirect(global struct agx_gs_setup_indirect_params *gsi,
enum mesa_prim mode, uint local_id)
{
global uint *in_draw = (global uint *)ia->draws;
global struct agx_geometry_params *p = gsi->geom;
global struct agx_ia_state *ia = gsi->ia;
/* Determine the (primitives, instances) grid size. */
uint vertex_count = in_draw[0];
uint instance_count = in_draw[1];
uint vertex_count = gsi->draw[0];
uint instance_count = gsi->draw[1];
ia->verts_per_instance = vertex_count;
@@ -478,22 +479,27 @@ libagx_gs_setup_indirect(global struct agx_geometry_params *p,
* indirect draw, the hardware would do this for us, but for software input
* assembly we need to do it ourselves.
*/
if (ia->index_buffer) {
ia->index_buffer += ((constant uint *)ia->draws)[2] * ia->index_size_B;
if (gsi->index_buffer) {
ia->index_buffer = gsi->index_buffer + gsi->draw[2] * gsi->index_size_B;
}
/* We may need to allocate VS and GS count buffers, do so now */
/* We need to allocate VS and GS count buffers, do so now */
global struct agx_geometry_state *state = p->state;
uint vertex_buffer_size =
libagx_tcs_in_size(vertex_count * instance_count, vs_outputs);
libagx_tcs_in_size(vertex_count * instance_count, gsi->vs_outputs);
p->count_buffer = (global uint *)(state->heap + state->heap_bottom);
state->heap_bottom +=
align(p->input_primitives * p->count_buffer_stride, 16);
*vertex_buffer = (uintptr_t)(state->heap + state->heap_bottom);
*(gsi->vertex_buffer) = (uintptr_t)(state->heap + state->heap_bottom);
state->heap_bottom += align(vertex_buffer_size, 4);
if (state->heap_bottom > 1024 * 1024 * 128) {
global uint *foo = (global uint *)(uintptr_t)0x1deadbeef;
*foo = 0x1234;
}
}
/*
@@ -567,12 +573,6 @@ libagx_prefix_sum(global uint *buffer, uint len, uint words, uint word)
}
}
bool
libagx_is_provoking_last(global struct agx_ia_state *ia)
{
return !ia->flatshade_first;
}
uintptr_t
libagx_vertex_output_address(uintptr_t buffer, uint64_t mask, uint vtx,
gl_varying_slot location)
src/asahi/lib/shaders/geometry.h
+54 -30
View File
@@ -24,16 +24,13 @@
/* Packed geometry state buffer */
struct agx_geometry_state {
/* Heap to allocate from, in either direction. By convention, the top is used
* for intra-draw allocations and the bottom is used for full-batch
* allocations. In the future we could use kernel support to improve this.
*/
/* Heap to allocate from. */
GLOBAL(uchar) heap;
uint32_t heap_bottom, heap_top, heap_size, padding;
uint32_t heap_bottom, heap_size;
} PACKED;
AGX_STATIC_ASSERT(sizeof(struct agx_geometry_state) == 6 * 4);
AGX_STATIC_ASSERT(sizeof(struct agx_geometry_state) == 4 * 4);
struct agx_ia_state {
struct agx_restart_unroll_params {
/* Heap to allocate from across draws */
GLOBAL(struct agx_geometry_state) heap;
@@ -46,36 +43,63 @@ struct agx_ia_state {
/* Input: indirect draw descriptor. Raw pointer since it's strided. */
uint64_t draws;
/* When unrolling primitive restart, output draw descriptors */
/* Output draw descriptors */
GLOBAL(uint) out_draws;
/* Input: maximum draw count, count is clamped to this */
uint32_t max_draws;
/* Primitive restart index */
uint32_t restart_index;
/* Input index buffer size in bytes */
uint32_t index_buffer_size_B;
/* Stride for the draw descriptor array */
uint32_t draw_stride;
/* Use first vertex as the provoking vertex for flat shading. We could stick
* this in the key, but meh, you're already hosed for perf on the unroll
* path.
*/
uint32_t flatshade_first;
} PACKED;
AGX_STATIC_ASSERT(sizeof(struct agx_restart_unroll_params) == 15 * 4);
struct agx_gs_setup_indirect_params {
/* Index buffer if present. */
CONST(uchar) index_buffer;
/* Indirect draw descriptor. */
CONST(uint) draw;
/* Pointer to be written with allocated vertex buffer */
GLOBAL(uintptr_t) vertex_buffer;
/* Output input assembly state */
GLOBAL(struct agx_ia_state) ia;
/* Output geometry parameters */
GLOBAL(struct agx_geometry_params) geom;
/* Vertex (TES) output mask for sizing the allocated buffer */
uint64_t vs_outputs;
/* The index size (1, 2, 4) or 0 if drawing without an index buffer. */
uint32_t index_size_B;
} PACKED;
AGX_STATIC_ASSERT(sizeof(struct agx_gs_setup_indirect_params) == 13 * 4);
struct agx_ia_state {
/* Index buffer if present. */
CONST(uchar) index_buffer;
/* Number of vertices per instance. Written by CPU for direct draw, indirect
* setup kernel for indirect. This is used for VS->GS and VS->TCS indexing.
*/
uint32_t verts_per_instance;
/* Input: maximum draw count, count is clamped to this */
uint32_t max_draws;
/* Primitive restart index, if unrolling */
uint32_t restart_index;
/* Input index buffer size in bytes, if unrolling */
uint32_t index_buffer_size_B;
/* Stride for the draw descrptor array */
uint32_t draw_stride;
/* When unrolling primitive restart, use first vertex as the provoking vertex
* for flat shading. We could stick this in the key, but meh, you're already
* hosed for perf on the unroll path.
*/
uint32_t flatshade_first;
/* The index size (1, 2, 4) or 0 if drawing without an index buffer. */
uint32_t index_size_B;
} PACKED;
AGX_STATIC_ASSERT(sizeof(struct agx_ia_state) == 17 * 4);
AGX_STATIC_ASSERT(sizeof(struct agx_ia_state) == 3 * 4);
struct agx_geometry_params {
/* Persistent (cross-draw) geometry state */
src/compiler/nir/nir_intrinsics.py
-3
View File
@@ -2024,9 +2024,6 @@ barrier("fence_helper_exit_agx")
# Pointer to the buffer passing outputs VS->TCS, VS->GS, or TES->GS linkage.
system_value("vs_output_buffer_agx", 1, bit_sizes=[64])
# Indirect for the above, used for indirect draws.
system_value("vs_output_buffer_ptr_agx", 1, bit_sizes=[64])
# Mask of VS->TCS, VS->GS, or TES->GS outputs. This is modelled as a sysval
# directly so it can be dynamic with shader objects or constant folded with
# pipelines (including GPL)
src/gallium/drivers/asahi/agx_nir_lower_sysvals.c
+10 -4
View File
@@ -126,6 +126,13 @@ load_texture_handle(nir_builder *b, nir_intrinsic_instr *intr, void *base)
nir_imul_imm(b, nir_u2u32(b, intr->src[0].ssa), AGX_TEXTURE_DESC_STRIDE));
}
static nir_def *
load_provoking_vtx(nir_builder *b)
{
struct agx_draw_uniforms *u = NULL;
return load_sysval_root(b, 1, 16, &u->provoking_vertex);
}
static nir_def *
lower_intrinsic(nir_builder *b, nir_intrinsic_instr *intr,
bool lower_draw_params)
@@ -178,8 +185,6 @@ lower_intrinsic(nir_builder *b, nir_intrinsic_instr *intr,
case nir_intrinsic_load_vs_output_buffer_agx:
return nir_load_global_constant(
b, load_sysval_root(b, 1, 64, &u->vertex_output_buffer_ptr), 8, 1, 64);
case nir_intrinsic_load_vs_output_buffer_ptr_agx:
return load_sysval_root(b, 1, 64, &u->vertex_output_buffer_ptr);
case nir_intrinsic_load_vs_outputs_agx:
return load_sysval_root(b, 1, 64, &u->vertex_outputs);
case nir_intrinsic_load_tess_param_buffer_agx:
@@ -199,8 +204,9 @@ lower_intrinsic(nir_builder *b, nir_intrinsic_instr *intr,
return load_sysval_root(
b, 1, 16, &u->uvs_index[nir_intrinsic_io_semantics(intr).location]);
case nir_intrinsic_load_is_first_fan_agx:
/* TODO: Plumb this so we can stop using geometry shaders for this case */
return nir_imm_false(b);
return nir_ieq_imm(b, load_provoking_vtx(b), 1);
case nir_intrinsic_load_provoking_last:
return nir_b2b32(b, nir_ieq_imm(b, load_provoking_vtx(b), 2));
default:
break;
}
src/gallium/drivers/asahi/agx_state.c
+239 -136
View File
@@ -1546,8 +1546,8 @@ agx_nir_lower_point_sprite_zw(nir_builder *b, nir_intrinsic_instr *intr,
static struct agx_compiled_shader *
agx_compile_nir(struct agx_device *dev, nir_shader *nir,
struct util_debug_callback *debug, enum pipe_shader_type stage,
bool terminal, bool secondary, unsigned cf_base,
BITSET_WORD *attrib_components_read)
bool internal_kernel, bool terminal, bool secondary,
unsigned cf_base, BITSET_WORD *attrib_components_read)
{
struct agx_compiled_shader *compiled = CALLOC_STRUCT(agx_compiled_shader);
compiled->stage = stage;
@@ -1574,7 +1574,9 @@ agx_compile_nir(struct agx_device *dev, nir_shader *nir,
nir->info.fs.uses_sample_shading)
key.fs.inside_sample_loop = true;
if (!secondary) {
if (internal_kernel) {
key.reserved_preamble = 8;
} else if (!secondary) {
NIR_PASS(_, nir, agx_nir_lower_sysvals, stage, true);
NIR_PASS(_, nir, agx_nir_layout_uniforms, compiled,
&key.reserved_preamble);
@@ -1597,9 +1599,11 @@ agx_compile_nir(struct agx_device *dev, nir_shader *nir,
return compiled;
}
static struct agx_compiled_shader *agx_build_meta_shader_internal(
struct agx_context *ctx, meta_shader_builder_t builder, void *data,
size_t data_size, bool prolog, bool epilog, unsigned cf_base);
static struct agx_compiled_shader *
agx_build_meta_shader_internal(struct agx_context *ctx,
meta_shader_builder_t builder, void *data,
size_t data_size, bool prolog, bool epilog,
unsigned cf_base, bool internal_kernel);
/* Does not take ownership of key. Clones if necessary. */
static struct agx_compiled_shader *
@@ -1691,8 +1695,8 @@ agx_compile_variant(struct agx_device *dev, struct pipe_context *pctx,
NIR_PASS(_, nir, agx_nir_lower_multisampled_image_store);
struct agx_compiled_shader *compiled = agx_compile_nir(
dev, nir, debug, so->type, so->type != PIPE_SHADER_FRAGMENT, false, 0,
attrib_components_read);
dev, nir, debug, so->type, false, so->type != PIPE_SHADER_FRAGMENT, false,
0, attrib_components_read);
if (so->type == PIPE_SHADER_FRAGMENT) {
/* XXX: don't replicate this all over the driver */
@@ -1707,14 +1711,14 @@ agx_compile_variant(struct agx_device *dev, struct pipe_context *pctx,
/* Compile auxiliary programs */
if (gs_count) {
compiled->gs_count =
agx_compile_nir(dev, gs_count, debug, so->type, true, false, 0, NULL);
compiled->gs_count = agx_compile_nir(dev, gs_count, debug, so->type,
false, true, false, 0, NULL);
compiled->gs_count->so = so;
}
if (pre_gs) {
compiled->pre_gs = agx_compile_nir(
dev, pre_gs, debug, PIPE_SHADER_COMPUTE, true, false, 0, NULL);
dev, pre_gs, debug, PIPE_SHADER_COMPUTE, false, true, false, 0, NULL);
}
if (gs_copy) {
@@ -1729,8 +1733,9 @@ agx_compile_variant(struct agx_device *dev, struct pipe_context *pctx,
struct agx_unlinked_uvs_layout uvs = {0};
NIR_PASS(_, gs_copy, agx_nir_lower_uvs, &uvs);
compiled->gs_copy = agx_compile_nir(
dev, gs_copy, debug, PIPE_SHADER_GEOMETRY, true, false, 0, NULL);
compiled->gs_copy =
agx_compile_nir(dev, gs_copy, debug, PIPE_SHADER_GEOMETRY, false, true,
false, 0, NULL);
compiled->gs_copy->so = so;
compiled->gs_copy->stage = so->type;
compiled->gs_copy->uvs = uvs;
@@ -2143,19 +2148,19 @@ asahi_fast_link(struct agx_context *ctx, struct agx_uncompiled_shader *so,
if (so->type == MESA_SHADER_FRAGMENT) {
prolog = agx_build_meta_shader_internal(
ctx, build_fs_prolog, &key->prolog.fs, sizeof(key->prolog.fs), true,
false, key->prolog.fs.cf_base);
false, key->prolog.fs.cf_base, false);
epilog =
agx_build_meta_shader_internal(ctx, agx_nir_fs_epilog, &key->epilog.fs,
sizeof(key->epilog.fs), false, true, 0);
epilog = agx_build_meta_shader_internal(
ctx, agx_nir_fs_epilog, &key->epilog.fs, sizeof(key->epilog.fs), false,
true, 0, false);
} else {
assert(so->type == MESA_SHADER_VERTEX ||
so->type == MESA_SHADER_TESS_EVAL);
prolog =
agx_build_meta_shader_internal(ctx, agx_nir_vs_prolog, &key->prolog.vs,
sizeof(key->prolog.vs), true, false, 0);
prolog = agx_build_meta_shader_internal(
ctx, agx_nir_vs_prolog, &key->prolog.vs, sizeof(key->prolog.vs), true,
false, 0, false);
}
/* Fast-link it all together */
@@ -2531,7 +2536,7 @@ static struct agx_compiled_shader *
agx_build_meta_shader_internal(struct agx_context *ctx,
meta_shader_builder_t builder, void *data,
size_t data_size, bool prolog, bool epilog,
unsigned cf_base)
unsigned cf_base, bool internal_kernel)
{
/* Build the meta shader key */
size_t total_key_size = sizeof(struct agx_generic_meta_key) + data_size;
@@ -2561,7 +2566,7 @@ agx_build_meta_shader_internal(struct agx_context *ctx,
agx_preprocess_nir(b.shader, dev->libagx);
struct agx_compiled_shader *shader = agx_compile_nir(
dev, b.shader, NULL, PIPE_SHADER_COMPUTE,
dev, b.shader, NULL, PIPE_SHADER_COMPUTE, internal_kernel,
!prolog && !(b.shader->info.stage == MESA_SHADER_FRAGMENT &&
b.shader->info.fs.uses_sample_shading),
prolog || epilog, cf_base, NULL);
@@ -2584,7 +2589,7 @@ agx_build_meta_shader(struct agx_context *ctx, meta_shader_builder_t builder,
void *data, size_t data_size)
{
return agx_build_meta_shader_internal(ctx, builder, data, data_size, false,
false, 0);
false, 0, false);
}
static unsigned
@@ -2894,6 +2899,30 @@ agx_update_descriptors(struct agx_batch *batch, struct agx_compiled_shader *cs)
agx_pool_upload_aligned(&batch->pool, unif, sizeof(*unif), 16);
}
static void
agx_usc_immediates(struct agx_usc_builder *b, struct agx_batch *batch,
struct agx_compiled_shader *cs)
{
unsigned constant_push_ranges =
DIV_ROUND_UP(cs->b.info.immediate_size_16, 64);
if (cs->b.info.immediate_size_16) {
/* XXX: do ahead of time */
uint64_t ptr =
agx_pool_upload_aligned(&batch->pool, cs->b.info.immediates,
cs->b.info.immediate_size_16 * 2, 64);
for (unsigned range = 0; range < constant_push_ranges; ++range) {
unsigned offset = 64 * range;
assert(offset < cs->b.info.immediate_size_16);
agx_usc_uniform(b, cs->b.info.immediate_base_uniform + offset,
MIN2(64, cs->b.info.immediate_size_16 - offset),
ptr + (offset * 2));
}
}
}
static uint32_t
agx_build_pipeline(struct agx_batch *batch, struct agx_compiled_shader *cs,
struct agx_linked_shader *linked,
@@ -2947,21 +2976,7 @@ agx_build_pipeline(struct agx_batch *batch, struct agx_compiled_shader *cs,
table_ptr + cs->push[i].offset);
}
if (cs->b.info.immediate_size_16) {
/* XXX: do ahead of time */
uint64_t ptr =
agx_pool_upload_aligned(&batch->pool, cs->b.info.immediates,
cs->b.info.immediate_size_16 * 2, 64);
for (unsigned range = 0; range < constant_push_ranges; ++range) {
unsigned offset = 64 * range;
assert(offset < cs->b.info.immediate_size_16);
agx_usc_uniform(&b, cs->b.info.immediate_base_uniform + offset,
MIN2(64, cs->b.info.immediate_size_16 - offset),
ptr + (offset * 2));
}
}
agx_usc_immediates(&b, batch, cs);
uint32_t max_scratch_size =
MAX2(cs->b.info.scratch_size, cs->b.info.preamble_scratch_size);
@@ -3037,6 +3052,66 @@ agx_build_pipeline(struct agx_batch *batch, struct agx_compiled_shader *cs,
return t.gpu;
}
static uint32_t
agx_build_internal_usc(struct agx_batch *batch, struct agx_compiled_shader *cs,
void *data, size_t data_size)
{
size_t usc_size = agx_usc_size(12);
struct agx_ptr t =
agx_pool_alloc_aligned(&batch->pipeline_pool, usc_size, 64);
struct agx_usc_builder b = agx_usc_builder(t.cpu, usc_size);
uint64_t ptr = agx_pool_upload(&batch->pool, data, data_size);
agx_usc_uniform(&b, 0, 4, agx_pool_upload(&batch->pool, &ptr, 8));
agx_usc_immediates(&b, batch, cs);
assert(cs->b.info.scratch_size == 0 && "internal kernels don't spill");
assert(cs->b.info.preamble_scratch_size == 0 && "internal doesn't spill");
unsigned local_size = cs->b.info.local_size;
agx_usc_pack(&b, SHARED, cfg) {
cfg.layout = AGX_SHARED_LAYOUT_VERTEX_COMPUTE;
cfg.bytes_per_threadgroup = local_size > 0 ? local_size : 65536;
cfg.uses_shared_memory = local_size > 0;
}
agx_usc_pack(&b, SHADER, cfg) {
cfg.code = (cs->bo->ptr.gpu + cs->b.info.main_offset);
cfg.unk_2 = 3;
}
agx_usc_pack(&b, REGISTERS, cfg) {
cfg.register_count = cs->b.info.nr_gprs;
cfg.spill_size = 0;
}
if (cs->b.info.has_preamble) {
agx_usc_pack(&b, PRESHADER, cfg) {
cfg.code = cs->bo->ptr.gpu + cs->b.info.preamble_offset;
}
} else {
agx_usc_pack(&b, NO_PRESHADER, cfg)
;
}
return t.gpu;
}
void
agx_launch_with_data(struct agx_batch *batch, const struct pipe_grid_info *info,
meta_shader_builder_t builder, void *key, size_t key_size,
void *data, size_t data_size)
{
struct agx_compiled_shader *cs = agx_build_meta_shader_internal(
batch->ctx, builder, key, key_size, false, false, 0, true);
uint32_t usc = agx_build_internal_usc(batch, cs, data, data_size);
agx_launch_internal(batch, info, cs, PIPE_SHADER_COMPUTE, usc);
}
struct asahi_bg_eot
agx_build_bg_eot(struct agx_batch *batch, bool store, bool partial_render)
{
@@ -3879,35 +3954,6 @@ agx_batch_geometry_state(struct agx_batch *batch)
return batch->geometry_state;
}
static void
agx_upload_ia_params(struct agx_batch *batch, const struct pipe_draw_info *info,
const struct pipe_draw_indirect_info *indirect,
uint32_t count, uint64_t input_index_buffer,
size_t index_buffer_size_B, uint64_t unroll_output)
{
struct agx_ia_state ia = {
.heap = agx_batch_geometry_state(batch),
.index_buffer = input_index_buffer,
.index_size_B = info->index_size,
.out_draws = unroll_output,
.restart_index = info->restart_index,
.index_buffer_size_B = index_buffer_size_B,
.flatshade_first = batch->ctx->rast->base.flatshade_first,
};
if (indirect) {
struct agx_resource *rsrc = agx_resource(indirect->buffer);
agx_batch_reads(batch, rsrc);
ia.draws = rsrc->bo->ptr.gpu + indirect->offset;
} else {
ia.verts_per_instance = count;
}
batch->uniforms.input_assembly =
agx_pool_upload_aligned(&batch->pool, &ia, sizeof(ia), 8);
}
static uint64_t
agx_batch_geometry_params(struct agx_batch *batch, uint64_t input_index_buffer,
size_t index_buffer_size_B,
@@ -3915,8 +3961,13 @@ agx_batch_geometry_params(struct agx_batch *batch, uint64_t input_index_buffer,
const struct pipe_draw_start_count_bias *draw,
const struct pipe_draw_indirect_info *indirect)
{
agx_upload_ia_params(batch, info, indirect, draw ? draw->count : 0,
input_index_buffer, index_buffer_size_B, 0);
struct agx_ia_state ia = {
.index_buffer = input_index_buffer,
.verts_per_instance = draw ? draw->count : 0,
};
batch->uniforms.input_assembly =
agx_pool_upload_aligned(&batch->pool, &ia, sizeof(ia), 8);
struct agx_geometry_params params = {
.state = agx_batch_geometry_state(batch),
@@ -4061,10 +4112,29 @@ agx_launch_gs_prerast(struct agx_batch *batch,
.grid = {1, 1, 1},
};
agx_launch(batch, &grid_setup,
agx_build_meta_shader(ctx, agx_nir_gs_setup_indirect, &key,
sizeof(key)),
NULL, PIPE_SHADER_COMPUTE);
uint64_t ib = 0;
size_t ib_extent = 0;
if (info->index_size) {
ib = agx_index_buffer_ptr(batch, info, indirect ? NULL : draws,
&ib_extent);
}
struct agx_resource *rsrc = agx_resource(indirect->buffer);
agx_batch_reads(batch, rsrc);
struct agx_gs_setup_indirect_params gsi = {
.index_buffer = ib,
.draw = rsrc->bo->ptr.gpu + indirect->offset,
.vertex_buffer = batch->uniforms.vertex_output_buffer_ptr,
.ia = batch->uniforms.input_assembly,
.geom = batch->uniforms.geometry_params,
.vs_outputs = batch->uniforms.vertex_outputs,
.index_size_B = info->index_size,
};
agx_launch_with_data(batch, &grid_setup, agx_nir_gs_setup_indirect, &key,
sizeof(key), &gsi, sizeof(gsi));
/* Wrap the pool allocation in a fake resource for meta-Gallium use */
assert(batch->geom_params_bo != NULL);
@@ -4175,7 +4245,18 @@ agx_draw_without_restart(struct agx_batch *batch,
&batch->pool, 5 * sizeof(uint32_t) * indirect->draw_count, 4,
&out_draws_rsrc.bo);
agx_upload_ia_params(batch, info, indirect, 0, ib, ib_extent, out_draws.gpu);
struct agx_resource *indirect_rsrc = agx_resource(indirect->buffer);
agx_batch_reads(batch, indirect_rsrc);
struct agx_restart_unroll_params unroll = {
.heap = agx_batch_geometry_state(batch),
.index_buffer = ib,
.out_draws = out_draws.gpu,
.restart_index = info->restart_index,
.index_buffer_size_B = ib_extent,
.flatshade_first = batch->ctx->rast->base.flatshade_first,
.draws = indirect_rsrc->bo->ptr.gpu + indirect->offset,
};
/* Unroll the index buffer for each draw */
const struct pipe_grid_info grid_setup = {
@@ -4183,10 +4264,8 @@ agx_draw_without_restart(struct agx_batch *batch,
.grid = {indirect->draw_count, 1, 1},
};
agx_launch(
batch, &grid_setup,
agx_build_meta_shader(ctx, agx_nir_unroll_restart, &key, sizeof(key)),
NULL, PIPE_SHADER_COMPUTE);
agx_launch_with_data(batch, &grid_setup, agx_nir_unroll_restart, &key,
sizeof(key), &unroll, sizeof(unroll));
/* Now draw the results without restart */
struct pipe_draw_info new_info = {
@@ -4492,8 +4571,14 @@ agx_draw_patches(struct agx_context *ctx, const struct pipe_draw_info *info,
if (info->index_size)
ib = agx_index_buffer_ptr(batch, info, draws, &ib_extent);
agx_upload_ia_params(batch, info, indirect, draws ? draws->count : 0, ib,
ib_extent, 0);
struct agx_ia_state ia = {
.index_buffer = ib,
.verts_per_instance = draws ? draws->count : 0,
};
batch->uniforms.input_assembly =
agx_pool_upload_aligned(&batch->pool, &ia, sizeof(ia), 8);
agx_upload_draw_params(batch, indirect, draws, info);
/* Setup parameters */
@@ -4956,9 +5041,13 @@ agx_draw_vbo(struct pipe_context *pctx, const struct pipe_draw_info *info,
}
if (IS_DIRTY(RS)) {
struct pipe_rasterizer_state *rs = &ctx->rast->base;
batch->uniforms.fixed_point_size =
ctx->rast->base.point_size_per_vertex ? 0.0
: ctx->rast->base.point_size;
rs->point_size_per_vertex ? 0.0 : rs->point_size;
/* TODO: tri fans */
batch->uniforms.provoking_vertex = !rs->flatshade_first ? 2 : 0;
}
if (IS_DIRTY(QUERY)) {
@@ -5162,58 +5251,13 @@ agx_texture_barrier(struct pipe_context *pipe, unsigned flags)
}
void
agx_launch(struct agx_batch *batch, const struct pipe_grid_info *info,
struct agx_compiled_shader *cs, struct agx_linked_shader *linked,
enum pipe_shader_type stage)
agx_launch_internal(struct agx_batch *batch, const struct pipe_grid_info *info,
struct agx_compiled_shader *cs, enum pipe_shader_type stage,
uint32_t usc)
{
struct agx_context *ctx = batch->ctx;
struct agx_device *dev = agx_device(ctx->base.screen);
/* To implement load_num_workgroups, the number of workgroups needs to be
* available in GPU memory. This is either the indirect buffer, or just a
* buffer we upload ourselves if not indirect.
*/
if (info->indirect) {
struct agx_resource *indirect = agx_resource(info->indirect);
agx_batch_reads(batch, indirect);
batch->uniforms.tables[AGX_SYSVAL_TABLE_GRID] =
indirect->bo->ptr.gpu + info->indirect_offset;
} else {
static_assert(sizeof(info->grid) == 12,
"matches indirect dispatch buffer");
batch->uniforms.tables[AGX_SYSVAL_TABLE_GRID] = agx_pool_upload_aligned(
&batch->pool, info->grid, sizeof(info->grid), 4);
}
util_dynarray_foreach(&ctx->global_buffers, struct pipe_resource *, res) {
if (!*res)
continue;
struct agx_resource *buffer = agx_resource(*res);
agx_batch_writes(batch, buffer, 0);
batch->incoherent_writes = true;
}
agx_batch_add_bo(batch, cs->bo);
agx_update_descriptors(batch, cs);
agx_upload_uniforms(batch);
// TODO: This is broken.
size_t subgroups_per_core = 0;
#if 0
if (!info->indirect) {
size_t subgroups_per_workgroup =
DIV_ROUND_UP(info->block[0] * info->block[1] * info->block[2], 32);
subgroups_per_core =
local_workgroups *
DIV_ROUND_UP(info->grid[0] * info->grid[1] * info->grid[2],
ctx->scratch_cs.num_cores);
}
#endif
/* TODO: Ensure space if we allow multiple kernels in a batch */
uint8_t *out = batch->cdm.current;
@@ -5231,9 +5275,7 @@ agx_launch(struct agx_batch *batch, const struct pipe_grid_info *info,
}
agx_push(out, CDM_LAUNCH_WORD_1, cfg) {
cfg.pipeline =
agx_build_pipeline(batch, cs, linked, PIPE_SHADER_COMPUTE,
info->variable_shared_mem, subgroups_per_core);
cfg.pipeline = usc;
}
/* Added in G14X */
@@ -5243,10 +5285,12 @@ agx_launch(struct agx_batch *batch, const struct pipe_grid_info *info,
}
if (info->indirect) {
struct agx_resource *indirect = agx_resource(info->indirect);
uint64_t addr = indirect->bo->ptr.gpu + info->indirect_offset;
agx_push(out, CDM_INDIRECT, cfg) {
cfg.address_hi = batch->uniforms.tables[AGX_SYSVAL_TABLE_GRID] >> 32;
cfg.address_lo =
batch->uniforms.tables[AGX_SYSVAL_TABLE_GRID] & BITFIELD64_MASK(32);
cfg.address_hi = addr >> 32;
cfg.address_lo = addr & BITFIELD64_MASK(32);
}
} else {
uint32_t size[3];
@@ -5316,6 +5360,65 @@ agx_launch(struct agx_batch *batch, const struct pipe_grid_info *info,
"Failed to reserve sufficient space in encoder");
}
void
agx_launch(struct agx_batch *batch, const struct pipe_grid_info *info,
struct agx_compiled_shader *cs, struct agx_linked_shader *linked,
enum pipe_shader_type stage)
{
struct agx_context *ctx = batch->ctx;
/* To implement load_num_workgroups, the number of workgroups needs to be
* available in GPU memory. This is either the indirect buffer, or just a
* buffer we upload ourselves if not indirect.
*/
if (info->indirect) {
struct agx_resource *indirect = agx_resource(info->indirect);
agx_batch_reads(batch, indirect);
batch->uniforms.tables[AGX_SYSVAL_TABLE_GRID] =
indirect->bo->ptr.gpu + info->indirect_offset;
} else {
static_assert(sizeof(info->grid) == 12,
"matches indirect dispatch buffer");
batch->uniforms.tables[AGX_SYSVAL_TABLE_GRID] = agx_pool_upload_aligned(
&batch->pool, info->grid, sizeof(info->grid), 4);
}
util_dynarray_foreach(&ctx->global_buffers, struct pipe_resource *, res) {
if (!*res)
continue;
struct agx_resource *buffer = agx_resource(*res);
agx_batch_writes(batch, buffer, 0);
batch->incoherent_writes = true;
}
agx_batch_add_bo(batch, cs->bo);
agx_update_descriptors(batch, cs);
agx_upload_uniforms(batch);
// TODO: This is broken.
size_t subgroups_per_core = 0;
#if 0
if (!info->indirect) {
size_t subgroups_per_workgroup =
DIV_ROUND_UP(info->block[0] * info->block[1] * info->block[2], 32);
subgroups_per_core =
local_workgroups *
DIV_ROUND_UP(info->grid[0] * info->grid[1] * info->grid[2],
ctx->scratch_cs.num_cores);
}
#endif
uint32_t usc =
agx_build_pipeline(batch, cs, linked, PIPE_SHADER_COMPUTE,
info->variable_shared_mem, subgroups_per_core);
agx_launch_internal(batch, info, cs, stage, usc);
}
static void
agx_launch_grid(struct pipe_context *pipe, const struct pipe_grid_info *info)
{
src/gallium/drivers/asahi/agx_state.h
+24 -10
View File
@@ -166,6 +166,9 @@ struct PACKED agx_draw_uniforms {
/* ~0/0 boolean whether the epilog lacks any discard instrction */
uint16_t no_epilog_discard;
/* Provoking vertex: 0, 1, 2 */
uint16_t provoking_vertex;
/* Mapping from varying slots written by the last vertex stage to UVS
* indices. This mapping must be compatible with the fragment shader.
*/
@@ -759,6 +762,27 @@ agx_context(struct pipe_context *pctx)
}
struct agx_linked_shader;
typedef void (*meta_shader_builder_t)(struct nir_builder *b, const void *key);
void agx_init_meta_shaders(struct agx_context *ctx);
void agx_destroy_meta_shaders(struct agx_context *ctx);
struct agx_compiled_shader *agx_build_meta_shader(struct agx_context *ctx,
meta_shader_builder_t builder,
void *data, size_t data_size);
void agx_launch_with_data(struct agx_batch *batch,
const struct pipe_grid_info *info,
meta_shader_builder_t builder, void *key,
size_t key_size, void *data, size_t data_size);
void agx_launch_internal(struct agx_batch *batch,
const struct pipe_grid_info *info,
struct agx_compiled_shader *cs,
enum pipe_shader_type stage, uint32_t usc);
void agx_launch(struct agx_batch *batch, const struct pipe_grid_info *info,
struct agx_compiled_shader *cs,
struct agx_linked_shader *linked, enum pipe_shader_type stage);
@@ -1149,13 +1173,3 @@ agx_texture_buffer_size_el(enum pipe_format format, uint32_t size)
return MIN2(AGX_TEXTURE_BUFFER_MAX_SIZE, size / blocksize);
}
typedef void (*meta_shader_builder_t)(struct nir_builder *b, const void *key);
void agx_init_meta_shaders(struct agx_context *ctx);
void agx_destroy_meta_shaders(struct agx_context *ctx);
struct agx_compiled_shader *agx_build_meta_shader(struct agx_context *ctx,
meta_shader_builder_t builder,
void *data, size_t data_size);