spirv: Implement non-Multiview parts of SPV_NV_mesh_shader
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10600>
This commit is contained in:
parent
10a03e30cf
commit
b34f9740ca
|
@ -70,6 +70,7 @@ struct spirv_supported_capabilities {
|
|||
bool kernel_image;
|
||||
bool kernel_image_read_write;
|
||||
bool literal_sampler;
|
||||
bool mesh_shading_nv;
|
||||
bool min_lod;
|
||||
bool multiview;
|
||||
bool physical_storage_buffer_address;
|
||||
|
|
|
@ -61,6 +61,10 @@ stage_to_enum(char *stage)
|
|||
return MESA_SHADER_COMPUTE;
|
||||
else if (!strcmp(stage, "kernel"))
|
||||
return MESA_SHADER_KERNEL;
|
||||
else if (!strcmp(stage, "task"))
|
||||
return MESA_SHADER_TASK;
|
||||
else if (!strcmp(stage, "mesh"))
|
||||
return MESA_SHADER_MESH;
|
||||
else
|
||||
return MESA_SHADER_NONE;
|
||||
}
|
||||
|
@ -74,7 +78,7 @@ print_usage(char *exec_name, FILE *f)
|
|||
" -h --help Print this help.\n"
|
||||
" -s, --stage <stage> Specify the shader stage. Valid stages are:\n"
|
||||
" vertex, tess-ctrl, tess-eval, geometry, fragment,\n"
|
||||
" compute, and kernel (OpenCL-style compute).\n"
|
||||
" task, mesh, compute, and kernel (OpenCL-style compute).\n"
|
||||
" -e, --entry <name> Specify the entry-point name.\n"
|
||||
" -g, --opengl Use OpenGL environment instead of Vulkan for\n"
|
||||
" graphics stages.\n"
|
||||
|
|
|
@ -1084,6 +1084,8 @@ struct_member_decoration_cb(struct vtn_builder *b,
|
|||
break;
|
||||
|
||||
case SpvDecorationPatch:
|
||||
case SpvDecorationPerPrimitiveNV:
|
||||
case SpvDecorationPerTaskNV:
|
||||
break;
|
||||
|
||||
case SpvDecorationSpecId:
|
||||
|
@ -1128,6 +1130,11 @@ struct_member_decoration_cb(struct vtn_builder *b,
|
|||
/* User semantic decorations can safely be ignored by the driver. */
|
||||
break;
|
||||
|
||||
case SpvDecorationPerViewNV:
|
||||
/* TODO(mesh): Handle multiview. */
|
||||
vtn_warn("Mesh multiview not yet supported. Needed for decoration PerViewNV.");
|
||||
break;
|
||||
|
||||
default:
|
||||
vtn_fail_with_decoration("Unhandled decoration", dec->decoration);
|
||||
}
|
||||
|
@ -2216,8 +2223,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
|
|||
}
|
||||
|
||||
/* Now that we have the value, update the workgroup size if needed */
|
||||
if (b->entry_point_stage == MESA_SHADER_COMPUTE ||
|
||||
b->entry_point_stage == MESA_SHADER_KERNEL)
|
||||
if (gl_shader_stage_uses_workgroup(b->entry_point_stage))
|
||||
vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb,
|
||||
NULL);
|
||||
}
|
||||
|
@ -4154,8 +4160,12 @@ vtn_handle_barrier(struct vtn_builder *b, SpvOp opcode,
|
|||
* variables performed by any invocation executed prior to a
|
||||
* OpControlBarrier will be visible to any other invocation after
|
||||
* return from that OpControlBarrier."
|
||||
*
|
||||
* The same applies to VK_NV_mesh_shader.
|
||||
*/
|
||||
if (b->nb.shader->info.stage == MESA_SHADER_TESS_CTRL) {
|
||||
if (b->nb.shader->info.stage == MESA_SHADER_TESS_CTRL ||
|
||||
b->nb.shader->info.stage == MESA_SHADER_TASK ||
|
||||
b->nb.shader->info.stage == MESA_SHADER_MESH) {
|
||||
memory_semantics &= ~(SpvMemorySemanticsAcquireMask |
|
||||
SpvMemorySemanticsReleaseMask |
|
||||
SpvMemorySemanticsAcquireReleaseMask |
|
||||
|
@ -4190,10 +4200,12 @@ gl_primitive_from_spv_execution_mode(struct vtn_builder *b,
|
|||
case SpvExecutionModeOutputPoints:
|
||||
return 0; /* GL_POINTS */
|
||||
case SpvExecutionModeInputLines:
|
||||
case SpvExecutionModeOutputLinesNV:
|
||||
return 1; /* GL_LINES */
|
||||
case SpvExecutionModeInputLinesAdjacency:
|
||||
return 0x000A; /* GL_LINE_STRIP_ADJACENCY_ARB */
|
||||
case SpvExecutionModeTriangles:
|
||||
case SpvExecutionModeOutputTrianglesNV:
|
||||
return 4; /* GL_TRIANGLES */
|
||||
case SpvExecutionModeInputTrianglesAdjacency:
|
||||
return 0x000C; /* GL_TRIANGLES_ADJACENCY_ARB */
|
||||
|
@ -4262,6 +4274,10 @@ stage_for_execution_model(struct vtn_builder *b, SpvExecutionModel model)
|
|||
return MESA_SHADER_INTERSECTION;
|
||||
case SpvExecutionModelCallableKHR:
|
||||
return MESA_SHADER_CALLABLE;
|
||||
case SpvExecutionModelTaskNV:
|
||||
return MESA_SHADER_TASK;
|
||||
case SpvExecutionModelMeshNV:
|
||||
return MESA_SHADER_MESH;
|
||||
default:
|
||||
vtn_fail("Unsupported execution model: %s (%u)",
|
||||
spirv_executionmodel_to_string(model), model);
|
||||
|
@ -4695,6 +4711,10 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
|
|||
spv_check_supported(float64_atomic_min_max, cap);
|
||||
break;
|
||||
|
||||
case SpvCapabilityMeshShadingNV:
|
||||
spv_check_supported(mesh_shading_nv, cap);
|
||||
break;
|
||||
|
||||
default:
|
||||
vtn_fail("Unhandled capability: %s (%u)",
|
||||
spirv_capability_to_string(cap), cap);
|
||||
|
@ -4867,19 +4887,32 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
|
|||
break;
|
||||
|
||||
case SpvExecutionModeLocalSize:
|
||||
vtn_assert(gl_shader_stage_is_compute(b->shader->info.stage));
|
||||
b->shader->info.workgroup_size[0] = mode->operands[0];
|
||||
b->shader->info.workgroup_size[1] = mode->operands[1];
|
||||
b->shader->info.workgroup_size[2] = mode->operands[2];
|
||||
if (gl_shader_stage_uses_workgroup(b->shader->info.stage)) {
|
||||
b->shader->info.workgroup_size[0] = mode->operands[0];
|
||||
b->shader->info.workgroup_size[1] = mode->operands[1];
|
||||
b->shader->info.workgroup_size[2] = mode->operands[2];
|
||||
} else {
|
||||
vtn_fail("Execution mode LocalSize not supported in stage %s",
|
||||
_mesa_shader_stage_to_string(b->shader->info.stage));
|
||||
}
|
||||
break;
|
||||
|
||||
case SpvExecutionModeOutputVertices:
|
||||
if (b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
|
||||
b->shader->info.stage == MESA_SHADER_TESS_EVAL) {
|
||||
switch (b->shader->info.stage) {
|
||||
case MESA_SHADER_TESS_CTRL:
|
||||
case MESA_SHADER_TESS_EVAL:
|
||||
b->shader->info.tess.tcs_vertices_out = mode->operands[0];
|
||||
} else {
|
||||
vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
|
||||
break;
|
||||
case MESA_SHADER_GEOMETRY:
|
||||
b->shader->info.gs.vertices_out = mode->operands[0];
|
||||
break;
|
||||
case MESA_SHADER_MESH:
|
||||
b->shader->info.mesh.max_vertices_out = mode->operands[0];
|
||||
break;
|
||||
default:
|
||||
vtn_fail("Execution mode OutputVertices not supported in stage %s",
|
||||
_mesa_shader_stage_to_string(b->shader->info.stage));
|
||||
break;
|
||||
}
|
||||
break;
|
||||
|
||||
|
@ -4903,7 +4936,37 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
|
|||
}
|
||||
break;
|
||||
|
||||
case SpvExecutionModeOutputPoints:
|
||||
case SpvExecutionModeOutputPrimitivesNV:
|
||||
vtn_assert(b->shader->info.stage == MESA_SHADER_MESH);
|
||||
b->shader->info.mesh.max_primitives_out = mode->operands[0];
|
||||
break;
|
||||
|
||||
case SpvExecutionModeOutputLinesNV:
|
||||
case SpvExecutionModeOutputTrianglesNV:
|
||||
vtn_assert(b->shader->info.stage == MESA_SHADER_MESH);
|
||||
b->shader->info.mesh.primitive_type =
|
||||
gl_primitive_from_spv_execution_mode(b, mode->exec_mode);
|
||||
break;
|
||||
|
||||
case SpvExecutionModeOutputPoints: {
|
||||
const unsigned primitive =
|
||||
gl_primitive_from_spv_execution_mode(b, mode->exec_mode);
|
||||
|
||||
switch (b->shader->info.stage) {
|
||||
case MESA_SHADER_GEOMETRY:
|
||||
b->shader->info.gs.output_primitive = primitive;
|
||||
break;
|
||||
case MESA_SHADER_MESH:
|
||||
b->shader->info.mesh.primitive_type = primitive;
|
||||
break;
|
||||
default:
|
||||
vtn_fail("Execution mode OutputPoints not supported in stage %s",
|
||||
_mesa_shader_stage_to_string(b->shader->info.stage));
|
||||
break;
|
||||
}
|
||||
break;
|
||||
}
|
||||
|
||||
case SpvExecutionModeOutputLineStrip:
|
||||
case SpvExecutionModeOutputTriangleStrip:
|
||||
vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
|
||||
|
@ -5087,9 +5150,14 @@ vtn_handle_execution_mode_id(struct vtn_builder *b, struct vtn_value *entry_poin
|
|||
|
||||
switch (mode->exec_mode) {
|
||||
case SpvExecutionModeLocalSizeId:
|
||||
b->shader->info.workgroup_size[0] = vtn_constant_uint(b, mode->operands[0]);
|
||||
b->shader->info.workgroup_size[1] = vtn_constant_uint(b, mode->operands[1]);
|
||||
b->shader->info.workgroup_size[2] = vtn_constant_uint(b, mode->operands[2]);
|
||||
if (gl_shader_stage_uses_workgroup(b->shader->info.stage)) {
|
||||
b->shader->info.workgroup_size[0] = vtn_constant_uint(b, mode->operands[0]);
|
||||
b->shader->info.workgroup_size[1] = vtn_constant_uint(b, mode->operands[1]);
|
||||
b->shader->info.workgroup_size[2] = vtn_constant_uint(b, mode->operands[2]);
|
||||
} else {
|
||||
vtn_fail("Execution mode LocalSizeId not supported in stage %s",
|
||||
_mesa_shader_stage_to_string(b->shader->info.stage));
|
||||
}
|
||||
break;
|
||||
|
||||
case SpvExecutionModeLocalSizeHintId:
|
||||
|
@ -5393,6 +5461,58 @@ vtn_handle_ray_intrinsic(struct vtn_builder *b, SpvOp opcode,
|
|||
}
|
||||
}
|
||||
|
||||
static void
|
||||
vtn_handle_write_packed_primitive_indices(struct vtn_builder *b, SpvOp opcode,
|
||||
const uint32_t *w, unsigned count)
|
||||
{
|
||||
vtn_assert(opcode == SpvOpWritePackedPrimitiveIndices4x8NV);
|
||||
|
||||
/* TODO(mesh): Use or create a primitive that allow the unpacking to
|
||||
* happen in the backend. What we have here is functional but too
|
||||
* blunt.
|
||||
*/
|
||||
|
||||
struct vtn_type *offset_type = vtn_get_value_type(b, w[1]);
|
||||
vtn_fail_if(offset_type->base_type != vtn_base_type_scalar ||
|
||||
offset_type->type != glsl_uint_type(),
|
||||
"Index Offset type of OpWritePackedPrimitiveIndices4x8NV "
|
||||
"must be an OpTypeInt with 32-bit Width and 0 Signedness.");
|
||||
|
||||
struct vtn_type *packed_type = vtn_get_value_type(b, w[2]);
|
||||
vtn_fail_if(packed_type->base_type != vtn_base_type_scalar ||
|
||||
packed_type->type != glsl_uint_type(),
|
||||
"Packed Indices type of OpWritePackedPrimitiveIndices4x8NV "
|
||||
"must be an OpTypeInt with 32-bit Width and 0 Signedness.");
|
||||
|
||||
nir_deref_instr *indices = NULL;
|
||||
nir_foreach_variable_with_modes(var, b->nb.shader, nir_var_shader_out) {
|
||||
if (var->data.location == VARYING_SLOT_PRIMITIVE_INDICES) {
|
||||
indices = nir_build_deref_var(&b->nb, var);
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
/* TODO(mesh): It may be the case that the variable is not present in the
|
||||
* entry point interface list.
|
||||
*
|
||||
* See https://github.com/KhronosGroup/SPIRV-Registry/issues/104.
|
||||
*/
|
||||
vtn_fail_if(indices == NULL,
|
||||
"Missing output variable decorated with PrimitiveIndices builtin.");
|
||||
|
||||
nir_ssa_def *offset = vtn_get_nir_ssa(b, w[1]);
|
||||
nir_ssa_def *packed = vtn_get_nir_ssa(b, w[2]);
|
||||
nir_ssa_def *unpacked = nir_unpack_bits(&b->nb, packed, 8);
|
||||
for (int i = 0; i < 4; i++) {
|
||||
nir_deref_instr *offset_deref =
|
||||
nir_build_deref_array(&b->nb, indices,
|
||||
nir_iadd_imm(&b->nb, offset, i));
|
||||
nir_ssa_def *val = nir_u2u(&b->nb, nir_channel(&b->nb, unpacked, i), 32);
|
||||
|
||||
nir_store_deref(&b->nb, offset_deref, val, 0x1);
|
||||
}
|
||||
}
|
||||
|
||||
static bool
|
||||
vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
|
||||
const uint32_t *w, unsigned count)
|
||||
|
@ -5831,6 +5951,10 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
|
|||
vtn_handle_opencl_core_instruction(b, opcode, w, count);
|
||||
break;
|
||||
|
||||
case SpvOpWritePackedPrimitiveIndices4x8NV:
|
||||
vtn_handle_write_packed_primitive_indices(b, opcode, w, count);
|
||||
break;
|
||||
|
||||
default:
|
||||
vtn_fail_with_opcode("Unhandled opcode", opcode);
|
||||
}
|
||||
|
|
|
@ -787,15 +787,18 @@ vtn_get_builtin_location(struct vtn_builder *b,
|
|||
{
|
||||
switch (builtin) {
|
||||
case SpvBuiltInPosition:
|
||||
case SpvBuiltInPositionPerViewNV:
|
||||
*location = VARYING_SLOT_POS;
|
||||
break;
|
||||
case SpvBuiltInPointSize:
|
||||
*location = VARYING_SLOT_PSIZ;
|
||||
break;
|
||||
case SpvBuiltInClipDistance:
|
||||
*location = VARYING_SLOT_CLIP_DIST0; /* XXX CLIP_DIST1? */
|
||||
case SpvBuiltInClipDistancePerViewNV:
|
||||
*location = VARYING_SLOT_CLIP_DIST0;
|
||||
break;
|
||||
case SpvBuiltInCullDistance:
|
||||
case SpvBuiltInCullDistancePerViewNV:
|
||||
*location = VARYING_SLOT_CULL_DIST0;
|
||||
break;
|
||||
case SpvBuiltInVertexId:
|
||||
|
@ -840,7 +843,8 @@ vtn_get_builtin_location(struct vtn_builder *b,
|
|||
*mode = nir_var_shader_out;
|
||||
else if (b->options && b->options->caps.shader_viewport_index_layer &&
|
||||
(b->shader->info.stage == MESA_SHADER_VERTEX ||
|
||||
b->shader->info.stage == MESA_SHADER_TESS_EVAL))
|
||||
b->shader->info.stage == MESA_SHADER_TESS_EVAL ||
|
||||
b->shader->info.stage == MESA_SHADER_MESH))
|
||||
*mode = nir_var_shader_out;
|
||||
else
|
||||
vtn_fail("invalid stage for SpvBuiltInLayer");
|
||||
|
@ -851,7 +855,8 @@ vtn_get_builtin_location(struct vtn_builder *b,
|
|||
*mode = nir_var_shader_out;
|
||||
else if (b->options && b->options->caps.shader_viewport_index_layer &&
|
||||
(b->shader->info.stage == MESA_SHADER_VERTEX ||
|
||||
b->shader->info.stage == MESA_SHADER_TESS_EVAL))
|
||||
b->shader->info.stage == MESA_SHADER_TESS_EVAL ||
|
||||
b->shader->info.stage == MESA_SHADER_MESH))
|
||||
*mode = nir_var_shader_out;
|
||||
else if (b->shader->info.stage == MESA_SHADER_FRAGMENT)
|
||||
*mode = nir_var_shader_in;
|
||||
|
@ -1123,6 +1128,15 @@ vtn_get_builtin_location(struct vtn_builder *b,
|
|||
vtn_fail("invalid stage for SpvBuiltInPrimitiveShadingRateKHR");
|
||||
}
|
||||
break;
|
||||
case SpvBuiltInPrimitiveCountNV:
|
||||
*location = VARYING_SLOT_PRIMITIVE_COUNT;
|
||||
break;
|
||||
case SpvBuiltInPrimitiveIndicesNV:
|
||||
*location = VARYING_SLOT_PRIMITIVE_INDICES;
|
||||
break;
|
||||
case SpvBuiltInTaskCountNV:
|
||||
*location = VARYING_SLOT_TASK_COUNT;
|
||||
break;
|
||||
default:
|
||||
vtn_fail("Unsupported builtin: %s (%u)",
|
||||
spirv_builtin_to_string(builtin), builtin);
|
||||
|
@ -1276,18 +1290,64 @@ apply_var_decoration(struct vtn_builder *b,
|
|||
/* TODO: We should actually plumb alias information through NIR. */
|
||||
break;
|
||||
|
||||
case SpvDecorationPerPrimitiveNV:
|
||||
vtn_fail_if(
|
||||
!(b->shader->info.stage == MESA_SHADER_MESH && var_data->mode == nir_var_shader_out) &&
|
||||
!(b->shader->info.stage == MESA_SHADER_FRAGMENT && var_data->mode == nir_var_shader_in),
|
||||
"PerPrimitiveNV decoration only allowed for Mesh shader outputs or Fragment shader inputs");
|
||||
var_data->per_primitive = true;
|
||||
break;
|
||||
|
||||
case SpvDecorationPerTaskNV:
|
||||
vtn_fail_if(
|
||||
!(b->shader->info.stage == MESA_SHADER_TASK && var_data->mode == nir_var_shader_out) &&
|
||||
!(b->shader->info.stage == MESA_SHADER_MESH && var_data->mode == nir_var_shader_in),
|
||||
"PerTaskNV decoration only allowed for Task shader outputs or Mesh shader inputs");
|
||||
/* Don't set anything, because this decoration is implied by being a
|
||||
* non-builtin Task Output or Mesh Input.
|
||||
*/
|
||||
break;
|
||||
|
||||
case SpvDecorationPerViewNV:
|
||||
vtn_fail_if(b->shader->info.stage != MESA_SHADER_MESH,
|
||||
"PerViewNV decoration only allowed in Mesh shaders");
|
||||
var_data->per_view = true;
|
||||
break;
|
||||
|
||||
default:
|
||||
vtn_fail_with_decoration("Unhandled decoration", dec->decoration);
|
||||
}
|
||||
}
|
||||
|
||||
static void
|
||||
var_is_patch_cb(struct vtn_builder *b, struct vtn_value *val, int member,
|
||||
const struct vtn_decoration *dec, void *void_var)
|
||||
gather_var_kind_cb(struct vtn_builder *b, struct vtn_value *val, int member,
|
||||
const struct vtn_decoration *dec, void *void_var)
|
||||
{
|
||||
struct vtn_variable *vtn_var = void_var;
|
||||
if (dec->decoration == SpvDecorationPatch)
|
||||
switch (dec->decoration) {
|
||||
case SpvDecorationPatch:
|
||||
vtn_var->var->data.patch = true;
|
||||
break;
|
||||
case SpvDecorationPerPrimitiveNV:
|
||||
vtn_var->var->data.per_primitive = true;
|
||||
break;
|
||||
case SpvDecorationBuiltIn:
|
||||
if (b->shader->info.stage == MESA_SHADER_MESH) {
|
||||
SpvBuiltIn builtin = dec->operands[0];
|
||||
switch (builtin) {
|
||||
case SpvBuiltInPrimitiveIndicesNV:
|
||||
vtn_var->var->data.per_primitive = true;
|
||||
break;
|
||||
default:
|
||||
/* Nothing to do. */
|
||||
break;
|
||||
}
|
||||
}
|
||||
break;
|
||||
default:
|
||||
/* Nothing to do. */
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
static void
|
||||
|
@ -1878,12 +1938,12 @@ vtn_create_variable(struct vtn_builder *b, struct vtn_value *val,
|
|||
* it to be all or nothing, we'll call it patch if any of the members
|
||||
* are declared patch.
|
||||
*/
|
||||
vtn_foreach_decoration(b, val, var_is_patch_cb, var);
|
||||
vtn_foreach_decoration(b, val, gather_var_kind_cb, var);
|
||||
if (glsl_type_is_array(var->type->type) &&
|
||||
glsl_type_is_struct_or_ifc(without_array->type)) {
|
||||
vtn_foreach_decoration(b, vtn_value(b, without_array->id,
|
||||
vtn_value_type_type),
|
||||
var_is_patch_cb, var);
|
||||
gather_var_kind_cb, var);
|
||||
}
|
||||
|
||||
struct vtn_type *per_vertex_type = var->type;
|
||||
|
@ -1935,6 +1995,17 @@ vtn_create_variable(struct vtn_builder *b, struct vtn_value *val,
|
|||
vtn_foreach_decoration(b, vtn_value(b, per_vertex_type->id,
|
||||
vtn_value_type_type),
|
||||
var_decoration_cb, var);
|
||||
|
||||
/* PerTask I/O is always a single block without any Location, so
|
||||
* initialize the base_location of the block and let
|
||||
* assign_missing_member_locations() do the rest.
|
||||
*/
|
||||
if ((b->shader->info.stage == MESA_SHADER_TASK && var->mode == vtn_variable_mode_output) ||
|
||||
(b->shader->info.stage == MESA_SHADER_MESH && var->mode == vtn_variable_mode_input)) {
|
||||
if (var->type->block)
|
||||
var->base_location = VARYING_SLOT_VAR0;
|
||||
}
|
||||
|
||||
break;
|
||||
}
|
||||
|
||||
|
|
Loading…
Reference in New Issue