CJBok
6 years ago
14 changed files with 1643 additions and 10 deletions
-
11src/video_core/CMakeLists.txt
-
18src/video_core/renderer_vulkan/fixed_pipeline_state.cpp
-
10src/video_core/renderer_vulkan/fixed_pipeline_state.h
-
339src/video_core/renderer_vulkan/vk_compute_pass.cpp
-
77src/video_core/renderer_vulkan/vk_compute_pass.h
-
112src/video_core/renderer_vulkan/vk_compute_pipeline.cpp
-
66src/video_core/renderer_vulkan/vk_compute_pipeline.h
-
271src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp
-
90src/video_core/renderer_vulkan/vk_graphics_pipeline.h
-
395src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
-
200src/video_core/renderer_vulkan/vk_pipeline_cache.h
-
13src/video_core/renderer_vulkan/vk_rasterizer.h
-
34src/video_core/renderer_vulkan/vk_shader_util.cpp
-
17src/video_core/renderer_vulkan/vk_shader_util.h
@ -0,0 +1,339 @@ |
|||||
|
// Copyright 2019 yuzu Emulator Project
|
||||
|
// Licensed under GPLv2 or any later version
|
||||
|
// Refer to the license.txt file included.
|
||||
|
|
||||
|
#include <cstring>
|
||||
|
#include <memory>
|
||||
|
#include <optional>
|
||||
|
#include <utility>
|
||||
|
#include <vector>
|
||||
|
#include "common/alignment.h"
|
||||
|
#include "common/assert.h"
|
||||
|
#include "common/common_types.h"
|
||||
|
#include "video_core/renderer_vulkan/declarations.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_compute_pass.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_device.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_scheduler.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_staging_buffer_pool.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_update_descriptor.h"
|
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
namespace { |
||||
|
|
||||
|
// Quad array SPIR-V module. Generated from the "shaders/" directory, read the instructions there.
|
||||
|
constexpr u8 quad_array[] = { |
||||
|
0x03, 0x02, 0x23, 0x07, 0x00, 0x00, 0x01, 0x00, 0x07, 0x00, 0x08, 0x00, 0x54, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, 0x01, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x06, 0x00, |
||||
|
0x01, 0x00, 0x00, 0x00, 0x47, 0x4c, 0x53, 0x4c, 0x2e, 0x73, 0x74, 0x64, 0x2e, 0x34, 0x35, 0x30, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x0e, 0x00, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, |
||||
|
0x0f, 0x00, 0x06, 0x00, 0x05, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x6d, 0x61, 0x69, 0x6e, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x10, 0x00, 0x06, 0x00, 0x04, 0x00, 0x00, 0x00, |
||||
|
0x11, 0x00, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, |
||||
|
0x47, 0x00, 0x04, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x1c, 0x00, 0x00, 0x00, |
||||
|
0x47, 0x00, 0x04, 0x00, 0x13, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, |
||||
|
0x48, 0x00, 0x05, 0x00, 0x14, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x03, 0x00, 0x14, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, |
||||
|
0x47, 0x00, 0x04, 0x00, 0x16, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, |
||||
|
0x47, 0x00, 0x04, 0x00, 0x16, 0x00, 0x00, 0x00, 0x21, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, |
||||
|
0x48, 0x00, 0x05, 0x00, 0x29, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x03, 0x00, 0x29, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, |
||||
|
0x47, 0x00, 0x04, 0x00, 0x4a, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00, |
||||
|
0x13, 0x00, 0x02, 0x00, 0x02, 0x00, 0x00, 0x00, 0x21, 0x00, 0x03, 0x00, 0x03, 0x00, 0x00, 0x00, |
||||
|
0x02, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x07, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x17, 0x00, 0x04, 0x00, 0x09, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, |
||||
|
0x03, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x0a, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, |
||||
|
0x09, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, 0x0a, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, |
||||
|
0x01, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x0d, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x03, 0x00, 0x13, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, |
||||
|
0x1e, 0x00, 0x03, 0x00, 0x14, 0x00, 0x00, 0x00, 0x13, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, |
||||
|
0x15, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x14, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, |
||||
|
0x15, 0x00, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00, |
||||
|
0x18, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x14, 0x00, 0x02, 0x00, |
||||
|
0x1b, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x03, 0x00, 0x29, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, |
||||
|
0x20, 0x00, 0x04, 0x00, 0x2a, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x29, 0x00, 0x00, 0x00, |
||||
|
0x3b, 0x00, 0x04, 0x00, 0x2a, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, |
||||
|
0x2b, 0x00, 0x04, 0x00, 0x18, 0x00, 0x00, 0x00, 0x2c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, |
||||
|
0x20, 0x00, 0x04, 0x00, 0x2d, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, |
||||
|
0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x31, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, |
||||
|
0x1c, 0x00, 0x04, 0x00, 0x34, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, |
||||
|
0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x35, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, |
||||
|
0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x36, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, |
||||
|
0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x37, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, |
||||
|
0x2c, 0x00, 0x09, 0x00, 0x34, 0x00, 0x00, 0x00, 0x38, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, |
||||
|
0x35, 0x00, 0x00, 0x00, 0x36, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x36, 0x00, 0x00, 0x00, |
||||
|
0x37, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x3a, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, |
||||
|
0x34, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x44, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x18, 0x00, 0x00, 0x00, 0x47, 0x00, 0x00, 0x00, |
||||
|
0x01, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x49, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x04, 0x00, 0x00, 0x2c, 0x00, 0x06, 0x00, 0x09, 0x00, 0x00, 0x00, 0x4a, 0x00, 0x00, 0x00, |
||||
|
0x49, 0x00, 0x00, 0x00, 0x35, 0x00, 0x00, 0x00, 0x35, 0x00, 0x00, 0x00, 0x36, 0x00, 0x05, 0x00, |
||||
|
0x02, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, |
||||
|
0xf8, 0x00, 0x02, 0x00, 0x05, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, 0x3a, 0x00, 0x00, 0x00, |
||||
|
0x3b, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, 0xf9, 0x00, 0x02, 0x00, 0x4c, 0x00, 0x00, 0x00, |
||||
|
0xf8, 0x00, 0x02, 0x00, 0x4c, 0x00, 0x00, 0x00, 0xf6, 0x00, 0x04, 0x00, 0x4b, 0x00, 0x00, 0x00, |
||||
|
0x4e, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xf9, 0x00, 0x02, 0x00, 0x4d, 0x00, 0x00, 0x00, |
||||
|
0xf8, 0x00, 0x02, 0x00, 0x4d, 0x00, 0x00, 0x00, 0x41, 0x00, 0x05, 0x00, 0x0d, 0x00, 0x00, 0x00, |
||||
|
0x0e, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0x0e, 0x00, 0x00, 0x00, 0x84, 0x00, 0x05, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x12, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, |
||||
|
0x44, 0x00, 0x05, 0x00, 0x06, 0x00, 0x00, 0x00, 0x17, 0x00, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x7c, 0x00, 0x04, 0x00, 0x18, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00, |
||||
|
0x17, 0x00, 0x00, 0x00, 0x7c, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, |
||||
|
0x19, 0x00, 0x00, 0x00, 0xae, 0x00, 0x05, 0x00, 0x1b, 0x00, 0x00, 0x00, 0x1c, 0x00, 0x00, 0x00, |
||||
|
0x12, 0x00, 0x00, 0x00, 0x1a, 0x00, 0x00, 0x00, 0xf7, 0x00, 0x03, 0x00, 0x1e, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0xfa, 0x00, 0x04, 0x00, 0x1c, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x00, 0x00, |
||||
|
0x1e, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x1d, 0x00, 0x00, 0x00, 0xf9, 0x00, 0x02, 0x00, |
||||
|
0x4b, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x1e, 0x00, 0x00, 0x00, 0xf9, 0x00, 0x02, 0x00, |
||||
|
0x21, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x21, 0x00, 0x00, 0x00, 0xf5, 0x00, 0x07, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x53, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x00, 0x00, |
||||
|
0x48, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, 0xb0, 0x00, 0x05, 0x00, 0x1b, 0x00, 0x00, 0x00, |
||||
|
0x27, 0x00, 0x00, 0x00, 0x53, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0xf6, 0x00, 0x04, 0x00, |
||||
|
0x23, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xfa, 0x00, 0x04, 0x00, |
||||
|
0x27, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, |
||||
|
0x22, 0x00, 0x00, 0x00, 0x41, 0x00, 0x05, 0x00, 0x2d, 0x00, 0x00, 0x00, 0x2e, 0x00, 0x00, 0x00, |
||||
|
0x2b, 0x00, 0x00, 0x00, 0x2c, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, |
||||
|
0x2f, 0x00, 0x00, 0x00, 0x2e, 0x00, 0x00, 0x00, 0x84, 0x00, 0x05, 0x00, 0x06, 0x00, 0x00, 0x00, |
||||
|
0x32, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0x31, 0x00, 0x00, 0x00, 0x80, 0x00, 0x05, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x33, 0x00, 0x00, 0x00, 0x2f, 0x00, 0x00, 0x00, 0x32, 0x00, 0x00, 0x00, |
||||
|
0x3e, 0x00, 0x03, 0x00, 0x3b, 0x00, 0x00, 0x00, 0x38, 0x00, 0x00, 0x00, 0x41, 0x00, 0x05, 0x00, |
||||
|
0x07, 0x00, 0x00, 0x00, 0x3c, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x00, 0x00, 0x53, 0x00, 0x00, 0x00, |
||||
|
0x3d, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x00, 0x00, 0x3c, 0x00, 0x00, 0x00, |
||||
|
0x80, 0x00, 0x05, 0x00, 0x06, 0x00, 0x00, 0x00, 0x3e, 0x00, 0x00, 0x00, 0x33, 0x00, 0x00, 0x00, |
||||
|
0x3d, 0x00, 0x00, 0x00, 0x80, 0x00, 0x05, 0x00, 0x06, 0x00, 0x00, 0x00, 0x42, 0x00, 0x00, 0x00, |
||||
|
0x12, 0x00, 0x00, 0x00, 0x53, 0x00, 0x00, 0x00, 0x41, 0x00, 0x06, 0x00, 0x44, 0x00, 0x00, 0x00, |
||||
|
0x45, 0x00, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, 0x2c, 0x00, 0x00, 0x00, 0x42, 0x00, 0x00, 0x00, |
||||
|
0x3e, 0x00, 0x03, 0x00, 0x45, 0x00, 0x00, 0x00, 0x3e, 0x00, 0x00, 0x00, 0x80, 0x00, 0x05, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x48, 0x00, 0x00, 0x00, 0x53, 0x00, 0x00, 0x00, 0x47, 0x00, 0x00, 0x00, |
||||
|
0xf9, 0x00, 0x02, 0x00, 0x21, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x23, 0x00, 0x00, 0x00, |
||||
|
0xf9, 0x00, 0x02, 0x00, 0x4b, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x4e, 0x00, 0x00, 0x00, |
||||
|
0xf9, 0x00, 0x02, 0x00, 0x4c, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x4b, 0x00, 0x00, 0x00, |
||||
|
0xfd, 0x00, 0x01, 0x00, 0x38, 0x00, 0x01, 0x00}; |
||||
|
|
||||
|
// Uint8 SPIR-V module. Generated from the "shaders/" directory.
|
||||
|
constexpr u8 uint8_pass[] = { |
||||
|
0x03, 0x02, 0x23, 0x07, 0x00, 0x00, 0x01, 0x00, 0x07, 0x00, 0x08, 0x00, 0x2f, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, 0x01, 0x00, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, |
||||
|
0x51, 0x11, 0x00, 0x00, 0x11, 0x00, 0x02, 0x00, 0x61, 0x11, 0x00, 0x00, 0x0a, 0x00, 0x07, 0x00, |
||||
|
0x53, 0x50, 0x56, 0x5f, 0x4b, 0x48, 0x52, 0x5f, 0x31, 0x36, 0x62, 0x69, 0x74, 0x5f, 0x73, 0x74, |
||||
|
0x6f, 0x72, 0x61, 0x67, 0x65, 0x00, 0x00, 0x00, 0x0a, 0x00, 0x07, 0x00, 0x53, 0x50, 0x56, 0x5f, |
||||
|
0x4b, 0x48, 0x52, 0x5f, 0x38, 0x62, 0x69, 0x74, 0x5f, 0x73, 0x74, 0x6f, 0x72, 0x61, 0x67, 0x65, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x06, 0x00, 0x01, 0x00, 0x00, 0x00, 0x47, 0x4c, 0x53, 0x4c, |
||||
|
0x2e, 0x73, 0x74, 0x64, 0x2e, 0x34, 0x35, 0x30, 0x00, 0x00, 0x00, 0x00, 0x0e, 0x00, 0x03, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x06, 0x00, 0x05, 0x00, 0x00, 0x00, |
||||
|
0x04, 0x00, 0x00, 0x00, 0x6d, 0x61, 0x69, 0x6e, 0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, |
||||
|
0x10, 0x00, 0x06, 0x00, 0x04, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, |
||||
|
0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x0b, 0x00, 0x00, 0x00, |
||||
|
0x0b, 0x00, 0x00, 0x00, 0x1c, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x12, 0x00, 0x00, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x48, 0x00, 0x04, 0x00, 0x13, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x18, 0x00, 0x00, 0x00, 0x48, 0x00, 0x05, 0x00, 0x13, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x03, 0x00, |
||||
|
0x13, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x15, 0x00, 0x00, 0x00, |
||||
|
0x22, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x15, 0x00, 0x00, 0x00, |
||||
|
0x21, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x1f, 0x00, 0x00, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x48, 0x00, 0x04, 0x00, 0x20, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00, 0x48, 0x00, 0x05, 0x00, 0x20, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x03, 0x00, |
||||
|
0x20, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x22, 0x00, 0x00, 0x00, |
||||
|
0x22, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x22, 0x00, 0x00, 0x00, |
||||
|
0x21, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x47, 0x00, 0x04, 0x00, 0x2e, 0x00, 0x00, 0x00, |
||||
|
0x0b, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00, 0x13, 0x00, 0x02, 0x00, 0x02, 0x00, 0x00, 0x00, |
||||
|
0x21, 0x00, 0x03, 0x00, 0x03, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, |
||||
|
0x07, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x17, 0x00, 0x04, 0x00, |
||||
|
0x09, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, |
||||
|
0x0a, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, |
||||
|
0x0a, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, |
||||
|
0x0d, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00, |
||||
|
0x11, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x03, 0x00, |
||||
|
0x12, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x03, 0x00, 0x13, 0x00, 0x00, 0x00, |
||||
|
0x12, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x14, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, |
||||
|
0x13, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, 0x14, 0x00, 0x00, 0x00, 0x15, 0x00, 0x00, 0x00, |
||||
|
0x02, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00, 0x17, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00, |
||||
|
0x01, 0x00, 0x00, 0x00, 0x14, 0x00, 0x02, 0x00, 0x1a, 0x00, 0x00, 0x00, 0x15, 0x00, 0x04, 0x00, |
||||
|
0x1e, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x03, 0x00, |
||||
|
0x1f, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x00, 0x00, 0x1e, 0x00, 0x03, 0x00, 0x20, 0x00, 0x00, 0x00, |
||||
|
0x1f, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x21, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, |
||||
|
0x20, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, 0x21, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, |
||||
|
0x02, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x17, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x26, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, |
||||
|
0x11, 0x00, 0x00, 0x00, 0x20, 0x00, 0x04, 0x00, 0x2a, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, |
||||
|
0x1e, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x2c, 0x00, 0x00, 0x00, |
||||
|
0x00, 0x04, 0x00, 0x00, 0x2b, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x2d, 0x00, 0x00, 0x00, |
||||
|
0x01, 0x00, 0x00, 0x00, 0x2c, 0x00, 0x06, 0x00, 0x09, 0x00, 0x00, 0x00, 0x2e, 0x00, 0x00, 0x00, |
||||
|
0x2c, 0x00, 0x00, 0x00, 0x2d, 0x00, 0x00, 0x00, 0x2d, 0x00, 0x00, 0x00, 0x36, 0x00, 0x05, 0x00, |
||||
|
0x02, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, |
||||
|
0xf8, 0x00, 0x02, 0x00, 0x05, 0x00, 0x00, 0x00, 0x3b, 0x00, 0x04, 0x00, 0x07, 0x00, 0x00, 0x00, |
||||
|
0x08, 0x00, 0x00, 0x00, 0x07, 0x00, 0x00, 0x00, 0x41, 0x00, 0x05, 0x00, 0x0d, 0x00, 0x00, 0x00, |
||||
|
0x0e, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x0c, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0x0e, 0x00, 0x00, 0x00, 0x3e, 0x00, 0x03, 0x00, |
||||
|
0x08, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, |
||||
|
0x10, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x44, 0x00, 0x05, 0x00, 0x06, 0x00, 0x00, 0x00, |
||||
|
0x16, 0x00, 0x00, 0x00, 0x15, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x7c, 0x00, 0x04, 0x00, |
||||
|
0x17, 0x00, 0x00, 0x00, 0x18, 0x00, 0x00, 0x00, 0x16, 0x00, 0x00, 0x00, 0x7c, 0x00, 0x04, 0x00, |
||||
|
0x06, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00, 0x18, 0x00, 0x00, 0x00, 0xb0, 0x00, 0x05, 0x00, |
||||
|
0x1a, 0x00, 0x00, 0x00, 0x1b, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x19, 0x00, 0x00, 0x00, |
||||
|
0xf7, 0x00, 0x03, 0x00, 0x1d, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xfa, 0x00, 0x04, 0x00, |
||||
|
0x1b, 0x00, 0x00, 0x00, 0x1c, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, |
||||
|
0x1c, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x24, 0x00, 0x00, 0x00, |
||||
|
0x08, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, 0x06, 0x00, 0x00, 0x00, 0x25, 0x00, 0x00, 0x00, |
||||
|
0x08, 0x00, 0x00, 0x00, 0x41, 0x00, 0x06, 0x00, 0x26, 0x00, 0x00, 0x00, 0x27, 0x00, 0x00, 0x00, |
||||
|
0x15, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, 0x25, 0x00, 0x00, 0x00, 0x3d, 0x00, 0x04, 0x00, |
||||
|
0x11, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x27, 0x00, 0x00, 0x00, 0x71, 0x00, 0x04, 0x00, |
||||
|
0x1e, 0x00, 0x00, 0x00, 0x29, 0x00, 0x00, 0x00, 0x28, 0x00, 0x00, 0x00, 0x41, 0x00, 0x06, 0x00, |
||||
|
0x2a, 0x00, 0x00, 0x00, 0x2b, 0x00, 0x00, 0x00, 0x22, 0x00, 0x00, 0x00, 0x23, 0x00, 0x00, 0x00, |
||||
|
0x24, 0x00, 0x00, 0x00, 0x3e, 0x00, 0x03, 0x00, 0x2b, 0x00, 0x00, 0x00, 0x29, 0x00, 0x00, 0x00, |
||||
|
0xf9, 0x00, 0x02, 0x00, 0x1d, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x02, 0x00, 0x1d, 0x00, 0x00, 0x00, |
||||
|
0xfd, 0x00, 0x01, 0x00, 0x38, 0x00, 0x01, 0x00}; |
||||
|
|
||||
|
} // Anonymous namespace
|
||||
|
|
||||
|
VKComputePass::VKComputePass(const VKDevice& device, VKDescriptorPool& descriptor_pool, |
||||
|
const std::vector<vk::DescriptorSetLayoutBinding>& bindings, |
||||
|
const std::vector<vk::DescriptorUpdateTemplateEntry>& templates, |
||||
|
const std::vector<vk::PushConstantRange> push_constants, |
||||
|
std::size_t code_size, const u8* code) { |
||||
|
const auto dev = device.GetLogical(); |
||||
|
const auto& dld = device.GetDispatchLoader(); |
||||
|
|
||||
|
const vk::DescriptorSetLayoutCreateInfo descriptor_layout_ci( |
||||
|
{}, static_cast<u32>(bindings.size()), bindings.data()); |
||||
|
descriptor_set_layout = dev.createDescriptorSetLayoutUnique(descriptor_layout_ci, nullptr, dld); |
||||
|
|
||||
|
const vk::PipelineLayoutCreateInfo pipeline_layout_ci({}, 1, &*descriptor_set_layout, |
||||
|
static_cast<u32>(push_constants.size()), |
||||
|
push_constants.data()); |
||||
|
layout = dev.createPipelineLayoutUnique(pipeline_layout_ci, nullptr, dld); |
||||
|
|
||||
|
if (!templates.empty()) { |
||||
|
const vk::DescriptorUpdateTemplateCreateInfo template_ci( |
||||
|
{}, static_cast<u32>(templates.size()), templates.data(), |
||||
|
vk::DescriptorUpdateTemplateType::eDescriptorSet, *descriptor_set_layout, |
||||
|
vk::PipelineBindPoint::eGraphics, *layout, 0); |
||||
|
descriptor_template = dev.createDescriptorUpdateTemplateUnique(template_ci, nullptr, dld); |
||||
|
|
||||
|
descriptor_allocator.emplace(descriptor_pool, *descriptor_set_layout); |
||||
|
} |
||||
|
|
||||
|
auto code_copy = std::make_unique<u32[]>(code_size / sizeof(u32) + 1); |
||||
|
std::memcpy(code_copy.get(), code, code_size); |
||||
|
const vk::ShaderModuleCreateInfo module_ci({}, code_size, code_copy.get()); |
||||
|
module = dev.createShaderModuleUnique(module_ci, nullptr, dld); |
||||
|
|
||||
|
const vk::PipelineShaderStageCreateInfo stage_ci({}, vk::ShaderStageFlagBits::eCompute, *module, |
||||
|
"main", nullptr); |
||||
|
|
||||
|
const vk::ComputePipelineCreateInfo pipeline_ci({}, stage_ci, *layout, nullptr, 0); |
||||
|
pipeline = dev.createComputePipelineUnique(nullptr, pipeline_ci, nullptr, dld); |
||||
|
} |
||||
|
|
||||
|
VKComputePass::~VKComputePass() = default; |
||||
|
|
||||
|
vk::DescriptorSet VKComputePass::CommitDescriptorSet( |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue, VKFence& fence) { |
||||
|
if (!descriptor_template) { |
||||
|
return {}; |
||||
|
} |
||||
|
const auto set = descriptor_allocator->Commit(fence); |
||||
|
update_descriptor_queue.Send(*descriptor_template, set); |
||||
|
return set; |
||||
|
} |
||||
|
|
||||
|
QuadArrayPass::QuadArrayPass(const VKDevice& device, VKScheduler& scheduler, |
||||
|
VKDescriptorPool& descriptor_pool, |
||||
|
VKStagingBufferPool& staging_buffer_pool, |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue) |
||||
|
: VKComputePass(device, descriptor_pool, |
||||
|
{vk::DescriptorSetLayoutBinding(0, vk::DescriptorType::eStorageBuffer, 1, |
||||
|
vk::ShaderStageFlagBits::eCompute, nullptr)}, |
||||
|
{vk::DescriptorUpdateTemplateEntry(0, 0, 1, vk::DescriptorType::eStorageBuffer, |
||||
|
0, sizeof(DescriptorUpdateEntry))}, |
||||
|
{vk::PushConstantRange(vk::ShaderStageFlagBits::eCompute, 0, sizeof(u32))}, |
||||
|
std::size(quad_array), quad_array), |
||||
|
scheduler{scheduler}, staging_buffer_pool{staging_buffer_pool}, |
||||
|
update_descriptor_queue{update_descriptor_queue} {} |
||||
|
|
||||
|
QuadArrayPass::~QuadArrayPass() = default; |
||||
|
|
||||
|
std::pair<const vk::Buffer&, vk::DeviceSize> QuadArrayPass::Assemble(u32 num_vertices, u32 first) { |
||||
|
const u32 num_triangle_vertices = num_vertices * 6 / 4; |
||||
|
const std::size_t staging_size = num_triangle_vertices * sizeof(u32); |
||||
|
auto& buffer = staging_buffer_pool.GetUnusedBuffer(staging_size, false); |
||||
|
|
||||
|
update_descriptor_queue.Acquire(); |
||||
|
update_descriptor_queue.AddBuffer(&*buffer.handle, 0, staging_size); |
||||
|
const auto set = CommitDescriptorSet(update_descriptor_queue, scheduler.GetFence()); |
||||
|
|
||||
|
scheduler.RequestOutsideRenderPassOperationContext(); |
||||
|
|
||||
|
ASSERT(num_vertices % 4 == 0); |
||||
|
const u32 num_quads = num_vertices / 4; |
||||
|
scheduler.Record([layout = *layout, pipeline = *pipeline, buffer = *buffer.handle, num_quads, |
||||
|
first, set](auto cmdbuf, auto& dld) { |
||||
|
constexpr u32 dispatch_size = 1024; |
||||
|
cmdbuf.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline, dld); |
||||
|
cmdbuf.bindDescriptorSets(vk::PipelineBindPoint::eCompute, layout, 0, {set}, {}, dld); |
||||
|
cmdbuf.pushConstants(layout, vk::ShaderStageFlagBits::eCompute, 0, sizeof(first), &first, |
||||
|
dld); |
||||
|
cmdbuf.dispatch(Common::AlignUp(num_quads, dispatch_size) / dispatch_size, 1, 1, dld); |
||||
|
|
||||
|
const vk::BufferMemoryBarrier barrier( |
||||
|
vk::AccessFlagBits::eShaderWrite, vk::AccessFlagBits::eVertexAttributeRead, |
||||
|
VK_QUEUE_FAMILY_IGNORED, VK_QUEUE_FAMILY_IGNORED, buffer, 0, |
||||
|
static_cast<vk::DeviceSize>(num_quads) * 6 * sizeof(u32)); |
||||
|
cmdbuf.pipelineBarrier(vk::PipelineStageFlagBits::eComputeShader, |
||||
|
vk::PipelineStageFlagBits::eVertexInput, {}, {}, {barrier}, {}, dld); |
||||
|
}); |
||||
|
return {*buffer.handle, 0}; |
||||
|
} |
||||
|
|
||||
|
Uint8Pass::Uint8Pass(const VKDevice& device, VKScheduler& scheduler, |
||||
|
VKDescriptorPool& descriptor_pool, VKStagingBufferPool& staging_buffer_pool, |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue) |
||||
|
: VKComputePass(device, descriptor_pool, |
||||
|
{vk::DescriptorSetLayoutBinding(0, vk::DescriptorType::eStorageBuffer, 1, |
||||
|
vk::ShaderStageFlagBits::eCompute, nullptr), |
||||
|
vk::DescriptorSetLayoutBinding(1, vk::DescriptorType::eStorageBuffer, 1, |
||||
|
vk::ShaderStageFlagBits::eCompute, nullptr)}, |
||||
|
{vk::DescriptorUpdateTemplateEntry(0, 0, 2, vk::DescriptorType::eStorageBuffer, |
||||
|
0, sizeof(DescriptorUpdateEntry))}, |
||||
|
{}, std::size(uint8_pass), uint8_pass), |
||||
|
scheduler{scheduler}, staging_buffer_pool{staging_buffer_pool}, |
||||
|
update_descriptor_queue{update_descriptor_queue} {} |
||||
|
|
||||
|
Uint8Pass::~Uint8Pass() = default; |
||||
|
|
||||
|
std::pair<const vk::Buffer*, u64> Uint8Pass::Assemble(u32 num_vertices, vk::Buffer src_buffer, |
||||
|
u64 src_offset) { |
||||
|
const auto staging_size = static_cast<u32>(num_vertices * sizeof(u16)); |
||||
|
auto& buffer = staging_buffer_pool.GetUnusedBuffer(staging_size, false); |
||||
|
|
||||
|
update_descriptor_queue.Acquire(); |
||||
|
update_descriptor_queue.AddBuffer(&src_buffer, src_offset, num_vertices); |
||||
|
update_descriptor_queue.AddBuffer(&*buffer.handle, 0, staging_size); |
||||
|
const auto set = CommitDescriptorSet(update_descriptor_queue, scheduler.GetFence()); |
||||
|
|
||||
|
scheduler.RequestOutsideRenderPassOperationContext(); |
||||
|
scheduler.Record([layout = *layout, pipeline = *pipeline, buffer = *buffer.handle, set, |
||||
|
num_vertices](auto cmdbuf, auto& dld) { |
||||
|
constexpr u32 dispatch_size = 1024; |
||||
|
cmdbuf.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline, dld); |
||||
|
cmdbuf.bindDescriptorSets(vk::PipelineBindPoint::eCompute, layout, 0, {set}, {}, dld); |
||||
|
cmdbuf.dispatch(Common::AlignUp(num_vertices, dispatch_size) / dispatch_size, 1, 1, dld); |
||||
|
|
||||
|
const vk::BufferMemoryBarrier barrier( |
||||
|
vk::AccessFlagBits::eShaderWrite, vk::AccessFlagBits::eVertexAttributeRead, |
||||
|
VK_QUEUE_FAMILY_IGNORED, VK_QUEUE_FAMILY_IGNORED, buffer, 0, |
||||
|
static_cast<vk::DeviceSize>(num_vertices) * sizeof(u16)); |
||||
|
cmdbuf.pipelineBarrier(vk::PipelineStageFlagBits::eComputeShader, |
||||
|
vk::PipelineStageFlagBits::eVertexInput, {}, {}, {barrier}, {}, dld); |
||||
|
}); |
||||
|
return {&*buffer.handle, 0}; |
||||
|
} |
||||
|
|
||||
|
} // namespace Vulkan
|
||||
@ -0,0 +1,77 @@ |
|||||
|
// Copyright 2019 yuzu Emulator Project |
||||
|
// Licensed under GPLv2 or any later version |
||||
|
// Refer to the license.txt file included. |
||||
|
|
||||
|
#pragma once |
||||
|
|
||||
|
#include <optional> |
||||
|
#include <utility> |
||||
|
#include <vector> |
||||
|
#include "common/common_types.h" |
||||
|
#include "video_core/renderer_vulkan/declarations.h" |
||||
|
#include "video_core/renderer_vulkan/vk_descriptor_pool.h" |
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
class VKDevice; |
||||
|
class VKFence; |
||||
|
class VKScheduler; |
||||
|
class VKStagingBufferPool; |
||||
|
class VKUpdateDescriptorQueue; |
||||
|
|
||||
|
class VKComputePass { |
||||
|
public: |
||||
|
explicit VKComputePass(const VKDevice& device, VKDescriptorPool& descriptor_pool, |
||||
|
const std::vector<vk::DescriptorSetLayoutBinding>& bindings, |
||||
|
const std::vector<vk::DescriptorUpdateTemplateEntry>& templates, |
||||
|
const std::vector<vk::PushConstantRange> push_constants, |
||||
|
std::size_t code_size, const u8* code); |
||||
|
~VKComputePass(); |
||||
|
|
||||
|
protected: |
||||
|
vk::DescriptorSet CommitDescriptorSet(VKUpdateDescriptorQueue& update_descriptor_queue, |
||||
|
VKFence& fence); |
||||
|
|
||||
|
UniqueDescriptorUpdateTemplate descriptor_template; |
||||
|
UniquePipelineLayout layout; |
||||
|
UniquePipeline pipeline; |
||||
|
|
||||
|
private: |
||||
|
UniqueDescriptorSetLayout descriptor_set_layout; |
||||
|
std::optional<DescriptorAllocator> descriptor_allocator; |
||||
|
UniqueShaderModule module; |
||||
|
}; |
||||
|
|
||||
|
class QuadArrayPass final : public VKComputePass { |
||||
|
public: |
||||
|
explicit QuadArrayPass(const VKDevice& device, VKScheduler& scheduler, |
||||
|
VKDescriptorPool& descriptor_pool, |
||||
|
VKStagingBufferPool& staging_buffer_pool, |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue); |
||||
|
~QuadArrayPass(); |
||||
|
|
||||
|
std::pair<const vk::Buffer&, vk::DeviceSize> Assemble(u32 num_vertices, u32 first); |
||||
|
|
||||
|
private: |
||||
|
VKScheduler& scheduler; |
||||
|
VKStagingBufferPool& staging_buffer_pool; |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue; |
||||
|
}; |
||||
|
|
||||
|
class Uint8Pass final : public VKComputePass { |
||||
|
public: |
||||
|
explicit Uint8Pass(const VKDevice& device, VKScheduler& scheduler, |
||||
|
VKDescriptorPool& descriptor_pool, VKStagingBufferPool& staging_buffer_pool, |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue); |
||||
|
~Uint8Pass(); |
||||
|
|
||||
|
std::pair<const vk::Buffer*, u64> Assemble(u32 num_vertices, vk::Buffer src_buffer, |
||||
|
u64 src_offset); |
||||
|
|
||||
|
private: |
||||
|
VKScheduler& scheduler; |
||||
|
VKStagingBufferPool& staging_buffer_pool; |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue; |
||||
|
}; |
||||
|
|
||||
|
} // namespace Vulkan |
||||
@ -0,0 +1,112 @@ |
|||||
|
// Copyright 2019 yuzu Emulator Project
|
||||
|
// Licensed under GPLv2 or any later version
|
||||
|
// Refer to the license.txt file included.
|
||||
|
|
||||
|
#include <memory>
|
||||
|
#include <vector>
|
||||
|
|
||||
|
#include "video_core/renderer_vulkan/declarations.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_compute_pipeline.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_device.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_pipeline_cache.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_resource_manager.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_scheduler.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_shader_decompiler.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_update_descriptor.h"
|
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
VKComputePipeline::VKComputePipeline(const VKDevice& device, VKScheduler& scheduler, |
||||
|
VKDescriptorPool& descriptor_pool, |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue, |
||||
|
const SPIRVShader& shader) |
||||
|
: device{device}, scheduler{scheduler}, entries{shader.entries}, |
||||
|
descriptor_set_layout{CreateDescriptorSetLayout()}, |
||||
|
descriptor_allocator{descriptor_pool, *descriptor_set_layout}, |
||||
|
update_descriptor_queue{update_descriptor_queue}, layout{CreatePipelineLayout()}, |
||||
|
descriptor_template{CreateDescriptorUpdateTemplate()}, |
||||
|
shader_module{CreateShaderModule(shader.code)}, pipeline{CreatePipeline()} {} |
||||
|
|
||||
|
VKComputePipeline::~VKComputePipeline() = default; |
||||
|
|
||||
|
vk::DescriptorSet VKComputePipeline::CommitDescriptorSet() { |
||||
|
if (!descriptor_template) { |
||||
|
return {}; |
||||
|
} |
||||
|
const auto set = descriptor_allocator.Commit(scheduler.GetFence()); |
||||
|
update_descriptor_queue.Send(*descriptor_template, set); |
||||
|
return set; |
||||
|
} |
||||
|
|
||||
|
UniqueDescriptorSetLayout VKComputePipeline::CreateDescriptorSetLayout() const { |
||||
|
std::vector<vk::DescriptorSetLayoutBinding> bindings; |
||||
|
u32 binding = 0; |
||||
|
const auto AddBindings = [&](vk::DescriptorType descriptor_type, std::size_t num_entries) { |
||||
|
// TODO(Rodrigo): Maybe make individual bindings here?
|
||||
|
for (u32 bindpoint = 0; bindpoint < static_cast<u32>(num_entries); ++bindpoint) { |
||||
|
bindings.emplace_back(binding++, descriptor_type, 1, vk::ShaderStageFlagBits::eCompute, |
||||
|
nullptr); |
||||
|
} |
||||
|
}; |
||||
|
AddBindings(vk::DescriptorType::eUniformBuffer, entries.const_buffers.size()); |
||||
|
AddBindings(vk::DescriptorType::eStorageBuffer, entries.global_buffers.size()); |
||||
|
AddBindings(vk::DescriptorType::eUniformTexelBuffer, entries.texel_buffers.size()); |
||||
|
AddBindings(vk::DescriptorType::eCombinedImageSampler, entries.samplers.size()); |
||||
|
AddBindings(vk::DescriptorType::eStorageImage, entries.images.size()); |
||||
|
|
||||
|
const vk::DescriptorSetLayoutCreateInfo descriptor_set_layout_ci( |
||||
|
{}, static_cast<u32>(bindings.size()), bindings.data()); |
||||
|
|
||||
|
const auto dev = device.GetLogical(); |
||||
|
const auto& dld = device.GetDispatchLoader(); |
||||
|
return dev.createDescriptorSetLayoutUnique(descriptor_set_layout_ci, nullptr, dld); |
||||
|
} |
||||
|
|
||||
|
UniquePipelineLayout VKComputePipeline::CreatePipelineLayout() const { |
||||
|
const vk::PipelineLayoutCreateInfo layout_ci({}, 1, &*descriptor_set_layout, 0, nullptr); |
||||
|
const auto dev = device.GetLogical(); |
||||
|
return dev.createPipelineLayoutUnique(layout_ci, nullptr, device.GetDispatchLoader()); |
||||
|
} |
||||
|
|
||||
|
UniqueDescriptorUpdateTemplate VKComputePipeline::CreateDescriptorUpdateTemplate() const { |
||||
|
std::vector<vk::DescriptorUpdateTemplateEntry> template_entries; |
||||
|
u32 binding = 0; |
||||
|
u32 offset = 0; |
||||
|
FillDescriptorUpdateTemplateEntries(device, entries, binding, offset, template_entries); |
||||
|
if (template_entries.empty()) { |
||||
|
// If the shader doesn't use descriptor sets, skip template creation.
|
||||
|
return UniqueDescriptorUpdateTemplate{}; |
||||
|
} |
||||
|
|
||||
|
const vk::DescriptorUpdateTemplateCreateInfo template_ci( |
||||
|
{}, static_cast<u32>(template_entries.size()), template_entries.data(), |
||||
|
vk::DescriptorUpdateTemplateType::eDescriptorSet, *descriptor_set_layout, |
||||
|
vk::PipelineBindPoint::eGraphics, *layout, DESCRIPTOR_SET); |
||||
|
|
||||
|
const auto dev = device.GetLogical(); |
||||
|
const auto& dld = device.GetDispatchLoader(); |
||||
|
return dev.createDescriptorUpdateTemplateUnique(template_ci, nullptr, dld); |
||||
|
} |
||||
|
|
||||
|
UniqueShaderModule VKComputePipeline::CreateShaderModule(const std::vector<u32>& code) const { |
||||
|
const vk::ShaderModuleCreateInfo module_ci({}, code.size() * sizeof(u32), code.data()); |
||||
|
const auto dev = device.GetLogical(); |
||||
|
return dev.createShaderModuleUnique(module_ci, nullptr, device.GetDispatchLoader()); |
||||
|
} |
||||
|
|
||||
|
UniquePipeline VKComputePipeline::CreatePipeline() const { |
||||
|
vk::PipelineShaderStageCreateInfo shader_stage_ci({}, vk::ShaderStageFlagBits::eCompute, |
||||
|
*shader_module, "main", nullptr); |
||||
|
vk::PipelineShaderStageRequiredSubgroupSizeCreateInfoEXT subgroup_size_ci; |
||||
|
subgroup_size_ci.requiredSubgroupSize = GuestWarpSize; |
||||
|
if (entries.uses_warps && device.IsGuestWarpSizeSupported(vk::ShaderStageFlagBits::eCompute)) { |
||||
|
shader_stage_ci.pNext = &subgroup_size_ci; |
||||
|
} |
||||
|
|
||||
|
const vk::ComputePipelineCreateInfo create_info({}, shader_stage_ci, *layout, {}, 0); |
||||
|
const auto dev = device.GetLogical(); |
||||
|
return dev.createComputePipelineUnique({}, create_info, nullptr, device.GetDispatchLoader()); |
||||
|
} |
||||
|
|
||||
|
} // namespace Vulkan
|
||||
@ -0,0 +1,66 @@ |
|||||
|
// Copyright 2019 yuzu Emulator Project |
||||
|
// Licensed under GPLv2 or any later version |
||||
|
// Refer to the license.txt file included. |
||||
|
|
||||
|
#pragma once |
||||
|
|
||||
|
#include <memory> |
||||
|
|
||||
|
#include "common/common_types.h" |
||||
|
#include "video_core/renderer_vulkan/declarations.h" |
||||
|
#include "video_core/renderer_vulkan/vk_descriptor_pool.h" |
||||
|
#include "video_core/renderer_vulkan/vk_shader_decompiler.h" |
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
class VKDevice; |
||||
|
class VKScheduler; |
||||
|
class VKUpdateDescriptorQueue; |
||||
|
|
||||
|
class VKComputePipeline final { |
||||
|
public: |
||||
|
explicit VKComputePipeline(const VKDevice& device, VKScheduler& scheduler, |
||||
|
VKDescriptorPool& descriptor_pool, |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue, |
||||
|
const SPIRVShader& shader); |
||||
|
~VKComputePipeline(); |
||||
|
|
||||
|
vk::DescriptorSet CommitDescriptorSet(); |
||||
|
|
||||
|
vk::Pipeline GetHandle() const { |
||||
|
return *pipeline; |
||||
|
} |
||||
|
|
||||
|
vk::PipelineLayout GetLayout() const { |
||||
|
return *layout; |
||||
|
} |
||||
|
|
||||
|
const ShaderEntries& GetEntries() { |
||||
|
return entries; |
||||
|
} |
||||
|
|
||||
|
private: |
||||
|
UniqueDescriptorSetLayout CreateDescriptorSetLayout() const; |
||||
|
|
||||
|
UniquePipelineLayout CreatePipelineLayout() const; |
||||
|
|
||||
|
UniqueDescriptorUpdateTemplate CreateDescriptorUpdateTemplate() const; |
||||
|
|
||||
|
UniqueShaderModule CreateShaderModule(const std::vector<u32>& code) const; |
||||
|
|
||||
|
UniquePipeline CreatePipeline() const; |
||||
|
|
||||
|
const VKDevice& device; |
||||
|
VKScheduler& scheduler; |
||||
|
ShaderEntries entries; |
||||
|
|
||||
|
UniqueDescriptorSetLayout descriptor_set_layout; |
||||
|
DescriptorAllocator descriptor_allocator; |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue; |
||||
|
UniquePipelineLayout layout; |
||||
|
UniqueDescriptorUpdateTemplate descriptor_template; |
||||
|
UniqueShaderModule shader_module; |
||||
|
UniquePipeline pipeline; |
||||
|
}; |
||||
|
|
||||
|
} // namespace Vulkan |
||||
@ -0,0 +1,271 @@ |
|||||
|
// Copyright 2019 yuzu Emulator Project
|
||||
|
// Licensed under GPLv2 or any later version
|
||||
|
// Refer to the license.txt file included.
|
||||
|
|
||||
|
#include <vector>
|
||||
|
#include "common/assert.h"
|
||||
|
#include "common/common_types.h"
|
||||
|
#include "common/microprofile.h"
|
||||
|
#include "video_core/renderer_vulkan/declarations.h"
|
||||
|
#include "video_core/renderer_vulkan/fixed_pipeline_state.h"
|
||||
|
#include "video_core/renderer_vulkan/maxwell_to_vk.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_device.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_graphics_pipeline.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_pipeline_cache.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_renderpass_cache.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_scheduler.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_update_descriptor.h"
|
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
MICROPROFILE_DECLARE(Vulkan_PipelineCache); |
||||
|
|
||||
|
namespace { |
||||
|
|
||||
|
vk::StencilOpState GetStencilFaceState(const FixedPipelineState::StencilFace& face) { |
||||
|
return vk::StencilOpState(MaxwellToVK::StencilOp(face.action_stencil_fail), |
||||
|
MaxwellToVK::StencilOp(face.action_depth_pass), |
||||
|
MaxwellToVK::StencilOp(face.action_depth_fail), |
||||
|
MaxwellToVK::ComparisonOp(face.test_func), 0, 0, 0); |
||||
|
} |
||||
|
|
||||
|
bool SupportsPrimitiveRestart(vk::PrimitiveTopology topology) { |
||||
|
static constexpr std::array unsupported_topologies = { |
||||
|
vk::PrimitiveTopology::ePointList, |
||||
|
vk::PrimitiveTopology::eLineList, |
||||
|
vk::PrimitiveTopology::eTriangleList, |
||||
|
vk::PrimitiveTopology::eLineListWithAdjacency, |
||||
|
vk::PrimitiveTopology::eTriangleListWithAdjacency, |
||||
|
vk::PrimitiveTopology::ePatchList}; |
||||
|
return std::find(std::begin(unsupported_topologies), std::end(unsupported_topologies), |
||||
|
topology) == std::end(unsupported_topologies); |
||||
|
} |
||||
|
|
||||
|
} // Anonymous namespace
|
||||
|
|
||||
|
VKGraphicsPipeline::VKGraphicsPipeline(const VKDevice& device, VKScheduler& scheduler, |
||||
|
VKDescriptorPool& descriptor_pool, |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue, |
||||
|
VKRenderPassCache& renderpass_cache, |
||||
|
const GraphicsPipelineCacheKey& key, |
||||
|
const std::vector<vk::DescriptorSetLayoutBinding>& bindings, |
||||
|
const SPIRVProgram& program) |
||||
|
: device{device}, scheduler{scheduler}, fixed_state{key.fixed_state}, hash{key.Hash()}, |
||||
|
descriptor_set_layout{CreateDescriptorSetLayout(bindings)}, |
||||
|
descriptor_allocator{descriptor_pool, *descriptor_set_layout}, |
||||
|
update_descriptor_queue{update_descriptor_queue}, layout{CreatePipelineLayout()}, |
||||
|
descriptor_template{CreateDescriptorUpdateTemplate(program)}, modules{CreateShaderModules( |
||||
|
program)}, |
||||
|
renderpass{renderpass_cache.GetRenderPass(key.renderpass_params)}, pipeline{CreatePipeline( |
||||
|
key.renderpass_params, |
||||
|
program)} {} |
||||
|
|
||||
|
VKGraphicsPipeline::~VKGraphicsPipeline() = default; |
||||
|
|
||||
|
vk::DescriptorSet VKGraphicsPipeline::CommitDescriptorSet() { |
||||
|
if (!descriptor_template) { |
||||
|
return {}; |
||||
|
} |
||||
|
const auto set = descriptor_allocator.Commit(scheduler.GetFence()); |
||||
|
update_descriptor_queue.Send(*descriptor_template, set); |
||||
|
return set; |
||||
|
} |
||||
|
|
||||
|
UniqueDescriptorSetLayout VKGraphicsPipeline::CreateDescriptorSetLayout( |
||||
|
const std::vector<vk::DescriptorSetLayoutBinding>& bindings) const { |
||||
|
const vk::DescriptorSetLayoutCreateInfo descriptor_set_layout_ci( |
||||
|
{}, static_cast<u32>(bindings.size()), bindings.data()); |
||||
|
|
||||
|
const auto dev = device.GetLogical(); |
||||
|
const auto& dld = device.GetDispatchLoader(); |
||||
|
return dev.createDescriptorSetLayoutUnique(descriptor_set_layout_ci, nullptr, dld); |
||||
|
} |
||||
|
|
||||
|
UniquePipelineLayout VKGraphicsPipeline::CreatePipelineLayout() const { |
||||
|
const vk::PipelineLayoutCreateInfo pipeline_layout_ci({}, 1, &*descriptor_set_layout, 0, |
||||
|
nullptr); |
||||
|
const auto dev = device.GetLogical(); |
||||
|
const auto& dld = device.GetDispatchLoader(); |
||||
|
return dev.createPipelineLayoutUnique(pipeline_layout_ci, nullptr, dld); |
||||
|
} |
||||
|
|
||||
|
UniqueDescriptorUpdateTemplate VKGraphicsPipeline::CreateDescriptorUpdateTemplate( |
||||
|
const SPIRVProgram& program) const { |
||||
|
std::vector<vk::DescriptorUpdateTemplateEntry> template_entries; |
||||
|
u32 binding = 0; |
||||
|
u32 offset = 0; |
||||
|
for (const auto& stage : program) { |
||||
|
if (stage) { |
||||
|
FillDescriptorUpdateTemplateEntries(device, stage->entries, binding, offset, |
||||
|
template_entries); |
||||
|
} |
||||
|
} |
||||
|
if (template_entries.empty()) { |
||||
|
// If the shader doesn't use descriptor sets, skip template creation.
|
||||
|
return UniqueDescriptorUpdateTemplate{}; |
||||
|
} |
||||
|
|
||||
|
const vk::DescriptorUpdateTemplateCreateInfo template_ci( |
||||
|
{}, static_cast<u32>(template_entries.size()), template_entries.data(), |
||||
|
vk::DescriptorUpdateTemplateType::eDescriptorSet, *descriptor_set_layout, |
||||
|
vk::PipelineBindPoint::eGraphics, *layout, DESCRIPTOR_SET); |
||||
|
|
||||
|
const auto dev = device.GetLogical(); |
||||
|
const auto& dld = device.GetDispatchLoader(); |
||||
|
return dev.createDescriptorUpdateTemplateUnique(template_ci, nullptr, dld); |
||||
|
} |
||||
|
|
||||
|
std::vector<UniqueShaderModule> VKGraphicsPipeline::CreateShaderModules( |
||||
|
const SPIRVProgram& program) const { |
||||
|
std::vector<UniqueShaderModule> modules; |
||||
|
const auto dev = device.GetLogical(); |
||||
|
const auto& dld = device.GetDispatchLoader(); |
||||
|
for (std::size_t i = 0; i < Maxwell::MaxShaderStage; ++i) { |
||||
|
const auto& stage = program[i]; |
||||
|
if (!stage) { |
||||
|
continue; |
||||
|
} |
||||
|
const vk::ShaderModuleCreateInfo module_ci({}, stage->code.size() * sizeof(u32), |
||||
|
stage->code.data()); |
||||
|
modules.emplace_back(dev.createShaderModuleUnique(module_ci, nullptr, dld)); |
||||
|
} |
||||
|
return modules; |
||||
|
} |
||||
|
|
||||
|
UniquePipeline VKGraphicsPipeline::CreatePipeline(const RenderPassParams& renderpass_params, |
||||
|
const SPIRVProgram& program) const { |
||||
|
const auto& vi = fixed_state.vertex_input; |
||||
|
const auto& ia = fixed_state.input_assembly; |
||||
|
const auto& ds = fixed_state.depth_stencil; |
||||
|
const auto& cd = fixed_state.color_blending; |
||||
|
const auto& ts = fixed_state.tessellation; |
||||
|
const auto& rs = fixed_state.rasterizer; |
||||
|
|
||||
|
std::vector<vk::VertexInputBindingDescription> vertex_bindings; |
||||
|
std::vector<vk::VertexInputBindingDivisorDescriptionEXT> vertex_binding_divisors; |
||||
|
for (std::size_t i = 0; i < vi.num_bindings; ++i) { |
||||
|
const auto& binding = vi.bindings[i]; |
||||
|
const bool instanced = binding.divisor != 0; |
||||
|
const auto rate = instanced ? vk::VertexInputRate::eInstance : vk::VertexInputRate::eVertex; |
||||
|
vertex_bindings.emplace_back(binding.index, binding.stride, rate); |
||||
|
if (instanced) { |
||||
|
vertex_binding_divisors.emplace_back(binding.index, binding.divisor); |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
std::vector<vk::VertexInputAttributeDescription> vertex_attributes; |
||||
|
const auto& input_attributes = program[0]->entries.attributes; |
||||
|
for (std::size_t i = 0; i < vi.num_attributes; ++i) { |
||||
|
const auto& attribute = vi.attributes[i]; |
||||
|
if (input_attributes.find(attribute.index) == input_attributes.end()) { |
||||
|
// Skip attributes not used by the vertex shaders.
|
||||
|
continue; |
||||
|
} |
||||
|
vertex_attributes.emplace_back(attribute.index, attribute.buffer, |
||||
|
MaxwellToVK::VertexFormat(attribute.type, attribute.size), |
||||
|
attribute.offset); |
||||
|
} |
||||
|
|
||||
|
vk::PipelineVertexInputStateCreateInfo vertex_input_ci( |
||||
|
{}, static_cast<u32>(vertex_bindings.size()), vertex_bindings.data(), |
||||
|
static_cast<u32>(vertex_attributes.size()), vertex_attributes.data()); |
||||
|
|
||||
|
const vk::PipelineVertexInputDivisorStateCreateInfoEXT vertex_input_divisor_ci( |
||||
|
static_cast<u32>(vertex_binding_divisors.size()), vertex_binding_divisors.data()); |
||||
|
if (!vertex_binding_divisors.empty()) { |
||||
|
vertex_input_ci.pNext = &vertex_input_divisor_ci; |
||||
|
} |
||||
|
|
||||
|
const auto primitive_topology = MaxwellToVK::PrimitiveTopology(device, ia.topology); |
||||
|
const vk::PipelineInputAssemblyStateCreateInfo input_assembly_ci( |
||||
|
{}, primitive_topology, |
||||
|
ia.primitive_restart_enable && SupportsPrimitiveRestart(primitive_topology)); |
||||
|
|
||||
|
const vk::PipelineTessellationStateCreateInfo tessellation_ci({}, ts.patch_control_points); |
||||
|
|
||||
|
const vk::PipelineViewportStateCreateInfo viewport_ci({}, Maxwell::NumViewports, nullptr, |
||||
|
Maxwell::NumViewports, nullptr); |
||||
|
|
||||
|
// TODO(Rodrigo): Find out what's the default register value for front face
|
||||
|
const vk::PipelineRasterizationStateCreateInfo rasterizer_ci( |
||||
|
{}, rs.depth_clamp_enable, false, vk::PolygonMode::eFill, |
||||
|
rs.cull_enable ? MaxwellToVK::CullFace(rs.cull_face) : vk::CullModeFlagBits::eNone, |
||||
|
rs.cull_enable ? MaxwellToVK::FrontFace(rs.front_face) : vk::FrontFace::eCounterClockwise, |
||||
|
rs.depth_bias_enable, 0.0f, 0.0f, 0.0f, 1.0f); |
||||
|
|
||||
|
const vk::PipelineMultisampleStateCreateInfo multisampling_ci( |
||||
|
{}, vk::SampleCountFlagBits::e1, false, 0.0f, nullptr, false, false); |
||||
|
|
||||
|
const vk::CompareOp depth_test_compare = ds.depth_test_enable |
||||
|
? MaxwellToVK::ComparisonOp(ds.depth_test_function) |
||||
|
: vk::CompareOp::eAlways; |
||||
|
|
||||
|
const vk::PipelineDepthStencilStateCreateInfo depth_stencil_ci( |
||||
|
{}, ds.depth_test_enable, ds.depth_write_enable, depth_test_compare, ds.depth_bounds_enable, |
||||
|
ds.stencil_enable, GetStencilFaceState(ds.front_stencil), |
||||
|
GetStencilFaceState(ds.back_stencil), 0.0f, 0.0f); |
||||
|
|
||||
|
std::array<vk::PipelineColorBlendAttachmentState, Maxwell::NumRenderTargets> cb_attachments; |
||||
|
const std::size_t num_attachments = |
||||
|
std::min(cd.attachments_count, renderpass_params.color_attachments.size()); |
||||
|
for (std::size_t i = 0; i < num_attachments; ++i) { |
||||
|
constexpr std::array component_table{ |
||||
|
vk::ColorComponentFlagBits::eR, vk::ColorComponentFlagBits::eG, |
||||
|
vk::ColorComponentFlagBits::eB, vk::ColorComponentFlagBits::eA}; |
||||
|
const auto& blend = cd.attachments[i]; |
||||
|
|
||||
|
vk::ColorComponentFlags color_components{}; |
||||
|
for (std::size_t j = 0; j < component_table.size(); ++j) { |
||||
|
if (blend.components[j]) |
||||
|
color_components |= component_table[j]; |
||||
|
} |
||||
|
|
||||
|
cb_attachments[i] = vk::PipelineColorBlendAttachmentState( |
||||
|
blend.enable, MaxwellToVK::BlendFactor(blend.src_rgb_func), |
||||
|
MaxwellToVK::BlendFactor(blend.dst_rgb_func), |
||||
|
MaxwellToVK::BlendEquation(blend.rgb_equation), |
||||
|
MaxwellToVK::BlendFactor(blend.src_a_func), MaxwellToVK::BlendFactor(blend.dst_a_func), |
||||
|
MaxwellToVK::BlendEquation(blend.a_equation), color_components); |
||||
|
} |
||||
|
const vk::PipelineColorBlendStateCreateInfo color_blending_ci({}, false, vk::LogicOp::eCopy, |
||||
|
static_cast<u32>(num_attachments), |
||||
|
cb_attachments.data(), {}); |
||||
|
|
||||
|
constexpr std::array dynamic_states = { |
||||
|
vk::DynamicState::eViewport, vk::DynamicState::eScissor, |
||||
|
vk::DynamicState::eDepthBias, vk::DynamicState::eBlendConstants, |
||||
|
vk::DynamicState::eDepthBounds, vk::DynamicState::eStencilCompareMask, |
||||
|
vk::DynamicState::eStencilWriteMask, vk::DynamicState::eStencilReference}; |
||||
|
const vk::PipelineDynamicStateCreateInfo dynamic_state_ci( |
||||
|
{}, static_cast<u32>(dynamic_states.size()), dynamic_states.data()); |
||||
|
|
||||
|
vk::PipelineShaderStageRequiredSubgroupSizeCreateInfoEXT subgroup_size_ci; |
||||
|
subgroup_size_ci.requiredSubgroupSize = GuestWarpSize; |
||||
|
|
||||
|
std::vector<vk::PipelineShaderStageCreateInfo> shader_stages; |
||||
|
std::size_t module_index = 0; |
||||
|
for (std::size_t stage = 0; stage < Maxwell::MaxShaderStage; ++stage) { |
||||
|
if (!program[stage]) { |
||||
|
continue; |
||||
|
} |
||||
|
const auto stage_enum = static_cast<Tegra::Engines::ShaderType>(stage); |
||||
|
const auto vk_stage = MaxwellToVK::ShaderStage(stage_enum); |
||||
|
auto& stage_ci = shader_stages.emplace_back(vk::PipelineShaderStageCreateFlags{}, vk_stage, |
||||
|
*modules[module_index++], "main", nullptr); |
||||
|
if (program[stage]->entries.uses_warps && device.IsGuestWarpSizeSupported(vk_stage)) { |
||||
|
stage_ci.pNext = &subgroup_size_ci; |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
const vk::GraphicsPipelineCreateInfo create_info( |
||||
|
{}, static_cast<u32>(shader_stages.size()), shader_stages.data(), &vertex_input_ci, |
||||
|
&input_assembly_ci, &tessellation_ci, &viewport_ci, &rasterizer_ci, &multisampling_ci, |
||||
|
&depth_stencil_ci, &color_blending_ci, &dynamic_state_ci, *layout, renderpass, 0, {}, 0); |
||||
|
|
||||
|
const auto dev = device.GetLogical(); |
||||
|
const auto& dld = device.GetDispatchLoader(); |
||||
|
return dev.createGraphicsPipelineUnique(nullptr, create_info, nullptr, dld); |
||||
|
} |
||||
|
|
||||
|
} // namespace Vulkan
|
||||
@ -0,0 +1,90 @@ |
|||||
|
// Copyright 2019 yuzu Emulator Project |
||||
|
// Licensed under GPLv2 or any later version |
||||
|
// Refer to the license.txt file included. |
||||
|
|
||||
|
#pragma once |
||||
|
|
||||
|
#include <array> |
||||
|
#include <memory> |
||||
|
#include <optional> |
||||
|
#include <unordered_map> |
||||
|
#include <vector> |
||||
|
|
||||
|
#include "video_core/engines/maxwell_3d.h" |
||||
|
#include "video_core/renderer_vulkan/declarations.h" |
||||
|
#include "video_core/renderer_vulkan/fixed_pipeline_state.h" |
||||
|
#include "video_core/renderer_vulkan/vk_descriptor_pool.h" |
||||
|
#include "video_core/renderer_vulkan/vk_renderpass_cache.h" |
||||
|
#include "video_core/renderer_vulkan/vk_resource_manager.h" |
||||
|
#include "video_core/renderer_vulkan/vk_shader_decompiler.h" |
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
using Maxwell = Tegra::Engines::Maxwell3D::Regs; |
||||
|
|
||||
|
struct GraphicsPipelineCacheKey; |
||||
|
|
||||
|
class VKDescriptorPool; |
||||
|
class VKDevice; |
||||
|
class VKRenderPassCache; |
||||
|
class VKScheduler; |
||||
|
class VKUpdateDescriptorQueue; |
||||
|
|
||||
|
using SPIRVProgram = std::array<std::optional<SPIRVShader>, Maxwell::MaxShaderStage>; |
||||
|
|
||||
|
class VKGraphicsPipeline final { |
||||
|
public: |
||||
|
explicit VKGraphicsPipeline(const VKDevice& device, VKScheduler& scheduler, |
||||
|
VKDescriptorPool& descriptor_pool, |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue, |
||||
|
VKRenderPassCache& renderpass_cache, |
||||
|
const GraphicsPipelineCacheKey& key, |
||||
|
const std::vector<vk::DescriptorSetLayoutBinding>& bindings, |
||||
|
const SPIRVProgram& program); |
||||
|
~VKGraphicsPipeline(); |
||||
|
|
||||
|
vk::DescriptorSet CommitDescriptorSet(); |
||||
|
|
||||
|
vk::Pipeline GetHandle() const { |
||||
|
return *pipeline; |
||||
|
} |
||||
|
|
||||
|
vk::PipelineLayout GetLayout() const { |
||||
|
return *layout; |
||||
|
} |
||||
|
|
||||
|
vk::RenderPass GetRenderPass() const { |
||||
|
return renderpass; |
||||
|
} |
||||
|
|
||||
|
private: |
||||
|
UniqueDescriptorSetLayout CreateDescriptorSetLayout( |
||||
|
const std::vector<vk::DescriptorSetLayoutBinding>& bindings) const; |
||||
|
|
||||
|
UniquePipelineLayout CreatePipelineLayout() const; |
||||
|
|
||||
|
UniqueDescriptorUpdateTemplate CreateDescriptorUpdateTemplate( |
||||
|
const SPIRVProgram& program) const; |
||||
|
|
||||
|
std::vector<UniqueShaderModule> CreateShaderModules(const SPIRVProgram& program) const; |
||||
|
|
||||
|
UniquePipeline CreatePipeline(const RenderPassParams& renderpass_params, |
||||
|
const SPIRVProgram& program) const; |
||||
|
|
||||
|
const VKDevice& device; |
||||
|
VKScheduler& scheduler; |
||||
|
const FixedPipelineState fixed_state; |
||||
|
const u64 hash; |
||||
|
|
||||
|
UniqueDescriptorSetLayout descriptor_set_layout; |
||||
|
DescriptorAllocator descriptor_allocator; |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue; |
||||
|
UniquePipelineLayout layout; |
||||
|
UniqueDescriptorUpdateTemplate descriptor_template; |
||||
|
std::vector<UniqueShaderModule> modules; |
||||
|
|
||||
|
vk::RenderPass renderpass; |
||||
|
UniquePipeline pipeline; |
||||
|
}; |
||||
|
|
||||
|
} // namespace Vulkan |
||||
@ -0,0 +1,395 @@ |
|||||
|
// Copyright 2019 yuzu Emulator Project
|
||||
|
// Licensed under GPLv2 or any later version
|
||||
|
// Refer to the license.txt file included.
|
||||
|
|
||||
|
#include <algorithm>
|
||||
|
#include <cstddef>
|
||||
|
#include <memory>
|
||||
|
#include <vector>
|
||||
|
|
||||
|
#include "common/microprofile.h"
|
||||
|
#include "core/core.h"
|
||||
|
#include "core/memory.h"
|
||||
|
#include "video_core/engines/kepler_compute.h"
|
||||
|
#include "video_core/engines/maxwell_3d.h"
|
||||
|
#include "video_core/memory_manager.h"
|
||||
|
#include "video_core/renderer_vulkan/declarations.h"
|
||||
|
#include "video_core/renderer_vulkan/fixed_pipeline_state.h"
|
||||
|
#include "video_core/renderer_vulkan/maxwell_to_vk.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_compute_pipeline.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_descriptor_pool.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_device.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_graphics_pipeline.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_pipeline_cache.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_rasterizer.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_renderpass_cache.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_resource_manager.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_scheduler.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_update_descriptor.h"
|
||||
|
#include "video_core/shader/compiler_settings.h"
|
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
MICROPROFILE_DECLARE(Vulkan_PipelineCache); |
||||
|
|
||||
|
using Tegra::Engines::ShaderType; |
||||
|
|
||||
|
namespace { |
||||
|
|
||||
|
constexpr VideoCommon::Shader::CompilerSettings compiler_settings{ |
||||
|
VideoCommon::Shader::CompileDepth::FullDecompile}; |
||||
|
|
||||
|
/// Gets the address for the specified shader stage program
|
||||
|
GPUVAddr GetShaderAddress(Core::System& system, Maxwell::ShaderProgram program) { |
||||
|
const auto& gpu{system.GPU().Maxwell3D()}; |
||||
|
const auto& shader_config{gpu.regs.shader_config[static_cast<std::size_t>(program)]}; |
||||
|
return gpu.regs.code_address.CodeAddress() + shader_config.offset; |
||||
|
} |
||||
|
|
||||
|
/// Gets if the current instruction offset is a scheduler instruction
|
||||
|
constexpr bool IsSchedInstruction(std::size_t offset, std::size_t main_offset) { |
||||
|
// Sched instructions appear once every 4 instructions.
|
||||
|
constexpr std::size_t SchedPeriod = 4; |
||||
|
const std::size_t absolute_offset = offset - main_offset; |
||||
|
return (absolute_offset % SchedPeriod) == 0; |
||||
|
} |
||||
|
|
||||
|
/// Calculates the size of a program stream
|
||||
|
std::size_t CalculateProgramSize(const ProgramCode& program, bool is_compute) { |
||||
|
const std::size_t start_offset = is_compute ? 0 : 10; |
||||
|
// This is the encoded version of BRA that jumps to itself. All Nvidia
|
||||
|
// shaders end with one.
|
||||
|
constexpr u64 self_jumping_branch = 0xE2400FFFFF07000FULL; |
||||
|
constexpr u64 mask = 0xFFFFFFFFFF7FFFFFULL; |
||||
|
std::size_t offset = start_offset; |
||||
|
while (offset < program.size()) { |
||||
|
const u64 instruction = program[offset]; |
||||
|
if (!IsSchedInstruction(offset, start_offset)) { |
||||
|
if ((instruction & mask) == self_jumping_branch) { |
||||
|
// End on Maxwell's "nop" instruction
|
||||
|
break; |
||||
|
} |
||||
|
if (instruction == 0) { |
||||
|
break; |
||||
|
} |
||||
|
} |
||||
|
++offset; |
||||
|
} |
||||
|
// The last instruction is included in the program size
|
||||
|
return std::min(offset + 1, program.size()); |
||||
|
} |
||||
|
|
||||
|
/// Gets the shader program code from memory for the specified address
|
||||
|
ProgramCode GetShaderCode(Tegra::MemoryManager& memory_manager, const GPUVAddr gpu_addr, |
||||
|
const u8* host_ptr, bool is_compute) { |
||||
|
ProgramCode program_code(VideoCommon::Shader::MAX_PROGRAM_LENGTH); |
||||
|
ASSERT_OR_EXECUTE(host_ptr != nullptr, { |
||||
|
std::fill(program_code.begin(), program_code.end(), 0); |
||||
|
return program_code; |
||||
|
}); |
||||
|
memory_manager.ReadBlockUnsafe(gpu_addr, program_code.data(), |
||||
|
program_code.size() * sizeof(u64)); |
||||
|
program_code.resize(CalculateProgramSize(program_code, is_compute)); |
||||
|
return program_code; |
||||
|
} |
||||
|
|
||||
|
constexpr std::size_t GetStageFromProgram(std::size_t program) { |
||||
|
return program == 0 ? 0 : program - 1; |
||||
|
} |
||||
|
|
||||
|
constexpr ShaderType GetStageFromProgram(Maxwell::ShaderProgram program) { |
||||
|
return static_cast<ShaderType>(GetStageFromProgram(static_cast<std::size_t>(program))); |
||||
|
} |
||||
|
|
||||
|
ShaderType GetShaderType(Maxwell::ShaderProgram program) { |
||||
|
switch (program) { |
||||
|
case Maxwell::ShaderProgram::VertexB: |
||||
|
return ShaderType::Vertex; |
||||
|
case Maxwell::ShaderProgram::TesselationControl: |
||||
|
return ShaderType::TesselationControl; |
||||
|
case Maxwell::ShaderProgram::TesselationEval: |
||||
|
return ShaderType::TesselationEval; |
||||
|
case Maxwell::ShaderProgram::Geometry: |
||||
|
return ShaderType::Geometry; |
||||
|
case Maxwell::ShaderProgram::Fragment: |
||||
|
return ShaderType::Fragment; |
||||
|
default: |
||||
|
UNIMPLEMENTED_MSG("program={}", static_cast<u32>(program)); |
||||
|
return ShaderType::Vertex; |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
u32 FillDescriptorLayout(const ShaderEntries& entries, |
||||
|
std::vector<vk::DescriptorSetLayoutBinding>& bindings, |
||||
|
Maxwell::ShaderProgram program_type, u32 base_binding) { |
||||
|
const ShaderType stage = GetStageFromProgram(program_type); |
||||
|
const vk::ShaderStageFlags stage_flags = MaxwellToVK::ShaderStage(stage); |
||||
|
|
||||
|
u32 binding = base_binding; |
||||
|
const auto AddBindings = [&](vk::DescriptorType descriptor_type, std::size_t num_entries) { |
||||
|
for (std::size_t i = 0; i < num_entries; ++i) { |
||||
|
bindings.emplace_back(binding++, descriptor_type, 1, stage_flags, nullptr); |
||||
|
} |
||||
|
}; |
||||
|
AddBindings(vk::DescriptorType::eUniformBuffer, entries.const_buffers.size()); |
||||
|
AddBindings(vk::DescriptorType::eStorageBuffer, entries.global_buffers.size()); |
||||
|
AddBindings(vk::DescriptorType::eUniformTexelBuffer, entries.texel_buffers.size()); |
||||
|
AddBindings(vk::DescriptorType::eCombinedImageSampler, entries.samplers.size()); |
||||
|
AddBindings(vk::DescriptorType::eStorageImage, entries.images.size()); |
||||
|
return binding; |
||||
|
} |
||||
|
|
||||
|
} // Anonymous namespace
|
||||
|
|
||||
|
CachedShader::CachedShader(Core::System& system, Tegra::Engines::ShaderType stage, |
||||
|
GPUVAddr gpu_addr, VAddr cpu_addr, u8* host_ptr, |
||||
|
ProgramCode program_code, u32 main_offset) |
||||
|
: RasterizerCacheObject{host_ptr}, gpu_addr{gpu_addr}, cpu_addr{cpu_addr}, |
||||
|
program_code{std::move(program_code)}, locker{stage, GetEngine(system, stage)}, |
||||
|
shader_ir{this->program_code, main_offset, compiler_settings, locker}, |
||||
|
entries{GenerateShaderEntries(shader_ir)} {} |
||||
|
|
||||
|
CachedShader::~CachedShader() = default; |
||||
|
|
||||
|
Tegra::Engines::ConstBufferEngineInterface& CachedShader::GetEngine( |
||||
|
Core::System& system, Tegra::Engines::ShaderType stage) { |
||||
|
if (stage == Tegra::Engines::ShaderType::Compute) { |
||||
|
return system.GPU().KeplerCompute(); |
||||
|
} else { |
||||
|
return system.GPU().Maxwell3D(); |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
VKPipelineCache::VKPipelineCache(Core::System& system, RasterizerVulkan& rasterizer, |
||||
|
const VKDevice& device, VKScheduler& scheduler, |
||||
|
VKDescriptorPool& descriptor_pool, |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue) |
||||
|
: RasterizerCache{rasterizer}, system{system}, device{device}, scheduler{scheduler}, |
||||
|
descriptor_pool{descriptor_pool}, update_descriptor_queue{update_descriptor_queue}, |
||||
|
renderpass_cache(device) {} |
||||
|
|
||||
|
VKPipelineCache::~VKPipelineCache() = default; |
||||
|
|
||||
|
std::array<Shader, Maxwell::MaxShaderProgram> VKPipelineCache::GetShaders() { |
||||
|
const auto& gpu = system.GPU().Maxwell3D(); |
||||
|
auto& dirty = system.GPU().Maxwell3D().dirty.shaders; |
||||
|
if (!dirty) { |
||||
|
return last_shaders; |
||||
|
} |
||||
|
dirty = false; |
||||
|
|
||||
|
std::array<Shader, Maxwell::MaxShaderProgram> shaders; |
||||
|
for (std::size_t index = 0; index < Maxwell::MaxShaderProgram; ++index) { |
||||
|
const auto& shader_config = gpu.regs.shader_config[index]; |
||||
|
const auto program{static_cast<Maxwell::ShaderProgram>(index)}; |
||||
|
|
||||
|
// Skip stages that are not enabled
|
||||
|
if (!gpu.regs.IsShaderConfigEnabled(index)) { |
||||
|
continue; |
||||
|
} |
||||
|
|
||||
|
auto& memory_manager{system.GPU().MemoryManager()}; |
||||
|
const GPUVAddr program_addr{GetShaderAddress(system, program)}; |
||||
|
const auto host_ptr{memory_manager.GetPointer(program_addr)}; |
||||
|
auto shader = TryGet(host_ptr); |
||||
|
if (!shader) { |
||||
|
// No shader found - create a new one
|
||||
|
constexpr u32 stage_offset = 10; |
||||
|
const auto stage = static_cast<Tegra::Engines::ShaderType>(index == 0 ? 0 : index - 1); |
||||
|
auto code = GetShaderCode(memory_manager, program_addr, host_ptr, false); |
||||
|
|
||||
|
const std::optional cpu_addr = memory_manager.GpuToCpuAddress(program_addr); |
||||
|
ASSERT(cpu_addr); |
||||
|
|
||||
|
shader = std::make_shared<CachedShader>(system, stage, program_addr, *cpu_addr, |
||||
|
host_ptr, std::move(code), stage_offset); |
||||
|
Register(shader); |
||||
|
} |
||||
|
shaders[index] = std::move(shader); |
||||
|
} |
||||
|
return last_shaders = shaders; |
||||
|
} |
||||
|
|
||||
|
VKGraphicsPipeline& VKPipelineCache::GetGraphicsPipeline(const GraphicsPipelineCacheKey& key) { |
||||
|
MICROPROFILE_SCOPE(Vulkan_PipelineCache); |
||||
|
|
||||
|
if (last_graphics_pipeline && last_graphics_key == key) { |
||||
|
return *last_graphics_pipeline; |
||||
|
} |
||||
|
last_graphics_key = key; |
||||
|
|
||||
|
const auto [pair, is_cache_miss] = graphics_cache.try_emplace(key); |
||||
|
auto& entry = pair->second; |
||||
|
if (is_cache_miss) { |
||||
|
LOG_INFO(Render_Vulkan, "Compile 0x{:016X}", key.Hash()); |
||||
|
const auto [program, bindings] = DecompileShaders(key); |
||||
|
entry = std::make_unique<VKGraphicsPipeline>(device, scheduler, descriptor_pool, |
||||
|
update_descriptor_queue, renderpass_cache, key, |
||||
|
bindings, program); |
||||
|
} |
||||
|
return *(last_graphics_pipeline = entry.get()); |
||||
|
} |
||||
|
|
||||
|
VKComputePipeline& VKPipelineCache::GetComputePipeline(const ComputePipelineCacheKey& key) { |
||||
|
MICROPROFILE_SCOPE(Vulkan_PipelineCache); |
||||
|
|
||||
|
const auto [pair, is_cache_miss] = compute_cache.try_emplace(key); |
||||
|
auto& entry = pair->second; |
||||
|
if (!is_cache_miss) { |
||||
|
return *entry; |
||||
|
} |
||||
|
LOG_INFO(Render_Vulkan, "Compile 0x{:016X}", key.Hash()); |
||||
|
|
||||
|
auto& memory_manager = system.GPU().MemoryManager(); |
||||
|
const auto program_addr = key.shader; |
||||
|
const auto host_ptr = memory_manager.GetPointer(program_addr); |
||||
|
|
||||
|
auto shader = TryGet(host_ptr); |
||||
|
if (!shader) { |
||||
|
// No shader found - create a new one
|
||||
|
const auto cpu_addr = memory_manager.GpuToCpuAddress(program_addr); |
||||
|
ASSERT(cpu_addr); |
||||
|
|
||||
|
auto code = GetShaderCode(memory_manager, program_addr, host_ptr, true); |
||||
|
constexpr u32 kernel_main_offset = 0; |
||||
|
shader = std::make_shared<CachedShader>(system, Tegra::Engines::ShaderType::Compute, |
||||
|
program_addr, *cpu_addr, host_ptr, std::move(code), |
||||
|
kernel_main_offset); |
||||
|
Register(shader); |
||||
|
} |
||||
|
|
||||
|
Specialization specialization; |
||||
|
specialization.workgroup_size = key.workgroup_size; |
||||
|
specialization.shared_memory_size = key.shared_memory_size; |
||||
|
|
||||
|
const SPIRVShader spirv_shader{ |
||||
|
Decompile(device, shader->GetIR(), ShaderType::Compute, specialization), |
||||
|
shader->GetEntries()}; |
||||
|
entry = std::make_unique<VKComputePipeline>(device, scheduler, descriptor_pool, |
||||
|
update_descriptor_queue, spirv_shader); |
||||
|
return *entry; |
||||
|
} |
||||
|
|
||||
|
void VKPipelineCache::Unregister(const Shader& shader) { |
||||
|
bool finished = false; |
||||
|
const auto Finish = [&] { |
||||
|
// TODO(Rodrigo): Instead of finishing here, wait for the fences that use this pipeline and
|
||||
|
// flush.
|
||||
|
if (finished) { |
||||
|
return; |
||||
|
} |
||||
|
finished = true; |
||||
|
scheduler.Finish(); |
||||
|
}; |
||||
|
|
||||
|
const GPUVAddr invalidated_addr = shader->GetGpuAddr(); |
||||
|
for (auto it = graphics_cache.begin(); it != graphics_cache.end();) { |
||||
|
auto& entry = it->first; |
||||
|
if (std::find(entry.shaders.begin(), entry.shaders.end(), invalidated_addr) == |
||||
|
entry.shaders.end()) { |
||||
|
++it; |
||||
|
continue; |
||||
|
} |
||||
|
Finish(); |
||||
|
it = graphics_cache.erase(it); |
||||
|
} |
||||
|
for (auto it = compute_cache.begin(); it != compute_cache.end();) { |
||||
|
auto& entry = it->first; |
||||
|
if (entry.shader != invalidated_addr) { |
||||
|
++it; |
||||
|
continue; |
||||
|
} |
||||
|
Finish(); |
||||
|
it = compute_cache.erase(it); |
||||
|
} |
||||
|
|
||||
|
RasterizerCache::Unregister(shader); |
||||
|
} |
||||
|
|
||||
|
std::pair<SPIRVProgram, std::vector<vk::DescriptorSetLayoutBinding>> |
||||
|
VKPipelineCache::DecompileShaders(const GraphicsPipelineCacheKey& key) { |
||||
|
const auto& fixed_state = key.fixed_state; |
||||
|
auto& memory_manager = system.GPU().MemoryManager(); |
||||
|
const auto& gpu = system.GPU().Maxwell3D(); |
||||
|
|
||||
|
Specialization specialization; |
||||
|
specialization.primitive_topology = fixed_state.input_assembly.topology; |
||||
|
if (specialization.primitive_topology == Maxwell::PrimitiveTopology::Points) { |
||||
|
ASSERT(fixed_state.input_assembly.point_size != 0.0f); |
||||
|
specialization.point_size = fixed_state.input_assembly.point_size; |
||||
|
} |
||||
|
for (std::size_t i = 0; i < Maxwell::NumVertexAttributes; ++i) { |
||||
|
specialization.attribute_types[i] = fixed_state.vertex_input.attributes[i].type; |
||||
|
} |
||||
|
specialization.ndc_minus_one_to_one = fixed_state.rasterizer.ndc_minus_one_to_one; |
||||
|
specialization.tessellation.primitive = fixed_state.tessellation.primitive; |
||||
|
specialization.tessellation.spacing = fixed_state.tessellation.spacing; |
||||
|
specialization.tessellation.clockwise = fixed_state.tessellation.clockwise; |
||||
|
for (const auto& rt : key.renderpass_params.color_attachments) { |
||||
|
specialization.enabled_rendertargets.set(rt.index); |
||||
|
} |
||||
|
|
||||
|
SPIRVProgram program; |
||||
|
std::vector<vk::DescriptorSetLayoutBinding> bindings; |
||||
|
|
||||
|
for (std::size_t index = 0; index < Maxwell::MaxShaderProgram; ++index) { |
||||
|
const auto program_enum = static_cast<Maxwell::ShaderProgram>(index); |
||||
|
|
||||
|
// Skip stages that are not enabled
|
||||
|
if (!gpu.regs.IsShaderConfigEnabled(index)) { |
||||
|
continue; |
||||
|
} |
||||
|
|
||||
|
const GPUVAddr gpu_addr = GetShaderAddress(system, program_enum); |
||||
|
const auto host_ptr = memory_manager.GetPointer(gpu_addr); |
||||
|
const auto shader = TryGet(host_ptr); |
||||
|
ASSERT(shader); |
||||
|
|
||||
|
const std::size_t stage = index == 0 ? 0 : index - 1; // Stage indices are 0 - 5
|
||||
|
const auto program_type = GetShaderType(program_enum); |
||||
|
const auto& entries = shader->GetEntries(); |
||||
|
program[stage] = {Decompile(device, shader->GetIR(), program_type, specialization), |
||||
|
entries}; |
||||
|
|
||||
|
if (program_enum == Maxwell::ShaderProgram::VertexA) { |
||||
|
// VertexB was combined with VertexA, so we skip the VertexB iteration
|
||||
|
++index; |
||||
|
} |
||||
|
|
||||
|
const u32 old_binding = specialization.base_binding; |
||||
|
specialization.base_binding = |
||||
|
FillDescriptorLayout(entries, bindings, program_enum, specialization.base_binding); |
||||
|
ASSERT(old_binding + entries.NumBindings() == specialization.base_binding); |
||||
|
} |
||||
|
return {std::move(program), std::move(bindings)}; |
||||
|
} |
||||
|
|
||||
|
void FillDescriptorUpdateTemplateEntries( |
||||
|
const VKDevice& device, const ShaderEntries& entries, u32& binding, u32& offset, |
||||
|
std::vector<vk::DescriptorUpdateTemplateEntry>& template_entries) { |
||||
|
static constexpr auto entry_size = static_cast<u32>(sizeof(DescriptorUpdateEntry)); |
||||
|
const auto AddEntry = [&](vk::DescriptorType descriptor_type, std::size_t count_) { |
||||
|
const u32 count = static_cast<u32>(count_); |
||||
|
if (descriptor_type == vk::DescriptorType::eUniformTexelBuffer && |
||||
|
device.GetDriverID() == vk::DriverIdKHR::eNvidiaProprietary) { |
||||
|
// Nvidia has a bug where updating multiple uniform texels at once causes the driver to
|
||||
|
// crash.
|
||||
|
for (u32 i = 0; i < count; ++i) { |
||||
|
template_entries.emplace_back(binding + i, 0, 1, descriptor_type, |
||||
|
offset + i * entry_size, entry_size); |
||||
|
} |
||||
|
} else if (count != 0) { |
||||
|
template_entries.emplace_back(binding, 0, count, descriptor_type, offset, entry_size); |
||||
|
} |
||||
|
offset += count * entry_size; |
||||
|
binding += count; |
||||
|
}; |
||||
|
|
||||
|
AddEntry(vk::DescriptorType::eUniformBuffer, entries.const_buffers.size()); |
||||
|
AddEntry(vk::DescriptorType::eStorageBuffer, entries.global_buffers.size()); |
||||
|
AddEntry(vk::DescriptorType::eUniformTexelBuffer, entries.texel_buffers.size()); |
||||
|
AddEntry(vk::DescriptorType::eCombinedImageSampler, entries.samplers.size()); |
||||
|
AddEntry(vk::DescriptorType::eStorageImage, entries.images.size()); |
||||
|
} |
||||
|
|
||||
|
} // namespace Vulkan
|
||||
@ -0,0 +1,200 @@ |
|||||
|
// Copyright 2019 yuzu Emulator Project |
||||
|
// Licensed under GPLv2 or any later version |
||||
|
// Refer to the license.txt file included. |
||||
|
|
||||
|
#pragma once |
||||
|
|
||||
|
#include <array> |
||||
|
#include <cstddef> |
||||
|
#include <memory> |
||||
|
#include <tuple> |
||||
|
#include <type_traits> |
||||
|
#include <unordered_map> |
||||
|
#include <utility> |
||||
|
#include <vector> |
||||
|
|
||||
|
#include <boost/functional/hash.hpp> |
||||
|
|
||||
|
#include "common/common_types.h" |
||||
|
#include "video_core/engines/const_buffer_engine_interface.h" |
||||
|
#include "video_core/engines/maxwell_3d.h" |
||||
|
#include "video_core/rasterizer_cache.h" |
||||
|
#include "video_core/renderer_vulkan/declarations.h" |
||||
|
#include "video_core/renderer_vulkan/fixed_pipeline_state.h" |
||||
|
#include "video_core/renderer_vulkan/vk_graphics_pipeline.h" |
||||
|
#include "video_core/renderer_vulkan/vk_renderpass_cache.h" |
||||
|
#include "video_core/renderer_vulkan/vk_resource_manager.h" |
||||
|
#include "video_core/renderer_vulkan/vk_shader_decompiler.h" |
||||
|
#include "video_core/shader/const_buffer_locker.h" |
||||
|
#include "video_core/shader/shader_ir.h" |
||||
|
#include "video_core/surface.h" |
||||
|
|
||||
|
namespace Core { |
||||
|
class System; |
||||
|
} |
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
class RasterizerVulkan; |
||||
|
class VKComputePipeline; |
||||
|
class VKDescriptorPool; |
||||
|
class VKDevice; |
||||
|
class VKFence; |
||||
|
class VKScheduler; |
||||
|
class VKUpdateDescriptorQueue; |
||||
|
|
||||
|
class CachedShader; |
||||
|
using Shader = std::shared_ptr<CachedShader>; |
||||
|
using Maxwell = Tegra::Engines::Maxwell3D::Regs; |
||||
|
|
||||
|
using ProgramCode = std::vector<u64>; |
||||
|
|
||||
|
struct GraphicsPipelineCacheKey { |
||||
|
FixedPipelineState fixed_state; |
||||
|
std::array<GPUVAddr, Maxwell::MaxShaderProgram> shaders; |
||||
|
RenderPassParams renderpass_params; |
||||
|
|
||||
|
std::size_t Hash() const noexcept { |
||||
|
std::size_t hash = fixed_state.Hash(); |
||||
|
for (const auto& shader : shaders) { |
||||
|
boost::hash_combine(hash, shader); |
||||
|
} |
||||
|
boost::hash_combine(hash, renderpass_params.Hash()); |
||||
|
return hash; |
||||
|
} |
||||
|
|
||||
|
bool operator==(const GraphicsPipelineCacheKey& rhs) const noexcept { |
||||
|
return std::tie(fixed_state, shaders, renderpass_params) == |
||||
|
std::tie(rhs.fixed_state, rhs.shaders, rhs.renderpass_params); |
||||
|
} |
||||
|
}; |
||||
|
|
||||
|
struct ComputePipelineCacheKey { |
||||
|
GPUVAddr shader{}; |
||||
|
u32 shared_memory_size{}; |
||||
|
std::array<u32, 3> workgroup_size{}; |
||||
|
|
||||
|
std::size_t Hash() const noexcept { |
||||
|
return static_cast<std::size_t>(shader) ^ |
||||
|
((static_cast<std::size_t>(shared_memory_size) >> 7) << 40) ^ |
||||
|
static_cast<std::size_t>(workgroup_size[0]) ^ |
||||
|
(static_cast<std::size_t>(workgroup_size[1]) << 16) ^ |
||||
|
(static_cast<std::size_t>(workgroup_size[2]) << 24); |
||||
|
} |
||||
|
|
||||
|
bool operator==(const ComputePipelineCacheKey& rhs) const noexcept { |
||||
|
return std::tie(shader, shared_memory_size, workgroup_size) == |
||||
|
std::tie(rhs.shader, rhs.shared_memory_size, rhs.workgroup_size); |
||||
|
} |
||||
|
}; |
||||
|
|
||||
|
} // namespace Vulkan |
||||
|
|
||||
|
namespace std { |
||||
|
|
||||
|
template <> |
||||
|
struct hash<Vulkan::GraphicsPipelineCacheKey> { |
||||
|
std::size_t operator()(const Vulkan::GraphicsPipelineCacheKey& k) const noexcept { |
||||
|
return k.Hash(); |
||||
|
} |
||||
|
}; |
||||
|
|
||||
|
template <> |
||||
|
struct hash<Vulkan::ComputePipelineCacheKey> { |
||||
|
std::size_t operator()(const Vulkan::ComputePipelineCacheKey& k) const noexcept { |
||||
|
return k.Hash(); |
||||
|
} |
||||
|
}; |
||||
|
|
||||
|
} // namespace std |
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
class CachedShader final : public RasterizerCacheObject { |
||||
|
public: |
||||
|
explicit CachedShader(Core::System& system, Tegra::Engines::ShaderType stage, GPUVAddr gpu_addr, |
||||
|
VAddr cpu_addr, u8* host_ptr, ProgramCode program_code, u32 main_offset); |
||||
|
~CachedShader(); |
||||
|
|
||||
|
GPUVAddr GetGpuAddr() const { |
||||
|
return gpu_addr; |
||||
|
} |
||||
|
|
||||
|
VAddr GetCpuAddr() const override { |
||||
|
return cpu_addr; |
||||
|
} |
||||
|
|
||||
|
std::size_t GetSizeInBytes() const override { |
||||
|
return program_code.size() * sizeof(u64); |
||||
|
} |
||||
|
|
||||
|
VideoCommon::Shader::ShaderIR& GetIR() { |
||||
|
return shader_ir; |
||||
|
} |
||||
|
|
||||
|
const VideoCommon::Shader::ShaderIR& GetIR() const { |
||||
|
return shader_ir; |
||||
|
} |
||||
|
|
||||
|
const ShaderEntries& GetEntries() const { |
||||
|
return entries; |
||||
|
} |
||||
|
|
||||
|
private: |
||||
|
static Tegra::Engines::ConstBufferEngineInterface& GetEngine(Core::System& system, |
||||
|
Tegra::Engines::ShaderType stage); |
||||
|
|
||||
|
GPUVAddr gpu_addr{}; |
||||
|
VAddr cpu_addr{}; |
||||
|
ProgramCode program_code; |
||||
|
VideoCommon::Shader::ConstBufferLocker locker; |
||||
|
VideoCommon::Shader::ShaderIR shader_ir; |
||||
|
ShaderEntries entries; |
||||
|
}; |
||||
|
|
||||
|
class VKPipelineCache final : public RasterizerCache<Shader> { |
||||
|
public: |
||||
|
explicit VKPipelineCache(Core::System& system, RasterizerVulkan& rasterizer, |
||||
|
const VKDevice& device, VKScheduler& scheduler, |
||||
|
VKDescriptorPool& descriptor_pool, |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue); |
||||
|
~VKPipelineCache(); |
||||
|
|
||||
|
std::array<Shader, Maxwell::MaxShaderProgram> GetShaders(); |
||||
|
|
||||
|
VKGraphicsPipeline& GetGraphicsPipeline(const GraphicsPipelineCacheKey& key); |
||||
|
|
||||
|
VKComputePipeline& GetComputePipeline(const ComputePipelineCacheKey& key); |
||||
|
|
||||
|
protected: |
||||
|
void Unregister(const Shader& shader) override; |
||||
|
|
||||
|
void FlushObjectInner(const Shader& object) override {} |
||||
|
|
||||
|
private: |
||||
|
std::pair<SPIRVProgram, std::vector<vk::DescriptorSetLayoutBinding>> DecompileShaders( |
||||
|
const GraphicsPipelineCacheKey& key); |
||||
|
|
||||
|
Core::System& system; |
||||
|
const VKDevice& device; |
||||
|
VKScheduler& scheduler; |
||||
|
VKDescriptorPool& descriptor_pool; |
||||
|
VKUpdateDescriptorQueue& update_descriptor_queue; |
||||
|
|
||||
|
VKRenderPassCache renderpass_cache; |
||||
|
|
||||
|
std::array<Shader, Maxwell::MaxShaderProgram> last_shaders; |
||||
|
|
||||
|
GraphicsPipelineCacheKey last_graphics_key; |
||||
|
VKGraphicsPipeline* last_graphics_pipeline = nullptr; |
||||
|
|
||||
|
std::unordered_map<GraphicsPipelineCacheKey, std::unique_ptr<VKGraphicsPipeline>> |
||||
|
graphics_cache; |
||||
|
std::unordered_map<ComputePipelineCacheKey, std::unique_ptr<VKComputePipeline>> compute_cache; |
||||
|
}; |
||||
|
|
||||
|
void FillDescriptorUpdateTemplateEntries( |
||||
|
const VKDevice& device, const ShaderEntries& entries, u32& binding, u32& offset, |
||||
|
std::vector<vk::DescriptorUpdateTemplateEntry>& template_entries); |
||||
|
|
||||
|
} // namespace Vulkan |
||||
@ -0,0 +1,13 @@ |
|||||
|
// Copyright 2019 yuzu Emulator Project |
||||
|
// Licensed under GPLv2 or any later version |
||||
|
// Refer to the license.txt file included. |
||||
|
|
||||
|
#pragma once |
||||
|
|
||||
|
#include "video_core/rasterizer_interface.h" |
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
class RasterizerVulkan : public VideoCore::RasterizerInterface {}; |
||||
|
|
||||
|
} // namespace Vulkan |
||||
@ -0,0 +1,34 @@ |
|||||
|
// Copyright 2018 yuzu Emulator Project
|
||||
|
// Licensed under GPLv2 or any later version
|
||||
|
// Refer to the license.txt file included.
|
||||
|
|
||||
|
#include <cstring>
|
||||
|
#include <memory>
|
||||
|
#include <vector>
|
||||
|
#include "common/alignment.h"
|
||||
|
#include "common/assert.h"
|
||||
|
#include "common/common_types.h"
|
||||
|
#include "video_core/renderer_vulkan/declarations.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_device.h"
|
||||
|
#include "video_core/renderer_vulkan/vk_shader_util.h"
|
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
UniqueShaderModule BuildShader(const VKDevice& device, std::size_t code_size, const u8* code_data) { |
||||
|
// Avoid undefined behavior by copying to a staging allocation
|
||||
|
ASSERT(code_size % sizeof(u32) == 0); |
||||
|
const auto data = std::make_unique<u32[]>(code_size / sizeof(u32)); |
||||
|
std::memcpy(data.get(), code_data, code_size); |
||||
|
|
||||
|
const auto dev = device.GetLogical(); |
||||
|
const auto& dld = device.GetDispatchLoader(); |
||||
|
const vk::ShaderModuleCreateInfo shader_ci({}, code_size, data.get()); |
||||
|
vk::ShaderModule shader_module; |
||||
|
if (dev.createShaderModule(&shader_ci, nullptr, &shader_module, dld) != vk::Result::eSuccess) { |
||||
|
UNREACHABLE_MSG("Shader module failed to build!"); |
||||
|
} |
||||
|
|
||||
|
return UniqueShaderModule(shader_module, vk::ObjectDestroy(dev, nullptr, dld)); |
||||
|
} |
||||
|
|
||||
|
} // namespace Vulkan
|
||||
@ -0,0 +1,17 @@ |
|||||
|
// Copyright 2018 yuzu Emulator Project |
||||
|
// Licensed under GPLv2 or any later version |
||||
|
// Refer to the license.txt file included. |
||||
|
|
||||
|
#pragma once |
||||
|
|
||||
|
#include <vector> |
||||
|
#include "common/common_types.h" |
||||
|
#include "video_core/renderer_vulkan/declarations.h" |
||||
|
|
||||
|
namespace Vulkan { |
||||
|
|
||||
|
class VKDevice; |
||||
|
|
||||
|
UniqueShaderModule BuildShader(const VKDevice& device, std::size_t code_size, const u8* code_data); |
||||
|
|
||||
|
} // namespace Vulkan |
||||
Write
Preview
Loading…
Cancel
Save
Reference in new issue