vk: mesh programs with control topology 1 draw points, with the point sprite coordinate (grass shadow mesh on AMD, #302)

.github/workflows/source-regressions.yml

@@ -76,6 +76,14 @@ on:
       - tests/global_load_16/**
       - gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_memory.cpp
       - gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_instructions.h
+      - gpu/vk/eden/shader_recompiler/shader_info.h
+      - gpu/vk/eden/shader_recompiler/runtime_info.h
+      - gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
+      - gpu/vk/eden/shader_recompiler/frontend/maxwell/translate_program.cpp
+      - gpu/vk/eden/shader_recompiler/backend/spirv/spirv_emit_context.cpp
+      - gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv.cpp
+      - gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_mesh.cpp
+      - tests/test_mesh_point_topology.py
       - .github/workflows/source-regressions.yml
   pull_request:
     branches: [main]
@@ -152,6 +160,14 @@ on:
       - tests/global_load_16/**
       - gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_memory.cpp
       - gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_instructions.h
+      - gpu/vk/eden/shader_recompiler/shader_info.h
+      - gpu/vk/eden/shader_recompiler/runtime_info.h
+      - gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
+      - gpu/vk/eden/shader_recompiler/frontend/maxwell/translate_program.cpp
+      - gpu/vk/eden/shader_recompiler/backend/spirv/spirv_emit_context.cpp
+      - gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv.cpp
+      - gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_mesh.cpp
+      - tests/test_mesh_point_topology.py
       - .github/workflows/source-regressions.yml
   workflow_dispatch:
 
@@ -243,6 +259,8 @@ jobs:
         run: python -I -B tests/test_table_track.py
       - name: Check descriptor timing stays off
         run: python -I -B tests/test_desc_clock.py
+      - name: Check a mesh program with control topology 1 draws points
+        run: python -I -B tests/test_mesh_point_topology.py
       - name: Check the replay clock skips the performance counter
         if: runner.os == 'Linux'
         run: python -I -B tests/test_prof_clock.py

gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv.cpp

@@ -287,8 +287,9 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
         ctx.AddExecutionMode(main, spv::ExecutionMode::OutputVertices, program.mesh.max_vertices);
         ctx.AddExecutionMode(main, spv::ExecutionMode::OutputPrimitivesEXT,
                              program.mesh.max_primitives);
-        ctx.AddExecutionMode(main, program.mesh.topology == 1 ? spv::ExecutionMode::OutputLinesEXT
-                                                              : spv::ExecutionMode::OutputTrianglesEXT);
+        ctx.AddExecutionMode(main, program.mesh.Corners() == 1   ? spv::ExecutionMode::OutputPoints
+                                   : program.mesh.Corners() == 2 ? spv::ExecutionMode::OutputLinesEXT
+                                                                 : spv::ExecutionMode::OutputTrianglesEXT);
         break;
     case Stage::TessellationControl:
         execution_model = spv::ExecutionModel::TessellationControl;

gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp

@@ -307,6 +307,14 @@ Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr, Id vertex) {
     const u32 element{static_cast<u32>(attr) % 4};
     if (IR::IsGeneric(attr)) {
         const u32 index{IR::GenericAttributeIndex(attr)};
+        if (ctx.stage == Stage::Fragment && ctx.runtime_info.point_coord_generics.test(index)) {
+            // (hookshot) the point sprite's coordinate in place of the generic (RuntimeInfo::point_coord_generics)
+            if (element < 2) {
+                return ctx.OpLoad(ctx.F32[1],
+                                  ctx.OpAccessChain(ctx.input_f32, ctx.point_coord, ctx.Const(element)));
+            }
+            return ctx.Const(element == 3 ? 1.0f : 0.0f);
+        }
         const auto& generic{ctx.input_generics.at(index)};
         if (!ValidId(generic.id)) {
             // Attribute is disabled or varying component is not written

gpu/vk/eden/shader_recompiler/backend/spirv/emit_spirv_mesh.cpp

@@ -161,8 +161,8 @@ void MeshEpilogue(EmitContext& ctx) {
     for (u32 first = 0; first < mesh.max_primitives; first += threads) {
         const Id primitive{ctx.OpIAdd(ctx.U32[1], invocation, ctx.Const(first))};
         If(ctx, ctx.OpULessThan(ctx.U1, primitive, count), [&] {
-            // Vertex indices: map bytes 4 + 3 * primitive (lines: 4 + 2 * primitive).
-            const u32 corners{mesh.topology == 1 ? 2U : 3U};
+            // Vertex indices: map bytes 4 + 3 * primitive (lines: 4 + 2 * primitive, points: 4 + primitive).
+            const u32 corners{mesh.Corners()};
             const Id byte{ctx.OpIAdd(ctx.U32[1], ctx.OpIMul(ctx.U32[1], primitive, ctx.Const(corners)),
                                      ctx.Const(4U))};
             std::array<Id, 3> indices;
@@ -175,7 +175,9 @@ void MeshEpilogue(EmitContext& ctx) {
             }
             const Id output_indices{ctx.TypePointer(spv::StorageClass::Output, ctx.U32[corners])};
             ctx.OpStore(ctx.OpAccessChain(output_indices, ctx.mesh_primitive_indices, primitive),
-                        ctx.OpCompositeConstruct(ctx.U32[corners], std::span<const Id>(indices.data(), corners)));
+                        corners == 1 ? indices[0]   // (a point's index is a scalar)
+                                     : ctx.OpCompositeConstruct(ctx.U32[corners],
+                                                                std::span<const Id>(indices.data(), corners)));
             CopyLines(ctx, primitive, ctx.mesh_primitive_base, primitive_lines,
                       ctx.mesh_primitive_targets);
         });

gpu/vk/eden/shader_recompiler/backend/spirv/spirv_emit_context.cpp

@@ -1857,7 +1857,8 @@ void EmitContext::DefineInputs(const IR::Program& program) {
     if (loads[IR::Attribute::FrontFace]) {
         front_face = DefineInput(*this, U1, true, spv::BuiltIn::FrontFacing);
     }
-    if (loads[IR::Attribute::PointSpriteS] || loads[IR::Attribute::PointSpriteT]) {
+    if (loads[IR::Attribute::PointSpriteS] || loads[IR::Attribute::PointSpriteT] ||
+        (stage == Stage::Fragment && runtime_info.point_coord_generics.any())) {
         point_coord = DefineInput(*this, F32[2], true, spv::BuiltIn::PointCoord);
     }
     if (loads[IR::Attribute::TessellationEvaluationPointU] ||
@@ -1952,7 +1953,7 @@ void EmitContext::DefineMesh(const IR::Program& program) {
     mesh_attribute_words = (std::max)(
         mesh_primitive_base + Common::DivCeil(mesh_info.max_primitives, 32U) * primitive_lines * 32,
         1U);
-    mesh_map_bytes = 4 + (mesh_info.topology == 1 ? 2 : 3) * mesh_info.max_primitives;
+    mesh_map_bytes = 4 + mesh_info.Corners() * mesh_info.max_primitives;
     mesh_shared_u32 = TypePointer(spv::StorageClass::Workgroup, U32[1]);
     mesh_attributes = DefineVariable(*this, TypeArray(U32[1], Const(mesh_attribute_words)),
                                      std::nullopt, spv::StorageClass::Workgroup);
@@ -1966,10 +1967,12 @@ void EmitContext::DefineMesh(const IR::Program& program) {
 
     const Id vertex_count{Const(mesh_info.max_vertices)};
     const Id primitive_count{Const(mesh_info.max_primitives)};
-    const bool mesh_lines{mesh_info.topology == 1};
+    const u32 corners{mesh_info.Corners()};
     mesh_primitive_indices =
-        DefineVariable(*this, TypeArray(U32[mesh_lines ? 2 : 3], primitive_count),
-                       mesh_lines ? spv::BuiltIn::PrimitiveLineIndicesEXT : spv::BuiltIn::PrimitiveTriangleIndicesEXT,
+        DefineVariable(*this, TypeArray(U32[corners], primitive_count),
+                       corners == 1   ? spv::BuiltIn::PrimitivePointIndicesEXT
+                       : corners == 2 ? spv::BuiltIn::PrimitiveLineIndicesEXT
+                                      : spv::BuiltIn::PrimitiveTriangleIndicesEXT,
                        spv::StorageClass::Output);
     std::array<Id, IR::NUM_GENERICS> generics[2]{};
     Id point_size{}, primitive_layer{}, primitive_viewport{}, primitive_id_output{};

gpu/vk/eden/shader_recompiler/frontend/maxwell/translate_program.cpp

@@ -482,7 +482,8 @@ IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Blo
         const u32 threads{program.stage == Stage::Task ? init & 0x3f : (shader_b >> 20) & 0x3f};
         program.workgroup_size = {(std::max)(threads, 1U), 1, 1};
         if (program.stage == Stage::Mesh) {
-            if (program.mesh.topology != 4 && program.mesh.topology != 1) {   // (GL's numbers: 1 lines, 4 triangles)
+            if (program.mesh.topology != 4 && program.mesh.topology != 1 &&
+                program.mesh.topology != 0) {   // (GL's numbers: 0 points, 1 lines, 4 triangles)
                 throw NotImplementedException("Mesh output topology {} (SET_MESH_SHADER_A {:#010x})",
                                               program.mesh.topology, shader_a);
             }

gpu/vk/eden/shader_recompiler/runtime_info.h

@@ -88,6 +88,10 @@ struct RuntimeInfo {
     VaryingState previous_stage_stores;
     /// Fragment stage after a mesh program: generic inputs it writes per primitive.
     std::bitset<32> per_primitive_generics{};
+    /// Fragment stage after a mesh program that draws points: generic inputs the hardware's point sprite
+    /// replaces with the coordinate within the point (s, t, 0, 1), as NVN's do with the generics the mesh
+    /// program leaves unwritten.
+    std::bitset<32> point_coord_generics{};
     std::map<IR::Attribute, IR::Attribute> previous_stage_legacy_stores_mapping;
 
     bool convert_depth_mode{};

gpu/vk/eden/shader_recompiler/shader_info.h

@@ -262,9 +262,14 @@ struct MeshInfo {
     u32 payload_words{}; ///< Task output region / mesh input region, in 32-bit words
     u32 max_vertices{};
     u32 max_primitives{};
-    u32 topology{}; ///< SET_MESH_SHADER_A output topology (4 = triangles)
+    u32 topology{}; ///< SET_MESH_SHADER_A output topology (GL's numbers: 0 points, 1 lines, 4 triangles)
     boost::container::static_vector<u32, 136> vertex_lines;
     boost::container::static_vector<u32, 136> primitive_lines;
+
+    /// Vertex indices per primitive, one map byte each
+    u32 Corners() const {
+        return topology == 0 ? 1U : topology == 1 ? 2U : 3U;
+    }
 };
 
 struct Info {

gpu/vk/vk_shader.cpp

@@ -106,7 +106,7 @@ struct Code {
     uint32_t texture_bound = ~0u;   // the bank native TEX instructions name (~0: none; TEXS take the stage's own)
     bool isbe = false;              // an NVN2 task (stage 6) or mesh (stage 7) program: it writes the internal stage buffer
     // (task / mesh: control section +0x58 threads 15:0, payload bytes 31:16; +0x60 max primitives 15:0, max vertices
-    // 31:16; +0x5f the output topology, 3 triangles)
+    // 31:16; +0x5f the output topology, 1 points, 3 triangles)
     uint32_t mesh_threads = 0, payload_bytes = 0, max_vertices = 0, max_primitives = 0, topology = 0;
     bool header_only = false;       // a graphics header and no code (NVN2's per-primitive output map of a mesh program)
     uint8_t sph[sizeof(Shader::ProgramHeader)]{};
@@ -1199,7 +1199,11 @@ void miss_note(uint64_t key, const char* what, uint64_t a, uint64_t b, const Nvn
 }
 
 std::array<u32, 3> mesh_registers_of(const Code* task, const Code& mesh) {
-    const uint32_t topology = mesh.topology == 3 ? 4 : mesh.topology;   // (triangles)
+    // (the control section's topology in GL's numbers: 3 triangles; 1 points, not lines. OoT's four mesh programs with a 1 (the grass
+    // drawn into the 2400 x 2400 shadow map around Link) write one index byte per primitive, the primitive count = their vertex
+    // count, and a point size. Read as lines, each joined vertex 2i to 2i + 1 of its meshlet: a mesh of long thin dark stripes on
+    // the ground, the water and Link around him, on AMD; Hookshot issue 302)
+    const uint32_t topology = mesh.topology == 3 ? 4 : mesh.topology == 1 ? 0 : mesh.topology;
     const uint32_t lines = task ? (task->payload_bytes + 127) / 128 : 0;
     // (the larger of the two for both: the control section's pair is not always vertices and primitives in that order. The door's
     // says 126 and 64, where its code writes 64 vertices and up to 128 primitives; held to 64, a meshlet's other primitives were
@@ -1292,6 +1296,13 @@ void translate_mesh_into(Program& p, const NvnState& s, const Code* task, const
                 if (Shader::IR::IsGeneric(attribute))
                     frt.per_primitive_generics.set(Shader::IR::GenericAttributeIndex(attribute));
             }
+            // (points: a generic the fragment program reads and the mesh program never writes is the point sprite's coordinate,
+            // as the hardware replaces it. OoT's grass shadow reads its alpha mask at generic 5's (s, t): left at (0, 0, 0, 1),
+            // each point was a filled square in the shadow map instead of a tuft)
+            if (mesh_t->program.mesh.topology == 0)
+                for (size_t i = 0; i < Shader::IR::NUM_GENERICS; i++)
+                    if (fs->program.info.loads.Generic(i) && !mesh_t->program.info.stores.Generic(i))
+                        frt.point_coord_generics.set(i);
             for (int i = 0; i < 8; i++) frt.color_output_types[i] = static_cast<Shader::AttributeType>(color_types[i]);
             const std::vector<uint32_t> fspirv = Shader::Backend::SPIRV::EmitSPIRV(profile(), frt, fs->program, bindings);
             kept.push_back(fspirv);