summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--src/video_core/CMakeLists.txt11
-rw-r--r--src/video_core/renderer_vulkan/fixed_pipeline_state.cpp18
-rw-r--r--src/video_core/renderer_vulkan/fixed_pipeline_state.h10
-rw-r--r--src/video_core/renderer_vulkan/vk_compute_pass.cpp339
-rw-r--r--src/video_core/renderer_vulkan/vk_compute_pass.h77
-rw-r--r--src/video_core/renderer_vulkan/vk_compute_pipeline.cpp112
-rw-r--r--src/video_core/renderer_vulkan/vk_compute_pipeline.h66
-rw-r--r--src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp271
-rw-r--r--src/video_core/renderer_vulkan/vk_graphics_pipeline.h90
-rw-r--r--src/video_core/renderer_vulkan/vk_pipeline_cache.cpp395
-rw-r--r--src/video_core/renderer_vulkan/vk_pipeline_cache.h200
-rw-r--r--src/video_core/renderer_vulkan/vk_rasterizer.h13
-rw-r--r--src/video_core/renderer_vulkan/vk_shader_util.cpp34
-rw-r--r--src/video_core/renderer_vulkan/vk_shader_util.h17
14 files changed, 1643 insertions, 10 deletions
diff --git a/src/video_core/CMakeLists.txt b/src/video_core/CMakeLists.txt
index c80171fe6..142852082 100644
--- a/src/video_core/CMakeLists.txt
+++ b/src/video_core/CMakeLists.txt
@@ -155,14 +155,23 @@ if (ENABLE_VULKAN)
155 renderer_vulkan/maxwell_to_vk.h 155 renderer_vulkan/maxwell_to_vk.h
156 renderer_vulkan/vk_buffer_cache.cpp 156 renderer_vulkan/vk_buffer_cache.cpp
157 renderer_vulkan/vk_buffer_cache.h 157 renderer_vulkan/vk_buffer_cache.h
158 renderer_vulkan/vk_compute_pass.cpp
159 renderer_vulkan/vk_compute_pass.h
160 renderer_vulkan/vk_compute_pipeline.cpp
161 renderer_vulkan/vk_compute_pipeline.h
158 renderer_vulkan/vk_descriptor_pool.cpp 162 renderer_vulkan/vk_descriptor_pool.cpp
159 renderer_vulkan/vk_descriptor_pool.h 163 renderer_vulkan/vk_descriptor_pool.h
160 renderer_vulkan/vk_device.cpp 164 renderer_vulkan/vk_device.cpp
161 renderer_vulkan/vk_device.h 165 renderer_vulkan/vk_device.h
166 renderer_vulkan/vk_graphics_pipeline.cpp
167 renderer_vulkan/vk_graphics_pipeline.h
162 renderer_vulkan/vk_image.cpp 168 renderer_vulkan/vk_image.cpp
163 renderer_vulkan/vk_image.h 169 renderer_vulkan/vk_image.h
164 renderer_vulkan/vk_memory_manager.cpp 170 renderer_vulkan/vk_memory_manager.cpp
165 renderer_vulkan/vk_memory_manager.h 171 renderer_vulkan/vk_memory_manager.h
172 renderer_vulkan/vk_pipeline_cache.cpp
173 renderer_vulkan/vk_pipeline_cache.h
174 renderer_vulkan/vk_rasterizer.h
166 renderer_vulkan/vk_renderpass_cache.cpp 175 renderer_vulkan/vk_renderpass_cache.cpp
167 renderer_vulkan/vk_renderpass_cache.h 176 renderer_vulkan/vk_renderpass_cache.h
168 renderer_vulkan/vk_resource_manager.cpp 177 renderer_vulkan/vk_resource_manager.cpp
@@ -173,6 +182,8 @@ if (ENABLE_VULKAN)
173 renderer_vulkan/vk_scheduler.h 182 renderer_vulkan/vk_scheduler.h
174 renderer_vulkan/vk_shader_decompiler.cpp 183 renderer_vulkan/vk_shader_decompiler.cpp
175 renderer_vulkan/vk_shader_decompiler.h 184 renderer_vulkan/vk_shader_decompiler.h
185 renderer_vulkan/vk_shader_util.cpp
186 renderer_vulkan/vk_shader_util.h
176 renderer_vulkan/vk_staging_buffer_pool.cpp 187 renderer_vulkan/vk_staging_buffer_pool.cpp
177 renderer_vulkan/vk_staging_buffer_pool.h 188 renderer_vulkan/vk_staging_buffer_pool.h
178 renderer_vulkan/vk_stream_buffer.cpp 189 renderer_vulkan/vk_stream_buffer.cpp
diff --git a/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp b/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp
index 5a490f6ef..4e3ff231e 100644
--- a/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp
+++ b/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp
@@ -109,6 +109,9 @@ constexpr FixedPipelineState::Rasterizer GetRasterizerState(const Maxwell& regs)
109 const auto topology = static_cast<std::size_t>(regs.draw.topology.Value()); 109 const auto topology = static_cast<std::size_t>(regs.draw.topology.Value());
110 const bool depth_bias_enabled = enabled_lut[PolygonOffsetEnableLUT[topology]]; 110 const bool depth_bias_enabled = enabled_lut[PolygonOffsetEnableLUT[topology]];
111 111
112 const auto& clip = regs.view_volume_clip_control;
113 const bool depth_clamp_enabled = clip.depth_clamp_near == 1 || clip.depth_clamp_far == 1;
114
112 Maxwell::Cull::FrontFace front_face = regs.cull.front_face; 115 Maxwell::Cull::FrontFace front_face = regs.cull.front_face;
113 if (regs.screen_y_control.triangle_rast_flip != 0 && 116 if (regs.screen_y_control.triangle_rast_flip != 0 &&
114 regs.viewport_transform[0].scale_y > 0.0f) { 117 regs.viewport_transform[0].scale_y > 0.0f) {
@@ -119,8 +122,9 @@ constexpr FixedPipelineState::Rasterizer GetRasterizerState(const Maxwell& regs)
119 } 122 }
120 123
121 const bool gl_ndc = regs.depth_mode == Maxwell::DepthMode::MinusOneToOne; 124 const bool gl_ndc = regs.depth_mode == Maxwell::DepthMode::MinusOneToOne;
122 return FixedPipelineState::Rasterizer(regs.cull.enabled, depth_bias_enabled, gl_ndc, 125 return FixedPipelineState::Rasterizer(regs.cull.enabled, depth_bias_enabled,
123 regs.cull.cull_face, front_face); 126 depth_clamp_enabled, gl_ndc, regs.cull.cull_face,
127 front_face);
124} 128}
125 129
126} // Anonymous namespace 130} // Anonymous namespace
@@ -222,15 +226,17 @@ bool FixedPipelineState::Tessellation::operator==(const Tessellation& rhs) const
222std::size_t FixedPipelineState::Rasterizer::Hash() const noexcept { 226std::size_t FixedPipelineState::Rasterizer::Hash() const noexcept {
223 return static_cast<std::size_t>(cull_enable) ^ 227 return static_cast<std::size_t>(cull_enable) ^
224 (static_cast<std::size_t>(depth_bias_enable) << 1) ^ 228 (static_cast<std::size_t>(depth_bias_enable) << 1) ^
225 (static_cast<std::size_t>(ndc_minus_one_to_one) << 2) ^ 229 (static_cast<std::size_t>(depth_clamp_enable) << 2) ^
230 (static_cast<std::size_t>(ndc_minus_one_to_one) << 3) ^
226 (static_cast<std::size_t>(cull_face) << 24) ^ 231 (static_cast<std::size_t>(cull_face) << 24) ^
227 (static_cast<std::size_t>(front_face) << 48); 232 (static_cast<std::size_t>(front_face) << 48);
228} 233}
229 234
230bool FixedPipelineState::Rasterizer::operator==(const Rasterizer& rhs) const noexcept { 235bool FixedPipelineState::Rasterizer::operator==(const Rasterizer& rhs) const noexcept {
231 return std::tie(cull_enable, depth_bias_enable, ndc_minus_one_to_one, cull_face, front_face) == 236 return std::tie(cull_enable, depth_bias_enable, depth_clamp_enable, ndc_minus_one_to_one,
232 std::tie(rhs.cull_enable, rhs.depth_bias_enable, rhs.ndc_minus_one_to_one, rhs.cull_face, 237 cull_face, front_face) ==
233 rhs.front_face); 238 std::tie(rhs.cull_enable, rhs.depth_bias_enable, rhs.depth_clamp_enable,
239 rhs.ndc_minus_one_to_one, rhs.cull_face, rhs.front_face);
234} 240}
235 241
236std::size_t FixedPipelineState::DepthStencil::Hash() const noexcept { 242std::size_t FixedPipelineState::DepthStencil::Hash() const noexcept {
diff --git a/src/video_core/renderer_vulkan/fixed_pipeline_state.h b/src/video_core/renderer_vulkan/fixed_pipeline_state.h
index 04152c0d4..87056ef37 100644
--- a/src/video_core/renderer_vulkan/fixed_pipeline_state.h
+++ b/src/video_core/renderer_vulkan/fixed_pipeline_state.h
@@ -170,15 +170,17 @@ struct FixedPipelineState {
170 }; 170 };
171 171
172 struct Rasterizer { 172 struct Rasterizer {
173 constexpr Rasterizer(bool cull_enable, bool depth_bias_enable, bool ndc_minus_one_to_one, 173 constexpr Rasterizer(bool cull_enable, bool depth_bias_enable, bool depth_clamp_enable,
174 Maxwell::Cull::CullFace cull_face, Maxwell::Cull::FrontFace front_face) 174 bool ndc_minus_one_to_one, Maxwell::Cull::CullFace cull_face,
175 Maxwell::Cull::FrontFace front_face)
175 : cull_enable{cull_enable}, depth_bias_enable{depth_bias_enable}, 176 : cull_enable{cull_enable}, depth_bias_enable{depth_bias_enable},
176 ndc_minus_one_to_one{ndc_minus_one_to_one}, cull_face{cull_face}, front_face{ 177 depth_clamp_enable{depth_clamp_enable}, ndc_minus_one_to_one{ndc_minus_one_to_one},
177 front_face} {} 178 cull_face{cull_face}, front_face{front_face} {}
178 Rasterizer() = default; 179 Rasterizer() = default;
179 180
180 bool cull_enable; 181 bool cull_enable;
181 bool depth_bias_enable; 182 bool depth_bias_enable;
183 bool depth_clamp_enable;
182 bool ndc_minus_one_to_one; 184 bool ndc_minus_one_to_one;
183 Maxwell::Cull::CullFace cull_face; 185 Maxwell::Cull::CullFace cull_face;
184 Maxwell::Cull::FrontFace front_face; 186 Maxwell::Cull::FrontFace front_face;
diff --git a/src/video_core/renderer_vulkan/vk_compute_pass.cpp b/src/video_core/renderer_vulkan/vk_compute_pass.cpp
new file mode 100644
index 000000000..7bdda3d79
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_compute_pass.cpp
@@ -0,0 +1,339 @@
1// Copyright 2019 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <cstring>
6#include <memory>
7#include <optional>
8#include <utility>
9#include <vector>
10#include "common/alignment.h"
11#include "common/assert.h"
12#include "common/common_types.h"
13#include "video_core/renderer_vulkan/declarations.h"
14#include "video_core/renderer_vulkan/vk_compute_pass.h"
15#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
16#include "video_core/renderer_vulkan/vk_device.h"
17#include "video_core/renderer_vulkan/vk_scheduler.h"
18#include "video_core/renderer_vulkan/vk_staging_buffer_pool.h"
19#include "video_core/renderer_vulkan/vk_update_descriptor.h"
20
21namespace Vulkan {
22
23namespace {
24
25// Quad array SPIR-V module. Generated from the "shaders/" directory, read the instructions there.
26constexpr u8 quad_array[] = {
27 0x03, 0x02, 0x23, 0x07, 0x00, 0x00, 0x01, 0x00, 0x07, 0x00, 0x08, 0x00, 0x54, 0x00, 0x00, 0x00,
28 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, 0x01, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x06, 0x00,
29 0x01, 0x00, 0x00, 0x00, 0x47, 0x4c, 0x53, 0x4c, 0x2e, 0x73, 0x74, 0x64, 0x2e, 0x34, 0x35, 0x30,
30 0x00, 0x00, 0x00, 0x00, 0x0e, 0x00, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
31 0x0f, 0x00, 0x06, 0x00, 0x05, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x6d, 0x61, 0x69, 0x6e,
32 0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x10, 0x00, 0x06, 0x00, 0x04, 0x00, 0x00, 0x00,
33 0x11, 0x00, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
34 0x47, 0x00, 0x04, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x1c, 0x00, 0x00, 0x00,
35 0x47, 0x00, 0x04, 0x00, 0x13, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
36 0x48, 0x00, 0x05, 0x00, 0x14, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00,
37 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x03, 0x00, 0x14, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00,
38 0x47, 0x00, 0x04, 0x00, 0x16, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
39 0x47, 0x00, 0x04, 0x00, 0x16, 0x00, 0x00, 0x00, 0x21, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
40 0x48, 0x00, 0x05, 0x00, 0x29, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00,
41 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x03, 0x00, 0x29, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
42 0x47, 0x00, 0x04, 0x00, 0x4a, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00,
43 0x13, 0x00, 0x02, 0x00, 0x02, 0x00, 0x00, 0x00, 0x21, 0x00, 0x03, 0x00, 0x03, 0x00, 0x00, 0x00,
44 0x02, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00,
45 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x07, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00,
46 0x06, 0x00, 0x00, 0x00, 0x17, 0x00, 0x04, 0x00, 0x09, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00,
47 0x03, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x0a, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
48 0x09, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, 0x0a, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00,
49 0x01, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
50 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x0d, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
51 0x06, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
52 0x06, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x03, 0x00, 0x13, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00,
53 0x1e, 0x00, 0x03, 0x00, 0x14, 0x00, 0x00, 0x00, 0x13, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00,
54 0x15, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x14, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00,
55 0x15, 0x00, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00,
56 0x18, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x14, 0x00, 0x02, 0x00,
57 0x1b, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x03, 0x00, 0x29, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00,
58 0x20, 0x00, 0x04, 0x00, 0x2a, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x29, 0x00, 0x00, 0x00,
59 0x3b, 0x00, 0x04, 0x00, 0x2a, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00,
60 0x2b, 0x00, 0x04, 0x00, 0x18, 0x00, 0x00, 0x00, 0x2c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
61 0x20, 0x00, 0x04, 0x00, 0x2d, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00,
62 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x31, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
63 0x1c, 0x00, 0x04, 0x00, 0x34, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
64 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x35, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
65 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x36, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
66 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x37, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00,
67 0x2c, 0x00, 0x09, 0x00, 0x34, 0x00, 0x00, 0x00, 0x38, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00,
68 0x35, 0x00, 0x00, 0x00, 0x36, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x36, 0x00, 0x00, 0x00,
69 0x37, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x3a, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00,
70 0x34, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x44, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
71 0x06, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x18, 0x00, 0x00, 0x00, 0x47, 0x00, 0x00, 0x00,
72 0x01, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x49, 0x00, 0x00, 0x00,
73 0x00, 0x04, 0x00, 0x00, 0x2c, 0x00, 0x06, 0x00, 0x09, 0x00, 0x00, 0x00, 0x4a, 0x00, 0x00, 0x00,
74 0x49, 0x00, 0x00, 0x00, 0x35, 0x00, 0x00, 0x00, 0x35, 0x00, 0x00, 0x00, 0x36, 0x00, 0x05, 0x00,
75 0x02, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00,
76 0xf8, 0x00, 0x02, 0x00, 0x05, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, 0x3a, 0x00, 0x00, 0x00,
77 0x3b, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, 0xf9, 0x00, 0x02, 0x00, 0x4c, 0x00, 0x00, 0x00,
78 0xf8, 0x00, 0x02, 0x00, 0x4c, 0x00, 0x00, 0x00, 0xf6, 0x00, 0x04, 0x00, 0x4b, 0x00, 0x00, 0x00,
79 0x4e, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xf9, 0x00, 0x02, 0x00, 0x4d, 0x00, 0x00, 0x00,
80 0xf8, 0x00, 0x02, 0x00, 0x4d, 0x00, 0x00, 0x00, 0x41, 0x00, 0x05, 0x00, 0x0d, 0x00, 0x00, 0x00,
81 0x0e, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00,
82 0x06, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0x0e, 0x00, 0x00, 0x00, 0x84, 0x00, 0x05, 0x00,
83 0x06, 0x00, 0x00, 0x00, 0x12, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
84 0x44, 0x00, 0x05, 0x00, 0x06, 0x00, 0x00, 0x00, 0x17, 0x00, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00,
85 0x00, 0x00, 0x00, 0x00, 0x7c, 0x00, 0x04, 0x00, 0x18, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00,
86 0x17, 0x00, 0x00, 0x00, 0x7c, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00,
87 0x19, 0x00, 0x00, 0x00, 0xae, 0x00, 0x05, 0x00, 0x1b, 0x00, 0x00, 0x00, 0x1c, 0x00, 0x00, 0x00,
88 0x12, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, 0xf7, 0x00, 0x03, 0x00, 0x1e, 0x00, 0x00, 0x00,
89 0x00, 0x00, 0x00, 0x00, 0xfa, 0x00, 0x04, 0x00, 0x1c, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x00, 0x00,
90 0x1e, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x1d, 0x00, 0x00, 0x00, 0xf9, 0x00, 0x02, 0x00,
91 0x4b, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x1e, 0x00, 0x00, 0x00, 0xf9, 0x00, 0x02, 0x00,
92 0x21, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x21, 0x00, 0x00, 0x00, 0xf5, 0x00, 0x07, 0x00,
93 0x06, 0x00, 0x00, 0x00, 0x53, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x00, 0x00,
94 0x48, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, 0xb0, 0x00, 0x05, 0x00, 0x1b, 0x00, 0x00, 0x00,
95 0x27, 0x00, 0x00, 0x00, 0x53, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0xf6, 0x00, 0x04, 0x00,
96 0x23, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xfa, 0x00, 0x04, 0x00,
97 0x27, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00,
98 0x22, 0x00, 0x00, 0x00, 0x41, 0x00, 0x05, 0x00, 0x2d, 0x00, 0x00, 0x00, 0x2e, 0x00, 0x00, 0x00,
99 0x2b, 0x00, 0x00, 0x00, 0x2c, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00,
100 0x2f, 0x00, 0x00, 0x00, 0x2e, 0x00, 0x00, 0x00, 0x84, 0x00, 0x05, 0x00, 0x06, 0x00, 0x00, 0x00,
101 0x32, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0x31, 0x00, 0x00, 0x00, 0x80, 0x00, 0x05, 0x00,
102 0x06, 0x00, 0x00, 0x00, 0x33, 0x00, 0x00, 0x00, 0x2f, 0x00, 0x00, 0x00, 0x32, 0x00, 0x00, 0x00,
103 0x3e, 0x00, 0x03, 0x00, 0x3b, 0x00, 0x00, 0x00, 0x38, 0x00, 0x00, 0x00, 0x41, 0x00, 0x05, 0x00,
104 0x07, 0x00, 0x00, 0x00, 0x3c, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x00, 0x00, 0x53, 0x00, 0x00, 0x00,
105 0x3d, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x00, 0x00, 0x3c, 0x00, 0x00, 0x00,
106 0x80, 0x00, 0x05, 0x00, 0x06, 0x00, 0x00, 0x00, 0x3e, 0x00, 0x00, 0x00, 0x33, 0x00, 0x00, 0x00,
107 0x3d, 0x00, 0x00, 0x00, 0x80, 0x00, 0x05, 0x00, 0x06, 0x00, 0x00, 0x00, 0x42, 0x00, 0x00, 0x00,
108 0x12, 0x00, 0x00, 0x00, 0x53, 0x00, 0x00, 0x00, 0x41, 0x00, 0x06, 0x00, 0x44, 0x00, 0x00, 0x00,
109 0x45, 0x00, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, 0x2c, 0x00, 0x00, 0x00, 0x42, 0x00, 0x00, 0x00,
110 0x3e, 0x00, 0x03, 0x00, 0x45, 0x00, 0x00, 0x00, 0x3e, 0x00, 0x00, 0x00, 0x80, 0x00, 0x05, 0x00,
111 0x06, 0x00, 0x00, 0x00, 0x48, 0x00, 0x00, 0x00, 0x53, 0x00, 0x00, 0x00, 0x47, 0x00, 0x00, 0x00,
112 0xf9, 0x00, 0x02, 0x00, 0x21, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x23, 0x00, 0x00, 0x00,
113 0xf9, 0x00, 0x02, 0x00, 0x4b, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x4e, 0x00, 0x00, 0x00,
114 0xf9, 0x00, 0x02, 0x00, 0x4c, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x4b, 0x00, 0x00, 0x00,
115 0xfd, 0x00, 0x01, 0x00, 0x38, 0x00, 0x01, 0x00};
116
117// Uint8 SPIR-V module. Generated from the "shaders/" directory.
118constexpr u8 uint8_pass[] = {
119 0x03, 0x02, 0x23, 0x07, 0x00, 0x00, 0x01, 0x00, 0x07, 0x00, 0x08, 0x00, 0x2f, 0x00, 0x00, 0x00,
120 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, 0x01, 0x00, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00,
121 0x51, 0x11, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, 0x61, 0x11, 0x00, 0x00, 0x0a, 0x00, 0x07, 0x00,
122 0x53, 0x50, 0x56, 0x5f, 0x4b, 0x48, 0x52, 0x5f, 0x31, 0x36, 0x62, 0x69, 0x74, 0x5f, 0x73, 0x74,
123 0x6f, 0x72, 0x61, 0x67, 0x65, 0x00, 0x00, 0x00, 0x0a, 0x00, 0x07, 0x00, 0x53, 0x50, 0x56, 0x5f,
124 0x4b, 0x48, 0x52, 0x5f, 0x38, 0x62, 0x69, 0x74, 0x5f, 0x73, 0x74, 0x6f, 0x72, 0x61, 0x67, 0x65,
125 0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x06, 0x00, 0x01, 0x00, 0x00, 0x00, 0x47, 0x4c, 0x53, 0x4c,
126 0x2e, 0x73, 0x74, 0x64, 0x2e, 0x34, 0x35, 0x30, 0x00, 0x00, 0x00, 0x00, 0x0e, 0x00, 0x03, 0x00,
127 0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x06, 0x00, 0x05, 0x00, 0x00, 0x00,
128 0x04, 0x00, 0x00, 0x00, 0x6d, 0x61, 0x69, 0x6e, 0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00,
129 0x10, 0x00, 0x06, 0x00, 0x04, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00,
130 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x0b, 0x00, 0x00, 0x00,
131 0x0b, 0x00, 0x00, 0x00, 0x1c, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x12, 0x00, 0x00, 0x00,
132 0x06, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x48, 0x00, 0x04, 0x00, 0x13, 0x00, 0x00, 0x00,
133 0x00, 0x00, 0x00, 0x00, 0x18, 0x00, 0x00, 0x00, 0x48, 0x00, 0x05, 0x00, 0x13, 0x00, 0x00, 0x00,
134 0x00, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x03, 0x00,
135 0x13, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x15, 0x00, 0x00, 0x00,
136 0x22, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x15, 0x00, 0x00, 0x00,
137 0x21, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x1f, 0x00, 0x00, 0x00,
138 0x06, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x48, 0x00, 0x04, 0x00, 0x20, 0x00, 0x00, 0x00,
139 0x00, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00, 0x48, 0x00, 0x05, 0x00, 0x20, 0x00, 0x00, 0x00,
140 0x00, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x03, 0x00,
141 0x20, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x22, 0x00, 0x00, 0x00,
142 0x22, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x22, 0x00, 0x00, 0x00,
143 0x21, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x2e, 0x00, 0x00, 0x00,
144 0x0b, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00, 0x13, 0x00, 0x02, 0x00, 0x02, 0x00, 0x00, 0x00,
145 0x21, 0x00, 0x03, 0x00, 0x03, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00,
146 0x06, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00,
147 0x07, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x17, 0x00, 0x04, 0x00,
148 0x09, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00,
149 0x0a, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00,
150 0x0a, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00,
151 0x06, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00,
152 0x0d, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00,
153 0x11, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x03, 0x00,
154 0x12, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x03, 0x00, 0x13, 0x00, 0x00, 0x00,
155 0x12, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x14, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
156 0x13, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, 0x14, 0x00, 0x00, 0x00, 0x15, 0x00, 0x00, 0x00,
157 0x02, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00, 0x17, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00,
158 0x01, 0x00, 0x00, 0x00, 0x14, 0x00, 0x02, 0x00, 0x1a, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00,
159 0x1e, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x03, 0x00,
160 0x1f, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x03, 0x00, 0x20, 0x00, 0x00, 0x00,
161 0x1f, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x21, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
162 0x20, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, 0x21, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00,
163 0x02, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x17, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00,
164 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x26, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
165 0x11, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x2a, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
166 0x1e, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x2c, 0x00, 0x00, 0x00,
167 0x00, 0x04, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x2d, 0x00, 0x00, 0x00,
168 0x01, 0x00, 0x00, 0x00, 0x2c, 0x00, 0x06, 0x00, 0x09, 0x00, 0x00, 0x00, 0x2e, 0x00, 0x00, 0x00,
169 0x2c, 0x00, 0x00, 0x00, 0x2d, 0x00, 0x00, 0x00, 0x2d, 0x00, 0x00, 0x00, 0x36, 0x00, 0x05, 0x00,
170 0x02, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00,
171 0xf8, 0x00, 0x02, 0x00, 0x05, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, 0x07, 0x00, 0x00, 0x00,
172 0x08, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, 0x41, 0x00, 0x05, 0x00, 0x0d, 0x00, 0x00, 0x00,
173 0x0e, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00,
174 0x06, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0x0e, 0x00, 0x00, 0x00, 0x3e, 0x00, 0x03, 0x00,
175 0x08, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00,
176 0x10, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x44, 0x00, 0x05, 0x00, 0x06, 0x00, 0x00, 0x00,
177 0x16, 0x00, 0x00, 0x00, 0x15, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x7c, 0x00, 0x04, 0x00,
178 0x17, 0x00, 0x00, 0x00, 0x18, 0x00, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, 0x7c, 0x00, 0x04, 0x00,
179 0x06, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00, 0x18, 0x00, 0x00, 0x00, 0xb0, 0x00, 0x05, 0x00,
180 0x1a, 0x00, 0x00, 0x00, 0x1b, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00,
181 0xf7, 0x00, 0x03, 0x00, 0x1d, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xfa, 0x00, 0x04, 0x00,
182 0x1b, 0x00, 0x00, 0x00, 0x1c, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00,
183 0x1c, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x24, 0x00, 0x00, 0x00,
184 0x08, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x25, 0x00, 0x00, 0x00,
185 0x08, 0x00, 0x00, 0x00, 0x41, 0x00, 0x06, 0x00, 0x26, 0x00, 0x00, 0x00, 0x27, 0x00, 0x00, 0x00,
186 0x15, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, 0x25, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00,
187 0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x27, 0x00, 0x00, 0x00, 0x71, 0x00, 0x04, 0x00,
188 0x1e, 0x00, 0x00, 0x00, 0x29, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x41, 0x00, 0x06, 0x00,
189 0x2a, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00,
190 0x24, 0x00, 0x00, 0x00, 0x3e, 0x00, 0x03, 0x00, 0x2b, 0x00, 0x00, 0x00, 0x29, 0x00, 0x00, 0x00,
191 0xf9, 0x00, 0x02, 0x00, 0x1d, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x1d, 0x00, 0x00, 0x00,
192 0xfd, 0x00, 0x01, 0x00, 0x38, 0x00, 0x01, 0x00};
193
194} // Anonymous namespace
195
196VKComputePass::VKComputePass(const VKDevice& device, VKDescriptorPool& descriptor_pool,
197 const std::vector<vk::DescriptorSetLayoutBinding>& bindings,
198 const std::vector<vk::DescriptorUpdateTemplateEntry>& templates,
199 const std::vector<vk::PushConstantRange> push_constants,
200 std::size_t code_size, const u8* code) {
201 const auto dev = device.GetLogical();
202 const auto& dld = device.GetDispatchLoader();
203
204 const vk::DescriptorSetLayoutCreateInfo descriptor_layout_ci(
205 {}, static_cast<u32>(bindings.size()), bindings.data());
206 descriptor_set_layout = dev.createDescriptorSetLayoutUnique(descriptor_layout_ci, nullptr, dld);
207
208 const vk::PipelineLayoutCreateInfo pipeline_layout_ci({}, 1, &*descriptor_set_layout,
209 static_cast<u32>(push_constants.size()),
210 push_constants.data());
211 layout = dev.createPipelineLayoutUnique(pipeline_layout_ci, nullptr, dld);
212
213 if (!templates.empty()) {
214 const vk::DescriptorUpdateTemplateCreateInfo template_ci(
215 {}, static_cast<u32>(templates.size()), templates.data(),
216 vk::DescriptorUpdateTemplateType::eDescriptorSet, *descriptor_set_layout,
217 vk::PipelineBindPoint::eGraphics, *layout, 0);
218 descriptor_template = dev.createDescriptorUpdateTemplateUnique(template_ci, nullptr, dld);
219
220 descriptor_allocator.emplace(descriptor_pool, *descriptor_set_layout);
221 }
222
223 auto code_copy = std::make_unique<u32[]>(code_size / sizeof(u32) + 1);
224 std::memcpy(code_copy.get(), code, code_size);
225 const vk::ShaderModuleCreateInfo module_ci({}, code_size, code_copy.get());
226 module = dev.createShaderModuleUnique(module_ci, nullptr, dld);
227
228 const vk::PipelineShaderStageCreateInfo stage_ci({}, vk::ShaderStageFlagBits::eCompute, *module,
229 "main", nullptr);
230
231 const vk::ComputePipelineCreateInfo pipeline_ci({}, stage_ci, *layout, nullptr, 0);
232 pipeline = dev.createComputePipelineUnique(nullptr, pipeline_ci, nullptr, dld);
233}
234
235VKComputePass::~VKComputePass() = default;
236
237vk::DescriptorSet VKComputePass::CommitDescriptorSet(
238 VKUpdateDescriptorQueue& update_descriptor_queue, VKFence& fence) {
239 if (!descriptor_template) {
240 return {};
241 }
242 const auto set = descriptor_allocator->Commit(fence);
243 update_descriptor_queue.Send(*descriptor_template, set);
244 return set;
245}
246
247QuadArrayPass::QuadArrayPass(const VKDevice& device, VKScheduler& scheduler,
248 VKDescriptorPool& descriptor_pool,
249 VKStagingBufferPool& staging_buffer_pool,
250 VKUpdateDescriptorQueue& update_descriptor_queue)
251 : VKComputePass(device, descriptor_pool,
252 {vk::DescriptorSetLayoutBinding(0, vk::DescriptorType::eStorageBuffer, 1,
253 vk::ShaderStageFlagBits::eCompute, nullptr)},
254 {vk::DescriptorUpdateTemplateEntry(0, 0, 1, vk::DescriptorType::eStorageBuffer,
255 0, sizeof(DescriptorUpdateEntry))},
256 {vk::PushConstantRange(vk::ShaderStageFlagBits::eCompute, 0, sizeof(u32))},
257 std::size(quad_array), quad_array),
258 scheduler{scheduler}, staging_buffer_pool{staging_buffer_pool},
259 update_descriptor_queue{update_descriptor_queue} {}
260
261QuadArrayPass::~QuadArrayPass() = default;
262
263std::pair<const vk::Buffer&, vk::DeviceSize> QuadArrayPass::Assemble(u32 num_vertices, u32 first) {
264 const u32 num_triangle_vertices = num_vertices * 6 / 4;
265 const std::size_t staging_size = num_triangle_vertices * sizeof(u32);
266 auto& buffer = staging_buffer_pool.GetUnusedBuffer(staging_size, false);
267
268 update_descriptor_queue.Acquire();
269 update_descriptor_queue.AddBuffer(&*buffer.handle, 0, staging_size);
270 const auto set = CommitDescriptorSet(update_descriptor_queue, scheduler.GetFence());
271
272 scheduler.RequestOutsideRenderPassOperationContext();
273
274 ASSERT(num_vertices % 4 == 0);
275 const u32 num_quads = num_vertices / 4;
276 scheduler.Record([layout = *layout, pipeline = *pipeline, buffer = *buffer.handle, num_quads,
277 first, set](auto cmdbuf, auto& dld) {
278 constexpr u32 dispatch_size = 1024;
279 cmdbuf.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline, dld);
280 cmdbuf.bindDescriptorSets(vk::PipelineBindPoint::eCompute, layout, 0, {set}, {}, dld);
281 cmdbuf.pushConstants(layout, vk::ShaderStageFlagBits::eCompute, 0, sizeof(first), &first,
282 dld);
283 cmdbuf.dispatch(Common::AlignUp(num_quads, dispatch_size) / dispatch_size, 1, 1, dld);
284
285 const vk::BufferMemoryBarrier barrier(
286 vk::AccessFlagBits::eShaderWrite, vk::AccessFlagBits::eVertexAttributeRead,
287 VK_QUEUE_FAMILY_IGNORED, VK_QUEUE_FAMILY_IGNORED, buffer, 0,
288 static_cast<vk::DeviceSize>(num_quads) * 6 * sizeof(u32));
289 cmdbuf.pipelineBarrier(vk::PipelineStageFlagBits::eComputeShader,
290 vk::PipelineStageFlagBits::eVertexInput, {}, {}, {barrier}, {}, dld);
291 });
292 return {*buffer.handle, 0};
293}
294
295Uint8Pass::Uint8Pass(const VKDevice& device, VKScheduler& scheduler,
296 VKDescriptorPool& descriptor_pool, VKStagingBufferPool& staging_buffer_pool,
297 VKUpdateDescriptorQueue& update_descriptor_queue)
298 : VKComputePass(device, descriptor_pool,
299 {vk::DescriptorSetLayoutBinding(0, vk::DescriptorType::eStorageBuffer, 1,
300 vk::ShaderStageFlagBits::eCompute, nullptr),
301 vk::DescriptorSetLayoutBinding(1, vk::DescriptorType::eStorageBuffer, 1,
302 vk::ShaderStageFlagBits::eCompute, nullptr)},
303 {vk::DescriptorUpdateTemplateEntry(0, 0, 2, vk::DescriptorType::eStorageBuffer,
304 0, sizeof(DescriptorUpdateEntry))},
305 {}, std::size(uint8_pass), uint8_pass),
306 scheduler{scheduler}, staging_buffer_pool{staging_buffer_pool},
307 update_descriptor_queue{update_descriptor_queue} {}
308
309Uint8Pass::~Uint8Pass() = default;
310
311std::pair<const vk::Buffer*, u64> Uint8Pass::Assemble(u32 num_vertices, vk::Buffer src_buffer,
312 u64 src_offset) {
313 const auto staging_size = static_cast<u32>(num_vertices * sizeof(u16));
314 auto& buffer = staging_buffer_pool.GetUnusedBuffer(staging_size, false);
315
316 update_descriptor_queue.Acquire();
317 update_descriptor_queue.AddBuffer(&src_buffer, src_offset, num_vertices);
318 update_descriptor_queue.AddBuffer(&*buffer.handle, 0, staging_size);
319 const auto set = CommitDescriptorSet(update_descriptor_queue, scheduler.GetFence());
320
321 scheduler.RequestOutsideRenderPassOperationContext();
322 scheduler.Record([layout = *layout, pipeline = *pipeline, buffer = *buffer.handle, set,
323 num_vertices](auto cmdbuf, auto& dld) {
324 constexpr u32 dispatch_size = 1024;
325 cmdbuf.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline, dld);
326 cmdbuf.bindDescriptorSets(vk::PipelineBindPoint::eCompute, layout, 0, {set}, {}, dld);
327 cmdbuf.dispatch(Common::AlignUp(num_vertices, dispatch_size) / dispatch_size, 1, 1, dld);
328
329 const vk::BufferMemoryBarrier barrier(
330 vk::AccessFlagBits::eShaderWrite, vk::AccessFlagBits::eVertexAttributeRead,
331 VK_QUEUE_FAMILY_IGNORED, VK_QUEUE_FAMILY_IGNORED, buffer, 0,
332 static_cast<vk::DeviceSize>(num_vertices) * sizeof(u16));
333 cmdbuf.pipelineBarrier(vk::PipelineStageFlagBits::eComputeShader,
334 vk::PipelineStageFlagBits::eVertexInput, {}, {}, {barrier}, {}, dld);
335 });
336 return {&*buffer.handle, 0};
337}
338
339} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_compute_pass.h b/src/video_core/renderer_vulkan/vk_compute_pass.h
new file mode 100644
index 000000000..7057eb837
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_compute_pass.h
@@ -0,0 +1,77 @@
1// Copyright 2019 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <optional>
8#include <utility>
9#include <vector>
10#include "common/common_types.h"
11#include "video_core/renderer_vulkan/declarations.h"
12#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
13
14namespace Vulkan {
15
16class VKDevice;
17class VKFence;
18class VKScheduler;
19class VKStagingBufferPool;
20class VKUpdateDescriptorQueue;
21
22class VKComputePass {
23public:
24 explicit VKComputePass(const VKDevice& device, VKDescriptorPool& descriptor_pool,
25 const std::vector<vk::DescriptorSetLayoutBinding>& bindings,
26 const std::vector<vk::DescriptorUpdateTemplateEntry>& templates,
27 const std::vector<vk::PushConstantRange> push_constants,
28 std::size_t code_size, const u8* code);
29 ~VKComputePass();
30
31protected:
32 vk::DescriptorSet CommitDescriptorSet(VKUpdateDescriptorQueue& update_descriptor_queue,
33 VKFence& fence);
34
35 UniqueDescriptorUpdateTemplate descriptor_template;
36 UniquePipelineLayout layout;
37 UniquePipeline pipeline;
38
39private:
40 UniqueDescriptorSetLayout descriptor_set_layout;
41 std::optional<DescriptorAllocator> descriptor_allocator;
42 UniqueShaderModule module;
43};
44
45class QuadArrayPass final : public VKComputePass {
46public:
47 explicit QuadArrayPass(const VKDevice& device, VKScheduler& scheduler,
48 VKDescriptorPool& descriptor_pool,
49 VKStagingBufferPool& staging_buffer_pool,
50 VKUpdateDescriptorQueue& update_descriptor_queue);
51 ~QuadArrayPass();
52
53 std::pair<const vk::Buffer&, vk::DeviceSize> Assemble(u32 num_vertices, u32 first);
54
55private:
56 VKScheduler& scheduler;
57 VKStagingBufferPool& staging_buffer_pool;
58 VKUpdateDescriptorQueue& update_descriptor_queue;
59};
60
61class Uint8Pass final : public VKComputePass {
62public:
63 explicit Uint8Pass(const VKDevice& device, VKScheduler& scheduler,
64 VKDescriptorPool& descriptor_pool, VKStagingBufferPool& staging_buffer_pool,
65 VKUpdateDescriptorQueue& update_descriptor_queue);
66 ~Uint8Pass();
67
68 std::pair<const vk::Buffer*, u64> Assemble(u32 num_vertices, vk::Buffer src_buffer,
69 u64 src_offset);
70
71private:
72 VKScheduler& scheduler;
73 VKStagingBufferPool& staging_buffer_pool;
74 VKUpdateDescriptorQueue& update_descriptor_queue;
75};
76
77} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp b/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp
new file mode 100644
index 000000000..9d5b8de7a
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp
@@ -0,0 +1,112 @@
1// Copyright 2019 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <memory>
6#include <vector>
7
8#include "video_core/renderer_vulkan/declarations.h"
9#include "video_core/renderer_vulkan/vk_compute_pipeline.h"
10#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
11#include "video_core/renderer_vulkan/vk_device.h"
12#include "video_core/renderer_vulkan/vk_pipeline_cache.h"
13#include "video_core/renderer_vulkan/vk_resource_manager.h"
14#include "video_core/renderer_vulkan/vk_scheduler.h"
15#include "video_core/renderer_vulkan/vk_shader_decompiler.h"
16#include "video_core/renderer_vulkan/vk_update_descriptor.h"
17
18namespace Vulkan {
19
20VKComputePipeline::VKComputePipeline(const VKDevice& device, VKScheduler& scheduler,
21 VKDescriptorPool& descriptor_pool,
22 VKUpdateDescriptorQueue& update_descriptor_queue,
23 const SPIRVShader& shader)
24 : device{device}, scheduler{scheduler}, entries{shader.entries},
25 descriptor_set_layout{CreateDescriptorSetLayout()},
26 descriptor_allocator{descriptor_pool, *descriptor_set_layout},
27 update_descriptor_queue{update_descriptor_queue}, layout{CreatePipelineLayout()},
28 descriptor_template{CreateDescriptorUpdateTemplate()},
29 shader_module{CreateShaderModule(shader.code)}, pipeline{CreatePipeline()} {}
30
31VKComputePipeline::~VKComputePipeline() = default;
32
33vk::DescriptorSet VKComputePipeline::CommitDescriptorSet() {
34 if (!descriptor_template) {
35 return {};
36 }
37 const auto set = descriptor_allocator.Commit(scheduler.GetFence());
38 update_descriptor_queue.Send(*descriptor_template, set);
39 return set;
40}
41
42UniqueDescriptorSetLayout VKComputePipeline::CreateDescriptorSetLayout() const {
43 std::vector<vk::DescriptorSetLayoutBinding> bindings;
44 u32 binding = 0;
45 const auto AddBindings = [&](vk::DescriptorType descriptor_type, std::size_t num_entries) {
46 // TODO(Rodrigo): Maybe make individual bindings here?
47 for (u32 bindpoint = 0; bindpoint < static_cast<u32>(num_entries); ++bindpoint) {
48 bindings.emplace_back(binding++, descriptor_type, 1, vk::ShaderStageFlagBits::eCompute,
49 nullptr);
50 }
51 };
52 AddBindings(vk::DescriptorType::eUniformBuffer, entries.const_buffers.size());
53 AddBindings(vk::DescriptorType::eStorageBuffer, entries.global_buffers.size());
54 AddBindings(vk::DescriptorType::eUniformTexelBuffer, entries.texel_buffers.size());
55 AddBindings(vk::DescriptorType::eCombinedImageSampler, entries.samplers.size());
56 AddBindings(vk::DescriptorType::eStorageImage, entries.images.size());
57
58 const vk::DescriptorSetLayoutCreateInfo descriptor_set_layout_ci(
59 {}, static_cast<u32>(bindings.size()), bindings.data());
60
61 const auto dev = device.GetLogical();
62 const auto& dld = device.GetDispatchLoader();
63 return dev.createDescriptorSetLayoutUnique(descriptor_set_layout_ci, nullptr, dld);
64}
65
66UniquePipelineLayout VKComputePipeline::CreatePipelineLayout() const {
67 const vk::PipelineLayoutCreateInfo layout_ci({}, 1, &*descriptor_set_layout, 0, nullptr);
68 const auto dev = device.GetLogical();
69 return dev.createPipelineLayoutUnique(layout_ci, nullptr, device.GetDispatchLoader());
70}
71
72UniqueDescriptorUpdateTemplate VKComputePipeline::CreateDescriptorUpdateTemplate() const {
73 std::vector<vk::DescriptorUpdateTemplateEntry> template_entries;
74 u32 binding = 0;
75 u32 offset = 0;
76 FillDescriptorUpdateTemplateEntries(device, entries, binding, offset, template_entries);
77 if (template_entries.empty()) {
78 // If the shader doesn't use descriptor sets, skip template creation.
79 return UniqueDescriptorUpdateTemplate{};
80 }
81
82 const vk::DescriptorUpdateTemplateCreateInfo template_ci(
83 {}, static_cast<u32>(template_entries.size()), template_entries.data(),
84 vk::DescriptorUpdateTemplateType::eDescriptorSet, *descriptor_set_layout,
85 vk::PipelineBindPoint::eGraphics, *layout, DESCRIPTOR_SET);
86
87 const auto dev = device.GetLogical();
88 const auto& dld = device.GetDispatchLoader();
89 return dev.createDescriptorUpdateTemplateUnique(template_ci, nullptr, dld);
90}
91
92UniqueShaderModule VKComputePipeline::CreateShaderModule(const std::vector<u32>& code) const {
93 const vk::ShaderModuleCreateInfo module_ci({}, code.size() * sizeof(u32), code.data());
94 const auto dev = device.GetLogical();
95 return dev.createShaderModuleUnique(module_ci, nullptr, device.GetDispatchLoader());
96}
97
98UniquePipeline VKComputePipeline::CreatePipeline() const {
99 vk::PipelineShaderStageCreateInfo shader_stage_ci({}, vk::ShaderStageFlagBits::eCompute,
100 *shader_module, "main", nullptr);
101 vk::PipelineShaderStageRequiredSubgroupSizeCreateInfoEXT subgroup_size_ci;
102 subgroup_size_ci.requiredSubgroupSize = GuestWarpSize;
103 if (entries.uses_warps && device.IsGuestWarpSizeSupported(vk::ShaderStageFlagBits::eCompute)) {
104 shader_stage_ci.pNext = &subgroup_size_ci;
105 }
106
107 const vk::ComputePipelineCreateInfo create_info({}, shader_stage_ci, *layout, {}, 0);
108 const auto dev = device.GetLogical();
109 return dev.createComputePipelineUnique({}, create_info, nullptr, device.GetDispatchLoader());
110}
111
112} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_compute_pipeline.h b/src/video_core/renderer_vulkan/vk_compute_pipeline.h
new file mode 100644
index 000000000..22235c6c9
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_compute_pipeline.h
@@ -0,0 +1,66 @@
1// Copyright 2019 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <memory>
8
9#include "common/common_types.h"
10#include "video_core/renderer_vulkan/declarations.h"
11#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
12#include "video_core/renderer_vulkan/vk_shader_decompiler.h"
13
14namespace Vulkan {
15
16class VKDevice;
17class VKScheduler;
18class VKUpdateDescriptorQueue;
19
20class VKComputePipeline final {
21public:
22 explicit VKComputePipeline(const VKDevice& device, VKScheduler& scheduler,
23 VKDescriptorPool& descriptor_pool,
24 VKUpdateDescriptorQueue& update_descriptor_queue,
25 const SPIRVShader& shader);
26 ~VKComputePipeline();
27
28 vk::DescriptorSet CommitDescriptorSet();
29
30 vk::Pipeline GetHandle() const {
31 return *pipeline;
32 }
33
34 vk::PipelineLayout GetLayout() const {
35 return *layout;
36 }
37
38 const ShaderEntries& GetEntries() {
39 return entries;
40 }
41
42private:
43 UniqueDescriptorSetLayout CreateDescriptorSetLayout() const;
44
45 UniquePipelineLayout CreatePipelineLayout() const;
46
47 UniqueDescriptorUpdateTemplate CreateDescriptorUpdateTemplate() const;
48
49 UniqueShaderModule CreateShaderModule(const std::vector<u32>& code) const;
50
51 UniquePipeline CreatePipeline() const;
52
53 const VKDevice& device;
54 VKScheduler& scheduler;
55 ShaderEntries entries;
56
57 UniqueDescriptorSetLayout descriptor_set_layout;
58 DescriptorAllocator descriptor_allocator;
59 VKUpdateDescriptorQueue& update_descriptor_queue;
60 UniquePipelineLayout layout;
61 UniqueDescriptorUpdateTemplate descriptor_template;
62 UniqueShaderModule shader_module;
63 UniquePipeline pipeline;
64};
65
66} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp
new file mode 100644
index 000000000..2e0536bf6
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp
@@ -0,0 +1,271 @@
1// Copyright 2019 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <vector>
6#include "common/assert.h"
7#include "common/common_types.h"
8#include "common/microprofile.h"
9#include "video_core/renderer_vulkan/declarations.h"
10#include "video_core/renderer_vulkan/fixed_pipeline_state.h"
11#include "video_core/renderer_vulkan/maxwell_to_vk.h"
12#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
13#include "video_core/renderer_vulkan/vk_device.h"
14#include "video_core/renderer_vulkan/vk_graphics_pipeline.h"
15#include "video_core/renderer_vulkan/vk_pipeline_cache.h"
16#include "video_core/renderer_vulkan/vk_renderpass_cache.h"
17#include "video_core/renderer_vulkan/vk_scheduler.h"
18#include "video_core/renderer_vulkan/vk_update_descriptor.h"
19
20namespace Vulkan {
21
22MICROPROFILE_DECLARE(Vulkan_PipelineCache);
23
24namespace {
25
26vk::StencilOpState GetStencilFaceState(const FixedPipelineState::StencilFace& face) {
27 return vk::StencilOpState(MaxwellToVK::StencilOp(face.action_stencil_fail),
28 MaxwellToVK::StencilOp(face.action_depth_pass),
29 MaxwellToVK::StencilOp(face.action_depth_fail),
30 MaxwellToVK::ComparisonOp(face.test_func), 0, 0, 0);
31}
32
33bool SupportsPrimitiveRestart(vk::PrimitiveTopology topology) {
34 static constexpr std::array unsupported_topologies = {
35 vk::PrimitiveTopology::ePointList,
36 vk::PrimitiveTopology::eLineList,
37 vk::PrimitiveTopology::eTriangleList,
38 vk::PrimitiveTopology::eLineListWithAdjacency,
39 vk::PrimitiveTopology::eTriangleListWithAdjacency,
40 vk::PrimitiveTopology::ePatchList};
41 return std::find(std::begin(unsupported_topologies), std::end(unsupported_topologies),
42 topology) == std::end(unsupported_topologies);
43}
44
45} // Anonymous namespace
46
47VKGraphicsPipeline::VKGraphicsPipeline(const VKDevice& device, VKScheduler& scheduler,
48 VKDescriptorPool& descriptor_pool,
49 VKUpdateDescriptorQueue& update_descriptor_queue,
50 VKRenderPassCache& renderpass_cache,
51 const GraphicsPipelineCacheKey& key,
52 const std::vector<vk::DescriptorSetLayoutBinding>& bindings,
53 const SPIRVProgram& program)
54 : device{device}, scheduler{scheduler}, fixed_state{key.fixed_state}, hash{key.Hash()},
55 descriptor_set_layout{CreateDescriptorSetLayout(bindings)},
56 descriptor_allocator{descriptor_pool, *descriptor_set_layout},
57 update_descriptor_queue{update_descriptor_queue}, layout{CreatePipelineLayout()},
58 descriptor_template{CreateDescriptorUpdateTemplate(program)}, modules{CreateShaderModules(
59 program)},
60 renderpass{renderpass_cache.GetRenderPass(key.renderpass_params)}, pipeline{CreatePipeline(
61 key.renderpass_params,
62 program)} {}
63
64VKGraphicsPipeline::~VKGraphicsPipeline() = default;
65
66vk::DescriptorSet VKGraphicsPipeline::CommitDescriptorSet() {
67 if (!descriptor_template) {
68 return {};
69 }
70 const auto set = descriptor_allocator.Commit(scheduler.GetFence());
71 update_descriptor_queue.Send(*descriptor_template, set);
72 return set;
73}
74
75UniqueDescriptorSetLayout VKGraphicsPipeline::CreateDescriptorSetLayout(
76 const std::vector<vk::DescriptorSetLayoutBinding>& bindings) const {
77 const vk::DescriptorSetLayoutCreateInfo descriptor_set_layout_ci(
78 {}, static_cast<u32>(bindings.size()), bindings.data());
79
80 const auto dev = device.GetLogical();
81 const auto& dld = device.GetDispatchLoader();
82 return dev.createDescriptorSetLayoutUnique(descriptor_set_layout_ci, nullptr, dld);
83}
84
85UniquePipelineLayout VKGraphicsPipeline::CreatePipelineLayout() const {
86 const vk::PipelineLayoutCreateInfo pipeline_layout_ci({}, 1, &*descriptor_set_layout, 0,
87 nullptr);
88 const auto dev = device.GetLogical();
89 const auto& dld = device.GetDispatchLoader();
90 return dev.createPipelineLayoutUnique(pipeline_layout_ci, nullptr, dld);
91}
92
93UniqueDescriptorUpdateTemplate VKGraphicsPipeline::CreateDescriptorUpdateTemplate(
94 const SPIRVProgram& program) const {
95 std::vector<vk::DescriptorUpdateTemplateEntry> template_entries;
96 u32 binding = 0;
97 u32 offset = 0;
98 for (const auto& stage : program) {
99 if (stage) {
100 FillDescriptorUpdateTemplateEntries(device, stage->entries, binding, offset,
101 template_entries);
102 }
103 }
104 if (template_entries.empty()) {
105 // If the shader doesn't use descriptor sets, skip template creation.
106 return UniqueDescriptorUpdateTemplate{};
107 }
108
109 const vk::DescriptorUpdateTemplateCreateInfo template_ci(
110 {}, static_cast<u32>(template_entries.size()), template_entries.data(),
111 vk::DescriptorUpdateTemplateType::eDescriptorSet, *descriptor_set_layout,
112 vk::PipelineBindPoint::eGraphics, *layout, DESCRIPTOR_SET);
113
114 const auto dev = device.GetLogical();
115 const auto& dld = device.GetDispatchLoader();
116 return dev.createDescriptorUpdateTemplateUnique(template_ci, nullptr, dld);
117}
118
119std::vector<UniqueShaderModule> VKGraphicsPipeline::CreateShaderModules(
120 const SPIRVProgram& program) const {
121 std::vector<UniqueShaderModule> modules;
122 const auto dev = device.GetLogical();
123 const auto& dld = device.GetDispatchLoader();
124 for (std::size_t i = 0; i < Maxwell::MaxShaderStage; ++i) {
125 const auto& stage = program[i];
126 if (!stage) {
127 continue;
128 }
129 const vk::ShaderModuleCreateInfo module_ci({}, stage->code.size() * sizeof(u32),
130 stage->code.data());
131 modules.emplace_back(dev.createShaderModuleUnique(module_ci, nullptr, dld));
132 }
133 return modules;
134}
135
136UniquePipeline VKGraphicsPipeline::CreatePipeline(const RenderPassParams& renderpass_params,
137 const SPIRVProgram& program) const {
138 const auto& vi = fixed_state.vertex_input;
139 const auto& ia = fixed_state.input_assembly;
140 const auto& ds = fixed_state.depth_stencil;
141 const auto& cd = fixed_state.color_blending;
142 const auto& ts = fixed_state.tessellation;
143 const auto& rs = fixed_state.rasterizer;
144
145 std::vector<vk::VertexInputBindingDescription> vertex_bindings;
146 std::vector<vk::VertexInputBindingDivisorDescriptionEXT> vertex_binding_divisors;
147 for (std::size_t i = 0; i < vi.num_bindings; ++i) {
148 const auto& binding = vi.bindings[i];
149 const bool instanced = binding.divisor != 0;
150 const auto rate = instanced ? vk::VertexInputRate::eInstance : vk::VertexInputRate::eVertex;
151 vertex_bindings.emplace_back(binding.index, binding.stride, rate);
152 if (instanced) {
153 vertex_binding_divisors.emplace_back(binding.index, binding.divisor);
154 }
155 }
156
157 std::vector<vk::VertexInputAttributeDescription> vertex_attributes;
158 const auto& input_attributes = program[0]->entries.attributes;
159 for (std::size_t i = 0; i < vi.num_attributes; ++i) {
160 const auto& attribute = vi.attributes[i];
161 if (input_attributes.find(attribute.index) == input_attributes.end()) {
162 // Skip attributes not used by the vertex shaders.
163 continue;
164 }
165 vertex_attributes.emplace_back(attribute.index, attribute.buffer,
166 MaxwellToVK::VertexFormat(attribute.type, attribute.size),
167 attribute.offset);
168 }
169
170 vk::PipelineVertexInputStateCreateInfo vertex_input_ci(
171 {}, static_cast<u32>(vertex_bindings.size()), vertex_bindings.data(),
172 static_cast<u32>(vertex_attributes.size()), vertex_attributes.data());
173
174 const vk::PipelineVertexInputDivisorStateCreateInfoEXT vertex_input_divisor_ci(
175 static_cast<u32>(vertex_binding_divisors.size()), vertex_binding_divisors.data());
176 if (!vertex_binding_divisors.empty()) {
177 vertex_input_ci.pNext = &vertex_input_divisor_ci;
178 }
179
180 const auto primitive_topology = MaxwellToVK::PrimitiveTopology(device, ia.topology);
181 const vk::PipelineInputAssemblyStateCreateInfo input_assembly_ci(
182 {}, primitive_topology,
183 ia.primitive_restart_enable && SupportsPrimitiveRestart(primitive_topology));
184
185 const vk::PipelineTessellationStateCreateInfo tessellation_ci({}, ts.patch_control_points);
186
187 const vk::PipelineViewportStateCreateInfo viewport_ci({}, Maxwell::NumViewports, nullptr,
188 Maxwell::NumViewports, nullptr);
189
190 // TODO(Rodrigo): Find out what's the default register value for front face
191 const vk::PipelineRasterizationStateCreateInfo rasterizer_ci(
192 {}, rs.depth_clamp_enable, false, vk::PolygonMode::eFill,
193 rs.cull_enable ? MaxwellToVK::CullFace(rs.cull_face) : vk::CullModeFlagBits::eNone,
194 rs.cull_enable ? MaxwellToVK::FrontFace(rs.front_face) : vk::FrontFace::eCounterClockwise,
195 rs.depth_bias_enable, 0.0f, 0.0f, 0.0f, 1.0f);
196
197 const vk::PipelineMultisampleStateCreateInfo multisampling_ci(
198 {}, vk::SampleCountFlagBits::e1, false, 0.0f, nullptr, false, false);
199
200 const vk::CompareOp depth_test_compare = ds.depth_test_enable
201 ? MaxwellToVK::ComparisonOp(ds.depth_test_function)
202 : vk::CompareOp::eAlways;
203
204 const vk::PipelineDepthStencilStateCreateInfo depth_stencil_ci(
205 {}, ds.depth_test_enable, ds.depth_write_enable, depth_test_compare, ds.depth_bounds_enable,
206 ds.stencil_enable, GetStencilFaceState(ds.front_stencil),
207 GetStencilFaceState(ds.back_stencil), 0.0f, 0.0f);
208
209 std::array<vk::PipelineColorBlendAttachmentState, Maxwell::NumRenderTargets> cb_attachments;
210 const std::size_t num_attachments =
211 std::min(cd.attachments_count, renderpass_params.color_attachments.size());
212 for (std::size_t i = 0; i < num_attachments; ++i) {
213 constexpr std::array component_table{
214 vk::ColorComponentFlagBits::eR, vk::ColorComponentFlagBits::eG,
215 vk::ColorComponentFlagBits::eB, vk::ColorComponentFlagBits::eA};
216 const auto& blend = cd.attachments[i];
217
218 vk::ColorComponentFlags color_components{};
219 for (std::size_t j = 0; j < component_table.size(); ++j) {
220 if (blend.components[j])
221 color_components |= component_table[j];
222 }
223
224 cb_attachments[i] = vk::PipelineColorBlendAttachmentState(
225 blend.enable, MaxwellToVK::BlendFactor(blend.src_rgb_func),
226 MaxwellToVK::BlendFactor(blend.dst_rgb_func),
227 MaxwellToVK::BlendEquation(blend.rgb_equation),
228 MaxwellToVK::BlendFactor(blend.src_a_func), MaxwellToVK::BlendFactor(blend.dst_a_func),
229 MaxwellToVK::BlendEquation(blend.a_equation), color_components);
230 }
231 const vk::PipelineColorBlendStateCreateInfo color_blending_ci({}, false, vk::LogicOp::eCopy,
232 static_cast<u32>(num_attachments),
233 cb_attachments.data(), {});
234
235 constexpr std::array dynamic_states = {
236 vk::DynamicState::eViewport, vk::DynamicState::eScissor,
237 vk::DynamicState::eDepthBias, vk::DynamicState::eBlendConstants,
238 vk::DynamicState::eDepthBounds, vk::DynamicState::eStencilCompareMask,
239 vk::DynamicState::eStencilWriteMask, vk::DynamicState::eStencilReference};
240 const vk::PipelineDynamicStateCreateInfo dynamic_state_ci(
241 {}, static_cast<u32>(dynamic_states.size()), dynamic_states.data());
242
243 vk::PipelineShaderStageRequiredSubgroupSizeCreateInfoEXT subgroup_size_ci;
244 subgroup_size_ci.requiredSubgroupSize = GuestWarpSize;
245
246 std::vector<vk::PipelineShaderStageCreateInfo> shader_stages;
247 std::size_t module_index = 0;
248 for (std::size_t stage = 0; stage < Maxwell::MaxShaderStage; ++stage) {
249 if (!program[stage]) {
250 continue;
251 }
252 const auto stage_enum = static_cast<Tegra::Engines::ShaderType>(stage);
253 const auto vk_stage = MaxwellToVK::ShaderStage(stage_enum);
254 auto& stage_ci = shader_stages.emplace_back(vk::PipelineShaderStageCreateFlags{}, vk_stage,
255 *modules[module_index++], "main", nullptr);
256 if (program[stage]->entries.uses_warps && device.IsGuestWarpSizeSupported(vk_stage)) {
257 stage_ci.pNext = &subgroup_size_ci;
258 }
259 }
260
261 const vk::GraphicsPipelineCreateInfo create_info(
262 {}, static_cast<u32>(shader_stages.size()), shader_stages.data(), &vertex_input_ci,
263 &input_assembly_ci, &tessellation_ci, &viewport_ci, &rasterizer_ci, &multisampling_ci,
264 &depth_stencil_ci, &color_blending_ci, &dynamic_state_ci, *layout, renderpass, 0, {}, 0);
265
266 const auto dev = device.GetLogical();
267 const auto& dld = device.GetDispatchLoader();
268 return dev.createGraphicsPipelineUnique(nullptr, create_info, nullptr, dld);
269}
270
271} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_graphics_pipeline.h b/src/video_core/renderer_vulkan/vk_graphics_pipeline.h
new file mode 100644
index 000000000..4f5e4ea2d
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_graphics_pipeline.h
@@ -0,0 +1,90 @@
1// Copyright 2019 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <array>
8#include <memory>
9#include <optional>
10#include <unordered_map>
11#include <vector>
12
13#include "video_core/engines/maxwell_3d.h"
14#include "video_core/renderer_vulkan/declarations.h"
15#include "video_core/renderer_vulkan/fixed_pipeline_state.h"
16#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
17#include "video_core/renderer_vulkan/vk_renderpass_cache.h"
18#include "video_core/renderer_vulkan/vk_resource_manager.h"
19#include "video_core/renderer_vulkan/vk_shader_decompiler.h"
20
21namespace Vulkan {
22
23using Maxwell = Tegra::Engines::Maxwell3D::Regs;
24
25struct GraphicsPipelineCacheKey;
26
27class VKDescriptorPool;
28class VKDevice;
29class VKRenderPassCache;
30class VKScheduler;
31class VKUpdateDescriptorQueue;
32
33using SPIRVProgram = std::array<std::optional<SPIRVShader>, Maxwell::MaxShaderStage>;
34
35class VKGraphicsPipeline final {
36public:
37 explicit VKGraphicsPipeline(const VKDevice& device, VKScheduler& scheduler,
38 VKDescriptorPool& descriptor_pool,
39 VKUpdateDescriptorQueue& update_descriptor_queue,
40 VKRenderPassCache& renderpass_cache,
41 const GraphicsPipelineCacheKey& key,
42 const std::vector<vk::DescriptorSetLayoutBinding>& bindings,
43 const SPIRVProgram& program);
44 ~VKGraphicsPipeline();
45
46 vk::DescriptorSet CommitDescriptorSet();
47
48 vk::Pipeline GetHandle() const {
49 return *pipeline;
50 }
51
52 vk::PipelineLayout GetLayout() const {
53 return *layout;
54 }
55
56 vk::RenderPass GetRenderPass() const {
57 return renderpass;
58 }
59
60private:
61 UniqueDescriptorSetLayout CreateDescriptorSetLayout(
62 const std::vector<vk::DescriptorSetLayoutBinding>& bindings) const;
63
64 UniquePipelineLayout CreatePipelineLayout() const;
65
66 UniqueDescriptorUpdateTemplate CreateDescriptorUpdateTemplate(
67 const SPIRVProgram& program) const;
68
69 std::vector<UniqueShaderModule> CreateShaderModules(const SPIRVProgram& program) const;
70
71 UniquePipeline CreatePipeline(const RenderPassParams& renderpass_params,
72 const SPIRVProgram& program) const;
73
74 const VKDevice& device;
75 VKScheduler& scheduler;
76 const FixedPipelineState fixed_state;
77 const u64 hash;
78
79 UniqueDescriptorSetLayout descriptor_set_layout;
80 DescriptorAllocator descriptor_allocator;
81 VKUpdateDescriptorQueue& update_descriptor_queue;
82 UniquePipelineLayout layout;
83 UniqueDescriptorUpdateTemplate descriptor_template;
84 std::vector<UniqueShaderModule> modules;
85
86 vk::RenderPass renderpass;
87 UniquePipeline pipeline;
88};
89
90} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
new file mode 100644
index 000000000..48e23d4cd
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
@@ -0,0 +1,395 @@
1// Copyright 2019 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <algorithm>
6#include <cstddef>
7#include <memory>
8#include <vector>
9
10#include "common/microprofile.h"
11#include "core/core.h"
12#include "core/memory.h"
13#include "video_core/engines/kepler_compute.h"
14#include "video_core/engines/maxwell_3d.h"
15#include "video_core/memory_manager.h"
16#include "video_core/renderer_vulkan/declarations.h"
17#include "video_core/renderer_vulkan/fixed_pipeline_state.h"
18#include "video_core/renderer_vulkan/maxwell_to_vk.h"
19#include "video_core/renderer_vulkan/vk_compute_pipeline.h"
20#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
21#include "video_core/renderer_vulkan/vk_device.h"
22#include "video_core/renderer_vulkan/vk_graphics_pipeline.h"
23#include "video_core/renderer_vulkan/vk_pipeline_cache.h"
24#include "video_core/renderer_vulkan/vk_rasterizer.h"
25#include "video_core/renderer_vulkan/vk_renderpass_cache.h"
26#include "video_core/renderer_vulkan/vk_resource_manager.h"
27#include "video_core/renderer_vulkan/vk_scheduler.h"
28#include "video_core/renderer_vulkan/vk_update_descriptor.h"
29#include "video_core/shader/compiler_settings.h"
30
31namespace Vulkan {
32
33MICROPROFILE_DECLARE(Vulkan_PipelineCache);
34
35using Tegra::Engines::ShaderType;
36
37namespace {
38
39constexpr VideoCommon::Shader::CompilerSettings compiler_settings{
40 VideoCommon::Shader::CompileDepth::FullDecompile};
41
42/// Gets the address for the specified shader stage program
43GPUVAddr GetShaderAddress(Core::System& system, Maxwell::ShaderProgram program) {
44 const auto& gpu{system.GPU().Maxwell3D()};
45 const auto& shader_config{gpu.regs.shader_config[static_cast<std::size_t>(program)]};
46 return gpu.regs.code_address.CodeAddress() + shader_config.offset;
47}
48
49/// Gets if the current instruction offset is a scheduler instruction
50constexpr bool IsSchedInstruction(std::size_t offset, std::size_t main_offset) {
51 // Sched instructions appear once every 4 instructions.
52 constexpr std::size_t SchedPeriod = 4;
53 const std::size_t absolute_offset = offset - main_offset;
54 return (absolute_offset % SchedPeriod) == 0;
55}
56
57/// Calculates the size of a program stream
58std::size_t CalculateProgramSize(const ProgramCode& program, bool is_compute) {
59 const std::size_t start_offset = is_compute ? 0 : 10;
60 // This is the encoded version of BRA that jumps to itself. All Nvidia
61 // shaders end with one.
62 constexpr u64 self_jumping_branch = 0xE2400FFFFF07000FULL;
63 constexpr u64 mask = 0xFFFFFFFFFF7FFFFFULL;
64 std::size_t offset = start_offset;
65 while (offset < program.size()) {
66 const u64 instruction = program[offset];
67 if (!IsSchedInstruction(offset, start_offset)) {
68 if ((instruction & mask) == self_jumping_branch) {
69 // End on Maxwell's "nop" instruction
70 break;
71 }
72 if (instruction == 0) {
73 break;
74 }
75 }
76 ++offset;
77 }
78 // The last instruction is included in the program size
79 return std::min(offset + 1, program.size());
80}
81
82/// Gets the shader program code from memory for the specified address
83ProgramCode GetShaderCode(Tegra::MemoryManager& memory_manager, const GPUVAddr gpu_addr,
84 const u8* host_ptr, bool is_compute) {
85 ProgramCode program_code(VideoCommon::Shader::MAX_PROGRAM_LENGTH);
86 ASSERT_OR_EXECUTE(host_ptr != nullptr, {
87 std::fill(program_code.begin(), program_code.end(), 0);
88 return program_code;
89 });
90 memory_manager.ReadBlockUnsafe(gpu_addr, program_code.data(),
91 program_code.size() * sizeof(u64));
92 program_code.resize(CalculateProgramSize(program_code, is_compute));
93 return program_code;
94}
95
96constexpr std::size_t GetStageFromProgram(std::size_t program) {
97 return program == 0 ? 0 : program - 1;
98}
99
100constexpr ShaderType GetStageFromProgram(Maxwell::ShaderProgram program) {
101 return static_cast<ShaderType>(GetStageFromProgram(static_cast<std::size_t>(program)));
102}
103
104ShaderType GetShaderType(Maxwell::ShaderProgram program) {
105 switch (program) {
106 case Maxwell::ShaderProgram::VertexB:
107 return ShaderType::Vertex;
108 case Maxwell::ShaderProgram::TesselationControl:
109 return ShaderType::TesselationControl;
110 case Maxwell::ShaderProgram::TesselationEval:
111 return ShaderType::TesselationEval;
112 case Maxwell::ShaderProgram::Geometry:
113 return ShaderType::Geometry;
114 case Maxwell::ShaderProgram::Fragment:
115 return ShaderType::Fragment;
116 default:
117 UNIMPLEMENTED_MSG("program={}", static_cast<u32>(program));
118 return ShaderType::Vertex;
119 }
120}
121
122u32 FillDescriptorLayout(const ShaderEntries& entries,
123 std::vector<vk::DescriptorSetLayoutBinding>& bindings,
124 Maxwell::ShaderProgram program_type, u32 base_binding) {
125 const ShaderType stage = GetStageFromProgram(program_type);
126 const vk::ShaderStageFlags stage_flags = MaxwellToVK::ShaderStage(stage);
127
128 u32 binding = base_binding;
129 const auto AddBindings = [&](vk::DescriptorType descriptor_type, std::size_t num_entries) {
130 for (std::size_t i = 0; i < num_entries; ++i) {
131 bindings.emplace_back(binding++, descriptor_type, 1, stage_flags, nullptr);
132 }
133 };
134 AddBindings(vk::DescriptorType::eUniformBuffer, entries.const_buffers.size());
135 AddBindings(vk::DescriptorType::eStorageBuffer, entries.global_buffers.size());
136 AddBindings(vk::DescriptorType::eUniformTexelBuffer, entries.texel_buffers.size());
137 AddBindings(vk::DescriptorType::eCombinedImageSampler, entries.samplers.size());
138 AddBindings(vk::DescriptorType::eStorageImage, entries.images.size());
139 return binding;
140}
141
142} // Anonymous namespace
143
144CachedShader::CachedShader(Core::System& system, Tegra::Engines::ShaderType stage,
145 GPUVAddr gpu_addr, VAddr cpu_addr, u8* host_ptr,
146 ProgramCode program_code, u32 main_offset)
147 : RasterizerCacheObject{host_ptr}, gpu_addr{gpu_addr}, cpu_addr{cpu_addr},
148 program_code{std::move(program_code)}, locker{stage, GetEngine(system, stage)},
149 shader_ir{this->program_code, main_offset, compiler_settings, locker},
150 entries{GenerateShaderEntries(shader_ir)} {}
151
152CachedShader::~CachedShader() = default;
153
154Tegra::Engines::ConstBufferEngineInterface& CachedShader::GetEngine(
155 Core::System& system, Tegra::Engines::ShaderType stage) {
156 if (stage == Tegra::Engines::ShaderType::Compute) {
157 return system.GPU().KeplerCompute();
158 } else {
159 return system.GPU().Maxwell3D();
160 }
161}
162
163VKPipelineCache::VKPipelineCache(Core::System& system, RasterizerVulkan& rasterizer,
164 const VKDevice& device, VKScheduler& scheduler,
165 VKDescriptorPool& descriptor_pool,
166 VKUpdateDescriptorQueue& update_descriptor_queue)
167 : RasterizerCache{rasterizer}, system{system}, device{device}, scheduler{scheduler},
168 descriptor_pool{descriptor_pool}, update_descriptor_queue{update_descriptor_queue},
169 renderpass_cache(device) {}
170
171VKPipelineCache::~VKPipelineCache() = default;
172
173std::array<Shader, Maxwell::MaxShaderProgram> VKPipelineCache::GetShaders() {
174 const auto& gpu = system.GPU().Maxwell3D();
175 auto& dirty = system.GPU().Maxwell3D().dirty.shaders;
176 if (!dirty) {
177 return last_shaders;
178 }
179 dirty = false;
180
181 std::array<Shader, Maxwell::MaxShaderProgram> shaders;
182 for (std::size_t index = 0; index < Maxwell::MaxShaderProgram; ++index) {
183 const auto& shader_config = gpu.regs.shader_config[index];
184 const auto program{static_cast<Maxwell::ShaderProgram>(index)};
185
186 // Skip stages that are not enabled
187 if (!gpu.regs.IsShaderConfigEnabled(index)) {
188 continue;
189 }
190
191 auto& memory_manager{system.GPU().MemoryManager()};
192 const GPUVAddr program_addr{GetShaderAddress(system, program)};
193 const auto host_ptr{memory_manager.GetPointer(program_addr)};
194 auto shader = TryGet(host_ptr);
195 if (!shader) {
196 // No shader found - create a new one
197 constexpr u32 stage_offset = 10;
198 const auto stage = static_cast<Tegra::Engines::ShaderType>(index == 0 ? 0 : index - 1);
199 auto code = GetShaderCode(memory_manager, program_addr, host_ptr, false);
200
201 const std::optional cpu_addr = memory_manager.GpuToCpuAddress(program_addr);
202 ASSERT(cpu_addr);
203
204 shader = std::make_shared<CachedShader>(system, stage, program_addr, *cpu_addr,
205 host_ptr, std::move(code), stage_offset);
206 Register(shader);
207 }
208 shaders[index] = std::move(shader);
209 }
210 return last_shaders = shaders;
211}
212
213VKGraphicsPipeline& VKPipelineCache::GetGraphicsPipeline(const GraphicsPipelineCacheKey& key) {
214 MICROPROFILE_SCOPE(Vulkan_PipelineCache);
215
216 if (last_graphics_pipeline && last_graphics_key == key) {
217 return *last_graphics_pipeline;
218 }
219 last_graphics_key = key;
220
221 const auto [pair, is_cache_miss] = graphics_cache.try_emplace(key);
222 auto& entry = pair->second;
223 if (is_cache_miss) {
224 LOG_INFO(Render_Vulkan, "Compile 0x{:016X}", key.Hash());
225 const auto [program, bindings] = DecompileShaders(key);
226 entry = std::make_unique<VKGraphicsPipeline>(device, scheduler, descriptor_pool,
227 update_descriptor_queue, renderpass_cache, key,
228 bindings, program);
229 }
230 return *(last_graphics_pipeline = entry.get());
231}
232
233VKComputePipeline& VKPipelineCache::GetComputePipeline(const ComputePipelineCacheKey& key) {
234 MICROPROFILE_SCOPE(Vulkan_PipelineCache);
235
236 const auto [pair, is_cache_miss] = compute_cache.try_emplace(key);
237 auto& entry = pair->second;
238 if (!is_cache_miss) {
239 return *entry;
240 }
241 LOG_INFO(Render_Vulkan, "Compile 0x{:016X}", key.Hash());
242
243 auto& memory_manager = system.GPU().MemoryManager();
244 const auto program_addr = key.shader;
245 const auto host_ptr = memory_manager.GetPointer(program_addr);
246
247 auto shader = TryGet(host_ptr);
248 if (!shader) {
249 // No shader found - create a new one
250 const auto cpu_addr = memory_manager.GpuToCpuAddress(program_addr);
251 ASSERT(cpu_addr);
252
253 auto code = GetShaderCode(memory_manager, program_addr, host_ptr, true);
254 constexpr u32 kernel_main_offset = 0;
255 shader = std::make_shared<CachedShader>(system, Tegra::Engines::ShaderType::Compute,
256 program_addr, *cpu_addr, host_ptr, std::move(code),
257 kernel_main_offset);
258 Register(shader);
259 }
260
261 Specialization specialization;
262 specialization.workgroup_size = key.workgroup_size;
263 specialization.shared_memory_size = key.shared_memory_size;
264
265 const SPIRVShader spirv_shader{
266 Decompile(device, shader->GetIR(), ShaderType::Compute, specialization),
267 shader->GetEntries()};
268 entry = std::make_unique<VKComputePipeline>(device, scheduler, descriptor_pool,
269 update_descriptor_queue, spirv_shader);
270 return *entry;
271}
272
273void VKPipelineCache::Unregister(const Shader& shader) {
274 bool finished = false;
275 const auto Finish = [&] {
276 // TODO(Rodrigo): Instead of finishing here, wait for the fences that use this pipeline and
277 // flush.
278 if (finished) {
279 return;
280 }
281 finished = true;
282 scheduler.Finish();
283 };
284
285 const GPUVAddr invalidated_addr = shader->GetGpuAddr();
286 for (auto it = graphics_cache.begin(); it != graphics_cache.end();) {
287 auto& entry = it->first;
288 if (std::find(entry.shaders.begin(), entry.shaders.end(), invalidated_addr) ==
289 entry.shaders.end()) {
290 ++it;
291 continue;
292 }
293 Finish();
294 it = graphics_cache.erase(it);
295 }
296 for (auto it = compute_cache.begin(); it != compute_cache.end();) {
297 auto& entry = it->first;
298 if (entry.shader != invalidated_addr) {
299 ++it;
300 continue;
301 }
302 Finish();
303 it = compute_cache.erase(it);
304 }
305
306 RasterizerCache::Unregister(shader);
307}
308
309std::pair<SPIRVProgram, std::vector<vk::DescriptorSetLayoutBinding>>
310VKPipelineCache::DecompileShaders(const GraphicsPipelineCacheKey& key) {
311 const auto& fixed_state = key.fixed_state;
312 auto& memory_manager = system.GPU().MemoryManager();
313 const auto& gpu = system.GPU().Maxwell3D();
314
315 Specialization specialization;
316 specialization.primitive_topology = fixed_state.input_assembly.topology;
317 if (specialization.primitive_topology == Maxwell::PrimitiveTopology::Points) {
318 ASSERT(fixed_state.input_assembly.point_size != 0.0f);
319 specialization.point_size = fixed_state.input_assembly.point_size;
320 }
321 for (std::size_t i = 0; i < Maxwell::NumVertexAttributes; ++i) {
322 specialization.attribute_types[i] = fixed_state.vertex_input.attributes[i].type;
323 }
324 specialization.ndc_minus_one_to_one = fixed_state.rasterizer.ndc_minus_one_to_one;
325 specialization.tessellation.primitive = fixed_state.tessellation.primitive;
326 specialization.tessellation.spacing = fixed_state.tessellation.spacing;
327 specialization.tessellation.clockwise = fixed_state.tessellation.clockwise;
328 for (const auto& rt : key.renderpass_params.color_attachments) {
329 specialization.enabled_rendertargets.set(rt.index);
330 }
331
332 SPIRVProgram program;
333 std::vector<vk::DescriptorSetLayoutBinding> bindings;
334
335 for (std::size_t index = 0; index < Maxwell::MaxShaderProgram; ++index) {
336 const auto program_enum = static_cast<Maxwell::ShaderProgram>(index);
337
338 // Skip stages that are not enabled
339 if (!gpu.regs.IsShaderConfigEnabled(index)) {
340 continue;
341 }
342
343 const GPUVAddr gpu_addr = GetShaderAddress(system, program_enum);
344 const auto host_ptr = memory_manager.GetPointer(gpu_addr);
345 const auto shader = TryGet(host_ptr);
346 ASSERT(shader);
347
348 const std::size_t stage = index == 0 ? 0 : index - 1; // Stage indices are 0 - 5
349 const auto program_type = GetShaderType(program_enum);
350 const auto& entries = shader->GetEntries();
351 program[stage] = {Decompile(device, shader->GetIR(), program_type, specialization),
352 entries};
353
354 if (program_enum == Maxwell::ShaderProgram::VertexA) {
355 // VertexB was combined with VertexA, so we skip the VertexB iteration
356 ++index;
357 }
358
359 const u32 old_binding = specialization.base_binding;
360 specialization.base_binding =
361 FillDescriptorLayout(entries, bindings, program_enum, specialization.base_binding);
362 ASSERT(old_binding + entries.NumBindings() == specialization.base_binding);
363 }
364 return {std::move(program), std::move(bindings)};
365}
366
367void FillDescriptorUpdateTemplateEntries(
368 const VKDevice& device, const ShaderEntries& entries, u32& binding, u32& offset,
369 std::vector<vk::DescriptorUpdateTemplateEntry>& template_entries) {
370 static constexpr auto entry_size = static_cast<u32>(sizeof(DescriptorUpdateEntry));
371 const auto AddEntry = [&](vk::DescriptorType descriptor_type, std::size_t count_) {
372 const u32 count = static_cast<u32>(count_);
373 if (descriptor_type == vk::DescriptorType::eUniformTexelBuffer &&
374 device.GetDriverID() == vk::DriverIdKHR::eNvidiaProprietary) {
375 // Nvidia has a bug where updating multiple uniform texels at once causes the driver to
376 // crash.
377 for (u32 i = 0; i < count; ++i) {
378 template_entries.emplace_back(binding + i, 0, 1, descriptor_type,
379 offset + i * entry_size, entry_size);
380 }
381 } else if (count != 0) {
382 template_entries.emplace_back(binding, 0, count, descriptor_type, offset, entry_size);
383 }
384 offset += count * entry_size;
385 binding += count;
386 };
387
388 AddEntry(vk::DescriptorType::eUniformBuffer, entries.const_buffers.size());
389 AddEntry(vk::DescriptorType::eStorageBuffer, entries.global_buffers.size());
390 AddEntry(vk::DescriptorType::eUniformTexelBuffer, entries.texel_buffers.size());
391 AddEntry(vk::DescriptorType::eCombinedImageSampler, entries.samplers.size());
392 AddEntry(vk::DescriptorType::eStorageImage, entries.images.size());
393}
394
395} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.h b/src/video_core/renderer_vulkan/vk_pipeline_cache.h
new file mode 100644
index 000000000..8678fc9c3
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.h
@@ -0,0 +1,200 @@
1// Copyright 2019 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <array>
8#include <cstddef>
9#include <memory>
10#include <tuple>
11#include <type_traits>
12#include <unordered_map>
13#include <utility>
14#include <vector>
15
16#include <boost/functional/hash.hpp>
17
18#include "common/common_types.h"
19#include "video_core/engines/const_buffer_engine_interface.h"
20#include "video_core/engines/maxwell_3d.h"
21#include "video_core/rasterizer_cache.h"
22#include "video_core/renderer_vulkan/declarations.h"
23#include "video_core/renderer_vulkan/fixed_pipeline_state.h"
24#include "video_core/renderer_vulkan/vk_graphics_pipeline.h"
25#include "video_core/renderer_vulkan/vk_renderpass_cache.h"
26#include "video_core/renderer_vulkan/vk_resource_manager.h"
27#include "video_core/renderer_vulkan/vk_shader_decompiler.h"
28#include "video_core/shader/const_buffer_locker.h"
29#include "video_core/shader/shader_ir.h"
30#include "video_core/surface.h"
31
32namespace Core {
33class System;
34}
35
36namespace Vulkan {
37
38class RasterizerVulkan;
39class VKComputePipeline;
40class VKDescriptorPool;
41class VKDevice;
42class VKFence;
43class VKScheduler;
44class VKUpdateDescriptorQueue;
45
46class CachedShader;
47using Shader = std::shared_ptr<CachedShader>;
48using Maxwell = Tegra::Engines::Maxwell3D::Regs;
49
50using ProgramCode = std::vector<u64>;
51
52struct GraphicsPipelineCacheKey {
53 FixedPipelineState fixed_state;
54 std::array<GPUVAddr, Maxwell::MaxShaderProgram> shaders;
55 RenderPassParams renderpass_params;
56
57 std::size_t Hash() const noexcept {
58 std::size_t hash = fixed_state.Hash();
59 for (const auto& shader : shaders) {
60 boost::hash_combine(hash, shader);
61 }
62 boost::hash_combine(hash, renderpass_params.Hash());
63 return hash;
64 }
65
66 bool operator==(const GraphicsPipelineCacheKey& rhs) const noexcept {
67 return std::tie(fixed_state, shaders, renderpass_params) ==
68 std::tie(rhs.fixed_state, rhs.shaders, rhs.renderpass_params);
69 }
70};
71
72struct ComputePipelineCacheKey {
73 GPUVAddr shader{};
74 u32 shared_memory_size{};
75 std::array<u32, 3> workgroup_size{};
76
77 std::size_t Hash() const noexcept {
78 return static_cast<std::size_t>(shader) ^
79 ((static_cast<std::size_t>(shared_memory_size) >> 7) << 40) ^
80 static_cast<std::size_t>(workgroup_size[0]) ^
81 (static_cast<std::size_t>(workgroup_size[1]) << 16) ^
82 (static_cast<std::size_t>(workgroup_size[2]) << 24);
83 }
84
85 bool operator==(const ComputePipelineCacheKey& rhs) const noexcept {
86 return std::tie(shader, shared_memory_size, workgroup_size) ==
87 std::tie(rhs.shader, rhs.shared_memory_size, rhs.workgroup_size);
88 }
89};
90
91} // namespace Vulkan
92
93namespace std {
94
95template <>
96struct hash<Vulkan::GraphicsPipelineCacheKey> {
97 std::size_t operator()(const Vulkan::GraphicsPipelineCacheKey& k) const noexcept {
98 return k.Hash();
99 }
100};
101
102template <>
103struct hash<Vulkan::ComputePipelineCacheKey> {
104 std::size_t operator()(const Vulkan::ComputePipelineCacheKey& k) const noexcept {
105 return k.Hash();
106 }
107};
108
109} // namespace std
110
111namespace Vulkan {
112
113class CachedShader final : public RasterizerCacheObject {
114public:
115 explicit CachedShader(Core::System& system, Tegra::Engines::ShaderType stage, GPUVAddr gpu_addr,
116 VAddr cpu_addr, u8* host_ptr, ProgramCode program_code, u32 main_offset);
117 ~CachedShader();
118
119 GPUVAddr GetGpuAddr() const {
120 return gpu_addr;
121 }
122
123 VAddr GetCpuAddr() const override {
124 return cpu_addr;
125 }
126
127 std::size_t GetSizeInBytes() const override {
128 return program_code.size() * sizeof(u64);
129 }
130
131 VideoCommon::Shader::ShaderIR& GetIR() {
132 return shader_ir;
133 }
134
135 const VideoCommon::Shader::ShaderIR& GetIR() const {
136 return shader_ir;
137 }
138
139 const ShaderEntries& GetEntries() const {
140 return entries;
141 }
142
143private:
144 static Tegra::Engines::ConstBufferEngineInterface& GetEngine(Core::System& system,
145 Tegra::Engines::ShaderType stage);
146
147 GPUVAddr gpu_addr{};
148 VAddr cpu_addr{};
149 ProgramCode program_code;
150 VideoCommon::Shader::ConstBufferLocker locker;
151 VideoCommon::Shader::ShaderIR shader_ir;
152 ShaderEntries entries;
153};
154
155class VKPipelineCache final : public RasterizerCache<Shader> {
156public:
157 explicit VKPipelineCache(Core::System& system, RasterizerVulkan& rasterizer,
158 const VKDevice& device, VKScheduler& scheduler,
159 VKDescriptorPool& descriptor_pool,
160 VKUpdateDescriptorQueue& update_descriptor_queue);
161 ~VKPipelineCache();
162
163 std::array<Shader, Maxwell::MaxShaderProgram> GetShaders();
164
165 VKGraphicsPipeline& GetGraphicsPipeline(const GraphicsPipelineCacheKey& key);
166
167 VKComputePipeline& GetComputePipeline(const ComputePipelineCacheKey& key);
168
169protected:
170 void Unregister(const Shader& shader) override;
171
172 void FlushObjectInner(const Shader& object) override {}
173
174private:
175 std::pair<SPIRVProgram, std::vector<vk::DescriptorSetLayoutBinding>> DecompileShaders(
176 const GraphicsPipelineCacheKey& key);
177
178 Core::System& system;
179 const VKDevice& device;
180 VKScheduler& scheduler;
181 VKDescriptorPool& descriptor_pool;
182 VKUpdateDescriptorQueue& update_descriptor_queue;
183
184 VKRenderPassCache renderpass_cache;
185
186 std::array<Shader, Maxwell::MaxShaderProgram> last_shaders;
187
188 GraphicsPipelineCacheKey last_graphics_key;
189 VKGraphicsPipeline* last_graphics_pipeline = nullptr;
190
191 std::unordered_map<GraphicsPipelineCacheKey, std::unique_ptr<VKGraphicsPipeline>>
192 graphics_cache;
193 std::unordered_map<ComputePipelineCacheKey, std::unique_ptr<VKComputePipeline>> compute_cache;
194};
195
196void FillDescriptorUpdateTemplateEntries(
197 const VKDevice& device, const ShaderEntries& entries, u32& binding, u32& offset,
198 std::vector<vk::DescriptorUpdateTemplateEntry>& template_entries);
199
200} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.h b/src/video_core/renderer_vulkan/vk_rasterizer.h
new file mode 100644
index 000000000..fc324952b
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_rasterizer.h
@@ -0,0 +1,13 @@
1// Copyright 2019 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include "video_core/rasterizer_interface.h"
8
9namespace Vulkan {
10
11class RasterizerVulkan : public VideoCore::RasterizerInterface {};
12
13} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_shader_util.cpp b/src/video_core/renderer_vulkan/vk_shader_util.cpp
new file mode 100644
index 000000000..b97c4cb3d
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_shader_util.cpp
@@ -0,0 +1,34 @@
1// Copyright 2018 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <cstring>
6#include <memory>
7#include <vector>
8#include "common/alignment.h"
9#include "common/assert.h"
10#include "common/common_types.h"
11#include "video_core/renderer_vulkan/declarations.h"
12#include "video_core/renderer_vulkan/vk_device.h"
13#include "video_core/renderer_vulkan/vk_shader_util.h"
14
15namespace Vulkan {
16
17UniqueShaderModule BuildShader(const VKDevice& device, std::size_t code_size, const u8* code_data) {
18 // Avoid undefined behavior by copying to a staging allocation
19 ASSERT(code_size % sizeof(u32) == 0);
20 const auto data = std::make_unique<u32[]>(code_size / sizeof(u32));
21 std::memcpy(data.get(), code_data, code_size);
22
23 const auto dev = device.GetLogical();
24 const auto& dld = device.GetDispatchLoader();
25 const vk::ShaderModuleCreateInfo shader_ci({}, code_size, data.get());
26 vk::ShaderModule shader_module;
27 if (dev.createShaderModule(&shader_ci, nullptr, &shader_module, dld) != vk::Result::eSuccess) {
28 UNREACHABLE_MSG("Shader module failed to build!");
29 }
30
31 return UniqueShaderModule(shader_module, vk::ObjectDestroy(dev, nullptr, dld));
32}
33
34} // namespace Vulkan
diff --git a/src/video_core/renderer_vulkan/vk_shader_util.h b/src/video_core/renderer_vulkan/vk_shader_util.h
new file mode 100644
index 000000000..c06d65970
--- /dev/null
+++ b/src/video_core/renderer_vulkan/vk_shader_util.h
@@ -0,0 +1,17 @@
1// Copyright 2018 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <vector>
8#include "common/common_types.h"
9#include "video_core/renderer_vulkan/declarations.h"
10
11namespace Vulkan {
12
13class VKDevice;
14
15UniqueShaderModule BuildShader(const VKDevice& device, std::size_t code_size, const u8* code_data);
16
17} // namespace Vulkan
HIC SVNT DRACONES