diff --git a/src/asahi/lib/agx_nir_lower_gs.c b/src/asahi/lib/agx_nir_lower_gs.c index 5d3848c4c75..1da1b31371e 100644 --- a/src/asahi/lib/agx_nir_lower_gs.c +++ b/src/asahi/lib/agx_nir_lower_gs.c @@ -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); diff --git a/src/asahi/lib/shaders/geometry.cl b/src/asahi/lib/shaders/geometry.cl index e5bb5afaed9..a8ed851596d 100644 --- a/src/asahi/lib/shaders/geometry.cl +++ b/src/asahi/lib/shaders/geometry.cl @@ -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) diff --git a/src/asahi/lib/shaders/geometry.h b/src/asahi/lib/shaders/geometry.h index 01867267e38..27ecff6b865 100644 --- a/src/asahi/lib/shaders/geometry.h +++ b/src/asahi/lib/shaders/geometry.h @@ -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 */ diff --git a/src/compiler/nir/nir_intrinsics.py b/src/compiler/nir/nir_intrinsics.py index 4ebd319cabe..db8622cb1d9 100644 --- a/src/compiler/nir/nir_intrinsics.py +++ b/src/compiler/nir/nir_intrinsics.py @@ -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) diff --git a/src/gallium/drivers/asahi/agx_nir_lower_sysvals.c b/src/gallium/drivers/asahi/agx_nir_lower_sysvals.c index 93466e1ffe3..a2375923b08 100644 --- a/src/gallium/drivers/asahi/agx_nir_lower_sysvals.c +++ b/src/gallium/drivers/asahi/agx_nir_lower_sysvals.c @@ -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; } diff --git a/src/gallium/drivers/asahi/agx_state.c b/src/gallium/drivers/asahi/agx_state.c index 3af8bc582fa..300e90b940a 100644 --- a/src/gallium/drivers/asahi/agx_state.c +++ b/src/gallium/drivers/asahi/agx_state.c @@ -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) { diff --git a/src/gallium/drivers/asahi/agx_state.h b/src/gallium/drivers/asahi/agx_state.h index dc8f9c60577..76cd5eb7718 100644 --- a/src/gallium/drivers/asahi/agx_state.h +++ b/src/gallium/drivers/asahi/agx_state.h @@ -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);