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

asahi,vtn: precompile kernels

Browse Source 
switch libagx to the precompilation pipeline. see the big comment in the
previous commit for why we're doing this.

while doing so, we move some dispatch stuff. there was so much churn from
precompile that this avoids doing the churn twice. that new header will be used
for DGC down the road.

there's also a small vtn/bindgen patch in here to skip bindgen'ing entrypoints,
as that conflicts with the new dispatch macros. this is the sane behaviour, we
just need to do the full precomp switch across the tree at once.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32339>
This commit is contained in:
Alyssa Rosenzweig
2024-11-19 20:26:43 -04:00
committed by Marge Bot
parent e3001352ad
commit f4a3ba5302
46 changed files with 1523 additions and 1852 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/asahi/clc/asahi_clc.c
+162 -66
View File
@@ -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 <fcntl.h>
#include <inttypes.h>
#include <stdio.h>
#include <string.h>
#include <unistd.h>
#include "util/u_math.h"
#include "util/macros.h"
#include <sys/mman.h>
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 <stdint.h>\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;
}
src/asahi/compiler/agx_compile.c
+6 -3
View File
@@ -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
src/asahi/compiler/agx_compile.h
+50 -14
View File
@@ -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;
src/asahi/lib/agx_bg_eot.c
+109
View File
@@ -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;
}
src/asahi/lib/agx_bg_eot.h
+23
View File
@@ -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,
src/asahi/lib/agx_device.c
+18 -3
View File
@@ -15,6 +15,7 @@
#include "agx_scratch.h"
#include "decode.h"
#include "glsl_types.h"
#include "libagx_dgc.h"
#include "libagx_shaders.h"
#include <fcntl.h>
@@ -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);
src/asahi/lib/agx_device.h
+13 -2
View File
@@ -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
src/asahi/lib/agx_helpers.h
+23 -15
View File
@@ -8,22 +8,13 @@
#include <stdbool.h>
#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],
src/asahi/lib/agx_nir_lower_gs.c
-125
View File
@@ -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);
}
src/asahi/lib/agx_nir_lower_gs.h
-59
View File
@@ -7,7 +7,6 @@
#include <stdbool.h>
#include <stdint.h>
#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);
src/asahi/lib/agx_scratch.c
-15
View File
@@ -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)
{
src/asahi/lib/agx_scratch.h
-2
View File
@@ -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);
src/asahi/lib/agx_usc.h
+27 -91
View File
@@ -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);
}
}
src/asahi/lib/meson.build
+8 -3
View File
@@ -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],
src/asahi/libagx/compression.cl
+16 -9
View File
@@ -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;
}
}
src/asahi/libagx/compression.h
+2 -7
View File
@@ -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);
src/asahi/libagx/draws.cl
+9 -8
View File
@@ -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) {
src/asahi/libagx/draws.h
-14
View File
@@ -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;
};
src/asahi/libagx/geometry.cl
+117 -69
View File
@@ -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.
src/asahi/libagx/geometry.h
+22 -72
View File
@@ -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
src/asahi/libagx/helper.cl
+1 -1
View File
@@ -17,7 +17,7 @@ enum helper_op {
OP_END = 15,
};
void
KERNEL(1)
libagx_helper(void)
{
uint64_t arg =
src/asahi/libagx/libagx.h
+10 -4
View File
@@ -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)
{
src/asahi/libagx/libagx_dgc.h
+362
View File
@@ -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);
}
src/asahi/libagx/meson.build
+1 -1
View File
@@ -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 = []
src/asahi/libagx/query.cl
+68 -65
View File
@@ -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;
}
src/asahi/libagx/query.h
-48
View File
@@ -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;
src/asahi/libagx/tessellation.cl
+39 -31
View File
@@ -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;
}
src/asahi/libagx/tessellator.cl
+13 -7
View File
@@ -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);
src/asahi/libagx/tessellator.h
+1 -34
View File
@@ -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);
src/asahi/vulkan/hk_cmd_buffer.c
+14 -31
View File
@@ -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 */
src/asahi/vulkan/hk_cmd_buffer.h
+14 -64
View File
@@ -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);
src/asahi/vulkan/hk_cmd_dispatch.c
+41 -131
View File
@@ -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, &params, 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));
}
src/asahi/vulkan/hk_cmd_draw.c
+130 -291
View File
@@ -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, &params, 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, &params, 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, &params, sizeof(params), 8));
}
}
src/asahi/vulkan/hk_device.c
+1 -1
View File
@@ -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;
src/asahi/vulkan/hk_device.h
-1
View File
@@ -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;
src/asahi/vulkan/hk_query_pool.c
+5 -51
View File
@@ -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, &params, 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, &params, 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);
}
src/asahi/vulkan/hk_queue.c
+4 -3
View File
@@ -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);
}
}
src/asahi/vulkan/hk_shader.c
+2 -2
View File
@@ -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);
src/asahi/vulkan/hk_shader.h
-7
View File
@@ -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;
src/asahi/vulkan/meson.build
+1 -1
View File
@@ -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 : [
src/compiler/spirv/spirv_to_nir.c
+3 -1
View File
@@ -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))
src/gallium/drivers/asahi/agx_pipe.c
+18 -3
View File
@@ -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;
}
src/gallium/drivers/asahi/agx_query.c
+8 -71
View File
@@ -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 = &params,
};
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;
}
src/gallium/drivers/asahi/agx_state.c
+173 -365
View File
@@ -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
src/gallium/drivers/asahi/agx_state.h
+8 -60
View File
@@ -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
src/gallium/drivers/asahi/meson.build
+1 -1
View File
@@ -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(