From b34f9740cacb2c6567a24297fd8674e3977a851a Mon Sep 17 00:00:00 2001 From: Caio Marcelo de Oliveira Filho Date: Thu, 29 Apr 2021 15:06:29 -0700 Subject: [PATCH] spirv: Implement non-Multiview parts of SPV_NV_mesh_shader MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Reviewed-by: Timur Kristóf Part-of: --- src/compiler/shader_info.h | 1 + src/compiler/spirv/spirv2nir.c | 6 +- src/compiler/spirv/spirv_to_nir.c | 154 ++++++++++++++++++++++++++--- src/compiler/spirv/vtn_variables.c | 87 ++++++++++++++-- 4 files changed, 224 insertions(+), 24 deletions(-) diff --git a/src/compiler/shader_info.h b/src/compiler/shader_info.h index 15a66b0144a..51f8668f8fd 100644 --- a/src/compiler/shader_info.h +++ b/src/compiler/shader_info.h @@ -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; diff --git a/src/compiler/spirv/spirv2nir.c b/src/compiler/spirv/spirv2nir.c index fed803e92a2..c473bf2c8d7 100644 --- a/src/compiler/spirv/spirv2nir.c +++ b/src/compiler/spirv/spirv2nir.c @@ -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 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 Specify the entry-point name.\n" " -g, --opengl Use OpenGL environment instead of Vulkan for\n" " graphics stages.\n" diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c index a64039f9469..697af85388e 100644 --- a/src/compiler/spirv/spirv_to_nir.c +++ b/src/compiler/spirv/spirv_to_nir.c @@ -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); } diff --git a/src/compiler/spirv/vtn_variables.c b/src/compiler/spirv/vtn_variables.c index d281061a31f..fc1cdc19a63 100644 --- a/src/compiler/spirv/vtn_variables.c +++ b/src/compiler/spirv/vtn_variables.c @@ -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; }