diff --git a/src/asahi/clc/asahi_clc.c b/src/asahi/clc/asahi_clc.c index 0fb2999c749..a5c9dc8554e 100644 --- a/src/asahi/clc/asahi_clc.c +++ b/src/asahi/clc/asahi_clc.c @@ -5,20 +5,30 @@ */ #include "asahi/compiler/agx_compile.h" +#include "asahi/compiler/agx_nir.h" #include "compiler/glsl_types.h" #include "compiler/spirv/nir_spirv.h" #include "nir.h" #include "nir_builder.h" -#include "nir_serialize.h" +#include "nir_builder_opcodes.h" +#include "nir_intrinsics.h" +#include "nir_precompiled.h" +#include "shader_enums.h" #include #include #include #include #include -#include "util/u_math.h" +#include "util/macros.h" #include +const char *targets[] = {"g13g", "g13x"}; + +#define foreach_target(target) \ + for (const char **target = &targets[0]; \ + target < &targets[ARRAY_SIZE(targets)]; ++target) + static const struct spirv_to_nir_options spirv_options = { .environment = NIR_SPIRV_OPENCL, .shared_addr_format = nir_address_format_62bit_generic, @@ -36,6 +46,8 @@ optimize(nir_shader *nir) do { progress = false; + NIR_PASS(progress, nir, nir_split_var_copies); + NIR_PASS(progress, nir, nir_split_struct_vars, nir_var_function_temp); NIR_PASS(progress, nir, nir_lower_var_copies); NIR_PASS(progress, nir, nir_lower_vars_to_ssa); @@ -58,8 +70,6 @@ optimize(nir_shader *nir) NIR_PASS(progress, nir, nir_opt_shrink_vectors, true); NIR_PASS(progress, nir, nir_opt_loop_unroll); - NIR_PASS(progress, nir, nir_split_var_copies); - NIR_PASS(progress, nir, nir_split_struct_vars, nir_var_function_temp); } while (progress); } @@ -76,9 +86,14 @@ compile(void *memctx, const uint32_t *spirv, size_t spirv_size) nir_validate_ssa_dominance(nir, "after spirv_to_nir"); ralloc_steal(memctx, nir); + nir_fixup_is_exported(nir); + NIR_PASS(_, nir, nir_lower_system_values); NIR_PASS(_, nir, nir_lower_calls_to_builtins); + nir_lower_compute_system_values_options cs = {.global_id_is_32bit = true}; + NIR_PASS(_, nir, nir_lower_compute_system_values, &cs); + /* We have to lower away local constant initializers right before we * inline functions. That way they get properly initialized at the top * of the function and not at the top of its caller. @@ -90,6 +105,9 @@ compile(void *memctx, const uint32_t *spirv, size_t spirv_size) NIR_PASS(_, nir, nir_copy_prop); NIR_PASS(_, nir, nir_opt_deref); + /* We can't deal with constant data, get rid of it */ + nir_lower_constant_to_temp(nir); + /* We can go ahead and lower the rest of the constant initializers. We do * this here so that nir_remove_dead_variables and split_per_member_structs * below see the corresponding stores. @@ -124,13 +142,7 @@ compile(void *memctx, const uint32_t *spirv, size_t spirv_size) nir_var_shader_temp | nir_var_function_temp | nir_var_mem_shared | nir_var_mem_global | nir_var_mem_constant, glsl_get_cl_type_size_align); - if (nir->constant_data_size > 0) { - assert(nir->constant_data == NULL); - nir->constant_data = rzalloc_size(nir, nir->constant_data_size); - nir_gather_explicit_io_initializers(nir, nir->constant_data, - nir->constant_data_size, - nir_var_mem_constant); - } + assert(nir->constant_data_size == 0); NIR_PASS(_, nir, nir_lower_memcpy); @@ -148,47 +160,82 @@ compile(void *memctx, const uint32_t *spirv, size_t spirv_size) NIR_PASS(_, nir, nir_opt_if, 0); NIR_PASS(_, nir, nir_opt_idiv_const, 16); + NIR_PASS(_, nir, agx_nir_lower_texture_early, false /* support_lod_bias */); + NIR_PASS(_, nir, agx_nir_lower_texture); + NIR_PASS(_, nir, agx_nir_lower_multisampled_image_store); + optimize(nir); return nir; } static void -print_u32_data(FILE *fp, const char *prefix, const char *arr_name, - const uint32_t *data, size_t len) +print_shader(FILE *fp, const char *name, const char *suffix, uint32_t variant, + struct agx_shader_part *p) { - fprintf(fp, "static const uint32_t %s_%s[] = {", prefix, arr_name); - for (unsigned i = 0; i < (len / 4); i++) { - if (i % 4 == 0) - fprintf(fp, "\n "); + struct agx_precompiled_kernel_info info = agx_compact_kernel_info(&p->info); + size_t sz_B = sizeof(info) + p->info.binary_size; + size_t sz_el = DIV_ROUND_UP(sz_B, 4); + uint32_t *mem = calloc(sz_el, 4); - fprintf(fp, " 0x%08" PRIx32 ",", data[i]); + memcpy(mem, &info, sizeof(info)); + memcpy((uint8_t *)mem + sizeof(info), p->binary, p->info.binary_size); + + nir_precomp_print_blob(fp, name, suffix, variant, mem, sz_B); + free(mem); +} + +static bool +gather_atomic_info(nir_builder *b, nir_intrinsic_instr *intr, void *data) +{ + bool *any = data; + + switch (intr->intrinsic) { + case nir_intrinsic_global_atomic: + case nir_intrinsic_global_atomic_agx: + case nir_intrinsic_deref_atomic: + case nir_intrinsic_global_atomic_swap: + case nir_intrinsic_global_atomic_swap_agx: + case nir_intrinsic_deref_atomic_swap: + *any = true; + return false; + default: + return false; } +} - if (len % 4) { - const uint8_t *data_u8 = (const uint8_t *)data; - uint32_t last = 0; - unsigned last_offs = ROUND_DOWN_TO(len, 4); - for (unsigned i = 0; i < len % 4; ++i) { - last |= (uint32_t)data_u8[last_offs + i] << (i * 8); - } +/* G13X variants are only compiled when atomics are used */ +static const char * +remap_variant(nir_function *func, unsigned variant, const char *target) +{ + bool has_atomic = func->pass_flags & BITFIELD_BIT(variant); - fprintf(fp, " 0x%08" PRIx32 ",", last); - } + if (!has_atomic && !strcmp(target, "g13x")) + return "g13g"; + else + return target; +} - fprintf(fp, "\n};\n"); +static nir_def * +load_kernel_input(nir_builder *b, unsigned num_components, unsigned bit_size, + unsigned offset_B) +{ + assert((offset_B & 1) == 0 && "half-aligned"); + return nir_load_preamble(b, num_components, bit_size, .base = offset_B / 2); } int main(int argc, char **argv) { - if (argc != 3) { - fprintf(stderr, "Usage: %s [input spir-v] [output header]\n", argv[0]); + if (argc != 4) { + fprintf(stderr, "Usage: %s [input spir-v] [output header] [output C]\n", + argv[0]); return 1; } const char *infile = argv[1]; - const char *outfile = argv[2]; + const char *outh_file = argv[2]; + const char *outc_file = argv[3]; void *mem_ctx = ralloc_context(NULL); @@ -209,56 +256,105 @@ main(int argc, char **argv) return 1; } - FILE *fp = fopen(outfile, "w"); + FILE *fp_h = fopen(outh_file, "w"); + FILE *fp_c = fopen(outc_file, "w"); glsl_type_singleton_init_or_ref(); - fprintf(fp, "/*\n"); - fprintf(fp, " * Copyright The Asahi Linux Contributors\n"); - fprintf(fp, " * SPDX-License-Identifier: MIT\n"); - fprintf(fp, " *\n"); - fprintf(fp, " * Autogenerated file, do not edit\n"); - fprintf(fp, " */\n"); - fprintf(fp, " #include \n"); + nir_precomp_print_header(fp_c, fp_h, "The Asahi Linux Contributors", + "libagx_shaders.h"); - /* Compile SPIR-V to NIR */ nir_shader *nir = compile(mem_ctx, spirv_map, spirv_len); - { - nir_builder b = nir_builder_init_simple_shader( - MESA_SHADER_COMPUTE, &agx_nir_options, "Helper shader"); + nir_foreach_entrypoint(libfunc, nir) { + libfunc->pass_flags = 0; + struct nir_precomp_layout layout = nir_precomp_derive_layout(libfunc); + unsigned nr_vars = nir_precomp_nr_variants(libfunc); - nir_function *func = - nir_shader_get_function_for_name(nir, "libagx_helper"); + nir_precomp_print_layout_struct(fp_h, libfunc); - nir_call(&b, nir_function_clone(b.shader, func)); + for (unsigned v = 0; v < nr_vars; ++v) { + nir_shader *s = nir_precompiled_build_variant( + libfunc, v, &agx_nir_options, load_kernel_input); - struct agx_shader_part compiled; - struct agx_shader_key key = { - .libagx = nir, - .is_helper = true, - }; + agx_link_libagx(s, nir); - agx_preprocess_nir(b.shader, nir); - agx_compile_shader_nir(b.shader, &key, NULL, &compiled); + NIR_PASS(_, s, nir_lower_vars_to_explicit_types, nir_var_mem_shared, + glsl_get_cl_type_size_align); - print_u32_data(fp, "libagx_g13", "helper", compiled.binary, - compiled.info.binary_size); - free(compiled.binary); - ralloc_free(b.shader); + NIR_PASS(_, s, nir_lower_explicit_io, nir_var_mem_shared, + nir_address_format_62bit_generic); - /* Remove the NIR function, it's compiled, we don't need it at runtime */ - exec_node_remove(&func->node); + /* Unroll loops before lowering indirects */ + bool progress = false; + do { + progress = false; + NIR_PASS(progress, s, nir_opt_loop); + } while (progress); + + agx_preprocess_nir(s, NULL); + + bool has_atomic = false; + nir_shader_intrinsics_pass(s, gather_atomic_info, nir_metadata_all, + &has_atomic); + if (has_atomic) { + libfunc->pass_flags |= BITFIELD_BIT(v); + } + + foreach_target(target) + { + /* Skip unused variants */ + if (strcmp(*target, remap_variant(libfunc, v, *target))) + continue; + + struct agx_shader_part compiled; + bool is_helper = !strcmp(libfunc->name, "libagx_helper"); + struct agx_shader_key key = { + .libagx = nir, + .promote_constants = !is_helper, + .reserved_preamble = layout.size_B / 2, + .is_helper = is_helper, + }; + + if (has_atomic) { + key.dev.needs_g13x_coherency = + u_tristate_make(!strcmp(*target, "g13x")); + } + + nir_shader *clone = nir_shader_clone(NULL, s); + agx_compile_shader_nir(clone, &key, NULL, &compiled); + print_shader(fp_c, libfunc->name, *target, v, &compiled); + free(compiled.binary); + ralloc_free(clone); + + assert(compiled.info.scratch_size == 0 && + "internal shaders do not spill"); + + assert(compiled.info.preamble_scratch_size == 0 && + "internal shader preambles do not spill"); + } + + ralloc_free(s); + } } - /* Serialize NIR for embedding */ - struct blob blob; - blob_init(&blob); - nir_serialize(&blob, nir, true /* strip */); - print_u32_data(fp, "libagx", "nir", (const uint32_t *)blob.data, blob.size); - blob_finish(&blob); + nir_precomp_print_program_enum(fp_h, nir, "libagx"); + nir_precomp_print_dispatch_macros(fp_h, nir); + + /* For each target, generate a table mapping programs to binaries */ + foreach_target(target) + { + nir_precomp_print_extern_binary_map(fp_h, "libagx", *target); + nir_precomp_print_binary_map(fp_c, nir, "libagx", *target, remap_variant); + } + + /* Remove the NIR functions we compiled to binaries to save memory */ + nir_remove_entrypoints(nir); + + nir_precomp_print_nir(fp_c, fp_h, nir, "libagx", "nir"); glsl_type_singleton_decref(); - fclose(fp); + fclose(fp_c); + fclose(fp_h); ralloc_free(mem_ctx); return 0; } diff --git a/src/asahi/compiler/agx_compile.c b/src/asahi/compiler/agx_compile.c index 0379a9118e9..4f0af5fd29b 100644 --- a/src/asahi/compiler/agx_compile.c +++ b/src/asahi/compiler/agx_compile.c @@ -2977,7 +2977,7 @@ optimize_bounds(nir_builder *b, nir_intrinsic_instr *intr, void *data) } static void -agx_optimize_nir(nir_shader *nir, bool soft_fault, unsigned *preamble_size) +agx_optimize_nir(nir_shader *nir, bool soft_fault, uint16_t *preamble_size) { /* This runs only once up front since other optimizations don't affect it */ NIR_PASS(_, nir, nir_opt_shrink_stores, true); @@ -3075,8 +3075,11 @@ agx_optimize_nir(nir_shader *nir, bool soft_fault, unsigned *preamble_size) */ NIR_PASS(_, nir, agx_nir_lower_fminmax); - if (preamble_size && (!(agx_compiler_debug & AGX_DBG_NOPREAMBLE))) - NIR_PASS(_, nir, agx_nir_opt_preamble, preamble_size); + if (preamble_size && (!(agx_compiler_debug & AGX_DBG_NOPREAMBLE))) { + unsigned temp = *preamble_size; + NIR_PASS(_, nir, agx_nir_opt_preamble, &temp); + *preamble_size = temp; + } /* Forming preambles may dramatically reduce the instruction count * in certain blocks, causing some if-else statements to become diff --git a/src/asahi/compiler/agx_compile.h b/src/asahi/compiler/agx_compile.h index 44da67e7733..b7023242c36 100644 --- a/src/asahi/compiler/agx_compile.h +++ b/src/asahi/compiler/agx_compile.h @@ -65,6 +65,17 @@ struct agx_interp_info { }; static_assert(sizeof(struct agx_interp_info) == 16, "packed"); +struct agx_rodata { + /* Offset in the binary */ + uint32_t offset; + + /* Base uniform to map constants */ + uint16_t base_uniform; + + /* Number of 16-bit constants to map contiguously there */ + uint16_t size_16; +}; + struct agx_shader_info { enum pipe_shader_type stage; uint32_t binary_size; @@ -72,13 +83,13 @@ struct agx_shader_info { union agx_varyings varyings; /* Number of uniforms */ - unsigned push_count; + uint16_t push_count; /* Local memory allocation in bytes */ - unsigned local_size; + uint16_t local_size; /* Local imageblock allocation in bytes per thread */ - unsigned imageblock_stride; + uint16_t imageblock_stride; /* Scratch memory allocation in bytes for main/preamble respectively */ unsigned scratch_size, preamble_scratch_size; @@ -131,7 +142,7 @@ struct agx_shader_info { /* Number of 16-bit registers used by the main shader and preamble * respectively. */ - unsigned nr_gprs, nr_preamble_gprs; + uint16_t nr_gprs, nr_preamble_gprs; /* Output mask set during driver lowering */ uint64_t outputs; @@ -142,18 +153,43 @@ struct agx_shader_info { /* There may be constants in the binary. The driver must map these to uniform * registers as specified hre. */ - struct { - /* Offset in the binary */ - uint32_t offset; - - /* Base uniform to map constants */ - uint16_t base_uniform; - - /* Number of 16-bit constants to map contiguously there */ - uint16_t size_16; - } rodata; + struct agx_rodata rodata; }; +struct agx_precompiled_kernel_info { + uint32_t preamble_offset, main_offset; + uint32_t main_size, binary_size; + struct agx_rodata rodata; + uint16_t nr_gprs, nr_preamble_gprs; + uint16_t push_count; + uint16_t workgroup_size[3]; + uint16_t local_size; + uint16_t imageblock_stride; + bool uses_txf; +}; + +static inline struct agx_precompiled_kernel_info +agx_compact_kernel_info(struct agx_shader_info *info) +{ + assert(info->has_preamble == (info->nr_preamble_gprs > 0)); + + return (struct agx_precompiled_kernel_info){ + .preamble_offset = info->preamble_offset, + .main_offset = info->main_offset, + .main_size = info->main_size, + .binary_size = info->binary_size, + .rodata = info->rodata, + .nr_gprs = info->nr_gprs, + .nr_preamble_gprs = info->nr_preamble_gprs, + .push_count = info->push_count, + .workgroup_size = {info->workgroup_size[0], info->workgroup_size[1], + info->workgroup_size[2]}, + .local_size = info->local_size, + .imageblock_stride = info->imageblock_stride, + .uses_txf = info->uses_txf, + }; +} + struct agx_shader_part { struct agx_shader_info info; void *binary; diff --git a/src/asahi/lib/agx_bg_eot.c b/src/asahi/lib/agx_bg_eot.c index f3a52d52b56..b9057758484 100644 --- a/src/asahi/lib/agx_bg_eot.c +++ b/src/asahi/lib/agx_bg_eot.c @@ -4,14 +4,19 @@ */ #include "agx_bg_eot.h" +#include "util/simple_mtx.h" +#include "util/u_debug.h" #include "agx_compile.h" #include "agx_device.h" #include "agx_nir.h" #include "agx_nir_texture.h" #include "agx_tilebuffer.h" +#include "agx_usc.h" +#include "libagx_shaders.h" #include "nir.h" #include "nir_builder.h" #include "nir_intrinsics.h" +#include "pool.h" static bool lower_tex_handle_to_u0(nir_builder *b, nir_intrinsic_instr *intr, void *data) @@ -213,6 +218,7 @@ void agx_bg_eot_init(struct agx_bg_eot_cache *cache, struct agx_device *dev) { agx_pool_init(&cache->pool, dev, AGX_BO_EXEC | AGX_BO_LOW_VA, true); + simple_mtx_init(&cache->lock, mtx_plain); cache->ht = agx_bg_eot_key_table_create(NULL); cache->dev = dev; } @@ -222,6 +228,109 @@ agx_bg_eot_cleanup(struct agx_bg_eot_cache *cache) { agx_pool_cleanup(&cache->pool); _mesa_hash_table_destroy(cache->ht, NULL); + simple_mtx_destroy(&cache->lock); cache->ht = NULL; cache->dev = NULL; } + +static struct agx_precompiled_shader * +agx_get_precompiled_locked(struct agx_bg_eot_cache *cache, unsigned program) +{ + simple_mtx_assert_locked(&cache->lock); + + /* It is possible that, while waiting for the lock, another thread uploaded + * the shader. Check for that so we don't double-upload. + */ + if (cache->precomp[program]) + return cache->precomp[program]; + + /* Otherwise, we need to upload. */ + struct agx_precompiled_shader *p = + ralloc(cache->ht, struct agx_precompiled_shader); + + const uint32_t *bin = cache->dev->libagx_programs[program]; + const struct agx_precompiled_kernel_info *info = (void *)bin; + const void *binary = (const uint8_t *)bin + sizeof(*info); + + assert(info->main_offset == 0 || program != LIBAGX_HELPER); + + p->b.workgroup = + agx_workgroup(info->workgroup_size[0], info->workgroup_size[1], + info->workgroup_size[2]); + + p->ptr = agx_pool_upload_aligned_with_bo(&cache->pool, binary, + info->binary_size, 128, &p->bo); + + /* Bake launch */ + agx_pack(&p->b.launch, CDM_LAUNCH_WORD_0, cfg) { + cfg.sampler_state_register_count = 1; + cfg.uniform_register_count = info->push_count; + cfg.preshader_register_count = info->nr_preamble_gprs; + } + + /* Bake USC */ + struct agx_usc_builder b = + agx_usc_builder(p->b.usc.data, sizeof(p->b.usc.data)); + + agx_usc_immediates(&b, &info->rodata, p->ptr); + + if (info->uses_txf) + agx_usc_push_packed(&b, SAMPLER, cache->dev->txf_sampler); + + agx_usc_shared(&b, info->local_size, info->imageblock_stride, 0); + + agx_usc_pack(&b, SHADER, cfg) { + cfg.code = agx_usc_addr(cache->dev, p->ptr + info->main_offset); + cfg.unk_2 = 3; + } + + agx_usc_pack(&b, REGISTERS, cfg) { + cfg.register_count = info->nr_gprs; + cfg.spill_size = 0; + } + + if (info->nr_preamble_gprs) { + agx_usc_pack(&b, PRESHADER, cfg) { + cfg.code = agx_usc_addr(cache->dev, p->ptr + info->preamble_offset); + } + } else { + agx_usc_pack(&b, NO_PRESHADER, cfg) + ; + } + + p->b.usc.size = b.head - p->b.usc.data; + + /* We must only write to the cache once we are done compiling, since other + * threads may be reading the cache concurrently. Do this last. + */ + p_atomic_set(&cache->precomp[program], p); + return p; +} + +struct agx_precompiled_shader * +agx_get_precompiled(struct agx_bg_eot_cache *cache, unsigned program) +{ + /* Shaders are immutable once written, so if we atomically read a non-NULL + * shader, then we have a valid cached shader and are done. + */ + struct agx_precompiled_shader *ret = p_atomic_read(cache->precomp + program); + + if (ret != NULL) + return ret; + + /* Otherwise, take the lock and upload. */ + simple_mtx_lock(&cache->lock); + ret = agx_get_precompiled_locked(cache, program); + simple_mtx_unlock(&cache->lock); + + return ret; +} + +uint64_t +agx_helper_program(struct agx_bg_eot_cache *cache) +{ + struct agx_precompiled_shader *pc = + agx_get_precompiled(cache, LIBAGX_HELPER); + + return pc->ptr | 1; +} diff --git a/src/asahi/lib/agx_bg_eot.h b/src/asahi/lib/agx_bg_eot.h index 610625fbe36..0bffe214855 100644 --- a/src/asahi/lib/agx_bg_eot.h +++ b/src/asahi/lib/agx_bg_eot.h @@ -6,17 +6,40 @@ #pragma once #include "asahi/compiler/agx_compile.h" +#include "util/simple_mtx.h" #include "agx_tilebuffer.h" +#include "libagx_dgc.h" +#include "libagx_shaders.h" #include "pool.h" +struct agx_precompiled_shader { + struct agx_shader b; + struct agx_bo *bo; + uint64_t ptr; +}; + struct agx_bg_eot_cache { struct agx_device *dev; struct agx_pool pool; + simple_mtx_t lock; /* Map from agx_bg_eot_key to agx_bg_eot_shader */ struct hash_table *ht; + + struct agx_precompiled_shader *precomp[LIBAGX_NUM_PROGRAMS]; }; +/* + * Get a precompiled shader, uploading if necessary. This is thread-safe. + */ +struct agx_precompiled_shader * +agx_get_precompiled(struct agx_bg_eot_cache *cache, unsigned program); + +/* + * Get the address of the cached helper program. This is thread-safe. + */ +uint64_t agx_helper_program(struct agx_bg_eot_cache *cache); + enum agx_bg_eot_op { AGX_BG_EOT_NONE, AGX_BG_CLEAR, diff --git a/src/asahi/lib/agx_device.c b/src/asahi/lib/agx_device.c index e66e1f2d486..56d3d868c5d 100644 --- a/src/asahi/lib/agx_device.c +++ b/src/asahi/lib/agx_device.c @@ -15,6 +15,7 @@ #include "agx_scratch.h" #include "decode.h" #include "glsl_types.h" +#include "libagx_dgc.h" #include "libagx_shaders.h" #include @@ -570,10 +571,25 @@ agx_open_device(void *memctx, struct agx_device *dev) glsl_type_singleton_init_or_ref(); struct blob_reader blob; - blob_reader_init(&blob, (void *)libagx_nir, sizeof(libagx_nir)); + blob_reader_init(&blob, (void *)libagx_0_nir, sizeof(libagx_0_nir)); dev->libagx = nir_deserialize(memctx, &agx_nir_options, &blob); - dev->helper = agx_build_helper(dev); + if (agx_gather_device_key(dev).needs_g13x_coherency == U_TRISTATE_YES) { + dev->libagx_programs = libagx_g13x; + } else { + dev->libagx_programs = libagx_g13g; + } + + if (dev->params.gpu_generation >= 14 && dev->params.num_clusters_total > 1) { + dev->chip = AGX_CHIP_G14X; + } else if (dev->params.gpu_generation >= 14) { + dev->chip = AGX_CHIP_G14G; + } else if (dev->params.gpu_generation >= 13 && + dev->params.num_clusters_total > 1) { + dev->chip = AGX_CHIP_G13X; + } else { + dev->chip = AGX_CHIP_G13G; + } return true; } @@ -582,7 +598,6 @@ void agx_close_device(struct agx_device *dev) { ralloc_free((void *)dev->libagx); - agx_bo_unreference(dev, dev->helper); agx_bo_cache_evict_all(dev); util_sparse_array_finish(&dev->bo_map); agxdecode_destroy_context(dev->agxdecode); diff --git a/src/asahi/lib/agx_device.h b/src/asahi/lib/agx_device.h index dfb19eeabcf..3f0d762530d 100644 --- a/src/asahi/lib/agx_device.h +++ b/src/asahi/lib/agx_device.h @@ -13,8 +13,10 @@ #include "util/timespec.h" #include "util/vma.h" #include "agx_bo.h" +#include "agx_pack.h" #include "decode.h" #include "layout.h" +#include "libagx_dgc.h" #include "unstable_asahi_drm.h" #include "vdrm.h" @@ -90,6 +92,9 @@ struct agx_device { /* NIR library of AGX helpers/shaders. Immutable once created. */ const struct nir_shader *libagx; + /* Precompiled libagx binary table */ + const uint32_t **libagx_programs; + char name[64]; struct drm_asahi_params_global params; uint64_t next_global_id, last_global_id; @@ -142,9 +147,15 @@ struct agx_device { uint64_t hits, misses; } bo_cache; - struct agx_bo *helper; - struct agxdecode_ctx *agxdecode; + + /* Prepacked USC Sampler word to bind the txf sampler, used for + * precompiled shaders on both drivers. + */ + struct agx_usc_sampler_packed txf_sampler; + + /* Simplified device selection */ + enum agx_chip chip; }; static inline bool diff --git a/src/asahi/lib/agx_helpers.h b/src/asahi/lib/agx_helpers.h index d3e6c83f326..36298d368eb 100644 --- a/src/asahi/lib/agx_helpers.h +++ b/src/asahi/lib/agx_helpers.h @@ -8,22 +8,13 @@ #include #include "asahi/compiler/agx_compile.h" #include "asahi/layout/layout.h" -#include "asahi/libagx/compression.h" #include "agx_pack.h" #include "agx_ppp.h" +#include "libagx_shaders.h" #define AGX_MAX_OCCLUSION_QUERIES (32768) #define AGX_MAX_VIEWPORTS (16) -#define agx_push(ptr, T, cfg) \ - for (unsigned _loop = 0; _loop < 1; ++_loop, ptr += AGX_##T##_LENGTH) \ - agx_pack(ptr, T, cfg) - -#define agx_push_packed(ptr, src, T) \ - STATIC_ASSERT(sizeof(src) == AGX_##T##_LENGTH); \ - memcpy(ptr, &src, sizeof(src)); \ - ptr += sizeof(src); - static inline enum agx_sampler_states agx_translate_sampler_state_count(unsigned count, bool extended) { @@ -270,12 +261,12 @@ agx_calculate_vbo_clamp(uint64_t vbuf, uint64_t sink, enum pipe_format format, } } -static void -agx_fill_decompress_push(struct libagx_decompress_push *push, - struct ail_layout *layout, unsigned layer, - unsigned level, uint64_t ptr) +static struct libagx_decompress_args +agx_fill_decompress_args(struct ail_layout *layout, unsigned layer, + unsigned level, uint64_t ptr, uint64_t images) { - *push = (struct libagx_decompress_push){ + return (struct libagx_decompress_args){ + .images = images, .tile_uncompressed = ail_tile_mode_uncompressed(layout->format), .metadata = ptr + layout->metadata_offset_B + layout->level_offsets_compressed_B[level] + @@ -286,6 +277,23 @@ agx_fill_decompress_push(struct libagx_decompress_push *push, }; } +#undef libagx_decompress +#define libagx_decompress(context, grid, layout, layer, level, ptr, images) \ + libagx_decompress_struct( \ + context, grid, \ + agx_fill_decompress_args(layout, layer, level, ptr, images), \ + util_logbase2(layout->sample_count_sa)) + +#define libagx_tessellate(context, grid, prim, mode, state) \ + if (prim == TESS_PRIMITIVE_QUADS) { \ + libagx_tess_quad(context, grid, state, mode); \ + } else if (prim == TESS_PRIMITIVE_TRIANGLES) { \ + libagx_tess_tri(context, grid, state, mode); \ + } else { \ + assert(prim == TESS_PRIMITIVE_ISOLINES); \ + libagx_tess_isoline(context, grid, state, mode); \ + } + struct agx_border_packed; void agx_pack_border(struct agx_border_packed *out, const uint32_t in[4], diff --git a/src/asahi/lib/agx_nir_lower_gs.c b/src/asahi/lib/agx_nir_lower_gs.c index 8521255b75f..82e1302abf2 100644 --- a/src/asahi/lib/agx_nir_lower_gs.c +++ b/src/asahi/lib/agx_nir_lower_gs.c @@ -1498,128 +1498,3 @@ agx_nir_lower_vs_before_gs(struct nir_shader *vs, return progress; } - -void -agx_nir_prefix_sum_gs(nir_builder *b, const void *data) -{ - const unsigned *words = data; - - b->shader->info.workgroup_size[0] = 1024; - - libagx_prefix_sum(b, load_geometry_param(b, count_buffer), - load_geometry_param(b, input_primitives), - nir_imm_int(b, *words), - nir_channel(b, nir_load_workgroup_id(b), 0)); -} - -void -agx_nir_prefix_sum_tess(nir_builder *b, const void *data) -{ - b->shader->info.workgroup_size[0] = 1024; - libagx_prefix_sum_tess(b, nir_load_preamble(b, 1, 64, .base = 0)); -} - -void -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_preamble(b, 1, 64, .base = 0), - nir_imm_int(b, key->prim), - nir_channel(b, nir_load_local_invocation_id(b), 0)); -} - -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; - - 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); - - libagx_unroll_restart(b, ia, mode, draw, lane); -} - -void -agx_nir_tessellate(nir_builder *b, const void *data) -{ - const struct agx_tessellator_key *key = data; - b->shader->info.workgroup_size[0] = 64; - - nir_def *params = nir_load_preamble(b, 1, 64, .base = 0); - nir_def *patch = nir_channel(b, nir_load_global_invocation_id(b, 32), 0); - nir_def *mode = nir_imm_int(b, key->mode); - - if (key->prim == TESS_PRIMITIVE_ISOLINES) - libagx_tess_isoline(b, params, mode, patch); - else if (key->prim == TESS_PRIMITIVE_TRIANGLES) - libagx_tess_tri(b, params, mode, patch); - else if (key->prim == TESS_PRIMITIVE_QUADS) - libagx_tess_quad(b, params, mode, patch); - else - unreachable("invalid tess primitive"); -} - -void -agx_nir_tess_setup_indirect(nir_builder *b, const void *data) -{ - const struct agx_tess_setup_indirect_key *key = data; - - nir_def *params = nir_load_preamble(b, 1, 64, .base = 0); - nir_def *point_mode = nir_imm_bool(b, key->point_mode); - - libagx_tess_setup_indirect(b, params, point_mode); -} - -void -agx_nir_increment_statistic(nir_builder *b, const void *data) -{ - libagx_increment_statistic(b, nir_load_preamble(b, 1, 64, .base = 0)); -} - -void -agx_nir_increment_cs_invocations(nir_builder *b, const void *data) -{ - libagx_increment_cs_invocations(b, nir_load_preamble(b, 1, 64, .base = 0)); -} - -void -agx_nir_increment_ia_counters(nir_builder *b, const void *data) -{ - const struct agx_increment_ia_counters_key *key = data; - b->shader->info.workgroup_size[0] = key->index_size_B ? 1024 : 1; - - nir_def *params = nir_load_preamble(b, 1, 64, .base = 0); - nir_def *index_size_B = nir_imm_int(b, key->index_size_B); - nir_def *thread = nir_channel(b, nir_load_global_invocation_id(b, 32), 0); - - libagx_increment_ia_counters(b, params, index_size_B, thread); -} - -void -agx_nir_predicate_indirect(nir_builder *b, const void *data) -{ - const struct agx_predicate_indirect_key *key = data; - - nir_def *params = nir_load_preamble(b, 1, 64, .base = 0); - nir_def *indexed = nir_imm_bool(b, key->indexed); - nir_def *thread = nir_channel(b, nir_load_global_invocation_id(b, 32), 0); - - libagx_predicate_indirect(b, params, thread, indexed); -} - -void -agx_nir_decompress(nir_builder *b, const void *data) -{ - const struct agx_decompress_key *key = data; - - nir_def *params = nir_load_preamble(b, 1, 64, .base = 0); - nir_def *tile = nir_load_workgroup_id(b); - nir_def *local = nir_channel(b, nir_load_local_invocation_id(b), 0); - nir_def *samples = nir_imm_int(b, key->nr_samples); - - libagx_decompress(b, params, tile, local, samples); -} diff --git a/src/asahi/lib/agx_nir_lower_gs.h b/src/asahi/lib/agx_nir_lower_gs.h index 67220801fc3..6489aaace6b 100644 --- a/src/asahi/lib/agx_nir_lower_gs.h +++ b/src/asahi/lib/agx_nir_lower_gs.h @@ -7,7 +7,6 @@ #include #include -#include "libagx/tessellator.h" #include "nir.h" #include "shader_enums.h" @@ -34,36 +33,6 @@ bool agx_nir_lower_gs(struct nir_shader *gs, const struct nir_shader *libagx, struct nir_shader **gs_copy, struct nir_shader **pre_gs, enum mesa_prim *out_mode, unsigned *out_count_words); -void agx_nir_prefix_sum_gs(struct nir_builder *b, const void *data); - -void agx_nir_prefix_sum_tess(struct nir_builder *b, const void *data); - -struct agx_gs_setup_indirect_key { - enum mesa_prim prim; -}; - -void agx_nir_gs_setup_indirect(struct nir_builder *b, const void *key); - -struct agx_unroll_restart_key { - enum mesa_prim prim; -}; - -void agx_nir_unroll_restart(struct nir_builder *b, const void *key); - -struct agx_tessellator_key { - enum tess_primitive_mode prim : 8; - enum libagx_tess_mode mode : 8; - unsigned pad : 16; -}; -static_assert(sizeof(struct agx_tessellator_key) == 4, "padded"); - -struct agx_tess_setup_indirect_key { - bool point_mode; -}; -static_assert(sizeof(struct agx_tess_setup_indirect_key) == 1, "padded"); - -void agx_nir_tessellate(struct nir_builder *b, const void *key); - bool agx_nir_lower_tcs(struct nir_shader *tcs, const struct nir_shader *libagx); bool agx_nir_lower_tes(struct nir_shader *tes, const struct nir_shader *libagx, @@ -72,31 +41,3 @@ bool agx_nir_lower_tes(struct nir_shader *tes, const struct nir_shader *libagx, uint64_t agx_tcs_per_vertex_outputs(const struct nir_shader *nir); unsigned agx_tcs_output_stride(const struct nir_shader *nir); - -void agx_nir_tess_setup_indirect(struct nir_builder *b, const void *data); - -void agx_nir_increment_statistic(struct nir_builder *b, const void *data); - -void agx_nir_increment_cs_invocations(struct nir_builder *b, const void *data); - -struct agx_increment_ia_counters_key { - /* Implies primitive restart */ - uint8_t index_size_B; -}; -static_assert(sizeof(struct agx_increment_ia_counters_key) == 1, "padded"); - -void agx_nir_increment_ia_counters(struct nir_builder *b, const void *data); - -struct agx_predicate_indirect_key { - bool indexed; -}; -static_assert(sizeof(struct agx_predicate_indirect_key) == 1, "padded"); - -void agx_nir_predicate_indirect(struct nir_builder *b, const void *data); - -struct agx_decompress_key { - uint8_t nr_samples; -}; -static_assert(sizeof(struct agx_decompress_key) == 1, "padded"); - -void agx_nir_decompress(struct nir_builder *b, const void *data); diff --git a/src/asahi/lib/agx_scratch.c b/src/asahi/lib/agx_scratch.c index 24d8d482b1a..8b4135d9a78 100644 --- a/src/asahi/lib/agx_scratch.c +++ b/src/asahi/lib/agx_scratch.c @@ -25,21 +25,6 @@ struct spill_size { uint32_t count; }; -struct agx_bo * -agx_build_helper(struct agx_device *dev) -{ - struct agx_bo *bo = agx_bo_create( - dev, sizeof(libagx_g13_helper), 0, - AGX_BO_READONLY | AGX_BO_EXEC | AGX_BO_LOW_VA, "Helper shader"); - assert(bo); - memcpy(bo->map, libagx_g13_helper, sizeof(libagx_g13_helper)); - - if (dev->debug & AGX_DBG_SCRATCH) - fprintf(stderr, "Helper: 0x%" PRIx64 "\n", bo->va->addr); - - return bo; -} - static struct spill_size agx_scratch_get_spill_size(unsigned dwords) { diff --git a/src/asahi/lib/agx_scratch.h b/src/asahi/lib/agx_scratch.h index cbd1e3a0649..eab82f5d579 100644 --- a/src/asahi/lib/agx_scratch.h +++ b/src/asahi/lib/agx_scratch.h @@ -29,8 +29,6 @@ struct agx_scratch { #endif }; -struct agx_bo *agx_build_helper(struct agx_device *dev); - void agx_scratch_init(struct agx_device *dev, struct agx_scratch *scratch); void agx_scratch_fini(struct agx_scratch *scratch); void agx_scratch_debug_pre(struct agx_scratch *scratch); diff --git a/src/asahi/lib/agx_usc.h b/src/asahi/lib/agx_usc.h index eb2da3681fe..42ea85720b4 100644 --- a/src/asahi/lib/agx_usc.h +++ b/src/asahi/lib/agx_usc.h @@ -7,16 +7,8 @@ #include "asahi/genxml/agx_pack.h" #include "agx_compile.h" - -/* Opaque structure representing a USC program being constructed */ -struct agx_usc_builder { - uint8_t *head; - -#ifndef NDEBUG - uint8_t *begin; - size_t size; -#endif -}; +#include "libagx_dgc.h" +#include "shader_enums.h" static inline unsigned agx_usc_size(unsigned num_reg_bindings) @@ -36,66 +28,6 @@ agx_usc_size(unsigned num_reg_bindings) return size; } -static struct agx_usc_builder -agx_usc_builder(void *out, ASSERTED size_t size) -{ - return (struct agx_usc_builder){ - .head = out, - -#ifndef NDEBUG - .begin = out, - .size = size, -#endif - }; -} - -static bool -agx_usc_builder_validate(struct agx_usc_builder *b, size_t size) -{ -#ifndef NDEBUG - assert(((b->head - b->begin) + size) <= b->size); -#endif - - return true; -} - -#define agx_usc_pack(b, struct_name, template) \ - for (bool it = \ - agx_usc_builder_validate((b), AGX_USC_##struct_name##_LENGTH); \ - it; it = false, (b)->head += AGX_USC_##struct_name##_LENGTH) \ - agx_pack((b)->head, USC_##struct_name, template) - -#define agx_usc_push_blob(b, blob, length) \ - for (bool it = agx_usc_builder_validate((b), length); it; \ - it = false, (b)->head += length) \ - memcpy((b)->head, blob, length); - -#define agx_usc_push_packed(b, struct_name, packed) \ - agx_usc_push_blob(b, packed.opaque, AGX_USC_##struct_name##_LENGTH); - -static void -agx_usc_uniform(struct agx_usc_builder *b, unsigned start_halfs, - unsigned size_halfs, uint64_t buffer) -{ - assert((start_halfs + size_halfs) <= (1 << 9) && "uniform file overflow"); - assert(size_halfs <= 64 && "caller's responsibility to split"); - assert(size_halfs > 0 && "no empty uniforms"); - - if (start_halfs & BITFIELD_BIT(8)) { - agx_usc_pack(b, UNIFORM_HIGH, cfg) { - cfg.start_halfs = start_halfs & BITFIELD_MASK(8); - cfg.size_halfs = size_halfs; - cfg.buffer = buffer; - } - } else { - agx_usc_pack(b, UNIFORM, cfg) { - cfg.start_halfs = start_halfs; - cfg.size_halfs = size_halfs; - cfg.buffer = buffer; - } - } -} - static void agx_usc_shared_none(struct agx_usc_builder *b) { @@ -105,28 +37,23 @@ agx_usc_shared_none(struct agx_usc_builder *b) } } -static void -agx_usc_shared_non_fragment(struct agx_usc_builder *b, - struct agx_shader_info *info, - unsigned variable_shared_mem) +static inline void +agx_usc_shared(struct agx_usc_builder *b, uint16_t local_size, + uint16_t imageblock_stride, unsigned variable_shared_mem) { - if (info->stage == PIPE_SHADER_FRAGMENT) { - return; - } else if (info->stage == PIPE_SHADER_COMPUTE && info->imageblock_stride) { - assert(info->local_size == 0 && "we don't handle this interaction"); + if (imageblock_stride) { + assert(local_size == 0 && "we don't handle this interaction"); assert(variable_shared_mem == 0 && "we don't handle this interaction"); agx_usc_pack(b, SHARED, cfg) { cfg.layout = AGX_SHARED_LAYOUT_32X32; cfg.uses_shared_memory = true; cfg.sample_count = 1; - cfg.sample_stride_in_8_bytes = - DIV_ROUND_UP(info->imageblock_stride, 8); + cfg.sample_stride_in_8_bytes = DIV_ROUND_UP(imageblock_stride, 8); cfg.bytes_per_threadgroup = cfg.sample_stride_in_8_bytes * 8 * 32 * 32; } - } else if (info->stage == PIPE_SHADER_COMPUTE || - info->stage == PIPE_SHADER_TESS_CTRL) { - unsigned size = info->local_size + variable_shared_mem; + } else if (local_size || variable_shared_mem) { + unsigned size = local_size + variable_shared_mem; agx_usc_pack(b, SHARED, cfg) { cfg.layout = AGX_SHARED_LAYOUT_VERTEX_COMPUTE; @@ -139,17 +66,26 @@ agx_usc_shared_non_fragment(struct agx_usc_builder *b, } static inline void -agx_usc_immediates(struct agx_usc_builder *b, struct agx_shader_info *info, +agx_usc_immediates(struct agx_usc_builder *b, const struct agx_rodata *ro, uint64_t base_addr) { - uint16_t size_16 = info->rodata.size_16; - - for (unsigned range = 0; range < DIV_ROUND_UP(size_16, 64); ++range) { + for (unsigned range = 0; range < DIV_ROUND_UP(ro->size_16, 64); ++range) { unsigned offset = 64 * range; - assert(offset < size_16); + assert(offset < ro->size_16); - agx_usc_uniform(b, info->rodata.base_uniform + offset, - MIN2(64, size_16 - offset), - base_addr + info->rodata.offset + (offset * 2)); + agx_usc_uniform(b, ro->base_uniform + offset, + MIN2(64, ro->size_16 - offset), + base_addr + ro->offset + (offset * 2)); + } +} + +static void +agx_usc_shared_non_fragment(struct agx_usc_builder *b, + const struct agx_shader_info *info, + unsigned variable_shared_mem) +{ + if (info->stage != PIPE_SHADER_FRAGMENT) { + agx_usc_shared(b, info->local_size, info->imageblock_stride, + variable_shared_mem); } } diff --git a/src/asahi/lib/meson.build b/src/asahi/lib/meson.build index 126970ad450..4fd49603136 100644 --- a/src/asahi/lib/meson.build +++ b/src/asahi/lib/meson.build @@ -51,14 +51,19 @@ libasahi_decode_shared = shared_library( ) libagx_shaders = custom_target( - 'libagx_shaders.h', + 'libagx_shaders', input : libagx_spv, - output : 'libagx_shaders.h', - command : [prog_asahi_clc, libagx_spv, '@OUTPUT@'], + output : ['libagx_shaders.h', 'libagx_shaders.c'], + command : [prog_asahi_clc, libagx_spv, '@OUTPUT0@', '@OUTPUT1@'], env: ['MESA_SHADER_CACHE_DISABLE=true'], depends : [prog_asahi_clc], ) +idep_libagx_shaders_h = declare_dependency( + sources : [libagx_shaders], + include_directories : include_directories('.'), +) + libasahi_lib = static_library( 'asahi_lib', [libasahi_lib_files, libagx_shaders, agx_pack, sha1_h], diff --git a/src/asahi/libagx/compression.cl b/src/asahi/libagx/compression.cl index 9b5286f3319..aebc5effa20 100644 --- a/src/asahi/libagx/compression.cl +++ b/src/asahi/libagx/compression.cl @@ -22,7 +22,8 @@ * offset in bytes. Since the descriptors are all in the u0_u1 push, the former * is hardcoded and the latter is an offsetof. */ -#define HANDLE(field) (uint2)(0, offsetof(struct libagx_decompress_push, field)) +#define HANDLE(field) \ + (uint2)(0, offsetof(struct libagx_decompress_images, field)) /* * The metadata buffer is fully twiddled, so interleave the X/Y coordinate bits. @@ -70,17 +71,23 @@ sample_id(int4 c, uint samples) return c.y & 1; } -void -libagx_decompress(constant struct libagx_decompress_push *push, uint3 coord_tl, - uint local_id, uint samples) +KERNEL(32) +libagx_decompress(constant struct libagx_decompress_images *images, + global uint64_t *metadata, uint64_t tile_uncompressed, + uint32_t metadata_layer_stride_tl, uint16_t metadata_width_tl, + uint16_t metadata_height_tl, + uint log2_samples__3 /* 1x, 2x, 4x */) { + uint3 coord_tl = (uint3)(get_group_id(0), get_group_id(1), get_group_id(2)); + uint local_id = get_local_id(0); + uint samples = 1 << log2_samples__3; + /* Index into the metadata buffer */ - uint index_tl = - index_metadata(coord_tl, push->metadata_width_tl, - push->metadata_height_tl, push->metadata_layer_stride_tl); + uint index_tl = index_metadata(coord_tl, metadata_width_tl, + metadata_height_tl, metadata_layer_stride_tl); /* If the tile is already uncompressed, there's nothing to do. */ - if (push->metadata[index_tl] == push->tile_uncompressed) + if (metadata[index_tl] == tile_uncompressed) return; /* Tiles are 16x16 */ @@ -132,6 +139,6 @@ libagx_decompress(constant struct libagx_decompress_push *push, uint3 coord_tl, /* We've replaced the body buffer. Mark the tile as uncompressed. */ if (local_id == 0) { - push->metadata[index_tl] = push->tile_uncompressed; + metadata[index_tl] = tile_uncompressed; } } diff --git a/src/asahi/libagx/compression.h b/src/asahi/libagx/compression.h index b750b54b4ba..f22228b7df1 100644 --- a/src/asahi/libagx/compression.h +++ b/src/asahi/libagx/compression.h @@ -7,13 +7,8 @@ #pragma once -struct libagx_decompress_push { +struct libagx_decompress_images { struct agx_texture_packed compressed; struct agx_pbe_packed uncompressed; - GLOBAL(uint64_t) metadata; - uint64_t tile_uncompressed; - uint32_t metadata_layer_stride_tl; - uint16_t metadata_width_tl; - uint16_t metadata_height_tl; }; -AGX_STATIC_ASSERT(sizeof(struct libagx_decompress_push) == 72); +AGX_STATIC_ASSERT(sizeof(struct libagx_decompress_images) == 48); diff --git a/src/asahi/libagx/draws.cl b/src/asahi/libagx/draws.cl index 6f0d8e95644..64925c6ef63 100644 --- a/src/asahi/libagx/draws.cl +++ b/src/asahi/libagx/draws.cl @@ -3,20 +3,21 @@ * SPDX-License-Identifier: MIT */ #include "libagx.h" -#include "draws.h" /* * To implement drawIndirectCount generically, we dispatch a kernel to * clone-and-patch the indirect buffer, predicating out draws as appropriate. */ -void -libagx_predicate_indirect(constant struct libagx_predicate_indirect_push *push, - uint draw, bool indexed) +KERNEL(32) +libagx_predicate_indirect(global uint32_t *out, constant uint32_t *in, + constant uint32_t *draw_count, uint32_t stride_el, + uint indexed__2) { - uint words = indexed ? 5 : 4; - global uint *out = &push->out[draw * words]; - constant uint *in = &push->in[draw * push->stride_el]; - bool enabled = draw < *(push->draw_count); + uint draw = get_global_id(0); + uint words = indexed__2 ? 5 : 4; + bool enabled = draw < *draw_count; + out += draw * words; + in += draw * stride_el; /* Copy enabled draws, zero predicated draws. */ for (uint i = 0; i < words; ++i) { diff --git a/src/asahi/libagx/draws.h b/src/asahi/libagx/draws.h deleted file mode 100644 index d22a1912248..00000000000 --- a/src/asahi/libagx/draws.h +++ /dev/null @@ -1,14 +0,0 @@ -/* - * Copyright 2024 Valve Corporation - * SPDX-License-Identifier: MIT - */ -#include "libagx.h" - -#pragma once - -struct libagx_predicate_indirect_push { - GLOBAL(uint32_t) out; - CONSTANT(uint32_t) in; - CONSTANT(uint32_t) draw_count; - uint32_t stride_el; -}; diff --git a/src/asahi/libagx/geometry.cl b/src/asahi/libagx/geometry.cl index efb4459c189..e3a58ab1746 100644 --- a/src/asahi/libagx/geometry.cl +++ b/src/asahi/libagx/geometry.cl @@ -5,6 +5,7 @@ */ #include "geometry.h" +#include "query.h" #include "tessellator.h" /* Compatible with util/u_math.h */ @@ -228,7 +229,7 @@ vertex_id_for_topology(enum mesa_prim mode, bool flatshade_first, uint prim, } } -uint +static uint libagx_load_index_buffer_internal(uintptr_t index_buffer, uint32_t index_buffer_range_el, uint id, uint index_size) @@ -268,6 +269,61 @@ libagx_load_index_buffer(constant struct agx_ia_state *p, uint id, p->index_buffer, p->index_buffer_range_el, id, index_size); } +static void +increment_ia_counters(global uint32_t *ia_vertices, + global uint32_t *vs_invocations, uint count) +{ + if (ia_vertices) { + *ia_vertices += count; + } + + if (vs_invocations) { + *vs_invocations += count; + } +} + +KERNEL(1) +libagx_increment_ia(global uint32_t *ia_vertices, + global uint32_t *vs_invocations, constant uint32_t *draw) +{ + increment_ia_counters(ia_vertices, vs_invocations, draw[0] * draw[1]); +} + +KERNEL(1024) +libagx_increment_ia_restart(global uint32_t *ia_vertices, + global uint32_t *vs_invocations, + constant uint32_t *draw, uint64_t index_buffer, + uint32_t index_buffer_range_el, + uint32_t restart_index, uint32_t index_size_B) +{ + uint tid = get_global_id(0); + unsigned count = draw[0]; + local uint scratch; + + uint start = draw[2]; + uint partial = 0; + + /* Count non-restart indices */ + for (uint i = tid; i < count; i += 1024) { + uint index = libagx_load_index_buffer_internal( + index_buffer, index_buffer_range_el, start + i, index_size_B); + + if (index != restart_index) + partial++; + } + + /* Accumulate the partials across the workgroup */ + scratch = 0; + barrier(CLK_LOCAL_MEM_FENCE); + atomic_add(&scratch, partial); + barrier(CLK_LOCAL_MEM_FENCE); + + /* Elect a single thread from the workgroup to increment the counters */ + if (tid == 0) { + increment_ia_counters(ia_vertices, vs_invocations, scratch * draw[1]); + } +} + /* * Return the ID of the first thread in the workgroup where cond is true, or * 1024 if cond is false across the workgroup. @@ -284,26 +340,15 @@ first_true_thread_in_workgroup(bool cond, local uint *scratch) return (first_group * 32) + off; } -/* - * Allocate memory from the heap (thread-safe). Returns the offset into the - * heap. The allocation will be word-aligned. - */ -static inline uint -libagx_atomic_alloc(global struct agx_geometry_state *heap, uint size_B) -{ - return atomic_fetch_add((volatile atomic_uint *)(&heap->heap_bottom), - align(size_B, 8)); -} - /* * When unrolling the index buffer for a draw, we translate the old indirect * draws to new indirect draws. This routine allocates the new index buffer and * sets up most of the new draw descriptor. */ static global void * -setup_unroll_for_draw(constant struct agx_restart_unroll_params *p, - constant uint *in_draw, uint draw, enum mesa_prim mode, - uint index_size_B) +setup_unroll_for_draw(global struct agx_geometry_state *heap, + constant uint *in_draw, global uint *out, + 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 @@ -313,14 +358,13 @@ setup_unroll_for_draw(constant struct agx_restart_unroll_params *p, uint max_verts = max_prims * mesa_vertices_per_prim(mode); uint alloc_size = max_verts * index_size_B; - /* Allocate unrolled index buffer. Atomic since multiple threads may be - * running to handle multidraw in parallel. + /* Allocate unrolled index buffer. + * + * TODO: For multidraw, should be atomic. But multidraw+unroll isn't + * currently wired up in any driver. */ - global struct agx_geometry_state *heap = p->heap; - uint old_heap_bottom_B = libagx_atomic_alloc(p->heap, alloc_size); - - /* Regardless of the input stride, we use tightly packed output draws */ - global uint *out = &p->out_draws[5 * draw]; + uint old_heap_bottom_B = heap->heap_bottom; + heap->heap_bottom += align(alloc_size, 8); /* Setup most of the descriptor. Count will be determined after unroll. */ out[1] = in_draw[1]; /* instance count */ @@ -332,26 +376,24 @@ setup_unroll_for_draw(constant struct agx_restart_unroll_params *p, return (global uchar *)heap->heap + old_heap_bottom_B; } -kernel void -libagx_unroll_restart(constant struct agx_restart_unroll_params *p, - enum mesa_prim mode, uint draw, uint tid) +KERNEL(1024) +libagx_unroll_restart(global struct agx_geometry_state *heap, + uint64_t index_buffer, constant uint *in_draw, + global uint32_t *out_draw, uint64_t zero_sink, + uint32_t max_draws, uint32_t restart_index, + uint32_t index_buffer_size_el, + uint32_t index_size_log2__3, uint32_t flatshade_first, + uint mode__11) { - /* For an indirect multidraw, we are dispatched maxDraws times and - * terminate trailing invocations. - */ - if (p->count && draw >= *(p->count)) - return; - - constant uint *in_draw = - (constant uint *)(p->draws + (draw * p->draw_stride)); - + uint32_t index_size_B = 1 << index_size_log2__3; + enum mesa_prim mode = libagx_uncompact_prim(mode__11); + uint tid = get_local_id(0); uint count = in_draw[0]; local uintptr_t out_ptr; if (tid == 0) { - out_ptr = (uintptr_t)setup_unroll_for_draw(p, in_draw, draw, mode, - p->index_size_B); - + out_ptr = (uintptr_t)setup_unroll_for_draw(heap, in_draw, out_draw, mode, + index_size_B); } barrier(CLK_LOCAL_MEM_FENCE); @@ -360,15 +402,11 @@ libagx_unroll_restart(constant struct agx_restart_unroll_params *p, global uint8_t *out_8 = (global uint8_t *)out_ptr; uintptr_t in_ptr = (uintptr_t)(libagx_index_buffer( - p->index_buffer, p->index_buffer_size_el, in_draw[2], p->index_size_B, - p->zero_sink)); + index_buffer, index_buffer_size_el, in_draw[2], index_size_B, zero_sink)); local uint scratch[32]; uint out_prims = 0; - uint32_t restart_idx = p->restart_index; - bool flatshade_first = p->flatshade_first; - uint needle = 0; uint per_prim = mesa_vertices_per_prim(mode); while (needle < count) { @@ -377,8 +415,8 @@ libagx_unroll_restart(constant struct agx_restart_unroll_params *p, for (;;) { uint idx = next_restart + tid; bool restart = idx >= count || libagx_load_index_buffer_internal( - in_ptr, p->index_buffer_size_el, idx, - p->index_size_B) == restart_idx; + in_ptr, index_buffer_size_el, idx, + index_size_B) == restart_index; uint next_offs = first_true_thread_in_workgroup(restart, scratch); @@ -399,11 +437,11 @@ libagx_unroll_restart(constant struct agx_restart_unroll_params *p, uint x = ((out_prims_base + i) * per_prim) + vtx; uint y = libagx_load_index_buffer_internal( - in_ptr, p->index_buffer_size_el, offset, p->index_size_B); + in_ptr, index_buffer_size_el, offset, index_size_B); - if (p->index_size_B == 4) + if (index_size_B == 4) out_32[x] = y; - else if (p->index_size_B == 2) + else if (index_size_B == 2) out_16[x] = y; else out_8[x] = y; @@ -415,7 +453,7 @@ libagx_unroll_restart(constant struct agx_restart_unroll_params *p, } if (tid == 0) - p->out_draws[(5 * draw) + 0] = out_prims * per_prim; + out_draw[0] = out_prims * per_prim; } uint @@ -497,21 +535,25 @@ libagx_build_gs_draw(global struct agx_geometry_params *p, uint vertices, } } -void -libagx_gs_setup_indirect(global struct agx_gs_setup_indirect_params *gsi, - enum mesa_prim mode, uint local_id) +KERNEL(1) +libagx_gs_setup_indirect( + uint64_t index_buffer, constant uint *draw, + global uintptr_t *vertex_buffer /* output */, + global struct agx_ia_state *ia /* output */, + global struct agx_geometry_params *p /* output */, uint64_t zero_sink, + uint64_t vs_outputs /* Vertex (TES) output mask */, + uint32_t index_size_B /* 0 if no index bffer */, + uint32_t index_buffer_range_el, + uint32_t prim /* Input primitive type, enum mesa_prim */) { - 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 = gsi->draw[0]; - uint instance_count = gsi->draw[1]; + uint vertex_count = draw[0]; + uint instance_count = draw[1]; ia->verts_per_instance = vertex_count; /* Calculate number of primitives input into the GS */ - uint prim_per_instance = u_decomposed_prims_for_vertices(mode, vertex_count); + uint prim_per_instance = u_decomposed_prims_for_vertices(prim, vertex_count); p->input_primitives = prim_per_instance * instance_count; /* Invoke VS as (vertices, instances); GS as (primitives, instances) */ @@ -528,30 +570,29 @@ libagx_gs_setup_indirect(global struct agx_gs_setup_indirect_params *gsi, * indirect draw, the hardware would do this for us, but for software input * assembly we need to do it ourselves. */ - if (gsi->index_size_B) { - ia->index_buffer = - libagx_index_buffer(gsi->index_buffer, gsi->index_buffer_range_el, - gsi->draw[2], gsi->index_size_B, gsi->zero_sink); + if (index_size_B) { + ia->index_buffer = libagx_index_buffer( + index_buffer, index_buffer_range_el, draw[2], index_size_B, zero_sink); ia->index_buffer_range_el = - libagx_index_buffer_range_el(gsi->index_buffer_range_el, gsi->draw[2]); + libagx_index_buffer_range_el(index_buffer_range_el, draw[2]); } /* 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, gsi->vs_outputs); + libagx_tcs_in_size(vertex_count * instance_count, vs_outputs); p->count_buffer = (global uint *)(state->heap + state->heap_bottom); state->heap_bottom += align(p->input_primitives * p->count_buffer_stride, 16); p->input_buffer = (uintptr_t)(state->heap + state->heap_bottom); - *(gsi->vertex_buffer) = p->input_buffer; + *vertex_buffer = p->input_buffer; state->heap_bottom += align(vertex_buffer_size, 4); - p->input_mask = gsi->vs_outputs; + p->input_mask = vs_outputs; if (state->heap_bottom > state->heap_size) { global uint *foo = (global uint *)(uintptr_t)0x1deadbeef; @@ -598,8 +639,8 @@ libagx_work_group_scan_inclusive_add(uint x, local uint *scratch) return (uint2)(prefix, reduction); } -kernel void -libagx_prefix_sum(global uint *buffer, uint len, uint words, uint word) +KERNEL(1024) +_libagx_prefix_sum(global uint *buffer, uint len, uint words, uint word) { local uint scratch[32]; uint tid = get_local_id(0); @@ -630,10 +671,17 @@ libagx_prefix_sum(global uint *buffer, uint len, uint words, uint word) } } -kernel void +KERNEL(1024) +libagx_prefix_sum_geom(constant struct agx_geometry_params *p) +{ + _libagx_prefix_sum(p->count_buffer, p->input_primitives, + p->count_buffer_stride / 4, get_group_id(0)); +} + +KERNEL(1024) libagx_prefix_sum_tess(global struct libagx_tess_args *p) { - libagx_prefix_sum(p->counts, p->nr_patches, 1 /* words */, 0 /* word */); + _libagx_prefix_sum(p->counts, p->nr_patches, 1 /* words */, 0 /* word */); /* After prefix summing, we know the total # of indices, so allocate the * index buffer now. Elect a thread for the allocation. diff --git a/src/asahi/libagx/geometry.h b/src/asahi/libagx/geometry.h index 553a65924a0..fdb0538f8c8 100644 --- a/src/asahi/libagx/geometry.h +++ b/src/asahi/libagx/geometry.h @@ -32,78 +32,6 @@ struct agx_geometry_state { } PACKED; AGX_STATIC_ASSERT(sizeof(struct agx_geometry_state) == 4 * 4); -struct agx_restart_unroll_params { - /* Heap to allocate from across draws */ - GLOBAL(struct agx_geometry_state) heap; - - /* Input: index buffer if present. */ - uint64_t index_buffer; - - /* Input: draw count */ - CONST(uint) count; - - /* Input: indirect draw descriptor. Raw pointer since it's strided. */ - uint64_t draws; - - /* Output draw descriptors */ - GLOBAL(uint) out_draws; - - /* Pointer to zero */ - uint64_t zero_sink; - - /* 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 elements */ - uint32_t index_buffer_size_el; - - /* Size of an index in bytes */ - uint32_t index_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) == 18 * 4); - -struct agx_gs_setup_indirect_params { - /* Index buffer if present. */ - uint64_t 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; - - /* Pointer to zero */ - uint64_t zero_sink; - - /* 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; - - /* Size of the index buffer */ - uint32_t index_buffer_range_el; -} PACKED; -AGX_STATIC_ASSERT(sizeof(struct agx_gs_setup_indirect_params) == 16 * 4); - struct agx_ia_state { /* Index buffer if present. */ uint64_t index_buffer; @@ -306,4 +234,26 @@ libagx_tcs_out_stride(uint nr_patch_out, uint out_patch_size, /* In a tess eval shader, stride for hw vertex ID */ #define LIBAGX_TES_PATCH_ID_STRIDE 8192 +static uint +libagx_compact_prim(enum mesa_prim prim) +{ + AGX_STATIC_ASSERT(MESA_PRIM_QUAD_STRIP == MESA_PRIM_QUADS + 1); + AGX_STATIC_ASSERT(MESA_PRIM_POLYGON == MESA_PRIM_QUADS + 2); + +#ifndef __OPENCL_VERSION__ + assert(prim != MESA_PRIM_QUADS); + assert(prim != MESA_PRIM_QUAD_STRIP); + assert(prim != MESA_PRIM_POLYGON); + assert(prim != MESA_PRIM_PATCHES); +#endif + + return (prim >= MESA_PRIM_QUADS) ? (prim - 3) : prim; +} + +static enum mesa_prim +libagx_uncompact_prim(uint packed) +{ + return (packed >= MESA_PRIM_QUADS) ? (packed + 3) : packed; +} + #endif diff --git a/src/asahi/libagx/helper.cl b/src/asahi/libagx/helper.cl index b03839af132..ccb4fb4f5df 100644 --- a/src/asahi/libagx/helper.cl +++ b/src/asahi/libagx/helper.cl @@ -17,7 +17,7 @@ enum helper_op { OP_END = 15, }; -void +KERNEL(1) libagx_helper(void) { uint64_t arg = diff --git a/src/asahi/libagx/libagx.h b/src/asahi/libagx/libagx.h index 156b89ba803..2f39150fe2d 100644 --- a/src/asahi/libagx/libagx.h +++ b/src/asahi/libagx/libagx.h @@ -13,11 +13,19 @@ #define GLOBAL(type_) uint64_t #define CONSTANT(type_) uint64_t #define AGX_STATIC_ASSERT(_COND) static_assert(_COND, #_COND) +#define global_ +#define constant_ #else #pragma OPENCL EXTENSION cl_khr_fp16 : enable #define PACKED __attribute__((packed, aligned(4))) #define GLOBAL(type_) global type_ * #define CONSTANT(type_) constant type_ * +#define global_ global +#define constant_ constant +#define ASSERTED +#define assert(x) +#define BITFIELD_BIT(b) (1u << b) +#define BITFIELD_MASK(m) (((m) == 32) ? 0xffffffff : ((1u << (m)) - 1)) typedef ulong uint64_t; typedef uint uint32_t; @@ -49,10 +57,6 @@ void nir_bindless_image_store(uint2 handle, int4 coord, uint sample, uint image_array, uint format, uint access, uint src_type); -uint libagx_load_index_buffer_internal(uintptr_t index_buffer, - uint32_t index_buffer_range_el, uint id, - uint index_size); - /* I have no idea why CL doesn't have this */ uint nir_ballot(bool cond); @@ -62,6 +66,8 @@ uint nir_ballot(bool cond); #define AGX_STATIC_ASSERT(_COND) \ typedef char AGX_PASTE(static_assertion, __LINE__)[(_COND) ? 1 : -1] +#define KERNEL(x) __attribute__((reqd_work_group_size(x, 1, 1))) kernel void + static inline uint align(uint x, uint y) { diff --git a/src/asahi/libagx/libagx_dgc.h b/src/asahi/libagx/libagx_dgc.h new file mode 100644 index 00000000000..9595cb4268b --- /dev/null +++ b/src/asahi/libagx/libagx_dgc.h @@ -0,0 +1,362 @@ +/* + * Copyright 2024 Valve Corporation + * SPDX-License-Identifier: MIT + */ + +#pragma once + +#include "agx_pack.h" +#include "libagx.h" + +#ifdef __OPENCL_VERSION__ +#define NDEBUG +#define memcpy __builtin_memcpy +#define STATIC_ASSERT AGX_STATIC_ASSERT +#endif + +#define agx_push(ptr, T, cfg) \ + for (unsigned _loop = 0; _loop < 1; \ + ++_loop, ptr = (global_ void *)(((uintptr_t)ptr) + AGX_##T##_LENGTH)) \ + agx_pack(ptr, T, cfg) + +#define agx_push_packed(ptr, src, T) \ + STATIC_ASSERT(sizeof(src) == AGX_##T##_LENGTH); \ + memcpy(ptr, &src, sizeof(src)); \ + ptr = (global_ void *)(((uintptr_t)ptr) + sizeof(src)); + +struct agx_workgroup { + uint32_t x, y, z; +}; + +static inline struct agx_workgroup +agx_workgroup(uint32_t x, uint32_t y, uint32_t z) +{ + return (struct agx_workgroup){.x = x, .y = y, .z = z}; +} + +static inline unsigned +agx_workgroup_threads(struct agx_workgroup wg) +{ + return wg.x * wg.y * wg.z; +} + +struct agx_grid { + enum agx_cdm_mode mode; + union { + uint32_t count[3]; + uint64_t ptr; + }; +}; + +static struct agx_grid +agx_3d(uint32_t x, uint32_t y, uint32_t z) +{ + return (struct agx_grid){.mode = AGX_CDM_MODE_DIRECT, .count = {x, y, z}}; +} + +static struct agx_grid +agx_1d(uint32_t x) +{ + return agx_3d(x, 1, 1); +} + +static struct agx_grid +agx_grid_indirect(uint64_t ptr) +{ + return (struct agx_grid){.mode = AGX_CDM_MODE_INDIRECT_GLOBAL, .ptr = ptr}; +} + +static struct agx_grid +agx_grid_indirect_local(uint64_t ptr) +{ + return (struct agx_grid){.mode = AGX_CDM_MODE_INDIRECT_LOCAL, .ptr = ptr}; +} + +static inline bool +agx_is_indirect(struct agx_grid grid) +{ + return grid.mode != AGX_CDM_MODE_DIRECT; +} + +enum agx_chip { + AGX_CHIP_G13G, + AGX_CHIP_G13X, + AGX_CHIP_G14G, + AGX_CHIP_G14X, +}; + +static inline global_ uint32_t * +agx_cdm_launch(global_ uint32_t *out, enum agx_chip chip, struct agx_grid grid, + struct agx_workgroup wg, + struct agx_cdm_launch_word_0_packed launch, uint32_t usc) +{ +#ifndef __OPENCL_VERSION__ + struct agx_cdm_launch_word_0_packed mode; + agx_pack(&mode, CDM_LAUNCH_WORD_0, cfg) { + cfg.mode = grid.mode; + } + + agx_merge(launch, mode, CDM_LAUNCH_WORD_0); +#endif + + agx_push_packed(out, launch, CDM_LAUNCH_WORD_0); + + agx_push(out, CDM_LAUNCH_WORD_1, cfg) { + cfg.pipeline = usc; + } + + if (chip == AGX_CHIP_G14X) { + agx_push(out, CDM_UNK_G14X, cfg) + ; + } + + if (agx_is_indirect(grid)) { + agx_push(out, CDM_INDIRECT, cfg) { + cfg.address_hi = grid.ptr >> 32; + cfg.address_lo = grid.ptr; + } + } else { + agx_push(out, CDM_GLOBAL_SIZE, cfg) { + cfg.x = grid.count[0]; + cfg.y = grid.count[1]; + cfg.z = grid.count[2]; + } + } + + if (grid.mode != AGX_CDM_MODE_INDIRECT_LOCAL) { + agx_push(out, CDM_LOCAL_SIZE, cfg) { + cfg.x = wg.x; + cfg.y = wg.y; + cfg.z = wg.z; + } + } + + return out; +} + +static inline global_ uint32_t * +agx_cdm_barrier(global_ uint32_t *out, enum agx_chip chip) +{ + agx_push(out, CDM_BARRIER, cfg) { + cfg.unk_5 = true; + cfg.unk_6 = true; + cfg.unk_8 = true; + // cfg.unk_11 = true; + // cfg.unk_20 = true; + // cfg.unk_24 = true; if clustered? + if (chip == AGX_CHIP_G13X) { + cfg.unk_4 = true; + // cfg.unk_26 = true; + } + + /* With multiple launches in the same CDM stream, we can get cache + * coherency (? or sync?) issues. We hit this with blits, which need - in + * between dispatches - need the PBE cache to be flushed and the texture + * cache to be invalidated. Until we know what bits mean what exactly, + * let's just set these after every launch to be safe. We can revisit in + * the future when we figure out what the bits mean. + */ + cfg.unk_0 = true; + cfg.unk_1 = true; + cfg.unk_2 = true; + cfg.usc_cache_inval = true; + cfg.unk_4 = true; + cfg.unk_5 = true; + cfg.unk_6 = true; + cfg.unk_7 = true; + cfg.unk_8 = true; + cfg.unk_9 = true; + cfg.unk_10 = true; + cfg.unk_11 = true; + cfg.unk_12 = true; + cfg.unk_13 = true; + cfg.unk_14 = true; + cfg.unk_15 = true; + cfg.unk_16 = true; + cfg.unk_17 = true; + cfg.unk_18 = true; + cfg.unk_19 = true; + } + + return out; +} + +static inline global_ uint32_t * +agx_cdm_return(global_ uint32_t *out) +{ + agx_push(out, CDM_STREAM_RETURN, cfg) + ; + + return out; +} + +static inline global_ uint32_t * +agx_cdm_terminate(global_ uint32_t *out) +{ + agx_push(out, CDM_STREAM_TERMINATE, _) + ; + + return out; +} + +static inline global_ uint32_t * +agx_vdm_terminate(global_ uint32_t *out) +{ + agx_push(out, VDM_STREAM_TERMINATE, _) + ; + + return out; +} + +static inline global_ uint32_t * +agx_cdm_jump(global_ uint32_t *out, uint64_t target) +{ + agx_push(out, CDM_STREAM_LINK, cfg) { + cfg.target_lo = target & BITFIELD_MASK(32); + cfg.target_hi = target >> 32; + } + + return out; +} + +static inline global_ uint32_t * +agx_vdm_jump(global_ uint32_t *out, uint64_t target) +{ + agx_push(out, VDM_STREAM_LINK, cfg) { + cfg.target_lo = target & BITFIELD_MASK(32); + cfg.target_hi = target >> 32; + } + + return out; +} + +static inline global_ uint32_t * +agx_cdm_call(global_ uint32_t *out, uint64_t target) +{ + agx_push(out, CDM_STREAM_LINK, cfg) { + cfg.target_lo = target & BITFIELD_MASK(32); + cfg.target_hi = target >> 32; + cfg.with_return = true; + } + + return out; +} + +static inline global_ uint32_t * +agx_vdm_call(global_ uint32_t *out, uint64_t target) +{ + agx_push(out, VDM_STREAM_LINK, cfg) { + cfg.target_lo = target & BITFIELD_MASK(32); + cfg.target_hi = target >> 32; + cfg.with_return = true; + } + + return out; +} + +#define AGX_MAX_LINKED_USC_SIZE \ + (AGX_USC_PRESHADER_LENGTH + AGX_USC_FRAGMENT_PROPERTIES_LENGTH + \ + AGX_USC_REGISTERS_LENGTH + AGX_USC_SHADER_LENGTH + AGX_USC_SHARED_LENGTH + \ + AGX_USC_SAMPLER_LENGTH + (AGX_USC_UNIFORM_LENGTH * 9)) + +/* + * This data structure contains everything needed to dispatch a compute shader + * (and hopefully eventually graphics?). + * + * It is purely flat, no CPU pointers. That makes it suitable for sharing + * between CPU and GPU. The intention is that it is packed on the CPU side and + * then consumed on either host or device for dispatching work. + */ +struct agx_shader { + struct agx_cdm_launch_word_0_packed launch; + struct agx_workgroup workgroup; + + struct { + uint32_t size; + uint8_t data[AGX_MAX_LINKED_USC_SIZE]; + } usc; +}; + +/* Opaque structure representing a USC program being constructed */ +struct agx_usc_builder { + global_ uint8_t *head; + +#ifndef NDEBUG + uint8_t *begin; + size_t size; +#endif +} PACKED; + +#ifdef __OPENCL_VERSION__ +AGX_STATIC_ASSERT(sizeof(struct agx_usc_builder) == 8); +#endif + +static struct agx_usc_builder +agx_usc_builder(global_ void *out, ASSERTED size_t size) +{ + return (struct agx_usc_builder){ + .head = out, + +#ifndef NDEBUG + .begin = out, + .size = size, +#endif + }; +} + +static bool +agx_usc_builder_validate(struct agx_usc_builder *b, size_t size) +{ +#ifndef NDEBUG + assert(((b->head - b->begin) + size) <= b->size); +#endif + + return true; +} + +#define agx_usc_pack(b, struct_name, template) \ + for (bool it = \ + agx_usc_builder_validate((b), AGX_USC_##struct_name##_LENGTH); \ + it; it = false, (b)->head += AGX_USC_##struct_name##_LENGTH) \ + agx_pack((b)->head, USC_##struct_name, template) + +#define agx_usc_push_blob(b, blob, length) \ + for (bool it = agx_usc_builder_validate((b), length); it; \ + it = false, (b)->head += length) \ + memcpy((b)->head, blob, length); + +#define agx_usc_push_packed(b, struct_name, packed) \ + agx_usc_push_blob(b, packed.opaque, AGX_USC_##struct_name##_LENGTH); + +static void +agx_usc_uniform(struct agx_usc_builder *b, unsigned start_halfs, + unsigned size_halfs, uint64_t buffer) +{ + assert((start_halfs + size_halfs) <= (1 << 9) && "uniform file overflow"); + assert(size_halfs <= 64 && "caller's responsibility to split"); + assert(size_halfs > 0 && "no empty uniforms"); + + if (start_halfs & BITFIELD_BIT(8)) { + agx_usc_pack(b, UNIFORM_HIGH, cfg) { + cfg.start_halfs = start_halfs & BITFIELD_MASK(8); + cfg.size_halfs = size_halfs; + cfg.buffer = buffer; + } + } else { + agx_usc_pack(b, UNIFORM, cfg) { + cfg.start_halfs = start_halfs; + cfg.size_halfs = size_halfs; + cfg.buffer = buffer; + } + } +} + +static inline void +agx_usc_words_precomp(global_ uint32_t *out, constant_ struct agx_shader *s, + uint64_t data, unsigned data_size) +{ + /* Map the data directly as uniforms starting at u0 */ + struct agx_usc_builder b = agx_usc_builder(out, sizeof(s->usc.data)); + agx_usc_uniform(&b, 0, DIV_ROUND_UP(data_size, 2), data); + agx_usc_push_blob(&b, s->usc.data, s->usc.size); +} diff --git a/src/asahi/libagx/meson.build b/src/asahi/libagx/meson.build index d32c60b4e8b..4fd5e60be6a 100644 --- a/src/asahi/libagx/meson.build +++ b/src/asahi/libagx/meson.build @@ -5,7 +5,6 @@ libagx_shader_files = files( 'libagx.h', 'compression.cl', 'draws.cl', - 'draws.h', 'geometry.cl', 'geometry.h', 'query.cl', @@ -15,6 +14,7 @@ libagx_shader_files = files( 'tessellator.h', 'texture.cl', 'helper.cl', + 'libagx_dgc.h', ) prepended_input_args = [] diff --git a/src/asahi/libagx/query.cl b/src/asahi/libagx/query.cl index 3e31182be4e..065d3789326 100644 --- a/src/asahi/libagx/query.cl +++ b/src/asahi/libagx/query.cl @@ -19,31 +19,69 @@ write_query_result(uintptr_t dst_addr, int32_t idx, bool is_64, uint64_t result) } } -void -libagx_copy_query(constant struct libagx_copy_query_push *push, unsigned i) +/* TODO: Optimize workgroup size */ +KERNEL(1) +libagx_copy_query(global uint32_t *availability, global uint64_t *results, + global uint16_t *oq_index, uint64_t dst_addr, + uint64_t dst_stride, uint32_t first_query, uint16_t partial, + uint16_t _64, uint16_t with_availability, + uint16_t reports_per_query) { - uint64_t dst = push->dst_addr + (((uint64_t)i) * push->dst_stride); - uint32_t query = push->first_query + i; - bool available = push->availability[query]; + uint i = get_global_id(0); + uint64_t dst = dst_addr + (((uint64_t)i) * dst_stride); + uint32_t query = first_query + i; + bool available = availability[query]; - if (available || push->partial) { + if (available || partial) { /* For occlusion queries, results[] points to the device global heap. We * need to remap indices according to the query pool's allocation. */ - uint result_index = push->oq_index ? push->oq_index[query] : query; - uint idx = result_index * push->reports_per_query; + uint result_index = oq_index ? oq_index[query] : query; + uint idx = result_index * reports_per_query; - for (unsigned i = 0; i < push->reports_per_query; ++i) { - write_query_result(dst, i, push->_64, push->results[idx + i]); + for (unsigned i = 0; i < reports_per_query; ++i) { + write_query_result(dst, i, _64, results[idx + i]); } } - if (push->with_availability) { - write_query_result(dst, push->reports_per_query, push->_64, available); + if (with_availability) { + write_query_result(dst, reports_per_query, _64, available); } } -void +/* TODO: Share with Gallium... */ +enum pipe_query_value_type { + PIPE_QUERY_TYPE_I32, + PIPE_QUERY_TYPE_U32, + PIPE_QUERY_TYPE_I64, + PIPE_QUERY_TYPE_U64, +}; + +KERNEL(1) +libagx_copy_query_gl(global uint64_t *query, global uint64_t *dest, + ushort value_type, ushort bool_size) +{ + uint64_t value = *query; + + if (bool_size == 4) { + value = (uint32_t)value; + } + + if (bool_size) { + value = value != 0; + } + + if (value_type <= PIPE_QUERY_TYPE_U32) { + global uint32_t *dest32 = (global uint32_t *)dest; + bool u32 = (value_type == PIPE_QUERY_TYPE_U32); + + *dest32 = u32 ? convert_uint_sat(value) : convert_int_sat((int64_t)value); + } else { + *dest = value; + } +} + +KERNEL(4) libagx_copy_xfb_counters(constant struct libagx_xfb_counter_copy *push) { unsigned i = get_local_id(0); @@ -51,64 +89,29 @@ libagx_copy_xfb_counters(constant struct libagx_xfb_counter_copy *push) *(push->dest[i]) = push->src[i] ? *(push->src[i]) : 0; } -void -libagx_increment_statistic(constant struct libagx_increment_params *p) +KERNEL(1) +libagx_increment_statistic(global uint32_t *statistic, uint32_t delta) { - *(p->statistic) += p->delta; + *statistic += delta; } -void -libagx_increment_cs_invocations(constant struct libagx_cs_invocation_params *p) +KERNEL(1) +libagx_increment_cs_invocations(global uint *grid, global uint32_t *statistic, + uint32_t local_size_threads) { - *(p->statistic) += libagx_cs_invocations(p->local_size_threads, p->grid[0], - p->grid[1], p->grid[2]); + *statistic += + libagx_cs_invocations(local_size_threads, grid[0], grid[1], grid[2]); } -kernel void -libagx_increment_ia_counters(constant struct libagx_increment_ia_counters *p, - uint index_size_B, uint tid) -{ - unsigned count = p->draw[0]; - local uint scratch; - - if (index_size_B /* implies primitive restart */) { - uint start = p->draw[2]; - uint partial = 0; - - /* Count non-restart indices */ - for (uint i = tid; i < count; i += 1024) { - uint index = libagx_load_index_buffer_internal( - p->index_buffer, p->index_buffer_range_el, start + i, index_size_B); - - if (index != p->restart_index) - partial++; - } - - /* Accumulate the partials across the workgroup */ - scratch = 0; - barrier(CLK_LOCAL_MEM_FENCE); - atomic_add(&scratch, partial); - barrier(CLK_LOCAL_MEM_FENCE); - count = scratch; - - /* Elect a single thread from the workgroup to increment the counters */ - if (tid != 0) - return; - } - - count *= p->draw[1]; - - if (p->ia_vertices) { - *(p->ia_vertices) += count; - } - - if (p->vs_invocations) { - *(p->vs_invocations) += count; - } -} - -void -libagx_write_u32s(constant struct libagx_imm_write *p, uint id) +KERNEL(32) +libagx_write_u32s(constant struct libagx_imm_write *p) { + uint id = get_global_id(0); *(p[id].address) = p[id].value; } + +KERNEL(1) +libagx_write_u32(global uint32_t *address, uint32_t value) +{ + *address = value; +} diff --git a/src/asahi/libagx/query.h b/src/asahi/libagx/query.h index c136d9e092c..3d25999db69 100644 --- a/src/asahi/libagx/query.h +++ b/src/asahi/libagx/query.h @@ -8,45 +8,11 @@ #pragma once -struct libagx_copy_query_push { - GLOBAL(uint32_t) availability; - GLOBAL(uint64_t) results; - GLOBAL(uint16_t) oq_index; - uint64_t dst_addr; - uint64_t dst_stride; - uint32_t first_query; - - /* Flags. Could specialize the shader? */ - uint16_t partial; - uint16_t _64; - uint16_t with_availability; - uint16_t reports_per_query; -}; - struct libagx_xfb_counter_copy { GLOBAL(uint32_t) dest[4]; GLOBAL(uint32_t) src[4]; }; -struct libagx_increment_params { - /* Pointer to the invocation statistic */ - GLOBAL(uint32_t) statistic; - - /* Value to increment by */ - uint32_t delta; -}; - -struct libagx_cs_invocation_params { - /* Pointer to the indirect dispatch grid */ - GLOBAL(uint32_t) grid; - - /* Pointer to the compute shader invocation statistic */ - GLOBAL(uint32_t) statistic; - - /* Local workgroup size in threads */ - uint32_t local_size_threads; -}; - static inline uint32_t libagx_cs_invocations(uint32_t local_size_threads, uint32_t x, uint32_t y, uint32_t z) @@ -54,20 +20,6 @@ libagx_cs_invocations(uint32_t local_size_threads, uint32_t x, uint32_t y, return local_size_threads * x * y * z; } -struct libagx_increment_ia_counters { - /* Statistics */ - GLOBAL(uint32_t) ia_vertices; - GLOBAL(uint32_t) vs_invocations; - - /* Input draw */ - CONSTANT(uint32_t) draw; - - /* Index buffer */ - uint64_t index_buffer; - uint32_t index_buffer_range_el; - uint32_t restart_index; -}; - struct libagx_imm_write { GLOBAL(uint32_t) address; uint32_t value; diff --git a/src/asahi/libagx/tessellation.cl b/src/asahi/libagx/tessellation.cl index bc9910d0fae..34c919e8e14 100644 --- a/src/asahi/libagx/tessellation.cl +++ b/src/asahi/libagx/tessellation.cl @@ -143,15 +143,24 @@ libagx_tess_level_inner_default(constant struct libagx_tess_args *p) p->tess_level_inner_default[1]); } -void -libagx_tess_setup_indirect(global struct libagx_tess_args *p, bool point_mode) +KERNEL(1) +libagx_tess_setup_indirect( + global struct libagx_tess_args *p, + global uint32_t *grids /* output: VS then TCS then tess */, + global struct agx_ia_state *ia /* output */, global uint32_t *indirect, + global uint64_t *vertex_output_buffer_ptr, uint64_t in_index_buffer, + uint32_t in_index_buffer_range_el, uint32_t in_index_size_B, + uint64_t vertex_outputs /* bitfield */, + + /* Tess control invocation counter if active, else zero */ + global uint32_t *tcs_statistic) { - uint count = p->indirect[0], instance_count = p->indirect[1]; + uint count = indirect[0], instance_count = indirect[1]; unsigned in_patches = count / p->input_patch_size; /* TCS invocation counter increments once per-patch */ - if (p->tcs_statistic) { - *(p->tcs_statistic) += in_patches; + if (tcs_statistic) { + *tcs_statistic += in_patches; } size_t draw_stride = 5 * sizeof(uint32_t); @@ -168,7 +177,7 @@ libagx_tess_setup_indirect(global struct libagx_tess_args *p, bool point_mode) alloc += unrolled_patches * sizeof(uint32_t); uint vb_offs = alloc; - uint vb_size = libagx_tcs_in_size(count * instance_count, p->vertex_outputs); + uint vb_size = libagx_tcs_in_size(count * instance_count, vertex_outputs); alloc += vb_size; /* Allocate all patch calculations in one go */ @@ -180,10 +189,9 @@ libagx_tess_setup_indirect(global struct libagx_tess_args *p, bool point_mode) p->coord_allocs = (global uint *)(blob + patch_coord_offs); p->nr_patches = unrolled_patches; - *(p->vertex_output_buffer_ptr) = (uintptr_t)(blob + vb_offs); + *vertex_output_buffer_ptr = (uintptr_t)(blob + vb_offs); p->counts = (global uint32_t *)(blob + count_offs); - global struct agx_ia_state *ia = p->ia; ia->verts_per_instance = count; /* If indexing is enabled, the third word is the offset into the index buffer @@ -193,42 +201,42 @@ libagx_tess_setup_indirect(global struct libagx_tess_args *p, bool point_mode) * * XXX: Deduplicate? */ - if (p->in_index_size_B) { + if (in_index_size_B) { ia->index_buffer = - libagx_index_buffer(p->in_index_buffer, p->in_index_buffer_range_el, - p->indirect[2], p->in_index_size_B, 0); + libagx_index_buffer(in_index_buffer, in_index_buffer_range_el, + indirect[2], in_index_size_B, 0); - ia->index_buffer_range_el = libagx_index_buffer_range_el( - p->in_index_buffer_range_el, p->indirect[2]); + ia->index_buffer_range_el = + libagx_index_buffer_range_el(in_index_buffer_range_el, indirect[2]); } /* VS grid size */ - p->grids[0] = count; - p->grids[1] = instance_count; - p->grids[2] = 1; + grids[0] = count; + grids[1] = instance_count; + grids[2] = 1; /* VS workgroup size */ - p->grids[3] = 64; - p->grids[4] = 1; - p->grids[5] = 1; + grids[3] = 64; + grids[4] = 1; + grids[5] = 1; /* TCS grid size */ - p->grids[6] = in_patches * p->output_patch_size; - p->grids[7] = instance_count; - p->grids[8] = 1; + grids[6] = in_patches * p->output_patch_size; + grids[7] = instance_count; + grids[8] = 1; /* TCS workgroup size */ - p->grids[9] = p->output_patch_size; - p->grids[10] = 1; - p->grids[11] = 1; + grids[9] = p->output_patch_size; + grids[10] = 1; + grids[11] = 1; /* Tess grid size */ - p->grids[12] = unrolled_patches; - p->grids[13] = 1; - p->grids[14] = 1; + grids[12] = unrolled_patches; + grids[13] = 1; + grids[14] = 1; /* Tess workgroup size */ - p->grids[15] = 64; - p->grids[16] = 1; - p->grids[17] = 1; + grids[15] = 64; + grids[16] = 1; + grids[17] = 1; } diff --git a/src/asahi/libagx/tessellator.cl b/src/asahi/libagx/tessellator.cl index f34b3a4fb28..b7e1ef403a2 100644 --- a/src/asahi/libagx/tessellator.cl +++ b/src/asahi/libagx/tessellator.cl @@ -723,10 +723,12 @@ StitchTransition(private struct CHWTessellator *ctx, int baseIndexOffset, } } -void +KERNEL(64) libagx_tess_isoline(constant struct libagx_tess_args *p, - enum libagx_tess_mode mode, uint patch) + enum libagx_tess_mode mode__2) { + enum libagx_tess_mode mode = mode__2; + uint patch = get_global_id(0); enum libagx_tess_partitioning partitioning = p->partitioning; bool lineDensityOdd; @@ -822,10 +824,12 @@ libagx_tess_isoline(constant struct libagx_tess_args *p, } } -void -libagx_tess_tri(constant struct libagx_tess_args *p, enum libagx_tess_mode mode, - uint patch) +KERNEL(64) +libagx_tess_tri(constant struct libagx_tess_args *p, + enum libagx_tess_mode mode__2) { + enum libagx_tess_mode mode = mode__2; + uint patch = get_global_id(0); enum libagx_tess_partitioning partitioning = p->partitioning; global float *factors = tess_factors(p, patch); @@ -1151,10 +1155,12 @@ libagx_tess_tri(constant struct libagx_tess_args *p, enum libagx_tess_mode mode, } } -void +KERNEL(64) libagx_tess_quad(constant struct libagx_tess_args *p, - enum libagx_tess_mode mode, uint patch) + enum libagx_tess_mode mode__2) { + enum libagx_tess_mode mode = mode__2; + uint patch = get_global_id(0); enum libagx_tess_partitioning partitioning = p->partitioning; global float *factors = tess_factors(p, patch); diff --git a/src/asahi/libagx/tessellator.h b/src/asahi/libagx/tessellator.h index 3f07f612c0b..0eb920811be 100644 --- a/src/asahi/libagx/tessellator.h +++ b/src/asahi/libagx/tessellator.h @@ -59,44 +59,11 @@ struct libagx_tess_args { */ GLOBAL(uint32_t) statistic; - /* Address of the tess control invocation counter for implementing pipeline - * statistics, if active. Zero if inactive. Incremented by indirect tess - * setup kernel. - */ - GLOBAL(uint32_t) tcs_statistic; - - /* For indirect draws with tessellation, the grid sizes. VS then TCS then - * tess. Allocated by the CPU and written by the tessellation - * setup indirect kernel. - */ - GLOBAL(uint32_t) grids; - - /* For indirect draws, the output input assembly descriptor */ - GLOBAL(struct agx_ia_state) ia; - - /* For indirect draws, the indirect draw descriptor. */ - GLOBAL(uint32_t) indirect; - - /* For indirect draws, the allocation for the vertex buffer. - * - * TODO: We could move these fields to an indirect setup kernel, not sure if - * it's worth it though... - */ - GLOBAL(uint64_t) vertex_output_buffer_ptr; - - /* Yet more indirect draw garbage. I need a refactor. */ - uint64_t in_index_buffer; - uint32_t in_index_buffer_range_el; - uint32_t in_index_size_B; - /* When geom+tess used together, the buffer containing TES outputs (executed * as a hardware compute shader). */ uint64_t tes_buffer; - /* For indirect draws, the bitfield of VS outputs */ - uint64_t vertex_outputs; - /* Bitfield of TCS per-vertex outputs */ uint64_t tcs_per_vertex_outputs; @@ -137,4 +104,4 @@ struct libagx_tess_args { */ uint32_t ccw; } PACKED; -AGX_STATIC_ASSERT(sizeof(struct libagx_tess_args) == 51 * 4); +AGX_STATIC_ASSERT(sizeof(struct libagx_tess_args) == 35 * 4); diff --git a/src/asahi/vulkan/hk_cmd_buffer.c b/src/asahi/vulkan/hk_cmd_buffer.c index 34200375ad9..c2abbb90b36 100644 --- a/src/asahi/vulkan/hk_cmd_buffer.c +++ b/src/asahi/vulkan/hk_cmd_buffer.c @@ -22,6 +22,7 @@ #include "hk_physical_device.h" #include "hk_shader.h" +#include "libagx_dgc.h" #include "pool.h" #include "shader_enums.h" #include "vk_pipeline_layout.h" @@ -707,33 +708,21 @@ hk_upload_usc_words(struct hk_cmd_buffer *cmd, struct hk_shader *s, return agx_usc_addr(&dev->dev, t.gpu); } -/* Specialized variant of hk_upload_usc_words for internal dispatches that do - * not use any state except for some directly mapped uniforms. - */ -uint32_t -hk_upload_usc_words_kernel(struct hk_cmd_buffer *cmd, struct hk_shader *s, - void *data, size_t data_size) +void +hk_dispatch_precomp(struct hk_cs *cs, struct agx_grid grid, + enum libagx_program idx, void *data, size_t data_size) { - struct hk_device *dev = hk_cmd_buffer_device(cmd); + struct hk_device *dev = hk_cmd_buffer_device(cs->cmd); + struct agx_precompiled_shader *prog = agx_get_precompiled(&dev->bg_eot, idx); - assert(s->info.stage == MESA_SHADER_COMPUTE); - assert(s->b.info.scratch_size == 0 && "you shouldn't be spilling!"); - assert(s->b.info.preamble_scratch_size == 0 && "you shouldn't be spilling!"); + struct agx_ptr t = hk_pool_usc_alloc(cs->cmd, agx_usc_size(15), 64); + uint64_t uploaded_data = hk_pool_upload(cs->cmd, data, data_size, 4); - unsigned constant_push_ranges = DIV_ROUND_UP(s->b.info.rodata.size_16, 64); - size_t usc_size = agx_usc_size(constant_push_ranges + 7); - struct agx_ptr t = hk_pool_usc_alloc(cmd, usc_size, 64); - if (!t.cpu) - return 0; + agx_usc_words_precomp(t.cpu, &prog->b, uploaded_data, data_size); - struct agx_usc_builder b = agx_usc_builder(t.cpu, usc_size); - - /* Map the data directly as uniforms starting at u0 */ - agx_usc_uniform(&b, 0, DIV_ROUND_UP(data_size, 2), - hk_pool_upload(cmd, data, data_size, 4)); - - agx_usc_push_blob(&b, s->only_linked->usc.data, s->only_linked->usc.size); - return agx_usc_addr(&dev->dev, t.gpu); + hk_dispatch_with_usc_launch(dev, cs, prog->b.launch, + agx_usc_addr(&dev->dev, t.gpu), grid, + prog->b.workgroup); } void @@ -824,15 +813,9 @@ hk_ensure_cs_has_space(struct hk_cmd_buffer *cmd, struct hk_cs *cs, /* Jump from the old control stream to the new control stream */ if (vdm) { - agx_pack(cs->current, VDM_STREAM_LINK, cfg) { - cfg.target_lo = T.gpu & BITFIELD_MASK(32); - cfg.target_hi = T.gpu >> 32; - } + agx_vdm_jump(cs->current, T.gpu); } else { - agx_pack(cs->current, CDM_STREAM_LINK, cfg) { - cfg.target_lo = T.gpu & BITFIELD_MASK(32); - cfg.target_hi = T.gpu >> 32; - } + agx_cdm_jump(cs->current, T.gpu); } /* Swap out the control stream */ diff --git a/src/asahi/vulkan/hk_cmd_buffer.h b/src/asahi/vulkan/hk_cmd_buffer.h index cbba03e8241..6bc6b6931c4 100644 --- a/src/asahi/vulkan/hk_cmd_buffer.h +++ b/src/asahi/vulkan/hk_cmd_buffer.h @@ -15,6 +15,7 @@ #include "agx_pack.h" #include "agx_tilebuffer.h" #include "agx_uvs.h" +#include "libagx_dgc.h" #include "pool.h" #include "shader_enums.h" @@ -611,11 +612,7 @@ hk_cmd_buffer_end_compute_internal(struct hk_cmd_buffer *cmd, hk_dispatch_imm_writes(cmd, cs); } - void *map = cs->current; - agx_push(map, CDM_STREAM_TERMINATE, _) - ; - - cs->current = map; + cs->current = agx_cdm_terminate(cs->current); } *ptr = NULL; @@ -633,10 +630,6 @@ hk_cmd_buffer_end_graphics(struct hk_cmd_buffer *cmd) struct hk_cs *cs = cmd->current_cs.gfx; if (cs) { - void *map = cs->current; - agx_push(map, VDM_STREAM_TERMINATE, _) - ; - /* Scissor and depth bias arrays are staged to dynamic arrays on the CPU. * When we end the control stream, they're done growing and are ready for * upload. @@ -649,7 +642,7 @@ hk_cmd_buffer_end_graphics(struct hk_cmd_buffer *cmd) /* TODO: maybe free scissor/depth_bias now? */ - cmd->current_cs.gfx->current = map; + cmd->current_cs.gfx->current = agx_vdm_terminate(cs->current); cmd->current_cs.gfx = NULL; } @@ -736,54 +729,25 @@ void hk_reserve_scratch(struct hk_cmd_buffer *cmd, struct hk_cs *cs, uint32_t hk_upload_usc_words(struct hk_cmd_buffer *cmd, struct hk_shader *s, struct hk_linked_shader *linked); -uint32_t hk_upload_usc_words_kernel(struct hk_cmd_buffer *cmd, - struct hk_shader *s, void *data, - size_t data_size); - void hk_usc_upload_spilled_rt_descs(struct agx_usc_builder *b, struct hk_cmd_buffer *cmd); void hk_cdm_cache_flush(struct hk_device *dev, struct hk_cs *cs); -struct hk_grid { - bool indirect; - bool indirect_local; - union { - uint32_t count[3]; - uint64_t ptr; - }; -}; - -static struct hk_grid -hk_grid(uint32_t x, uint32_t y, uint32_t z) -{ - return (struct hk_grid){.indirect = false, .count = {x, y, z}}; -} - -static struct hk_grid -hk_grid_indirect(uint64_t ptr) -{ - return (struct hk_grid){.indirect = true, .ptr = ptr}; -} - -static struct hk_grid -hk_grid_indirect_local(uint64_t ptr) -{ - return (struct hk_grid){ - .indirect = true, - .indirect_local = true, - .ptr = ptr, - }; -} +void hk_dispatch_with_usc_launch(struct hk_device *dev, struct hk_cs *cs, + struct agx_cdm_launch_word_0_packed launch, + uint32_t usc, struct agx_grid grid, + struct agx_workgroup local_size); void hk_dispatch_with_usc(struct hk_device *dev, struct hk_cs *cs, struct agx_shader_info *info, uint32_t usc, - struct hk_grid grid, struct hk_grid local_size); + struct agx_grid grid, + struct agx_workgroup local_size); static inline void hk_dispatch_with_local_size(struct hk_cmd_buffer *cmd, struct hk_cs *cs, - struct hk_shader *s, struct hk_grid grid, - struct hk_grid local_size) + struct hk_shader *s, struct agx_grid grid, + struct agx_workgroup local_size) { struct hk_device *dev = hk_cmd_buffer_device(cmd); uint32_t usc = hk_upload_usc_words(cmd, s, s->only_linked); @@ -792,24 +756,10 @@ hk_dispatch_with_local_size(struct hk_cmd_buffer *cmd, struct hk_cs *cs, hk_dispatch_with_usc(dev, cs, &s->b.info, usc, grid, local_size); } -static inline void -hk_dispatch(struct hk_cmd_buffer *cmd, struct hk_cs *cs, struct hk_shader *s, - struct hk_grid grid) -{ - assert(s->info.stage == MESA_SHADER_COMPUTE); +void hk_dispatch_precomp(struct hk_cs *cs, struct agx_grid gird, + enum libagx_program idx, void *data, size_t data_size); - struct hk_grid local_size = - hk_grid(s->b.info.workgroup_size[0], s->b.info.workgroup_size[1], - s->b.info.workgroup_size[2]); - - if (!grid.indirect) { - grid.count[0] *= local_size.count[0]; - grid.count[1] *= local_size.count[1]; - grid.count[2] *= local_size.count[2]; - } - - hk_dispatch_with_local_size(cmd, cs, s, grid, local_size); -} +#define MESA_DISPATCH_PRECOMP hk_dispatch_precomp void hk_queue_write(struct hk_cmd_buffer *cmd, uint64_t address, uint32_t value, bool after_gfx); diff --git a/src/asahi/vulkan/hk_cmd_dispatch.c b/src/asahi/vulkan/hk_cmd_dispatch.c index 35fceb85bc9..d1131c99280 100644 --- a/src/asahi/vulkan/hk_cmd_dispatch.c +++ b/src/asahi/vulkan/hk_cmd_dispatch.c @@ -19,6 +19,8 @@ #include "hk_entrypoints.h" #include "hk_physical_device.h" #include "hk_shader.h" +#include "libagx_dgc.h" +#include "libagx_shaders.h" #include "pool.h" void @@ -47,130 +49,42 @@ hk_cdm_cache_flush(struct hk_device *dev, struct hk_cs *cs) assert(cs->current + AGX_CDM_BARRIER_LENGTH < cs->end && "caller must ensure space"); - uint8_t *out = cs->current; - - agx_push(out, CDM_BARRIER, cfg) { - cfg.unk_5 = true; - cfg.unk_6 = true; - cfg.unk_8 = true; - // cfg.unk_11 = true; - // cfg.unk_20 = true; - if (dev->dev.params.num_clusters_total > 1) { - // cfg.unk_24 = true; - if (dev->dev.params.gpu_generation == 13) { - cfg.unk_4 = true; - // cfg.unk_26 = true; - } - } - - /* With multiple launches in the same CDM stream, we can get cache - * coherency (? or sync?) issues. We hit this with blits, which need - in - * between dispatches - need the PBE cache to be flushed and the texture - * cache to be invalidated. Until we know what bits mean what exactly, - * let's just set these after every launch to be safe. We can revisit in - * the future when we figure out what the bits mean. - */ - cfg.unk_0 = true; - cfg.unk_1 = true; - cfg.unk_2 = true; - cfg.usc_cache_inval = true; - cfg.unk_4 = true; - cfg.unk_5 = true; - cfg.unk_6 = true; - cfg.unk_7 = true; - cfg.unk_8 = true; - cfg.unk_9 = true; - cfg.unk_10 = true; - cfg.unk_11 = true; - cfg.unk_12 = true; - cfg.unk_13 = true; - cfg.unk_14 = true; - cfg.unk_15 = true; - cfg.unk_16 = true; - cfg.unk_17 = true; - cfg.unk_18 = true; - cfg.unk_19 = true; - } - - cs->current = out; + cs->current = agx_cdm_barrier(cs->current, dev->dev.chip); cs->stats.flushes++; } -/* - * Enqueue workgroups to a given CDM control stream with a given prepared USC - * words. This does not interact with any global state, so it is suitable for - * internal dispatches that do not save/restore state. That can be simpler / - * lower overhead than vk_meta for special operations that logically operate - * as graphics. - */ +void +hk_dispatch_with_usc_launch(struct hk_device *dev, struct hk_cs *cs, + struct agx_cdm_launch_word_0_packed launch, + uint32_t usc, struct agx_grid grid, + struct agx_workgroup wg) +{ + assert(cs->current + 0x2000 < cs->end && "should have ensured space"); + cs->stats.cmds++; + + cs->current = + agx_cdm_launch(cs->current, dev->dev.chip, grid, wg, launch, usc); + + hk_cdm_cache_flush(dev, cs); +} + void hk_dispatch_with_usc(struct hk_device *dev, struct hk_cs *cs, struct agx_shader_info *info, uint32_t usc, - struct hk_grid grid, struct hk_grid local_size) + struct agx_grid grid, struct agx_workgroup local_size) { - assert(cs->current + 0x2000 < cs->end && "should have ensured space"); - uint8_t *out = cs->current; - - cs->stats.cmds++; - - agx_push(out, CDM_LAUNCH_WORD_0, cfg) { - if (grid.indirect) { - if (grid.indirect_local) - cfg.mode = AGX_CDM_MODE_INDIRECT_LOCAL; - else - cfg.mode = AGX_CDM_MODE_INDIRECT_GLOBAL; - } else { - cfg.mode = AGX_CDM_MODE_DIRECT; - } - - /* For now, always bind the txf sampler and nothing else */ + struct agx_cdm_launch_word_0_packed launch; + agx_pack(&launch, CDM_LAUNCH_WORD_0, cfg) { cfg.sampler_state_register_count = 1; - cfg.uniform_register_count = info->push_count; cfg.preshader_register_count = info->nr_preamble_gprs; } - agx_push(out, CDM_LAUNCH_WORD_1, cfg) { - cfg.pipeline = usc; - } - - /* Added in G14X */ - if (dev->dev.params.gpu_generation >= 14 && - dev->dev.params.num_clusters_total > 1) { - - agx_push(out, CDM_UNK_G14X, cfg) - ; - } - - assert(!local_size.indirect); - - if (grid.indirect) { - agx_push(out, CDM_INDIRECT, cfg) { - cfg.address_hi = grid.ptr >> 32; - cfg.address_lo = grid.ptr & BITFIELD64_MASK(32); - } - } else { - agx_push(out, CDM_GLOBAL_SIZE, cfg) { - cfg.x = grid.count[0]; - cfg.y = grid.count[1]; - cfg.z = grid.count[2]; - } - } - - if (!grid.indirect_local) { - agx_push(out, CDM_LOCAL_SIZE, cfg) { - cfg.x = local_size.count[0]; - cfg.y = local_size.count[1]; - cfg.z = local_size.count[2]; - } - } - - cs->current = out; - hk_cdm_cache_flush(dev, cs); + hk_dispatch_with_usc_launch(dev, cs, launch, usc, grid, local_size); } static void -dispatch(struct hk_cmd_buffer *cmd, struct hk_grid grid) +dispatch(struct hk_cmd_buffer *cmd, struct agx_grid grid) { struct hk_device *dev = hk_cmd_buffer_device(cmd); struct hk_shader *s = hk_only_variant(cmd->state.cs.shader); @@ -178,34 +92,30 @@ dispatch(struct hk_cmd_buffer *cmd, struct hk_grid grid) if (!cs) return; + struct agx_workgroup local_size = + agx_workgroup(s->b.info.workgroup_size[0], s->b.info.workgroup_size[1], + s->b.info.workgroup_size[2]); + uint64_t stat = hk_pipeline_stat_addr( cmd, VK_QUERY_PIPELINE_STATISTIC_COMPUTE_SHADER_INVOCATIONS_BIT); if (stat) { - uint32_t local_size_threads = (uint32_t)s->b.info.workgroup_size[0] * - (uint32_t)s->b.info.workgroup_size[1] * - (uint32_t)s->b.info.workgroup_size[2]; - - struct libagx_cs_invocation_params p = { - .grid = cmd->state.cs.descriptors.root.cs.group_count_addr, - .local_size_threads = local_size_threads, - .statistic = stat, - }; - - struct hk_shader *s = - hk_meta_kernel(dev, agx_nir_increment_cs_invocations, NULL, 0); - - uint64_t params = hk_pool_upload(cmd, &p, sizeof(p), 8); - uint32_t usc = - hk_upload_usc_words_kernel(cmd, s, ¶ms, sizeof(params)); - perf_debug(dev, "CS invocation statistic"); - hk_dispatch_with_usc(dev, cs, &s->b.info, usc, hk_grid(1, 1, 1), - hk_grid(1, 1, 1)); + uint64_t grid = cmd->state.cs.descriptors.root.cs.group_count_addr; + + libagx_increment_cs_invocations(cs, agx_1d(1), grid, stat, + agx_workgroup_threads(local_size)); } hk_ensure_cs_has_space(cmd, cs, 0x2000 /* TODO */); - hk_dispatch(cmd, cs, s, grid); + + if (!agx_is_indirect(grid)) { + grid.count[0] *= local_size.x; + grid.count[1] *= local_size.y; + grid.count[2] *= local_size.z; + } + + hk_dispatch_with_local_size(cmd, cs, s, grid, local_size); cs->stats.calls++; } @@ -236,7 +146,7 @@ hk_CmdDispatchBase(VkCommandBuffer commandBuffer, uint32_t baseGroupX, desc->root.cs.group_count_addr = hk_pool_upload(cmd, &group_count, sizeof(group_count), 8); - dispatch(cmd, hk_grid(groupCountX, groupCountY, groupCountZ)); + dispatch(cmd, agx_3d(groupCountX, groupCountY, groupCountZ)); } VKAPI_ATTR void VKAPI_CALL @@ -257,5 +167,5 @@ hk_CmdDispatchIndirect(VkCommandBuffer commandBuffer, VkBuffer _buffer, assert(dispatch_addr != 0); desc->root.cs.group_count_addr = dispatch_addr; - dispatch(cmd, hk_grid_indirect(dispatch_addr)); + dispatch(cmd, agx_grid_indirect(dispatch_addr)); } diff --git a/src/asahi/vulkan/hk_cmd_draw.c b/src/asahi/vulkan/hk_cmd_draw.c index c9ca3fda2be..35f4670333a 100644 --- a/src/asahi/vulkan/hk_cmd_draw.c +++ b/src/asahi/vulkan/hk_cmd_draw.c @@ -29,7 +29,7 @@ #include "hk_shader.h" #include "asahi/genxml/agx_pack.h" -#include "asahi/libagx/draws.h" +#include "asahi/libagx/compression.h" #include "asahi/libagx/geometry.h" #include "asahi/libagx/libagx.h" #include "asahi/libagx/query.h" @@ -41,6 +41,7 @@ #include "util/ralloc.h" #include "vulkan/vulkan_core.h" #include "layout.h" +#include "libagx_shaders.h" #include "nir.h" #include "nir_builder.h" #include "nir_lower_blend.h" @@ -70,7 +71,7 @@ #define HK_TEST_INDIRECTS (0) struct hk_draw { - struct hk_grid b; + struct agx_grid b; struct hk_addr_range index; bool indexed; uint32_t start; @@ -85,7 +86,7 @@ struct hk_draw { UNUSED static inline void print_draw(struct hk_draw d, FILE *fp) { - if (d.b.indirect) + if (agx_is_indirect(d.b)) fprintf(fp, "indirect (buffer %" PRIx64 "):", d.b.ptr); else fprintf(fp, "direct (%ux%u):", d.b.count[0], d.b.count[1]); @@ -113,7 +114,7 @@ print_draw(struct hk_draw d, FILE *fp) static struct hk_draw hk_draw_indirect(uint64_t ptr) { - return (struct hk_draw){.b = hk_grid_indirect(ptr)}; + return (struct hk_draw){.b = agx_grid_indirect(ptr)}; } static struct hk_draw @@ -121,7 +122,7 @@ hk_draw_indexed_indirect(uint64_t ptr, struct hk_addr_range index, enum agx_index_size index_size, bool restart) { return (struct hk_draw){ - .b = hk_grid_indirect(ptr), + .b = agx_grid_indirect(ptr), .index = index, .indexed = true, .index_size = index_size, @@ -460,7 +461,7 @@ hk_build_bg_eot(struct hk_cmd_buffer *cmd, const VkRenderingInfo *info, } if (uses_txf) { - agx_usc_push_packed(&b, SAMPLER, dev->rodata.txf_sampler); + agx_usc_push_packed(&b, SAMPLER, dev->dev.txf_sampler); } /* For attachmentless rendering, we don't know the sample count until @@ -878,31 +879,20 @@ hk_CmdBeginRendering(VkCommandBuffer commandBuffer, return; unsigned level = view->vk.base_mip_level; + unsigned layer = view->vk.base_array_layer; + uint64_t base = hk_image_base_address(image, image_plane); - struct agx_ptr data = - hk_pool_alloc(cmd, sizeof(struct libagx_decompress_push), 64); - struct libagx_decompress_push *push = data.cpu; - agx_fill_decompress_push( - push, layout, view->vk.base_array_layer, level, - hk_image_base_address(image, image_plane)); - - push->compressed = view->planes[plane].emrt_texture; - push->uncompressed = view->planes[plane].emrt_pbe; - - struct hk_grid grid = - hk_grid(ail_metadata_width_tl(layout, level) * 32, - ail_metadata_height_tl(layout, level), layer_count); - - struct agx_decompress_key key = { - .nr_samples = layout->sample_count_sa, + struct libagx_decompress_images imgs = { + .compressed = view->planes[plane].emrt_texture, + .uncompressed = view->planes[plane].emrt_pbe, }; - struct hk_shader *s = - hk_meta_kernel(dev, agx_nir_decompress, &key, sizeof(key)); + struct agx_grid grid = + agx_3d(ail_metadata_width_tl(layout, level) * 32, + ail_metadata_height_tl(layout, level), layer_count); - uint32_t usc = hk_upload_usc_words_kernel(cmd, s, &data.gpu, 8); - hk_dispatch_with_usc(dev, cs, &s->b.info, usc, grid, - hk_grid(32, 1, 1)); + libagx_decompress(cs, grid, layout, layer, level, base, + hk_pool_upload(cmd, &imgs, sizeof(imgs), 64)); } } } @@ -1099,48 +1089,11 @@ hk_geometry_state(struct hk_cmd_buffer *cmd) return dev->rodata.geometry_state; } -static uint64_t -hk_upload_gsi_params(struct hk_cmd_buffer *cmd, struct hk_draw draw) -{ - struct hk_device *dev = hk_cmd_buffer_device(cmd); - struct hk_descriptor_state *desc = &cmd->state.gfx.descriptors; - struct hk_graphics_state *gfx = &cmd->state.gfx; - struct hk_shader *vs = hk_bound_sw_vs_before_gs(gfx); - - unsigned index_size_B = - draw.indexed ? agx_index_size_to_B(draw.index_size) : 0; - - uint64_t vb; - if (cmd->state.gfx.shaders[MESA_SHADER_TESS_EVAL]) { - assert(index_size_B == 4); - - vb = desc->root.draw.tess_params + - offsetof(struct libagx_tess_args, tes_buffer); - } else { - vb = desc->root.root_desc_addr + - offsetof(struct hk_root_descriptor_table, draw.vertex_output_buffer); - } - - struct agx_gs_setup_indirect_params gsi = { - .index_buffer = draw.index.addr, - .index_size_B = index_size_B, - .index_buffer_range_el = draw.index.range / index_size_B, - .zero_sink = dev->rodata.zero_sink, - .draw = draw.b.ptr, - .vertex_buffer = vb, - .ia = desc->root.draw.input_assembly, - .geom = desc->root.draw.geometry_params, - .vs_outputs = vs->b.info.outputs, - }; - - return hk_pool_upload(cmd, &gsi, sizeof(gsi), 8); -} - static uint64_t hk_upload_ia_params(struct hk_cmd_buffer *cmd, struct hk_draw draw) { struct hk_device *dev = hk_cmd_buffer_device(cmd); - assert(!draw.b.indirect && "indirect params written by GPU"); + assert(!agx_is_indirect(draw.b) && "indirect params written by GPU"); struct agx_ia_state ia = {.verts_per_instance = draw.b.count[0]}; @@ -1199,7 +1152,7 @@ hk_upload_geometry_params(struct hk_cmd_buffer *cmd, struct hk_draw draw) /* XXX: We should deduplicate this logic */ bool restart = (draw.indexed && draw.restart); bool indirect = - draw.b.indirect || gfx->shaders[MESA_SHADER_TESS_EVAL] || restart; + agx_is_indirect(draw.b) || gfx->shaders[MESA_SHADER_TESS_EVAL] || restart; enum mesa_prim mode = hk_gs_in_prim(cmd); if (restart) { @@ -1239,10 +1192,9 @@ hk_upload_geometry_params(struct hk_cmd_buffer *cmd, struct hk_draw draw) /* Calculate input primitive count for direct draws, and allocate the vertex * & count buffers. GPU calculates and allocates for indirect draws. */ - unsigned count_buffer_stride = count->info.gs.count_words * 4; + params.count_buffer_stride = count->info.gs.count_words * 4; if (indirect) { - params.count_buffer_stride = count_buffer_stride; params.vs_grid[2] = params.gs_grid[2] = 1; } else { uint32_t verts = draw.b.count[0], instances = draw.b.count[1]; @@ -1253,7 +1205,7 @@ hk_upload_geometry_params(struct hk_cmd_buffer *cmd, struct hk_draw draw) params.primitives_log2 = util_logbase2_ceil(params.gs_grid[0]); params.input_primitives = params.gs_grid[0] * instances; - unsigned size = params.input_primitives * count_buffer_stride; + unsigned size = params.input_primitives * params.count_buffer_stride; if (size) { params.count_buffer = hk_pool_alloc(cmd, size, 4).gpu; } @@ -1306,7 +1258,7 @@ hk_upload_tess_params(struct hk_cmd_buffer *cmd, struct libagx_tess_args *out, /* heap is allocated by hk_geometry_state */ args.patch_coord_buffer = dev->heap->va->addr; - if (!draw.b.indirect) { + if (!agx_is_indirect(draw.b)) { unsigned in_patches = draw.b.count[0] / args.input_patch_size; unsigned unrolled_patches = in_patches * draw.b.count[1]; @@ -1332,31 +1284,11 @@ hk_upload_tess_params(struct hk_cmd_buffer *cmd, struct libagx_tess_args *out, args.out_draws = blob.gpu + draw_offs; args.counts = blob.gpu + count_offs; } else { - args.tcs_statistic = hk_pipeline_stat_addr( - cmd, - VK_QUERY_PIPELINE_STATISTIC_TESSELLATION_CONTROL_SHADER_PATCHES_BIT); - - args.indirect = draw.b.ptr; - /* Allocate 3x indirect global+local grids for VS/TCS/tess */ uint32_t grid_stride = sizeof(uint32_t) * 6; - args.grids = hk_pool_alloc(cmd, grid_stride * 3, 4).gpu; - gfx->tess.grids = args.grids; + gfx->tess.grids = hk_pool_alloc(cmd, grid_stride * 3, 4).gpu; - struct hk_shader *vs = hk_bound_sw_vs(gfx); - args.vertex_outputs = vs->b.info.outputs; - args.vertex_output_buffer_ptr = - gfx->root + - offsetof(struct hk_root_descriptor_table, draw.vertex_output_buffer); - args.ia = gfx->descriptors.root.draw.input_assembly; args.out_draws = hk_pool_alloc(cmd, draw_stride_B, 4).gpu; - - if (draw.indexed) { - args.in_index_buffer = draw.index.addr; - args.in_index_size_B = agx_index_size_to_B(draw.index_size); - args.in_index_buffer_range_el = - draw.index.range / args.in_index_size_B; - } } gfx->tess.out_draws = args.out_draws; @@ -1389,22 +1321,6 @@ hk_build_meta_shader_locked(struct hk_device *dev, struct hk_internal_key *key, .robustness = &rs, }; - /* We need to link libagx and assign shared before preprocessing, matching - * what the driver would otherwise produce. - */ - agx_link_libagx(b.shader, dev->dev.libagx); - - if (info.stage == MESA_SHADER_COMPUTE) { - NIR_PASS(_, b.shader, nir_lower_vars_to_explicit_types, - nir_var_mem_shared, glsl_get_cl_type_size_align); - - /* Commit to the layout so we don't clobber later */ - b.shader->info.shared_memory_explicit_layout = true; - - NIR_PASS(_, b.shader, nir_lower_explicit_io, nir_var_mem_shared, - nir_address_format_62bit_generic); - } - hk_preprocess_nir_internal(dev->vk.physical, b.shader); struct hk_api_shader *s; @@ -1453,7 +1369,7 @@ hk_draw_without_restart(struct hk_cmd_buffer *cmd, struct hk_cs *cs, perf_debug(dev, "Unrolling primitive restart due to GS/XFB"); /* The unroll kernel assumes an indirect draw. Synthesize one if needed */ - if (!draw.b.indirect) { + if (!agx_is_indirect(draw.b)) { uint32_t desc[5] = {draw.b.count[0], draw.b.count[1], draw.start, draw.index_bias, draw.start_instance}; @@ -1463,41 +1379,33 @@ hk_draw_without_restart(struct hk_cmd_buffer *cmd, struct hk_cs *cs, } /* Next, we unroll the index buffer used by the indirect draw */ - struct agx_unroll_restart_key key = { - .prim = vk_conv_topology(dyn->ia.primitive_topology), - }; + enum mesa_prim prim = vk_conv_topology(dyn->ia.primitive_topology); - uint32_t index_size_B = agx_index_size_to_B(draw.index_size); + assert(draw_count == 1 && "TODO: multidraw"); - struct agx_restart_unroll_params ia = { + struct libagx_unroll_restart_args ia = { .heap = hk_geometry_state(cmd), .index_buffer = draw.index.addr, - .count = hk_pool_upload(cmd, &draw_count, sizeof(uint32_t), 4), - .draws = draw.b.ptr, - .out_draws = hk_pool_alloc(cmd, 5 * sizeof(uint32_t) * draw_count, 4).gpu, + .in_draw = draw.b.ptr, + .out_draw = hk_pool_alloc(cmd, 5 * sizeof(uint32_t) * draw_count, 4).gpu, .max_draws = 1 /* TODO: MDI */, .restart_index = gfx->index.restart, - .index_buffer_size_el = draw.index.range / index_size_B, + .index_buffer_size_el = + draw.index.range / agx_translate_index_size(draw.index_size), .flatshade_first = dyn->rs.provoking_vertex == VK_PROVOKING_VERTEX_MODE_FIRST_VERTEX_EXT, .zero_sink = dev->rodata.zero_sink, - .index_size_B = index_size_B, }; - struct hk_shader *s = - hk_meta_kernel(dev, agx_nir_unroll_restart, &key, sizeof(key)); - - uint64_t params = hk_pool_upload(cmd, &ia, sizeof(ia), 8); - uint32_t usc = hk_upload_usc_words_kernel(cmd, s, ¶ms, sizeof(params)); - hk_dispatch_with_usc(dev, cs, &s->b.info, usc, - hk_grid(1024 * draw_count, 1, 1), hk_grid(1024, 1, 1)); + libagx_unroll_restart_struct(cs, agx_1d(1024 * draw_count), ia, + draw.index_size, libagx_compact_prim(prim)); struct hk_addr_range out_index = { .addr = dev->heap->va->addr, .range = dev->heap->size, }; - return hk_draw_indexed_indirect(ia.out_draws, out_index, draw.index_size, + return hk_draw_indexed_indirect(ia.out_draw, out_index, draw.index_size, false /* restart */); } @@ -1509,7 +1417,7 @@ hk_launch_gs_prerast(struct hk_cmd_buffer *cmd, struct hk_cs *cs, struct hk_graphics_state *gfx = &cmd->state.gfx; struct hk_descriptor_state *desc = &cmd->state.gfx.descriptors; struct hk_api_shader *gs = gfx->shaders[MESA_SHADER_GEOMETRY]; - struct hk_grid grid_vs, grid_gs; + struct agx_grid grid_vs, grid_gs; struct vk_dynamic_graphics_state *dyn = &cmd->vk.dynamic_graphics_state; bool rast_disc = dyn->rs.rasterizer_discard_enable; @@ -1521,6 +1429,7 @@ hk_launch_gs_prerast(struct hk_cmd_buffer *cmd, struct hk_cs *cs, struct hk_shader *count = hk_count_gs_variant(gs, rast_disc); struct hk_shader *pre_gs = hk_pre_gs_variant(gs, rast_disc); + uint64_t geometry_params = desc->root.draw.geometry_params; unsigned count_words = count->info.gs.count_words; if (false /* TODO */) @@ -1538,24 +1447,38 @@ hk_launch_gs_prerast(struct hk_cmd_buffer *cmd, struct hk_cs *cs, } /* Setup grids */ - if (draw.b.indirect) { - struct agx_gs_setup_indirect_key key = {.prim = mode}; + if (agx_is_indirect(draw.b)) { + struct libagx_gs_setup_indirect_args gsi = { + .index_buffer = draw.index.addr, + .zero_sink = dev->rodata.zero_sink, + .draw = draw.b.ptr, + .ia = desc->root.draw.input_assembly, + .p = desc->root.draw.geometry_params, + .vs_outputs = vs->b.info.outputs, + .prim = mode, + }; - struct hk_shader *gsi = - hk_meta_kernel(dev, agx_nir_gs_setup_indirect, &key, sizeof(key)); + if (cmd->state.gfx.shaders[MESA_SHADER_TESS_EVAL]) { + gsi.vertex_buffer = desc->root.draw.tess_params + + offsetof(struct libagx_tess_args, tes_buffer); + } else { + gsi.vertex_buffer = desc->root.root_desc_addr + + offsetof(struct hk_root_descriptor_table, + draw.vertex_output_buffer); + } - uint64_t push = hk_upload_gsi_params(cmd, draw); - uint32_t usc = hk_upload_usc_words_kernel(cmd, gsi, &push, sizeof(push)); + if (draw.indexed) { + gsi.index_size_B = agx_index_size_to_B(draw.index_size); + gsi.index_buffer_range_el = draw.index.range / gsi.index_size_B; + } - hk_dispatch_with_usc(dev, cs, &gsi->b.info, usc, hk_grid(1, 1, 1), - hk_grid(1, 1, 1)); + libagx_gs_setup_indirect_struct(cs, agx_1d(1), gsi); - uint64_t geometry_params = desc->root.draw.geometry_params; - grid_vs = hk_grid_indirect(geometry_params + - offsetof(struct agx_geometry_params, vs_grid)); + grid_vs = agx_grid_indirect( + geometry_params + offsetof(struct agx_geometry_params, vs_grid)); - grid_gs = hk_grid_indirect(geometry_params + - offsetof(struct agx_geometry_params, gs_grid)); + grid_gs = agx_grid_indirect( + geometry_params + offsetof(struct agx_geometry_params, gs_grid)); } else { grid_vs = grid_gs = draw.b; grid_gs.count[0] = u_decomposed_prims_for_vertices(mode, draw.b.count[0]); @@ -1568,29 +1491,22 @@ hk_launch_gs_prerast(struct hk_cmd_buffer *cmd, struct hk_cs *cs, vs->info.stage == MESA_SHADER_VERTEX ? gfx->linked[MESA_SHADER_VERTEX] : vs->only_linked), - grid_vs, hk_grid(1, 1, 1)); + grid_vs, agx_workgroup(1, 1, 1)); /* If we need counts, launch the count shader and prefix sum the results. */ if (count_words) { - hk_dispatch_with_local_size(cmd, cs, count, grid_gs, hk_grid(1, 1, 1)); + hk_dispatch_with_local_size(cmd, cs, count, grid_gs, + agx_workgroup(1, 1, 1)); - struct hk_api_shader *prefix_sum = hk_meta_shader( - dev, agx_nir_prefix_sum_gs, &count_words, sizeof(count_words)); - - /* XXX: hack */ - hk_only_variant(prefix_sum)->info.stage = MESA_SHADER_GEOMETRY; - - hk_dispatch_with_local_size(cmd, cs, hk_only_variant(prefix_sum), - hk_grid(1024 * count_words, 1, 1), - hk_grid(1024, 1, 1)); + libagx_prefix_sum_geom(cs, agx_1d(1024 * count_words), geometry_params); } /* Pre-GS shader */ - hk_dispatch_with_local_size(cmd, cs, pre_gs, hk_grid(1, 1, 1), - hk_grid(1, 1, 1)); + hk_dispatch_with_local_size(cmd, cs, pre_gs, agx_1d(1), + agx_workgroup(1, 1, 1)); /* Pre-rast geometry shader */ - hk_dispatch_with_local_size(cmd, cs, main, grid_gs, hk_grid(1, 1, 1)); + hk_dispatch_with_local_size(cmd, cs, main, grid_gs, agx_workgroup(1, 1, 1)); struct hk_addr_range range = (struct hk_addr_range){ .addr = dev->heap->va->addr, @@ -1607,7 +1523,7 @@ hk_launch_tess(struct hk_cmd_buffer *cmd, struct hk_cs *cs, struct hk_draw draw) { struct hk_device *dev = hk_cmd_buffer_device(cmd); struct hk_graphics_state *gfx = &cmd->state.gfx; - struct hk_grid grid_vs, grid_tcs, grid_tess; + struct agx_grid grid_vs, grid_tcs, grid_tess; struct hk_shader *vs = hk_bound_sw_vs(gfx); struct hk_shader *tcs = hk_only_variant(gfx->shaders[MESA_SHADER_TESS_CTRL]); @@ -1625,51 +1541,47 @@ hk_launch_tess(struct hk_cmd_buffer *cmd, struct hk_cs *cs, struct hk_draw draw) cmd, VK_QUERY_PIPELINE_STATISTIC_TESSELLATION_CONTROL_SHADER_PATCHES_BIT); /* Setup grids */ - if (draw.b.indirect) { + if (agx_is_indirect(draw.b)) { perf_debug(dev, "Indirect tessellation"); - struct agx_tess_setup_indirect_key key = { - .point_mode = info.points, + struct libagx_tess_setup_indirect_args args = { + .p = state, + .grids = gfx->tess.grids, + .indirect = draw.b.ptr, + .ia = gfx->descriptors.root.draw.input_assembly, + .vertex_outputs = vs->b.info.outputs, + .vertex_output_buffer_ptr = + gfx->root + offsetof(struct hk_root_descriptor_table, + draw.vertex_output_buffer), + .tcs_statistic = hk_pipeline_stat_addr( + cmd, + VK_QUERY_PIPELINE_STATISTIC_TESSELLATION_CONTROL_SHADER_PATCHES_BIT), }; - struct hk_shader *tsi = - hk_meta_kernel(dev, agx_nir_tess_setup_indirect, &key, sizeof(key)); + if (draw.indexed) { + args.in_index_buffer = draw.index.addr; + args.in_index_size_B = agx_index_size_to_B(draw.index_size); + args.in_index_buffer_range_el = + draw.index.range / args.in_index_size_B; + } - /* TODO */ - uint32_t usc = - hk_upload_usc_words_kernel(cmd, tsi, &state, sizeof(state)); - - hk_dispatch_with_usc(dev, cs, &tsi->b.info, usc, hk_grid(1, 1, 1), - hk_grid(1, 1, 1)); + libagx_tess_setup_indirect_struct(cs, agx_1d(1), args); uint32_t grid_stride = sizeof(uint32_t) * 6; - grid_vs = hk_grid_indirect_local(gfx->tess.grids + 0 * grid_stride); - grid_tcs = hk_grid_indirect_local(gfx->tess.grids + 1 * grid_stride); - grid_tess = hk_grid_indirect_local(gfx->tess.grids + 2 * grid_stride); + grid_vs = agx_grid_indirect_local(gfx->tess.grids + 0 * grid_stride); + grid_tcs = agx_grid_indirect_local(gfx->tess.grids + 1 * grid_stride); + grid_tess = agx_grid_indirect_local(gfx->tess.grids + 2 * grid_stride); } else { uint32_t patches = draw.b.count[0] / input_patch_size; grid_vs = grid_tcs = draw.b; grid_tcs.count[0] = patches * tcs->info.tess.tcs_output_patch_size; - grid_tess = hk_grid(patches * draw.b.count[1], 1, 1); + grid_tess = agx_1d(patches * draw.b.count[1]); /* TCS invocation counter increments once per-patch */ if (tcs_stat) { perf_debug(dev, "Direct TCS statistic"); - - struct libagx_increment_params args = { - .statistic = tcs_stat, - .delta = patches, - }; - - struct hk_shader *s = - hk_meta_kernel(dev, agx_nir_increment_statistic, NULL, 0); - - uint64_t push = hk_pool_upload(cmd, &args, sizeof(args), 8); - uint32_t usc = hk_upload_usc_words_kernel(cmd, s, &push, sizeof(push)); - - hk_dispatch_with_usc(dev, cs, &s->b.info, usc, hk_grid(1, 1, 1), - hk_grid(1, 1, 1)); + libagx_increment_statistic(cs, agx_1d(1), tcs_stat, patches); } } @@ -1680,51 +1592,17 @@ hk_launch_tess(struct hk_cmd_buffer *cmd, struct hk_cs *cs, struct hk_draw draw) hk_dispatch_with_usc( dev, cs, &vs->b.info, hk_upload_usc_words(cmd, vs, gfx->linked[MESA_SHADER_VERTEX]), grid_vs, - hk_grid(64, 1, 1)); + agx_workgroup(64, 1, 1)); hk_dispatch_with_usc( dev, cs, &tcs->b.info, hk_upload_usc_words(cmd, tcs, tcs->only_linked), - grid_tcs, hk_grid(tcs->info.tess.tcs_output_patch_size, 1, 1)); + grid_tcs, agx_workgroup(tcs->info.tess.tcs_output_patch_size, 1, 1)); - struct agx_tessellator_key key = { - .prim = info.mode, - }; - - /* Generate counts */ - key.mode = LIBAGX_TESS_MODE_COUNT; - { - struct hk_shader *tess = - hk_meta_kernel(dev, agx_nir_tessellate, &key, sizeof(key)); - - hk_dispatch_with_usc( - dev, cs, &tess->b.info, - hk_upload_usc_words_kernel(cmd, tess, &state, sizeof(state)), - grid_tess, hk_grid(64, 1, 1)); - } - - /* Prefix sum counts, allocating index buffer space. */ - { - struct hk_shader *sum = - hk_meta_kernel(dev, agx_nir_prefix_sum_tess, NULL, 0); - - hk_dispatch_with_usc( - dev, cs, &sum->b.info, - hk_upload_usc_words_kernel(cmd, sum, &state, sizeof(state)), - hk_grid(1024, 1, 1), hk_grid(1024, 1, 1)); - } - - key.mode = LIBAGX_TESS_MODE_WITH_COUNTS; - - /* Now we can tessellate */ - { - struct hk_shader *tess = - hk_meta_kernel(dev, agx_nir_tessellate, &key, sizeof(key)); - - hk_dispatch_with_usc( - dev, cs, &tess->b.info, - hk_upload_usc_words_kernel(cmd, tess, &state, sizeof(state)), - grid_tess, hk_grid(64, 1, 1)); - } + /* First generate counts, then prefix sum them, and then tessellate. */ + libagx_tessellate(cs, grid_tess, info.mode, LIBAGX_TESS_MODE_COUNT, state); + libagx_prefix_sum_tess(cs, agx_1d(1024), state); + libagx_tessellate(cs, grid_tess, info.mode, LIBAGX_TESS_MODE_WITH_COUNTS, + state); struct hk_addr_range range = (struct hk_addr_range){ .addr = dev->heap->va->addr, @@ -2915,7 +2793,7 @@ hk_flush_dynamic_state(struct hk_cmd_buffer *cmd, struct hk_cs *cs, * avoid keying to indirectness. */ if (gfx->linked[MESA_SHADER_VERTEX]->b.uses_base_param) { - if (draw.b.indirect) { + if (agx_is_indirect(draw.b)) { gfx->draw_params = draw.b.ptr; if (draw.indexed) { @@ -3049,7 +2927,7 @@ hk_flush_dynamic_state(struct hk_cmd_buffer *cmd, struct hk_cs *cs, /* XXX: We should deduplicate this logic */ bool restart = (draw.indexed && draw.restart); - bool indirect = draw.b.indirect || restart; + bool indirect = agx_is_indirect(draw.b) || restart; desc->root.draw.input_assembly = indirect ? hk_pool_alloc(cmd, sizeof(struct agx_ia_state), 4).gpu @@ -3208,7 +3086,7 @@ hk_needs_index_robustness(struct hk_cmd_buffer *cmd, struct hk_draw draw) dev->vk.enabled_features.pipelineRobustness)) return false; - if (draw.b.indirect) { + if (agx_is_indirect(draw.b)) { return true; } else { uint32_t range_B = @@ -3477,38 +3355,24 @@ hk_ia_update(struct hk_cmd_buffer *cmd, struct hk_cs *cs, struct hk_draw draw, struct hk_device *dev = hk_cmd_buffer_device(cmd); perf_debug(dev, "Input assembly counters"); - struct agx_increment_ia_counters_key key = { - .index_size_B = draw.restart ? agx_index_size_to_B(draw.index_size) : 0, - }; - uint64_t draw_ptr; - if (draw.b.indirect) { + if (agx_is_indirect(draw.b)) { draw_ptr = draw.b.ptr; } else { uint32_t desc[] = {draw.b.count[0], draw.b.count[1], 0}; draw_ptr = hk_pool_upload(cmd, &desc, sizeof(desc), 4); } - struct libagx_increment_ia_counters args = { - .ia_vertices = ia_vertices, - .vs_invocations = vs_invocations, - .restart_index = cmd->state.gfx.index.restart, - .draw = draw_ptr, - .index_buffer = draw.index.addr, - .index_buffer_range_el = - key.index_size_B ? (draw.index.range / key.index_size_B) : 0, - }; + if (draw.restart && draw.indexed) { + uint32_t index_size_B = agx_index_size_to_B(draw.index_size); + uint32_t index_range_el = draw.index.range / index_size_B; - uint64_t wg_size = key.index_size_B ? 1024 : 1; - - struct hk_shader *s = - hk_meta_kernel(dev, agx_nir_increment_ia_counters, &key, sizeof(key)); - - uint64_t push = hk_pool_upload(cmd, &args, sizeof(args), 8); - uint32_t usc = hk_upload_usc_words_kernel(cmd, s, &push, sizeof(push)); - - hk_dispatch_with_usc(dev, cs, &s->b.info, usc, hk_grid(wg_size, 1, 1), - hk_grid(wg_size, 1, 1)); + libagx_increment_ia_restart(cs, agx_1d(1024), ia_vertices, vs_invocations, + draw_ptr, draw.index.addr, index_range_el, + cmd->state.gfx.index.restart, index_size_B); + } else { + libagx_increment_ia(cs, agx_1d(1), ia_vertices, vs_invocations, draw_ptr); + } } static void @@ -3518,7 +3382,8 @@ hk_draw(struct hk_cmd_buffer *cmd, uint16_t draw_id, struct hk_draw draw_) &cmd->vk.dynamic_graphics_state; /* Filter trivial draws so we don't need to worry about null index buffers */ - if (!draw_.b.indirect && (draw_.b.count[0] == 0 || draw_.b.count[1] == 0)) + if (!agx_is_indirect(draw_.b) && + (draw_.b.count[0] == 0 || draw_.b.count[1] == 0)) return; draw_.restart = dyn->ia.primitive_restart_enable; @@ -3581,13 +3446,13 @@ hk_draw(struct hk_cmd_buffer *cmd, uint16_t draw_id, struct hk_draw draw_) } uint64_t ib = draw.index.addr; - if (draw.indexed && !draw.b.indirect) + if (draw.indexed && !agx_is_indirect(draw.b)) ib += (draw.start << draw.index_size); agx_push(out, INDEX_LIST, cfg) { cfg.primitive = cmd->state.gfx.topology; - if (draw.b.indirect) { + if (agx_is_indirect(draw.b)) { cfg.indirect_buffer_present = true; } else { cfg.instance_count_present = true; @@ -3611,7 +3476,7 @@ hk_draw(struct hk_cmd_buffer *cmd, uint16_t draw_id, struct hk_draw draw_) } } - if (draw.b.indirect) { + if (agx_is_indirect(draw.b)) { agx_push(out, INDEX_LIST_INDIRECT_BUFFER, cfg) { cfg.address_hi = draw.b.ptr >> 32; cfg.address_lo = draw.b.ptr & BITFIELD_MASK(32); @@ -3659,7 +3524,7 @@ hk_CmdDraw(VkCommandBuffer commandBuffer, uint32_t vertexCount, draw = hk_draw_indirect(hk_pool_upload(cmd, data, sizeof(data), 4)); } else { draw = (struct hk_draw){ - .b = hk_grid(vertexCount, instanceCount, 1), + .b = agx_3d(vertexCount, instanceCount, 1), .start = firstVertex, .start_instance = firstInstance, }; @@ -3678,7 +3543,7 @@ hk_CmdDrawMultiEXT(VkCommandBuffer commandBuffer, uint32_t drawCount, for (unsigned i = 0; i < drawCount; ++i) { struct hk_draw draw = { - .b = hk_grid(pVertexInfo->vertexCount, instanceCount, 1), + .b = agx_3d(pVertexInfo->vertexCount, instanceCount, 1), .start = pVertexInfo->firstVertex, .start_instance = firstInstance, }; @@ -3706,7 +3571,7 @@ hk_draw_indexed(VkCommandBuffer commandBuffer, uint16_t draw_id, draw = hk_draw_indexed_indirect(addr, cmd->state.gfx.index.buffer, 0, 0); } else { draw = (struct hk_draw){ - .b = hk_grid(indexCount, instanceCount, 1), + .b = agx_3d(indexCount, instanceCount, 1), .indexed = true, .index = cmd->state.gfx.index.buffer, .start = firstIndex, @@ -3841,10 +3706,6 @@ hk_draw_indirect_count(VkCommandBuffer commandBuffer, VkBuffer _buffer, VK_FROM_HANDLE(hk_buffer, count_buffer, countBuffer); struct hk_device *dev = hk_cmd_buffer_device(cmd); - struct agx_predicate_indirect_key key = {.indexed = indexed}; - struct hk_shader *s = - hk_meta_kernel(dev, agx_nir_predicate_indirect, &key, sizeof(key)); - perf_debug(dev, "Draw indirect count"); struct hk_cs *cs = @@ -3858,19 +3719,11 @@ hk_draw_indirect_count(VkCommandBuffer commandBuffer, VkBuffer _buffer, size_t out_stride = sizeof(uint32_t) * (indexed ? 5 : 4); uint64_t patched = hk_pool_alloc(cmd, out_stride * maxDrawCount, 4).gpu; + uint64_t in = hk_buffer_address(buffer, offset); + uint64_t count_addr = hk_buffer_address(count_buffer, countBufferOffset); - struct libagx_predicate_indirect_push push = { - .in = hk_buffer_address(buffer, offset), - .out = patched, - .draw_count = hk_buffer_address(count_buffer, countBufferOffset), - .stride_el = stride / 4, - }; - - uint64_t push_ = hk_pool_upload(cmd, &push, sizeof(push), 8); - uint32_t usc = hk_upload_usc_words_kernel(cmd, s, &push_, sizeof(push_)); - - hk_dispatch_with_usc(dev, cs, &s->b.info, usc, hk_grid(maxDrawCount, 1, 1), - hk_grid(1, 1, 1)); + libagx_predicate_indirect(cs, agx_1d(maxDrawCount), patched, in, count_addr, + stride / 4, indexed); if (indexed) { hk_draw_indexed_indirect_inner(commandBuffer, patched, maxDrawCount, @@ -3930,14 +3783,6 @@ hk_CmdBindTransformFeedbackBuffersEXT(VkCommandBuffer commandBuffer, } } -static void -hk_libagx_copy_xfb_counters(nir_builder *b, const void *key) -{ - b->shader->info.workgroup_size_variable = true; - - libagx_copy_xfb_counters(b, nir_load_preamble(b, 1, 64)); -} - static void hk_begin_end_xfb(VkCommandBuffer commandBuffer, uint32_t firstCounterBuffer, uint32_t counterBufferCount, const VkBuffer *pCounterBuffers, @@ -3996,14 +3841,8 @@ hk_begin_end_xfb(VkCommandBuffer commandBuffer, uint32_t firstCounterBuffer, if (copies > 0) { perf_debug(dev, "XFB counter copy"); - struct hk_shader *s = - hk_meta_kernel(dev, hk_libagx_copy_xfb_counters, NULL, 0); - - uint64_t push = hk_pool_upload(cmd, ¶ms, sizeof(params), 8); - uint32_t usc = hk_upload_usc_words_kernel(cmd, s, &push, sizeof(push)); - - hk_dispatch_with_usc(dev, cs, &s->b.info, usc, hk_grid(copies, 1, 1), - hk_grid(copies, 1, 1)); + libagx_copy_xfb_counters(cs, agx_1d(copies), + hk_pool_upload(cmd, ¶ms, sizeof(params), 8)); } } diff --git a/src/asahi/vulkan/hk_device.c b/src/asahi/vulkan/hk_device.c index b647eb73683..15066abf148 100644 --- a/src/asahi/vulkan/hk_device.c +++ b/src/asahi/vulkan/hk_device.c @@ -62,7 +62,7 @@ hk_upload_rodata(struct hk_device *dev) uint32_t offs = 0; offs = align(offs, 8); - agx_pack(&dev->rodata.txf_sampler, USC_SAMPLER, cfg) { + agx_pack(&dev->dev.txf_sampler, USC_SAMPLER, cfg) { cfg.start = 0; cfg.count = 1; cfg.buffer = dev->rodata.bo->va->addr + offs; diff --git a/src/asahi/vulkan/hk_device.h b/src/asahi/vulkan/hk_device.h index 08545826d09..1faf4709f8d 100644 --- a/src/asahi/vulkan/hk_device.h +++ b/src/asahi/vulkan/hk_device.h @@ -83,7 +83,6 @@ struct hk_device { struct { struct agx_bo *bo; - struct agx_usc_sampler_packed txf_sampler; struct agx_usc_uniform_packed image_heap; uint64_t null_sink, zero_sink; uint64_t geometry_state; diff --git a/src/asahi/vulkan/hk_query_pool.c b/src/asahi/vulkan/hk_query_pool.c index 08ecc564428..290c8e758ca 100644 --- a/src/asahi/vulkan/hk_query_pool.c +++ b/src/asahi/vulkan/hk_query_pool.c @@ -14,6 +14,7 @@ #include "hk_entrypoints.h" #include "hk_shader.h" +#include "libagx_shaders.h" #include "vk_common_entrypoints.h" #include "asahi/lib/agx_bo.h" @@ -186,32 +187,6 @@ hk_query_report_map(struct hk_device *dev, struct hk_query_pool *pool, } } -struct hk_write_params { - uint64_t address; - uint32_t value; -}; - -static void -hk_nir_write_u32(nir_builder *b, UNUSED const void *key) -{ - nir_def *addr = nir_load_preamble( - b, 1, 64, .base = offsetof(struct hk_write_params, address) / 2); - - nir_def *value = nir_load_preamble( - b, 1, 32, .base = offsetof(struct hk_write_params, value) / 2); - - nir_store_global(b, addr, 4, value, nir_component_mask(1)); -} - -static void -hk_nir_write_u32s(nir_builder *b, const void *data) -{ - nir_def *params = nir_load_preamble(b, 1, 64, .base = 0); - nir_def *id = nir_channel(b, nir_load_global_invocation_id(b, 32), 0); - - libagx_write_u32s(b, params, id); -} - void hk_dispatch_imm_writes(struct hk_cmd_buffer *cmd, struct hk_cs *cs) { @@ -229,17 +204,14 @@ hk_dispatch_imm_writes(struct hk_cmd_buffer *cmd, struct hk_cs *cs) perf_debug(dev, "Queued writes"); - struct hk_shader *s = hk_meta_kernel(dev, hk_nir_write_u32s, NULL, 0); uint64_t params = hk_pool_upload(cmd, cs->imm_writes.data, cs->imm_writes.size, 16); - uint32_t usc = hk_upload_usc_words_kernel(cmd, s, ¶ms, sizeof(params)); uint32_t count = util_dynarray_num_elements(&cs->imm_writes, struct libagx_imm_write); assert(count > 0); - hk_dispatch_with_usc(dev, cs, &s->b.info, usc, hk_grid(count, 1, 1), - hk_grid(32, 1, 1)); + libagx_write_u32s(cs, agx_1d(count), params); } void @@ -276,13 +248,7 @@ hk_queue_write(struct hk_cmd_buffer *cmd, uint64_t address, uint32_t value, hk_cdm_cache_flush(dev, cs); perf_debug(dev, "Queued write"); - - struct hk_shader *s = hk_meta_kernel(dev, hk_nir_write_u32, NULL, 0); - struct hk_write_params params = {.address = address, .value = value}; - uint32_t usc = hk_upload_usc_words_kernel(cmd, s, ¶ms, sizeof(params)); - - hk_dispatch_with_usc(dev, cs, &s->b.info, usc, hk_grid(1, 1, 1), - hk_grid(1, 1, 1)); + libagx_write_u32(cs, agx_1d(1), address, value); } /** @@ -588,13 +554,6 @@ hk_GetQueryPoolResults(VkDevice device, VkQueryPool queryPool, return status; } -static void -hk_nir_copy_query(nir_builder *b, UNUSED const void *key) -{ - nir_def *id = nir_channel(b, nir_load_workgroup_id(b), 0); - libagx_copy_query(b, nir_load_preamble(b, 1, 64), id); -} - VKAPI_ATTR void VKAPI_CALL hk_CmdCopyQueryPoolResults(VkCommandBuffer commandBuffer, VkQueryPool queryPool, uint32_t firstQuery, uint32_t queryCount, @@ -613,7 +572,7 @@ hk_CmdCopyQueryPoolResults(VkCommandBuffer commandBuffer, VkQueryPool queryPool, perf_debug(dev, "Query pool copy"); hk_ensure_cs_has_space(cmd, cs, 0x2000 /* TODO */); - const struct libagx_copy_query_push info = { + struct libagx_copy_query_args info = { .availability = pool->bo->va->addr, .results = pool->oq_queries ? dev->occlusion_queries.bo->va->addr : pool->bo->va->addr + pool->query_start, @@ -629,10 +588,5 @@ hk_CmdCopyQueryPoolResults(VkCommandBuffer commandBuffer, VkQueryPool queryPool, .with_availability = flags & VK_QUERY_RESULT_WITH_AVAILABILITY_BIT, }; - uint64_t push = hk_pool_upload(cmd, &info, sizeof(info), 8); - - struct hk_shader *s = hk_meta_kernel(dev, hk_nir_copy_query, NULL, 0); - uint32_t usc = hk_upload_usc_words_kernel(cmd, s, &push, sizeof(push)); - hk_dispatch_with_usc(dev, cs, &s->b.info, usc, hk_grid(queryCount, 1, 1), - hk_grid(1, 1, 1)); + libagx_copy_query_struct(cs, agx_1d(queryCount), info); } diff --git a/src/asahi/vulkan/hk_queue.c b/src/asahi/vulkan/hk_queue.c index 0662558089e..ba61fe1707b 100644 --- a/src/asahi/vulkan/hk_queue.c +++ b/src/asahi/vulkan/hk_queue.c @@ -9,6 +9,7 @@ */ #include "hk_queue.h" +#include "agx_bg_eot.h" #include "agx_bo.h" #include "agx_device.h" #include "agx_pack.h" @@ -89,7 +90,7 @@ asahi_fill_cdm_command(struct hk_device *dev, struct hk_cs *cs, if (cs->scratch.cs.main || cs->scratch.cs.preamble) { cmd->helper_arg = dev->scratch.cs.buf->va->addr; cmd->helper_cfg = cs->scratch.cs.preamble ? (1 << 16) : 0; - cmd->helper_program = dev->dev.helper->va->addr | 1; + cmd->helper_program = agx_helper_program(&dev->bg_eot); } } @@ -230,13 +231,13 @@ asahi_fill_vdm_command(struct hk_device *dev, struct hk_cs *cs, c->flags |= ASAHI_RENDER_VERTEX_SPILLS; c->vertex_helper_arg = dev->scratch.vs.buf->va->addr; c->vertex_helper_cfg = cs->scratch.vs.preamble ? (1 << 16) : 0; - c->vertex_helper_program = dev->dev.helper->va->addr | 1; + c->vertex_helper_program = agx_helper_program(&dev->bg_eot); } if (cs->scratch.fs.main || cs->scratch.fs.preamble) { c->fragment_helper_arg = dev->scratch.fs.buf->va->addr; c->fragment_helper_cfg = cs->scratch.fs.preamble ? (1 << 16) : 0; - c->fragment_helper_program = dev->dev.helper->va->addr | 1; + c->fragment_helper_program = agx_helper_program(&dev->bg_eot); } } diff --git a/src/asahi/vulkan/hk_shader.c b/src/asahi/vulkan/hk_shader.c index 208a63ae434..d81947506ce 100644 --- a/src/asahi/vulkan/hk_shader.c +++ b/src/asahi/vulkan/hk_shader.c @@ -1479,13 +1479,13 @@ hk_fast_link(struct hk_device *dev, bool fragment, struct hk_shader *main, struct agx_usc_builder b = agx_usc_builder(s->usc.data, sizeof(s->usc.data)); if (main && main->b.info.rodata.size_16) { - agx_usc_immediates(&b, &main->b.info, main->bo->va->addr); + agx_usc_immediates(&b, &main->b.info.rodata, main->bo->va->addr); } agx_usc_push_packed(&b, UNIFORM, dev->rodata.image_heap); if (s->b.uses_txf) - agx_usc_push_packed(&b, SAMPLER, dev->rodata.txf_sampler); + agx_usc_push_packed(&b, SAMPLER, dev->dev.txf_sampler); agx_usc_shared_non_fragment(&b, &main->b.info, 0); agx_usc_push_packed(&b, SHADER, s->b.shader); diff --git a/src/asahi/vulkan/hk_shader.h b/src/asahi/vulkan/hk_shader.h index 1a04b41e79e..da2181f0104 100644 --- a/src/asahi/vulkan/hk_shader.h +++ b/src/asahi/vulkan/hk_shader.h @@ -383,13 +383,6 @@ struct hk_api_shader *hk_meta_shader(struct hk_device *dev, hk_internal_builder_t builder, void *data, size_t data_size); -static inline struct hk_shader * -hk_meta_kernel(struct hk_device *dev, hk_internal_builder_t builder, void *data, - size_t data_size) -{ - return hk_only_variant(hk_meta_shader(dev, builder, data, data_size)); -} - struct hk_passthrough_gs_key { /* Bit mask of outputs written by the VS/TES, to be passed through */ uint64_t outputs; diff --git a/src/asahi/vulkan/meson.build b/src/asahi/vulkan/meson.build index 38801dcf275..22b7360fec6 100644 --- a/src/asahi/vulkan/meson.build +++ b/src/asahi/vulkan/meson.build @@ -70,6 +70,7 @@ hk_deps = [ idep_vulkan_wsi_headers, idep_agx_pack, idep_libagx, + idep_libagx_shaders_h, ] libhk = static_library( @@ -77,7 +78,6 @@ libhk = static_library( [ hk_entrypoints, hk_files, - libagx_shaders, sha1_h, ], include_directories : [ diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c index 28d4c50fefb..9145486c4a9 100644 --- a/src/compiler/spirv/spirv_to_nir.c +++ b/src/compiler/spirv/spirv_to_nir.c @@ -7350,8 +7350,10 @@ spirv_library_to_nir_builder(FILE *fp, const uint32_t *words, size_t word_count, fprintf(fp, "#include \"compiler/nir/nir_builder.h\"\n\n"); + nir_fixup_is_exported(b->shader); + vtn_foreach_function(func, &b->functions) { - if (func->linkage != SpvLinkageTypeExport) + if (!func->nir_func->is_exported || func->nir_func->is_entrypoint) continue; if (!func_to_nir_builder(fp, func)) diff --git a/src/gallium/drivers/asahi/agx_pipe.c b/src/gallium/drivers/asahi/agx_pipe.c index 5f03bd3ca52..4196ec43a89 100644 --- a/src/gallium/drivers/asahi/agx_pipe.c +++ b/src/gallium/drivers/asahi/agx_pipe.c @@ -43,6 +43,7 @@ #include "util/u_upload_mgr.h" #include "util/xmlconfig.h" #include "agx_bg_eot.h" +#include "agx_bo.h" #include "agx_device.h" #include "agx_disk_cache.h" #include "agx_fence.h" @@ -1519,12 +1520,12 @@ agx_cmdbuf(struct agx_device *dev, struct drm_asahi_cmd_render *c, c->flags |= ASAHI_RENDER_VERTEX_SPILLS; c->vertex_helper_arg = batch->ctx->scratch_vs.buf->va->addr; c->vertex_helper_cfg = batch->vs_preamble_scratch << 16; - c->vertex_helper_program = dev->helper->va->addr | 1; + c->vertex_helper_program = agx_helper_program(&batch->ctx->bg_eot); } if (batch->fs_scratch) { c->fragment_helper_arg = batch->ctx->scratch_fs.buf->va->addr; c->fragment_helper_cfg = batch->fs_preamble_scratch << 16; - c->fragment_helper_program = dev->helper->va->addr | 1; + c->fragment_helper_program = agx_helper_program(&batch->ctx->bg_eot); } } @@ -1646,7 +1647,7 @@ agx_flush_compute(struct agx_context *ctx, struct agx_batch *batch, cmdbuf->helper_arg = ctx->scratch_cs.buf->va->addr; cmdbuf->helper_cfg = batch->cs_preamble_scratch << 16; // cmdbuf->helper_cfg |= 0x40; - cmdbuf->helper_program = dev->helper->va->addr | 1; + cmdbuf->helper_program = agx_helper_program(&batch->ctx->bg_eot); } } @@ -2738,5 +2739,19 @@ agx_screen_create(int fd, struct renderonly *ro, agx_disk_cache_init(agx_screen); + /* TODO: Refactor readonly data? */ + { + struct agx_bo *bo = + agx_bo_create(&agx_screen->dev, 16384, 0, 0, "Rodata"); + + agx_pack_txf_sampler((struct agx_sampler_packed *)bo->map); + + agx_pack(&agx_screen->dev.txf_sampler, USC_SAMPLER, cfg) { + cfg.start = 0; + cfg.count = 1; + cfg.buffer = bo->va->addr; + } + } + return screen; } diff --git a/src/gallium/drivers/asahi/agx_query.c b/src/gallium/drivers/asahi/agx_query.c index 200aef54659..af29e6e3202 100644 --- a/src/gallium/drivers/asahi/agx_query.c +++ b/src/gallium/drivers/asahi/agx_query.c @@ -11,15 +11,11 @@ #include "util/ralloc.h" #include "util/u_dump.h" #include "util/u_inlines.h" -#include "util/u_prim.h" #include "agx_bo.h" #include "agx_device.h" #include "agx_state.h" -#include "nir.h" -#include "nir_builder.h" -#include "nir_builder_opcodes.h" -#include "pool.h" -#include "shader_enums.h" +#include "libagx.h" +#include "libagx_shaders.h" static bool is_occlusion(struct agx_query *query) @@ -471,44 +467,6 @@ agx_get_query_result_resource_cpu(struct agx_context *ctx, result_type_size(result_type), &result.u64); } -struct query_copy_key { - enum pipe_query_value_type result; - enum query_copy_type query; -}; - -static void -agx_nir_query_copy(nir_builder *b, const void *key_) -{ - const struct query_copy_key *key = key_; - b->shader->info.num_ubos = 1; - - nir_def *params = - nir_load_ubo(b, 2, 64, nir_imm_int(b, 0), nir_imm_int(b, 0), - .align_mul = 8, .range = 8); - - nir_def *value = - nir_load_global_constant(b, nir_channel(b, params, 0), 8, 1, 64); - - if (key->query == QUERY_COPY_BOOL32 || key->query == QUERY_COPY_BOOL64) { - if (key->query == QUERY_COPY_BOOL32) - value = nir_u2u32(b, value); - - value = nir_u2u64(b, nir_ine_imm(b, value, 0)); - } - - if (key->result == PIPE_QUERY_TYPE_U32) { - value = - nir_u2u32(b, nir_umin(b, value, nir_imm_int64(b, u_uintN_max(32)))); - } else if (key->result == PIPE_QUERY_TYPE_I32) { - value = - nir_u2u32(b, nir_iclamp(b, value, nir_imm_int64(b, u_intN_min(32)), - nir_imm_int64(b, u_intN_max(32)))); - } - - nir_store_global(b, nir_channel(b, params, 1), result_type_size(key->result), - value, nir_component_mask(1)); -} - static bool agx_get_query_result_resource_gpu(struct agx_context *ctx, struct agx_query *query, @@ -529,42 +487,21 @@ agx_get_query_result_resource_gpu(struct agx_context *ctx, flush_query_writers(ctx, query, util_str_query_type(query->type, true)); struct agx_resource *rsrc = agx_resource(prsrc); - - struct query_copy_key key = { - .result = result_type, - .query = classify_query_type(query->type), - }; - - struct agx_compiled_shader *cs = - agx_build_meta_shader(ctx, agx_nir_query_copy, &key, sizeof(key)); + enum query_copy_type copy_type = classify_query_type(query->type); struct agx_batch *batch = agx_get_compute_batch(ctx); agx_batch_init_state(batch); agx_dirty_all(ctx); - /* Save cb */ - struct agx_stage *stage = &ctx->stage[PIPE_SHADER_COMPUTE]; - struct pipe_constant_buffer saved_cb = {NULL}; - pipe_resource_reference(&saved_cb.buffer, stage->cb[0].buffer); - memcpy(&saved_cb, &stage->cb[0], sizeof(struct pipe_constant_buffer)); - /* Set params */ - uint64_t params[2] = {query->ptr.gpu, rsrc->bo->va->addr + offset}; agx_batch_writes_range(batch, rsrc, offset, result_type_size(result_type)); - struct pipe_constant_buffer cb = { - .buffer_size = sizeof(params), - .user_buffer = ¶ms, - }; - ctx->base.set_constant_buffer(&ctx->base, PIPE_SHADER_COMPUTE, 0, false, - &cb); + unsigned bool_size = copy_type == QUERY_COPY_BOOL64 ? 8 + : copy_type == QUERY_COPY_BOOL32 ? 4 + : 0; - struct agx_grid grid = agx_grid_direct(1, 1, 1, 1, 1, 1); - agx_launch(batch, &grid, cs, NULL, PIPE_SHADER_COMPUTE, 0); - - /* take_ownership=true so do not unreference */ - ctx->base.set_constant_buffer(&ctx->base, PIPE_SHADER_COMPUTE, 0, true, - &saved_cb); + libagx_copy_query_gl(batch, agx_1d(1), query->ptr.gpu, + rsrc->bo->va->addr + offset, result_type, bool_size); return true; } diff --git a/src/gallium/drivers/asahi/agx_state.c b/src/gallium/drivers/asahi/agx_state.c index 2bd92eb1e4f..71e139976f4 100644 --- a/src/gallium/drivers/asahi/agx_state.c +++ b/src/gallium/drivers/asahi/agx_state.c @@ -59,6 +59,9 @@ #include "agx_nir_lower_gs.h" #include "agx_nir_lower_vbo.h" #include "agx_tilebuffer.h" +#include "libagx.h" +#include "libagx_dgc.h" +#include "libagx_shaders.h" #include "nir_builder.h" #include "nir_builder_opcodes.h" #include "nir_intrinsics.h" @@ -2989,7 +2992,7 @@ agx_build_pipeline(struct agx_batch *batch, struct agx_compiled_shader *cs, } if (cs->bo) { - agx_usc_immediates(&b, &cs->b.info, cs->bo->va->addr); + agx_usc_immediates(&b, &cs->b.info.rodata, cs->bo->va->addr); } uint32_t max_scratch_size = @@ -3060,91 +3063,46 @@ agx_build_pipeline(struct agx_batch *batch, struct agx_compiled_shader *cs, return agx_usc_addr(dev, t.gpu); } -static uint32_t -agx_build_internal_usc(struct agx_batch *batch, struct agx_shader_info *info, - uint64_t addr, uint64_t data) -{ - struct agx_device *dev = agx_device(batch->ctx->base.screen); - bool needs_sampler = info->uses_txf; - size_t usc_size = agx_usc_size(12 + (needs_sampler ? 1 : 0)); - - 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); - - agx_usc_uniform(&b, 0, 4, agx_pool_upload(&batch->pool, &data, 8)); - agx_usc_immediates(&b, info, addr); - - if (needs_sampler) { - /* TODO: deduplicate */ - struct agx_ptr t = agx_pool_alloc_aligned( - &batch->pool, sizeof(struct agx_sampler_packed), 64); - - agx_pack_txf_sampler((struct agx_sampler_packed *)t.cpu); - - agx_usc_pack(&b, SAMPLER, cfg) { - cfg.start = 0; - cfg.count = 1; - cfg.buffer = t.gpu; - } - } - - assert(info->scratch_size == 0 && "internal kernels don't spill"); - assert(info->preamble_scratch_size == 0 && "internal doesn't spill"); - - unsigned local_size = 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 = agx_usc_addr(dev, addr + info->main_offset); - cfg.unk_2 = 3; - } - - agx_usc_pack(&b, REGISTERS, cfg) { - cfg.register_count = info->nr_gprs; - cfg.spill_size = 0; - } - - if (info->has_preamble) { - agx_usc_pack(&b, PRESHADER, cfg) { - cfg.code = agx_usc_addr(dev, addr + info->preamble_offset); - } - } else { - agx_usc_pack(&b, NO_PRESHADER, cfg) - ; - } - - return agx_usc_addr(dev, t.gpu); -} - static void -agx_launch_with_uploaded_data(struct agx_batch *batch, - const struct agx_grid *grid, - meta_shader_builder_t builder, void *key, - size_t key_size, uint64_t data) +agx_launch_internal(struct agx_batch *batch, struct agx_grid grid, + struct agx_workgroup wg, + struct agx_cdm_launch_word_0_packed launch, + enum pipe_shader_type stage, uint32_t usc) { - struct agx_compiled_shader *cs = agx_build_meta_shader_internal( - batch->ctx, builder, key, key_size, false, false, 0, true); + struct agx_context *ctx = batch->ctx; + struct agx_device *dev = agx_device(ctx->base.screen); - uint32_t usc = - agx_build_internal_usc(batch, &cs->b.info, cs->bo->va->addr, data); + /* TODO: Ensure space if we allow multiple kernels in a batch */ + uint32_t *out = (uint32_t *)batch->cdm.current; - agx_launch_internal(batch, grid, cs, &cs->b.info, PIPE_SHADER_COMPUTE, usc); + out = agx_cdm_launch(out, dev->chip, grid, wg, launch, usc); + out = agx_cdm_barrier(out, dev->chip); + + batch->cdm.current = (void *)out; + assert(batch->cdm.current <= batch->cdm.end && + "Failed to reserve sufficient space in encoder"); } void -agx_launch_with_data(struct agx_batch *batch, const struct agx_grid *grid, - meta_shader_builder_t builder, void *key, size_t key_size, - void *data, size_t data_size) +agx_launch_precomp(struct agx_batch *batch, struct agx_grid grid, + enum libagx_program program, void *args, size_t arg_size) { - uint64_t upload = agx_pool_upload_aligned(&batch->pool, data, data_size, 4); - agx_launch_with_uploaded_data(batch, grid, builder, key, key_size, upload); + struct agx_device *dev = agx_device(batch->ctx->base.screen); + struct agx_precompiled_shader *cs = + agx_get_precompiled(&batch->ctx->bg_eot, program); + + struct agx_ptr t = + agx_pool_alloc_aligned(&batch->pipeline_pool, agx_usc_size(15), 64); + + uint64_t uploaded_data = + agx_pool_upload_aligned(&batch->pool, args, arg_size, 4); + + uint32_t usc = agx_usc_addr(dev, t.gpu); + agx_usc_words_precomp(t.cpu, &cs->b, uploaded_data, arg_size); + + agx_batch_add_bo(batch, cs->bo); + agx_launch_internal(batch, grid, cs->b.workgroup, cs->b.launch, + PIPE_SHADER_COMPUTE, usc); } struct asahi_bg_eot @@ -3935,33 +3893,27 @@ agx_ia_update(struct agx_batch *batch, const struct pipe_draw_info *info, struct agx_context *ctx = batch->ctx; struct agx_device *dev = agx_device(ctx->base.screen); - struct agx_increment_ia_counters_key key = { - .index_size_B = info->primitive_restart ? info->index_size : 0, - }; - - struct libagx_increment_ia_counters args = { - .ia_vertices = agx_get_query_address( - batch, ctx->pipeline_statistics[PIPE_STAT_QUERY_IA_VERTICES]), - - .vs_invocations = agx_get_query_address( - batch, ctx->pipeline_statistics[PIPE_STAT_QUERY_VS_INVOCATIONS]), - - .restart_index = info->restart_index, - .index_buffer = ib, - .index_buffer_range_el = ib_range_el, - .draw = draw, - }; - - uint64_t wg_size = key.index_size_B ? 1024 : 1; - struct agx_grid grid = agx_grid_direct(wg_size, 1, 1, wg_size, 1, 1); - if (!batch->cdm.bo) { batch->cdm = agx_encoder_allocate(batch, dev); } - perf_debug(dev, "Input assembly counters"); - agx_launch_with_data(batch, &grid, agx_nir_increment_ia_counters, &key, - sizeof(key), &args, sizeof(args)); + uint64_t ia_vertices = agx_get_query_address( + batch, ctx->pipeline_statistics[PIPE_STAT_QUERY_IA_VERTICES]); + + uint64_t vs_invocations = agx_get_query_address( + batch, ctx->pipeline_statistics[PIPE_STAT_QUERY_VS_INVOCATIONS]); + + if (info->primitive_restart) { + perf_debug(dev, "Input assembly counters with primitive restart"); + + libagx_increment_ia_restart(batch, agx_1d(1024), ia_vertices, + vs_invocations, draw, ib, ib_range_el, + info->restart_index, info->index_size); + } else { + perf_debug(dev, "Input assembly counters"); + + libagx_increment_ia(batch, agx_1d(1), ia_vertices, vs_invocations, draw); + } } static uint64_t @@ -4056,12 +4008,11 @@ agx_batch_geometry_params(struct agx_batch *batch, uint64_t input_index_buffer, /* Calculate input primitive count for direct draws, and allocate the vertex * & count buffers. GPU calculates and allocates for indirect draws. */ - unsigned count_buffer_stride = batch->ctx->gs->gs_count_words * 4; batch->uniforms.vertex_outputs = batch->ctx->vs->b.info.outputs; params.input_mask = batch->uniforms.vertex_outputs; + params.count_buffer_stride = batch->ctx->gs->gs_count_words * 4; if (indirect) { - params.count_buffer_stride = count_buffer_stride; batch->uniforms.vertex_output_buffer_ptr = agx_pool_alloc_aligned(&batch->pool, 8, 8).gpu; @@ -4077,7 +4028,7 @@ agx_batch_geometry_params(struct agx_batch *batch, uint64_t input_index_buffer, unsigned vb_size = libagx_tcs_in_size(draw->count * info->instance_count, batch->uniforms.vertex_outputs); - unsigned size = params.input_primitives * count_buffer_stride; + unsigned size = params.input_primitives * params.count_buffer_stride; if (size) { params.count_buffer = @@ -4139,14 +4090,12 @@ agx_launch_gs_prerast(struct agx_batch *batch, assert(!info->primitive_restart && "should have been lowered"); + uint64_t gp = batch->uniforms.geometry_params; struct agx_grid grid_vs, grid_gs; + struct agx_workgroup wg; /* Setup grids */ if (indirect) { - struct agx_gs_setup_indirect_key key = { - .prim = info->mode, - }; - uint64_t ib = 0; size_t ib_extent = 0; @@ -4155,61 +4104,53 @@ agx_launch_gs_prerast(struct agx_batch *batch, &ib_extent); } - struct agx_gs_setup_indirect_params gsi = { + struct libagx_gs_setup_indirect_args gsi = { .index_buffer = ib, .index_buffer_range_el = ib_extent / info->index_size, .draw = agx_indirect_buffer_ptr(batch, indirect), .vertex_buffer = batch->uniforms.vertex_output_buffer_ptr, .ia = batch->uniforms.input_assembly, - .geom = batch->uniforms.geometry_params, + .p = batch->uniforms.geometry_params, .vs_outputs = batch->uniforms.vertex_outputs, .index_size_B = info->index_size, + .prim = info->mode, }; - const struct agx_grid grid_setup = agx_grid_direct(1, 1, 1, 1, 1, 1); - agx_launch_with_data(batch, &grid_setup, agx_nir_gs_setup_indirect, &key, - sizeof(key), &gsi, sizeof(gsi)); + libagx_gs_setup_indirect_struct(batch, agx_1d(1), gsi); - uint64_t gp = batch->uniforms.geometry_params; - - grid_vs = agx_grid_indirect( - gp + offsetof(struct agx_geometry_params, vs_grid), 1, 1, 1); - - grid_gs = agx_grid_indirect( - gp + offsetof(struct agx_geometry_params, gs_grid), 1, 1, 1); - } else { + wg = agx_workgroup(1, 1, 1); grid_vs = - agx_grid_direct(draws->count, info->instance_count, 1, 64, 1, 1); + agx_grid_indirect(gp + offsetof(struct agx_geometry_params, vs_grid)); - grid_gs = agx_grid_direct( - u_decomposed_prims_for_vertices(info->mode, draws->count), - info->instance_count, 1, 64, 1, 1); + grid_gs = + agx_grid_indirect(gp + offsetof(struct agx_geometry_params, gs_grid)); + } else { + wg = agx_workgroup(64, 1, 1); + grid_vs = agx_3d(draws->count, info->instance_count, 1); + + grid_gs = + agx_3d(u_decomposed_prims_for_vertices(info->mode, draws->count), + info->instance_count, 1); } /* Launch the vertex shader first */ - agx_launch(batch, &grid_vs, ctx->vs, ctx->linked.vs, ctx->vs->stage, 0); + agx_launch(batch, grid_vs, wg, ctx->vs, ctx->linked.vs, ctx->vs->stage, 0); /* If there is a count shader, launch it and prefix sum the results. */ if (gs->gs_count) { perf_debug(dev, "Geometry shader count"); - agx_launch(batch, &grid_gs, gs->gs_count, NULL, PIPE_SHADER_GEOMETRY, 0); + agx_launch(batch, grid_gs, wg, gs->gs_count, NULL, PIPE_SHADER_GEOMETRY, + 0); - unsigned words = gs->gs_count_words; - struct agx_grid grid = - agx_grid_direct(1024 * gs->gs_count_words, 1, 1, 1024, 1, 1); - - agx_launch(batch, &grid, - agx_build_meta_shader(ctx, agx_nir_prefix_sum_gs, &words, - sizeof(words)), - NULL, PIPE_SHADER_COMPUTE, 0); + libagx_prefix_sum_geom(batch, agx_1d(1024 * gs->gs_count_words), gp); } /* Pre-GS shader */ - struct agx_grid grid = agx_grid_direct(1, 1, 1, 1, 1, 1); - agx_launch(batch, &grid, gs->pre_gs, NULL, PIPE_SHADER_COMPUTE, 0); + agx_launch(batch, agx_1d(1), agx_workgroup(1, 1, 1), gs->pre_gs, NULL, + PIPE_SHADER_COMPUTE, 0); /* Pre-rast geometry shader */ - agx_launch(batch, &grid_gs, gs, NULL, PIPE_SHADER_GEOMETRY, 0); + agx_launch(batch, grid_gs, wg, gs, NULL, PIPE_SHADER_GEOMETRY, 0); } static void @@ -4255,33 +4196,26 @@ agx_draw_without_restart(struct agx_batch *batch, if (!batch->cdm.bo) batch->cdm = agx_encoder_allocate(batch, dev); - struct agx_unroll_restart_key key = { - .prim = info->mode, - }; - /* Allocate output indirect draw descriptors. This is exact. */ struct agx_resource out_draws_rsrc = {0}; struct agx_ptr out_draws = agx_pool_alloc_aligned_with_bo( &batch->pool, 5 * sizeof(uint32_t) * indirect->draw_count, 4, &out_draws_rsrc.bo); - struct agx_restart_unroll_params unroll = { + struct libagx_unroll_restart_args unroll = { .heap = agx_batch_geometry_state(batch), .index_buffer = ib, - .out_draws = out_draws.gpu, + .out_draw = out_draws.gpu, .restart_index = info->restart_index, .index_buffer_size_el = ib_extent / info->index_size, .flatshade_first = batch->ctx->rast->base.flatshade_first, - .draws = agx_indirect_buffer_ptr(batch, indirect), - .index_size_B = info->index_size, + .in_draw = agx_indirect_buffer_ptr(batch, indirect), }; /* Unroll the index buffer for each draw */ - const struct agx_grid grid_setup = - agx_grid_direct(1024 * indirect->draw_count, 1, 1, 1024, 1, 1); - - agx_launch_with_data(batch, &grid_setup, agx_nir_unroll_restart, &key, - sizeof(key), &unroll, sizeof(unroll)); + libagx_unroll_restart_struct(batch, agx_1d(1024 * indirect->draw_count), + unroll, util_logbase2(info->index_size), + libagx_compact_prim(info->mode)); /* Now draw the results without restart */ struct pipe_draw_info new_info = { @@ -4642,7 +4576,6 @@ agx_draw_patches(struct agx_context *ctx, const struct pipe_draw_info *info, sizeof(ctx->default_inner_level)); struct agx_grid vs_grid, tcs_grid, tess_grid; - unsigned tess_wg_size = 64; agx_upload_vbos(batch); agx_update_vs(batch, info->index_size); @@ -4696,34 +4629,12 @@ agx_draw_patches(struct agx_context *ctx, const struct pipe_draw_info *info, batch->uniforms.vertex_output_buffer_ptr = agx_pool_upload(&batch->pool, &addr, 8); - vs_grid = - agx_grid_direct(draws->count, info->instance_count, 1, 64, 1, 1); + vs_grid = agx_3d(draws->count, info->instance_count, 1); + tcs_grid = agx_3d(in_patches * tcs->tess.output_patch_size, + info->instance_count, 1); - tcs_grid = agx_grid_direct(in_patches * tcs->tess.output_patch_size, - info->instance_count, 1, - tcs->tess.output_patch_size, 1, 1); - - tess_grid = agx_grid_direct(unrolled_patches, 1, 1, tess_wg_size, 1, 1); + tess_grid = agx_1d(unrolled_patches); } else if (indirect) { - args.tcs_statistic = agx_get_query_address( - batch, ctx->pipeline_statistics[PIPE_STAT_QUERY_HS_INVOCATIONS]); - - args.indirect = agx_indirect_buffer_ptr(batch, indirect); - - /* Allocate 3x indirect global+local grids for VS/TCS/tess */ - uint32_t grid_stride = sizeof(uint32_t) * 6; - args.grids = agx_pool_alloc_aligned(&batch->pool, grid_stride * 3, 4).gpu; - - vs_grid = agx_grid_indirect_local(args.grids + 0 * grid_stride); - tcs_grid = agx_grid_indirect_local(args.grids + 1 * grid_stride); - tess_grid = agx_grid_indirect_local(args.grids + 2 * grid_stride); - - args.vertex_outputs = ctx->vs->b.info.outputs; - args.vertex_output_buffer_ptr = - agx_pool_alloc_aligned(&batch->pool, 8, 8).gpu; - - batch->uniforms.vertex_output_buffer_ptr = args.vertex_output_buffer_ptr; - args.out_draws = agx_pool_alloc_aligned_with_bo(&batch->pool, draw_stride, 4, &draw_bo) .gpu; @@ -4733,43 +4644,45 @@ agx_draw_patches(struct agx_context *ctx, const struct pipe_draw_info *info, agx_pool_upload_aligned(&batch->pool, &args, sizeof(args), 4); if (indirect) { - const struct agx_grid indirect_grid = agx_grid_direct(1, 1, 1, 1, 1, 1); - struct agx_tess_setup_indirect_key indirect_key = { - .point_mode = point_mode, - }; + uint32_t grid_stride = sizeof(uint32_t) * 6; - agx_launch_with_uploaded_data(batch, &indirect_grid, - agx_nir_tess_setup_indirect, &indirect_key, - sizeof(indirect_key), state); + uint64_t vertex_out_ptr = agx_pool_alloc_aligned(&batch->pool, 8, 8).gpu; + uint64_t indirect_ptr = agx_indirect_buffer_ptr(batch, indirect); + + uint64_t tcs_statistic = agx_get_query_address( + batch, ctx->pipeline_statistics[PIPE_STAT_QUERY_HS_INVOCATIONS]); + + /* Allocate 3x indirect global+local grids for VS/TCS/tess */ + uint64_t grids = + agx_pool_alloc_aligned(&batch->pool, grid_stride * 3, 4).gpu; + + libagx_tess_setup_indirect(batch, agx_1d(1), state, grids, + 0 /* XXX: IA */, indirect_ptr, vertex_out_ptr, + 0, 0, 0 /* XXX: Index buffer */, + ctx->vs->b.info.outputs, tcs_statistic); + + batch->uniforms.vertex_output_buffer_ptr = vertex_out_ptr; + + vs_grid = agx_grid_indirect_local(grids + 0 * grid_stride); + tcs_grid = agx_grid_indirect_local(grids + 1 * grid_stride); + tess_grid = agx_grid_indirect_local(grids + 2 * grid_stride); } batch->uniforms.tess_params = state; - agx_launch(batch, &vs_grid, ctx->vs, ctx->linked.vs, PIPE_SHADER_VERTEX, 0); - agx_launch(batch, &tcs_grid, ctx->tcs, NULL, PIPE_SHADER_TESS_CTRL, 0); + agx_launch(batch, vs_grid, agx_workgroup(64, 1, 1), ctx->vs, ctx->linked.vs, + PIPE_SHADER_VERTEX, 0); + + agx_launch(batch, tcs_grid, agx_workgroup(tcs->tess.output_patch_size, 1, 1), + ctx->tcs, NULL, PIPE_SHADER_TESS_CTRL, 0); + batch->uniforms.vertex_output_buffer_ptr = 0; - struct agx_tessellator_key key = { - .prim = mode, - }; - - /* Generate counts */ - key.mode = LIBAGX_TESS_MODE_COUNT; - agx_launch_with_uploaded_data(batch, &tess_grid, agx_nir_tessellate, &key, - sizeof(key), state); - - /* Prefix sum counts, allocating index buffer space. */ - const struct agx_grid prefix_sum_grid = - agx_grid_direct(1024, 1, 1, 1024, 1, 1); - - agx_launch_with_uploaded_data(batch, &prefix_sum_grid, - agx_nir_prefix_sum_tess, NULL, 0, state); - - key.mode = LIBAGX_TESS_MODE_WITH_COUNTS; - - /* Now we can tessellate */ - agx_launch_with_uploaded_data(batch, &tess_grid, agx_nir_tessellate, &key, - sizeof(key), state); + /* Generate counts, then prefix sum them, then finally tessellate. */ + libagx_tessellate(batch, tess_grid, mode, LIBAGX_TESS_MODE_COUNT, state); + libagx_prefix_sum_tess(batch, agx_1d(1024), state); + libagx_tessellate(batch, tess_grid, mode, LIBAGX_TESS_MODE_WITH_COUNTS, + state); /* Face culling state needs to be specialized for tess */ ctx->dirty |= AGX_DIRTY_RS; @@ -4960,7 +4873,6 @@ agx_draw_vbo(struct pipe_context *pctx, const struct pipe_draw_info *info, } struct agx_batch *batch = agx_get_batch(ctx); - uint64_t ib = 0; size_t ib_extent = 0; @@ -5302,113 +5214,10 @@ agx_texture_barrier(struct pipe_context *pipe, unsigned flags) } void -agx_launch_internal(struct agx_batch *batch, const struct agx_grid *grid, - struct agx_compiled_shader *cs, - const struct agx_shader_info *info, - enum pipe_shader_type stage, uint32_t usc) -{ - struct agx_context *ctx = batch->ctx; - struct agx_device *dev = agx_device(ctx->base.screen); - - if (cs) - agx_batch_add_bo(batch, cs->bo); - - /* TODO: Ensure space if we allow multiple kernels in a batch */ - uint8_t *out = batch->cdm.current; - - agx_push(out, CDM_LAUNCH_WORD_0, cfg) { - cfg.mode = grid->mode; - cfg.uniform_register_count = info->push_count; - cfg.preshader_register_count = info->nr_preamble_gprs; - cfg.texture_state_register_count = - cs ? agx_nr_tex_descriptors(batch, cs) : 0; - cfg.sampler_state_register_count = - translate_sampler_state_count(ctx, stage); - } - - agx_push(out, CDM_LAUNCH_WORD_1, cfg) { - cfg.pipeline = usc; - } - - /* Added in G14X */ - if (dev->params.gpu_generation >= 14 && dev->params.num_clusters_total > 1) { - agx_push(out, CDM_UNK_G14X, cfg) - ; - } - - if (grid->mode == AGX_CDM_MODE_DIRECT) { - agx_push(out, CDM_GLOBAL_SIZE, cfg) { - cfg.x = grid->global[0]; - cfg.y = grid->global[1]; - cfg.z = grid->global[2]; - } - } else { - agx_push(out, CDM_INDIRECT, cfg) { - cfg.address_hi = grid->indirect >> 32; - cfg.address_lo = grid->indirect & BITFIELD64_MASK(32); - } - } - - if (grid->mode != AGX_CDM_MODE_INDIRECT_LOCAL) { - agx_push(out, CDM_LOCAL_SIZE, cfg) { - cfg.x = grid->local[0]; - cfg.y = grid->local[1]; - cfg.z = grid->local[2]; - } - } - - agx_push(out, CDM_BARRIER, cfg) { - cfg.unk_5 = true; - cfg.unk_6 = true; - cfg.unk_8 = true; - // cfg.unk_11 = true; - // cfg.unk_20 = true; - if (dev->params.num_clusters_total > 1) { - // cfg.unk_24 = true; - if (dev->params.gpu_generation == 13) { - cfg.unk_4 = true; - // cfg.unk_26 = true; - } - } - - /* With multiple launches in the same CDM stream, we can get cache - * coherency (? or sync?) issues. We hit this with blits, which need - in - * between dispatches - need the PBE cache to be flushed and the texture - * cache to be invalidated. Until we know what bits mean what exactly, - * let's just set these after every launch to be safe. We can revisit in - * the future when we figure out what the bits mean. - */ - cfg.unk_0 = true; - cfg.unk_1 = true; - cfg.unk_2 = true; - cfg.usc_cache_inval = true; - cfg.unk_4 = true; - cfg.unk_5 = true; - cfg.unk_6 = true; - cfg.unk_7 = true; - cfg.unk_8 = true; - cfg.unk_9 = true; - cfg.unk_10 = true; - cfg.unk_11 = true; - cfg.unk_12 = true; - cfg.unk_13 = true; - cfg.unk_14 = true; - cfg.unk_15 = true; - cfg.unk_16 = true; - cfg.unk_17 = true; - cfg.unk_18 = true; - cfg.unk_19 = true; - } - - batch->cdm.current = out; - assert(batch->cdm.current <= batch->cdm.end && - "Failed to reserve sufficient space in encoder"); -} - -void -agx_launch(struct agx_batch *batch, const struct agx_grid *grid, - struct agx_compiled_shader *cs, struct agx_linked_shader *linked, - enum pipe_shader_type stage, unsigned variable_shared_mem) +agx_launch(struct agx_batch *batch, struct agx_grid grid, + struct agx_workgroup wg, struct agx_compiled_shader *cs, + struct agx_linked_shader *linked, enum pipe_shader_type stage, + unsigned variable_shared_mem) { struct agx_context *ctx = batch->ctx; @@ -5416,17 +5225,17 @@ agx_launch(struct agx_batch *batch, const struct agx_grid *grid, * available in GPU memory. This is either the indirect buffer, or just a * buffer we upload ourselves if not indirect. */ - if (grid->mode == AGX_CDM_MODE_DIRECT) { + if (grid.mode == AGX_CDM_MODE_DIRECT) { uint32_t groups[3] = { - grid->global[0] / grid->local[0], - grid->global[1] / grid->local[1], - grid->global[2] / grid->local[2], + grid.count[0] / wg.x, + grid.count[1] / wg.y, + grid.count[2] / wg.z, }; batch->uniforms.tables[AGX_SYSVAL_TABLE_GRID] = agx_pool_upload_aligned(&batch->pool, groups, sizeof(groups), 4); } else { - batch->uniforms.tables[AGX_SYSVAL_TABLE_GRID] = grid->indirect; + batch->uniforms.tables[AGX_SYSVAL_TABLE_GRID] = grid.ptr; } util_dynarray_foreach(&ctx->global_buffers, struct pipe_resource *, res) { @@ -5457,7 +5266,20 @@ agx_launch(struct agx_batch *batch, const struct agx_grid *grid, uint32_t usc = agx_build_pipeline(batch, cs, linked, PIPE_SHADER_COMPUTE, variable_shared_mem, subgroups_per_core); - agx_launch_internal(batch, grid, cs, &cs->b.info, stage, usc); + if (cs) + agx_batch_add_bo(batch, cs->bo); + + struct agx_cdm_launch_word_0_packed launch; + agx_pack(&launch, CDM_LAUNCH_WORD_0, cfg) { + cfg.uniform_register_count = cs->b.info.push_count; + cfg.preshader_register_count = cs->b.info.nr_preamble_gprs; + cfg.texture_state_register_count = + cs ? agx_nr_tex_descriptors(batch, cs) : 0; + cfg.sampler_state_register_count = + translate_sampler_state_count(ctx, stage); + } + + agx_launch_internal(batch, grid, wg, launch, stage, usc); } static void @@ -5481,25 +5303,23 @@ agx_launch_grid(struct pipe_context *pipe, const struct pipe_grid_info *info) * * TODO: Can we use the hardware counter for this? */ - if (ctx->pipeline_statistics[PIPE_STAT_QUERY_CS_INVOCATIONS]) { - unsigned blocksize = info->block[0] * info->block[1] * info->block[2]; + struct agx_query *statistic = + ctx->pipeline_statistics[PIPE_STAT_QUERY_CS_INVOCATIONS]; + struct agx_workgroup wg = + agx_workgroup(info->block[0], info->block[1], info->block[2]); + + if (statistic) { if (indirect) { - struct libagx_cs_invocation_params p = { - .grid = indirect, - .local_size_threads = blocksize, - .statistic = agx_get_query_address( - batch, ctx->pipeline_statistics[PIPE_STAT_QUERY_CS_INVOCATIONS]), - }; + uint64_t addr = agx_get_query_address(batch, statistic); - const struct agx_grid g = agx_grid_direct(1, 1, 1, 1, 1, 1); - agx_launch_with_data(batch, &g, agx_nir_increment_cs_invocations, NULL, - 0, &p, sizeof(p)); + libagx_increment_cs_invocations(batch, agx_1d(1), indirect, addr, + agx_workgroup_threads(wg)); } else { agx_query_increment_cpu( - ctx, ctx->pipeline_statistics[PIPE_STAT_QUERY_CS_INVOCATIONS], - libagx_cs_invocations(blocksize, info->grid[0], info->grid[1], - info->grid[2])); + ctx, statistic, + libagx_cs_invocations(agx_workgroup_threads(wg), info->grid[0], + info->grid[1], info->grid[2])); } } @@ -5514,25 +5334,19 @@ agx_launch_grid(struct pipe_context *pipe, const struct pipe_grid_info *info) struct agx_compiled_shader *cs = _mesa_hash_table_next_entry(uncompiled->variants, NULL)->data; - struct agx_grid grid = { - .local[0] = info->block[0], - .local[1] = info->block[1], - .local[2] = info->block[2], - }; - + struct agx_grid grid; if (indirect) { - grid.mode = AGX_CDM_MODE_INDIRECT_GLOBAL; - grid.indirect = indirect; + grid = agx_grid_indirect(indirect); } else { - grid.mode = AGX_CDM_MODE_DIRECT; + grid = agx_3d(0, 0, 0); for (unsigned d = 0; d < 3; ++d) { - grid.global[d] = ((info->grid[d] - 1) * info->block[d]) + - (info->last_block[d] ?: info->block[d]); + grid.count[d] = ((info->grid[d] - 1) * info->block[d]) + + (info->last_block[d] ?: info->block[d]); } } - agx_launch(batch, &grid, cs, NULL, PIPE_SHADER_COMPUTE, + agx_launch(batch, grid, wg, cs, NULL, PIPE_SHADER_COMPUTE, info->variable_shared_mem); /* TODO: Dirty tracking? */ @@ -5612,30 +5426,24 @@ agx_decompress_inplace(struct agx_batch *batch, struct pipe_surface *surf, if (!batch->cdm.bo) batch->cdm = agx_encoder_allocate(batch, dev); - struct agx_ptr data = agx_pool_alloc_aligned( - &batch->pool, sizeof(struct libagx_decompress_push), 64); - struct libagx_decompress_push *push = data.cpu; - agx_fill_decompress_push(push, layout, surf->u.tex.first_layer, level, - agx_map_texture_gpu(rsrc, 0)); + struct agx_ptr images = agx_pool_alloc_aligned( + &batch->pool, sizeof(struct libagx_decompress_images), 64); + struct libagx_decompress_images *img = images.cpu; struct pipe_sampler_view sampler_view = sampler_view_for_surface(surf); sampler_view.target = PIPE_TEXTURE_2D_ARRAY; struct pipe_image_view view = image_view_for_surface(surf); - agx_pack_texture(&push->compressed, rsrc, surf->format, &sampler_view); - agx_batch_upload_pbe(batch, &push->uncompressed, &view, false, true, true, + agx_pack_texture(&img->compressed, rsrc, surf->format, &sampler_view); + agx_batch_upload_pbe(batch, &img->uncompressed, &view, false, true, true, true); - struct agx_grid grid = agx_grid_direct( - ail_metadata_width_tl(layout, level) * 32, - ail_metadata_height_tl(layout, level), - surf->u.tex.last_layer - surf->u.tex.first_layer + 1, 32, 1, 1); + struct agx_grid grid = + agx_3d(ail_metadata_width_tl(layout, level) * 32, + ail_metadata_height_tl(layout, level), + surf->u.tex.last_layer - surf->u.tex.first_layer + 1); - struct agx_decompress_key key = { - .nr_samples = layout->sample_count_sa, - }; - - agx_launch_with_uploaded_data(batch, &grid, agx_nir_decompress, &key, - sizeof(key), data.gpu); + libagx_decompress(batch, grid, layout, surf->u.tex.first_layer, level, + agx_map_texture_gpu(rsrc, 0), images.gpu); } void diff --git a/src/gallium/drivers/asahi/agx_state.h b/src/gallium/drivers/asahi/agx_state.h index acbc5c0fddf..4a7c255dfe5 100644 --- a/src/gallium/drivers/asahi/agx_state.h +++ b/src/gallium/drivers/asahi/agx_state.h @@ -790,69 +790,17 @@ struct agx_compiled_shader *agx_build_meta_shader(struct agx_context *ctx, meta_shader_builder_t builder, void *data, size_t data_size); -struct agx_grid { - /* Tag for the union */ - enum agx_cdm_mode mode; - - /* If mode != INDIRECT_LOCAL, the local size */ - uint32_t local[3]; - - union { - /* If mode == DIRECT, the global size. This is *not* multiplied by the - * local size, differing from the API definition but matching AGX. - */ - uint32_t global[3]; - - /* Address of the indirect buffer if mode != DIRECT */ - uint64_t indirect; - }; -}; - -static inline const struct agx_grid -agx_grid_direct(uint32_t global_x, uint32_t global_y, uint32_t global_z, - uint32_t local_x, uint32_t local_y, uint32_t local_z) -{ - return (struct agx_grid){ - .mode = AGX_CDM_MODE_DIRECT, - .global = {global_x, global_y, global_z}, - .local = {local_x, local_y, local_z}, - }; -} - -static inline const struct agx_grid -agx_grid_indirect(uint64_t indirect, uint32_t local_x, uint32_t local_y, - uint32_t local_z) -{ - return (struct agx_grid){ - .mode = AGX_CDM_MODE_INDIRECT_GLOBAL, - .local = {local_x, local_y, local_z}, - .indirect = indirect, - }; -} - -static inline const struct agx_grid -agx_grid_indirect_local(uint64_t indirect) -{ - return (struct agx_grid){ - .mode = AGX_CDM_MODE_INDIRECT_LOCAL, - .indirect = indirect, - }; -} - -void agx_launch_with_data(struct agx_batch *batch, const struct agx_grid *grid, - 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 agx_grid *grid, - struct agx_compiled_shader *cs, - const struct agx_shader_info *info, - enum pipe_shader_type stage, uint32_t usc); - -void agx_launch(struct agx_batch *batch, const struct agx_grid *grid, - struct agx_compiled_shader *cs, +void agx_launch(struct agx_batch *batch, struct agx_grid grid, + struct agx_workgroup wg, struct agx_compiled_shader *cs, struct agx_linked_shader *linked, enum pipe_shader_type stage, unsigned variable_shared_mem); +void agx_launch_precomp(struct agx_batch *batch, struct agx_grid grid, + enum libagx_program program, void *args, + size_t arg_size); + +#define MESA_DISPATCH_PRECOMP agx_launch_precomp + void agx_init_query_functions(struct pipe_context *ctx); void diff --git a/src/gallium/drivers/asahi/meson.build b/src/gallium/drivers/asahi/meson.build index 7ce039f5c99..efa317c216d 100644 --- a/src/gallium/drivers/asahi/meson.build +++ b/src/gallium/drivers/asahi/meson.build @@ -23,7 +23,7 @@ libasahi = static_library( include_directories : [inc_gallium_aux, inc_gallium, inc_include, inc_src, inc_asahi, inc_virtio_gpu, inc_virtio_vdrm], c_args : [c_msvc_compat_args], gnu_symbol_visibility : 'hidden', - dependencies : [idep_nir, idep_mesautil, idep_agx_pack, dep_libdrm, idep_mesaclc, idep_libagx], + dependencies : [idep_nir, idep_mesautil, idep_agx_pack, dep_libdrm, idep_mesaclc, idep_libagx, idep_libagx_shaders_h], ) driver_asahi = declare_dependency(