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

aco: Add Mesh and Task shader stages. 

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13580>
This commit is contained in:
Timur Kristóf 2021-10-21 11:33:10 +02:00 committed by Marge Bot
parent 7ae425c5d4
commit 6766e6a985
4 changed files with 41 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/amd/compiler/aco_instruction_selection.cpp
View File

@ -4761,7 +4761,7 @@ void
visit_store_output(isel_context* ctx, nir_intrinsic_instr* instr)
{
if (ctx->stage == vertex_vs || ctx->stage == tess_eval_vs || ctx->stage == fragment_fs ||
ctx->stage == vertex_ngg || ctx->stage == tess_eval_ngg ||
ctx->stage == vertex_ngg || ctx->stage == tess_eval_ngg || ctx->stage == mesh_ngg ||
(ctx->stage == vertex_tess_control_hs && ctx->shader->info.stage == MESA_SHADER_VERTEX) ||
ctx->shader->info.stage == MESA_SHADER_GEOMETRY) {
bool stored_to_temps = store_output_to_temps(ctx, instr);

27
src/amd/compiler/aco_instruction_selection_setup.cpp
View File

@ -328,6 +328,16 @@ setup_tes_variables(isel_context* ctx, nir_shader* nir)
}
}
void
setup_ms_variables(isel_context* ctx, nir_shader* nir)
{
setup_vs_output_info(ctx, nir, &ctx->program->info->ms.outinfo);
ctx->program->config->lds_size =
DIV_ROUND_UP(nir->info.shared_size, ctx->program->dev.lds_encoding_granule);
assert((ctx->program->config->lds_size * ctx->program->dev.lds_encoding_granule) < (32 * 1024));
}
void
setup_variables(isel_context* ctx, nir_shader* nir)
{
@ -335,7 +345,8 @@ setup_variables(isel_context* ctx, nir_shader* nir)
case MESA_SHADER_FRAGMENT: {
break;
}
case MESA_SHADER_COMPUTE: {
case MESA_SHADER_COMPUTE:
case MESA_SHADER_TASK: {
ctx->program->config->lds_size =
DIV_ROUND_UP(nir->info.shared_size, ctx->program->dev.lds_encoding_granule);
break;
@ -355,6 +366,10 @@ setup_variables(isel_context* ctx, nir_shader* nir)
setup_tes_variables(ctx, nir);
break;
}
case MESA_SHADER_MESH: {
setup_ms_variables(ctx, nir);
break;
}
default: unreachable("Unhandled shader stage.");
}
@ -835,6 +850,8 @@ setup_isel_context(Program* program, unsigned shader_count, struct nir_shader* c
break;
case MESA_SHADER_FRAGMENT: sw_stage = sw_stage | SWStage::FS; break;
case MESA_SHADER_COMPUTE: sw_stage = sw_stage | SWStage::CS; break;
case MESA_SHADER_TASK: sw_stage = sw_stage | SWStage::TS; break;
case MESA_SHADER_MESH: sw_stage = sw_stage | SWStage::MS; break;
default: unreachable("Shader stage not implemented");
}
}
@ -855,6 +872,10 @@ setup_isel_context(Program* program, unsigned shader_count, struct nir_shader* c
hw_stage = HWStage::CS;
else if (sw_stage == SWStage::GSCopy)
hw_stage = HWStage::VS;
else if (sw_stage == SWStage::TS)
hw_stage = HWStage::CS; /* Task shaders are implemented with compute shaders. */
else if (sw_stage == SWStage::MS)
hw_stage = HWStage::NGG; /* Mesh shaders only work on NGG and on GFX10.3+. */
else if (sw_stage == SWStage::VS_GS && gfx9_plus && !ngg)
hw_stage = HWStage::GS; /* GFX6-9: VS+GS merged into a GS (and GFX10/legacy) */
else if (sw_stage == SWStage::VS_GS && ngg)
@ -890,6 +911,10 @@ setup_isel_context(Program* program, unsigned shader_count, struct nir_shader* c
program->workgroup_size = program->info->workgroup_size;
assert(program->workgroup_size);
/* Mesh shading only works on GFX10.3+. */
ASSERTED bool mesh_shading = ctx.stage.has(SWStage::TS) || ctx.stage.has(SWStage::MS);
assert(!mesh_shading || ctx.program->chip_class >= GFX10_3);
if (ctx.stage == tess_control_hs)
setup_tcs_info(&ctx, shaders[0], NULL);
else if (ctx.stage == vertex_tess_control_hs)

15
src/amd/compiler/aco_ir.h
View File

@ -1904,7 +1904,7 @@ struct Block {
/*
* Shader stages as provided in Vulkan by the application. Contrast this to HWStage.
*/
enum class SWStage : uint8_t {
enum class SWStage : uint16_t {
None = 0,
VS = 1 << 0, /* Vertex Shader */
GS = 1 << 1, /* Geometry Shader */
@ -1912,7 +1912,9 @@ enum class SWStage : uint8_t {
TES = 1 << 3, /* Tessellation Evaluation aka Domain Shader */
FS = 1 << 4, /* Fragment aka Pixel Shader */
CS = 1 << 5, /* Compute Shader */
GSCopy = 1 << 6, /* GS Copy Shader (internal) */
TS = 1 << 6, /* Task Shader */
MS = 1 << 7, /* Mesh Shader */
GSCopy = 1 << 8, /* GS Copy Shader (internal) */
/* Stage combinations merged to run on a single HWStage */
VS_GS = VS | GS,
@ -1923,7 +1925,7 @@ enum class SWStage : uint8_t {
constexpr SWStage
operator|(SWStage a, SWStage b)
{
return static_cast<SWStage>(static_cast<uint8_t>(a) | static_cast<uint8_t>(b));
return static_cast<SWStage>(static_cast<uint16_t>(a) | static_cast<uint16_t>(b));
}
/*
@ -1956,10 +1958,10 @@ struct Stage {
/* Check if the given SWStage is included */
constexpr bool has(SWStage stage) const
{
return (static_cast<uint8_t>(sw) & static_cast<uint8_t>(stage));
return (static_cast<uint16_t>(sw) & static_cast<uint16_t>(stage));
}
unsigned num_sw_stages() const { return util_bitcount(static_cast<uint8_t>(sw)); }
unsigned num_sw_stages() const { return util_bitcount(static_cast<uint16_t>(sw)); }
constexpr bool operator==(const Stage& other) const { return sw == other.sw && hw == other.hw; }
@ -1978,6 +1980,9 @@ static constexpr Stage fragment_fs(HWStage::FS, SWStage::FS);
static constexpr Stage compute_cs(HWStage::CS, SWStage::CS);
static constexpr Stage tess_eval_vs(HWStage::VS, SWStage::TES);
static constexpr Stage gs_copy_vs(HWStage::VS, SWStage::GSCopy);
/* Mesh shading pipeline */
static constexpr Stage task_cs(HWStage::CS, SWStage::TS);
static constexpr Stage mesh_ngg(HWStage::NGG, SWStage::MS);
/* GFX10/NGG */
static constexpr Stage vertex_ngg(HWStage::NGG, SWStage::VS);
static constexpr Stage vertex_geometry_ngg(HWStage::NGG, SWStage::VS_GS);

4
src/amd/compiler/aco_print_ir.cpp
View File

@ -795,6 +795,10 @@ print_stage(Stage stage, FILE* output)
fprintf(output, "vertex_geometry_ngg");
else if (stage == tess_eval_geometry_ngg)
fprintf(output, "tess_eval_geometry_ngg");
else if (stage == mesh_ngg)
fprintf(output, "mesh_ngg");
else if (stage == task_cs)
fprintf(output, "task_cs");
else
fprintf(output, "unknown");