radv: implement PS epilogs
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com> Reviewed-by: Timur Kristóf <timur.kristof@gmail.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17485>
This commit is contained in:
parent
270cc39648
commit
5ee5c73d2d
|
@ -658,6 +658,8 @@ RADV driver environment variables
|
||||||
force all allocated buffers to be referenced in submissions
|
force all allocated buffers to be referenced in submissions
|
||||||
``checkir``
|
``checkir``
|
||||||
validate the LLVM IR before LLVM compiles the shader
|
validate the LLVM IR before LLVM compiles the shader
|
||||||
|
``epilogs``
|
||||||
|
dump fragment shader epilogs
|
||||||
``forcecompress``
|
``forcecompress``
|
||||||
Enables DCC,FMASK,CMASK,HTILE in situations where the driver supports it
|
Enables DCC,FMASK,CMASK,HTILE in situations where the driver supports it
|
||||||
but normally does not deem it beneficial.
|
but normally does not deem it beneficial.
|
||||||
|
|
|
@ -66,6 +66,7 @@ enum {
|
||||||
RADV_DEBUG_DUMP_PROLOGS = 1ull << 35,
|
RADV_DEBUG_DUMP_PROLOGS = 1ull << 35,
|
||||||
RADV_DEBUG_NO_DMA_BLIT = 1ull << 36,
|
RADV_DEBUG_NO_DMA_BLIT = 1ull << 36,
|
||||||
RADV_DEBUG_SPLIT_FMA = 1ull << 37,
|
RADV_DEBUG_SPLIT_FMA = 1ull << 37,
|
||||||
|
RADV_DEBUG_DUMP_EPILOGS = 1ull << 38,
|
||||||
};
|
};
|
||||||
|
|
||||||
enum {
|
enum {
|
||||||
|
|
|
@ -967,6 +967,7 @@ static const struct debug_control radv_debug_options[] = {
|
||||||
{"nonggc", RADV_DEBUG_NO_NGGC},
|
{"nonggc", RADV_DEBUG_NO_NGGC},
|
||||||
{"prologs", RADV_DEBUG_DUMP_PROLOGS},
|
{"prologs", RADV_DEBUG_DUMP_PROLOGS},
|
||||||
{"nodma", RADV_DEBUG_NO_DMA_BLIT},
|
{"nodma", RADV_DEBUG_NO_DMA_BLIT},
|
||||||
|
{"epilogs", RADV_DEBUG_DUMP_EPILOGS},
|
||||||
{NULL, 0}};
|
{NULL, 0}};
|
||||||
|
|
||||||
const char *
|
const char *
|
||||||
|
|
|
@ -2444,6 +2444,55 @@ radv_create_vs_prolog(struct radv_device *device, const struct radv_vs_prolog_ke
|
||||||
return prolog;
|
return prolog;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
struct radv_shader_part *
|
||||||
|
radv_create_ps_epilog(struct radv_device *device, const struct radv_ps_epilog_key *key)
|
||||||
|
{
|
||||||
|
struct radv_shader_args args = {0};
|
||||||
|
struct radv_nir_compiler_options options = {0};
|
||||||
|
options.family = device->physical_device->rad_info.family;
|
||||||
|
options.gfx_level = device->physical_device->rad_info.gfx_level;
|
||||||
|
options.address32_hi = device->physical_device->rad_info.address32_hi;
|
||||||
|
options.dump_shader = device->instance->debug_flags & RADV_DEBUG_DUMP_EPILOGS;
|
||||||
|
options.record_ir = device->instance->debug_flags & RADV_DEBUG_HANG;
|
||||||
|
options.dump_preoptir = device->instance->debug_flags & RADV_DEBUG_DUMP_EPILOGS;
|
||||||
|
options.dump_shader = device->instance->debug_flags & RADV_DEBUG_DUMP_EPILOGS;
|
||||||
|
|
||||||
|
struct radv_shader_info info = {0};
|
||||||
|
info.wave_size = key->wave32 ? 32 : 64;
|
||||||
|
info.workgroup_size = 64;
|
||||||
|
|
||||||
|
radv_declare_ps_epilog_args(device->physical_device->rad_info.gfx_level, key, &args);
|
||||||
|
|
||||||
|
#ifdef LLVM_AVAILABLE
|
||||||
|
if (options.dump_shader || options.record_ir)
|
||||||
|
ac_init_llvm_once();
|
||||||
|
#endif
|
||||||
|
|
||||||
|
struct radv_shader_part_binary *binary = NULL;
|
||||||
|
struct aco_shader_info ac_info;
|
||||||
|
struct aco_ps_epilog_key ac_key;
|
||||||
|
struct aco_compiler_options ac_opts;
|
||||||
|
radv_aco_convert_shader_info(&ac_info, &info);
|
||||||
|
radv_aco_convert_opts(&ac_opts, &options);
|
||||||
|
radv_aco_convert_ps_epilog_key(&ac_key, key);
|
||||||
|
aco_compile_ps_epilog(&ac_opts, &ac_info, &ac_key, &args, &radv_aco_build_shader_part,
|
||||||
|
(void **)&binary);
|
||||||
|
struct radv_shader_part *epilog = upload_shader_part(device, binary, info.wave_size);
|
||||||
|
if (epilog) {
|
||||||
|
epilog->disasm_string =
|
||||||
|
binary->disasm_size ? strdup((const char *)(binary->data + binary->code_size)) : NULL;
|
||||||
|
}
|
||||||
|
|
||||||
|
free(binary);
|
||||||
|
|
||||||
|
if (epilog && options.dump_shader) {
|
||||||
|
fprintf(stderr, "Fragment epilog");
|
||||||
|
fprintf(stderr, "\ndisasm:\n%s\n", epilog->disasm_string);
|
||||||
|
}
|
||||||
|
|
||||||
|
return epilog;
|
||||||
|
}
|
||||||
|
|
||||||
void
|
void
|
||||||
radv_shader_destroy(struct radv_device *device, struct radv_shader *shader)
|
radv_shader_destroy(struct radv_device *device, struct radv_shader *shader)
|
||||||
{
|
{
|
||||||
|
|
|
@ -596,6 +596,9 @@ void radv_trap_handler_shader_destroy(struct radv_device *device,
|
||||||
struct radv_shader_part *radv_create_vs_prolog(struct radv_device *device,
|
struct radv_shader_part *radv_create_vs_prolog(struct radv_device *device,
|
||||||
const struct radv_vs_prolog_key *key);
|
const struct radv_vs_prolog_key *key);
|
||||||
|
|
||||||
|
struct radv_shader_part *radv_create_ps_epilog(struct radv_device *device,
|
||||||
|
const struct radv_ps_epilog_key *key);
|
||||||
|
|
||||||
void radv_shader_destroy(struct radv_device *device, struct radv_shader *shader);
|
void radv_shader_destroy(struct radv_device *device, struct radv_shader *shader);
|
||||||
|
|
||||||
void radv_shader_part_destroy(struct radv_device *device, struct radv_shader_part *shader_part);
|
void radv_shader_part_destroy(struct radv_device *device, struct radv_shader_part *shader_part);
|
||||||
|
|
|
@ -450,6 +450,12 @@ declare_ps_input_vgprs(const struct radv_shader_info *info, struct radv_shader_a
|
||||||
vgpr_arg++;
|
vgpr_arg++;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
if (info->ps.has_epilog) {
|
||||||
|
/* FIXME: Ensure the main shader doesn't have less VGPRs than the epilog */
|
||||||
|
for (unsigned i = 0; i < MAX_RTS; i++)
|
||||||
|
ac_add_arg(&args->ac, AC_ARG_VGPR, 4, AC_ARG_INT, NULL);
|
||||||
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
static void
|
static void
|
||||||
|
|
Loading…
Reference in New Issue