diff options
Diffstat (limited to 'src/video_core')
52 files changed, 1043 insertions, 797 deletions
diff --git a/src/video_core/CMakeLists.txt b/src/video_core/CMakeLists.txt index 8ede4ba9b..ff53282c9 100644 --- a/src/video_core/CMakeLists.txt +++ b/src/video_core/CMakeLists.txt @@ -124,6 +124,8 @@ add_library(video_core STATIC shader/decode.cpp shader/expr.cpp shader/expr.h + shader/memory_util.cpp + shader/memory_util.h shader/node_helper.cpp shader/node_helper.h shader/node.h diff --git a/src/video_core/dma_pusher.cpp b/src/video_core/dma_pusher.cpp index 324dafdcd..16311f05e 100644 --- a/src/video_core/dma_pusher.cpp +++ b/src/video_core/dma_pusher.cpp @@ -71,16 +71,22 @@ bool DmaPusher::Step() { gpu.MemoryManager().ReadBlockUnsafe(dma_get, command_headers.data(), command_list_header.size * sizeof(u32)); - for (const CommandHeader& command_header : command_headers) { - - // now, see if we're in the middle of a command - if (dma_state.length_pending) { - // Second word of long non-inc methods command - method count - dma_state.length_pending = 0; - dma_state.method_count = command_header.method_count_; - } else if (dma_state.method_count) { + for (std::size_t index = 0; index < command_headers.size();) { + const CommandHeader& command_header = command_headers[index]; + + if (dma_state.method_count) { // Data word of methods command - CallMethod(command_header.argument); + if (dma_state.non_incrementing) { + const u32 max_write = static_cast<u32>( + std::min<std::size_t>(index + dma_state.method_count, command_headers.size()) - + index); + CallMultiMethod(&command_header.argument, max_write); + dma_state.method_count -= max_write; + index += max_write; + continue; + } else { + CallMethod(command_header.argument); + } if (!dma_state.non_incrementing) { dma_state.method++; @@ -120,6 +126,7 @@ bool DmaPusher::Step() { break; } } + index++; } if (!non_main) { @@ -140,4 +147,9 @@ void DmaPusher::CallMethod(u32 argument) const { gpu.CallMethod({dma_state.method, argument, dma_state.subchannel, dma_state.method_count}); } +void DmaPusher::CallMultiMethod(const u32* base_start, u32 num_methods) const { + gpu.CallMultiMethod(dma_state.method, dma_state.subchannel, base_start, num_methods, + dma_state.method_count); +} + } // namespace Tegra diff --git a/src/video_core/dma_pusher.h b/src/video_core/dma_pusher.h index d6188614a..6cef71306 100644 --- a/src/video_core/dma_pusher.h +++ b/src/video_core/dma_pusher.h @@ -75,6 +75,7 @@ private: void SetState(const CommandHeader& command_header); void CallMethod(u32 argument) const; + void CallMultiMethod(const u32* base_start, u32 num_methods) const; std::vector<CommandHeader> command_headers; ///< Buffer for list of commands fetched at once diff --git a/src/video_core/engines/fermi_2d.cpp b/src/video_core/engines/fermi_2d.cpp index bace6affb..8a47614d2 100644 --- a/src/video_core/engines/fermi_2d.cpp +++ b/src/video_core/engines/fermi_2d.cpp @@ -28,6 +28,12 @@ void Fermi2D::CallMethod(const GPU::MethodCall& method_call) { } } +void Fermi2D::CallMultiMethod(u32 method, const u32* base_start, u32 amount, u32 methods_pending) { + for (std::size_t i = 0; i < amount; i++) { + CallMethod({method, base_start[i], 0, methods_pending - static_cast<u32>(i)}); + } +} + static std::pair<u32, u32> DelimitLine(u32 src_1, u32 src_2, u32 dst_1, u32 dst_2, u32 src_line) { const u32 line_a = src_2 - src_1; const u32 line_b = dst_2 - dst_1; diff --git a/src/video_core/engines/fermi_2d.h b/src/video_core/engines/fermi_2d.h index dba342c70..939a5966d 100644 --- a/src/video_core/engines/fermi_2d.h +++ b/src/video_core/engines/fermi_2d.h @@ -39,6 +39,9 @@ public: /// Write the value to the register identified by method. void CallMethod(const GPU::MethodCall& method_call); + /// Write multiple values to the register identified by method. + void CallMultiMethod(u32 method, const u32* base_start, u32 amount, u32 methods_pending); + enum class Origin : u32 { Center = 0, Corner = 1, diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 368c75a66..00a12175f 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -51,6 +51,13 @@ void KeplerCompute::CallMethod(const GPU::MethodCall& method_call) { } } +void KeplerCompute::CallMultiMethod(u32 method, const u32* base_start, u32 amount, + u32 methods_pending) { + for (std::size_t i = 0; i < amount; i++) { + CallMethod({method, base_start[i], 0, methods_pending - static_cast<u32>(i)}); + } +} + Texture::FullTextureInfo KeplerCompute::GetTexture(std::size_t offset) const { const std::bitset<8> cbuf_mask = launch_description.const_buffer_enable_mask.Value(); ASSERT(cbuf_mask[regs.tex_cb_index]); diff --git a/src/video_core/engines/kepler_compute.h b/src/video_core/engines/kepler_compute.h index eeb79c56f..fe55fdfd0 100644 --- a/src/video_core/engines/kepler_compute.h +++ b/src/video_core/engines/kepler_compute.h @@ -202,6 +202,9 @@ public: /// Write the value to the register identified by method. void CallMethod(const GPU::MethodCall& method_call); + /// Write multiple values to the register identified by method. + void CallMultiMethod(u32 method, const u32* base_start, u32 amount, u32 methods_pending); + Texture::FullTextureInfo GetTexture(std::size_t offset) const; /// Given a texture handle, returns the TSC and TIC entries. diff --git a/src/video_core/engines/kepler_memory.cpp b/src/video_core/engines/kepler_memory.cpp index 597872e43..586ff15dc 100644 --- a/src/video_core/engines/kepler_memory.cpp +++ b/src/video_core/engines/kepler_memory.cpp @@ -41,4 +41,11 @@ void KeplerMemory::CallMethod(const GPU::MethodCall& method_call) { } } +void KeplerMemory::CallMultiMethod(u32 method, const u32* base_start, u32 amount, + u32 methods_pending) { + for (std::size_t i = 0; i < amount; i++) { + CallMethod({method, base_start[i], 0, methods_pending - static_cast<u32>(i)}); + } +} + } // namespace Tegra::Engines diff --git a/src/video_core/engines/kepler_memory.h b/src/video_core/engines/kepler_memory.h index 396fb6e86..bb26fb030 100644 --- a/src/video_core/engines/kepler_memory.h +++ b/src/video_core/engines/kepler_memory.h @@ -40,6 +40,9 @@ public: /// Write the value to the register identified by method. void CallMethod(const GPU::MethodCall& method_call); + /// Write multiple values to the register identified by method. + void CallMultiMethod(u32 method, const u32* base_start, u32 amount, u32 methods_pending); + struct Regs { static constexpr size_t NUM_REGS = 0x7F; diff --git a/src/video_core/engines/maxwell_3d.cpp b/src/video_core/engines/maxwell_3d.cpp index 2824ed707..39e3b66a2 100644 --- a/src/video_core/engines/maxwell_3d.cpp +++ b/src/video_core/engines/maxwell_3d.cpp @@ -280,6 +280,58 @@ void Maxwell3D::CallMethod(const GPU::MethodCall& method_call) { } } +void Maxwell3D::CallMultiMethod(u32 method, const u32* base_start, u32 amount, + u32 methods_pending) { + // Methods after 0xE00 are special, they're actually triggers for some microcode that was + // uploaded to the GPU during initialization. + if (method >= MacroRegistersStart) { + // We're trying to execute a macro + if (executing_macro == 0) { + // A macro call must begin by writing the macro method's register, not its argument. + ASSERT_MSG((method % 2) == 0, + "Can't start macro execution by writing to the ARGS register"); + executing_macro = method; + } + + for (std::size_t i = 0; i < amount; i++) { + macro_params.push_back(base_start[i]); + } + + // Call the macro when there are no more parameters in the command buffer + if (amount == methods_pending) { + CallMacroMethod(executing_macro, macro_params.size(), macro_params.data()); + macro_params.clear(); + } + return; + } + switch (method) { + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[0]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[1]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[2]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[3]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[4]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[5]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[6]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[7]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[8]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[9]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[10]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[11]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[12]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[13]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[14]): + case MAXWELL3D_REG_INDEX(const_buffer.cb_data[15]): { + ProcessCBMultiData(method, base_start, amount); + break; + } + default: { + for (std::size_t i = 0; i < amount; i++) { + CallMethod({method, base_start[i], 0, methods_pending - static_cast<u32>(i)}); + } + } + } +} + void Maxwell3D::StepInstance(const MMEDrawMode expected_mode, const u32 count) { if (mme_draw.current_mode == MMEDrawMode::Undefined) { if (mme_draw.gl_begin_consume) { @@ -570,6 +622,28 @@ void Maxwell3D::StartCBData(u32 method) { ProcessCBData(regs.const_buffer.cb_data[cb_data_state.id]); } +void Maxwell3D::ProcessCBMultiData(u32 method, const u32* start_base, u32 amount) { + if (cb_data_state.current != method) { + if (cb_data_state.current != null_cb_data) { + FinishCBData(); + } + constexpr u32 first_cb_data = MAXWELL3D_REG_INDEX(const_buffer.cb_data[0]); + cb_data_state.start_pos = regs.const_buffer.cb_pos; + cb_data_state.id = method - first_cb_data; + cb_data_state.current = method; + cb_data_state.counter = 0; + } + const std::size_t id = cb_data_state.id; + const std::size_t size = amount; + std::size_t i = 0; + for (; i < size; i++) { + cb_data_state.buffer[id][cb_data_state.counter] = start_base[i]; + cb_data_state.counter++; + } + // Increment the current buffer position. + regs.const_buffer.cb_pos = regs.const_buffer.cb_pos + 4 * amount; +} + void Maxwell3D::FinishCBData() { // Write the input value to the current const buffer at the current position. const GPUVAddr buffer_address = regs.const_buffer.BufferAddress(); diff --git a/src/video_core/engines/maxwell_3d.h b/src/video_core/engines/maxwell_3d.h index 59d5752d2..5e522e0d2 100644 --- a/src/video_core/engines/maxwell_3d.h +++ b/src/video_core/engines/maxwell_3d.h @@ -1179,6 +1179,7 @@ public: BitField<0, 1, u32> depth_range_0_1; BitField<3, 1, u32> depth_clamp_near; BitField<4, 1, u32> depth_clamp_far; + BitField<11, 1, u32> depth_clamp_disabled; } view_volume_clip_control; INSERT_UNION_PADDING_WORDS(0x1F); @@ -1259,7 +1260,8 @@ public: GPUVAddr LimitAddress() const { return static_cast<GPUVAddr>((static_cast<GPUVAddr>(limit_high) << 32) | - limit_low); + limit_low) + + 1; } } vertex_array_limit[NumVertexArrays]; @@ -1358,6 +1360,9 @@ public: /// Write the value to the register identified by method. void CallMethod(const GPU::MethodCall& method_call); + /// Write multiple values to the register identified by method. + void CallMultiMethod(u32 method, const u32* base_start, u32 amount, u32 methods_pending); + /// Write the value to the register identified by method. void CallMethodFromMME(const GPU::MethodCall& method_call); @@ -1511,6 +1516,7 @@ private: /// Handles a write to the CB_DATA[i] register. void StartCBData(u32 method); void ProcessCBData(u32 value); + void ProcessCBMultiData(u32 method, const u32* start_base, u32 amount); void FinishCBData(); /// Handles a write to the CB_BIND register. diff --git a/src/video_core/engines/maxwell_dma.cpp b/src/video_core/engines/maxwell_dma.cpp index 3bfed6ab8..6630005b0 100644 --- a/src/video_core/engines/maxwell_dma.cpp +++ b/src/video_core/engines/maxwell_dma.cpp @@ -36,6 +36,13 @@ void MaxwellDMA::CallMethod(const GPU::MethodCall& method_call) { #undef MAXWELLDMA_REG_INDEX } +void MaxwellDMA::CallMultiMethod(u32 method, const u32* base_start, u32 amount, + u32 methods_pending) { + for (std::size_t i = 0; i < amount; i++) { + CallMethod({method, base_start[i], 0, methods_pending - static_cast<u32>(i)}); + } +} + void MaxwellDMA::HandleCopy() { LOG_TRACE(HW_GPU, "Requested a DMA copy"); diff --git a/src/video_core/engines/maxwell_dma.h b/src/video_core/engines/maxwell_dma.h index 4f40d1d1f..c43ed8194 100644 --- a/src/video_core/engines/maxwell_dma.h +++ b/src/video_core/engines/maxwell_dma.h @@ -35,6 +35,9 @@ public: /// Write the value to the register identified by method. void CallMethod(const GPU::MethodCall& method_call); + /// Write multiple values to the register identified by method. + void CallMultiMethod(u32 method, const u32* base_start, u32 amount, u32 methods_pending); + struct Regs { static constexpr std::size_t NUM_REGS = 0x1D6; diff --git a/src/video_core/engines/shader_bytecode.h b/src/video_core/engines/shader_bytecode.h index 7231597d4..8dae754d4 100644 --- a/src/video_core/engines/shader_bytecode.h +++ b/src/video_core/engines/shader_bytecode.h @@ -655,6 +655,7 @@ union Instruction { } constexpr Instruction(u64 value) : value{value} {} + constexpr Instruction(const Instruction& instr) : value(instr.value) {} BitField<0, 8, Register> gpr0; BitField<8, 8, Register> gpr8; @@ -813,15 +814,17 @@ union Instruction { } alu_integer; union { + BitField<43, 1, u64> x; + } iadd; + + union { BitField<39, 1, u64> ftz; BitField<32, 1, u64> saturate; BitField<49, 2, HalfMerge> merge; - BitField<43, 1, u64> negate_a; BitField<44, 1, u64> abs_a; BitField<47, 2, HalfType> type_a; - BitField<31, 1, u64> negate_b; BitField<30, 1, u64> abs_b; BitField<28, 2, HalfType> type_b; diff --git a/src/video_core/fence_manager.h b/src/video_core/fence_manager.h index dabd1588c..8b2a6a42c 100644 --- a/src/video_core/fence_manager.h +++ b/src/video_core/fence_manager.h @@ -88,7 +88,8 @@ public: } PopAsyncFlushes(); if (current_fence->IsSemaphore()) { - memory_manager.Write<u32>(current_fence->GetAddress(), current_fence->GetPayload()); + memory_manager.template Write<u32>(current_fence->GetAddress(), + current_fence->GetPayload()); } else { gpu.IncrementSyncPoint(current_fence->GetPayload()); } @@ -134,7 +135,8 @@ private: } PopAsyncFlushes(); if (current_fence->IsSemaphore()) { - memory_manager.Write<u32>(current_fence->GetAddress(), current_fence->GetPayload()); + memory_manager.template Write<u32>(current_fence->GetAddress(), + current_fence->GetPayload()); } else { gpu.IncrementSyncPoint(current_fence->GetPayload()); } diff --git a/src/video_core/gpu.cpp b/src/video_core/gpu.cpp index 3b7572d61..b87fd873d 100644 --- a/src/video_core/gpu.cpp +++ b/src/video_core/gpu.cpp @@ -9,6 +9,7 @@ #include "core/core_timing_util.h" #include "core/frontend/emu_window.h" #include "core/memory.h" +#include "core/settings.h" #include "video_core/engines/fermi_2d.h" #include "video_core/engines/kepler_compute.h" #include "video_core/engines/kepler_memory.h" @@ -154,7 +155,10 @@ u64 GPU::GetTicks() const { constexpr u64 gpu_ticks_den = 625; const u64 cpu_ticks = system.CoreTiming().GetTicks(); - const u64 nanoseconds = Core::Timing::CyclesToNs(cpu_ticks).count(); + u64 nanoseconds = Core::Timing::CyclesToNs(cpu_ticks).count(); + if (Settings::values.use_fast_gpu_time) { + nanoseconds /= 256; + } const u64 nanoseconds_num = nanoseconds / gpu_ticks_den; const u64 nanoseconds_rem = nanoseconds % gpu_ticks_den; return nanoseconds_num * gpu_ticks_num + (nanoseconds_rem * gpu_ticks_num) / gpu_ticks_den; @@ -209,16 +213,32 @@ void GPU::CallMethod(const MethodCall& method_call) { ASSERT(method_call.subchannel < bound_engines.size()); - if (ExecuteMethodOnEngine(method_call)) { + if (ExecuteMethodOnEngine(method_call.method)) { CallEngineMethod(method_call); } else { CallPullerMethod(method_call); } } -bool GPU::ExecuteMethodOnEngine(const MethodCall& method_call) { - const auto method = static_cast<BufferMethods>(method_call.method); - return method >= BufferMethods::NonPullerMethods; +void GPU::CallMultiMethod(u32 method, u32 subchannel, const u32* base_start, u32 amount, + u32 methods_pending) { + LOG_TRACE(HW_GPU, "Processing method {:08X} on subchannel {}", method, subchannel); + + ASSERT(subchannel < bound_engines.size()); + + if (ExecuteMethodOnEngine(method)) { + CallEngineMultiMethod(method, subchannel, base_start, amount, methods_pending); + } else { + for (std::size_t i = 0; i < amount; i++) { + CallPullerMethod( + {method, base_start[i], subchannel, methods_pending - static_cast<u32>(i)}); + } + } +} + +bool GPU::ExecuteMethodOnEngine(u32 method) { + const auto buffer_method = static_cast<BufferMethods>(method); + return buffer_method >= BufferMethods::NonPullerMethods; } void GPU::CallPullerMethod(const MethodCall& method_call) { @@ -298,6 +318,31 @@ void GPU::CallEngineMethod(const MethodCall& method_call) { } } +void GPU::CallEngineMultiMethod(u32 method, u32 subchannel, const u32* base_start, u32 amount, + u32 methods_pending) { + const EngineID engine = bound_engines[subchannel]; + + switch (engine) { + case EngineID::FERMI_TWOD_A: + fermi_2d->CallMultiMethod(method, base_start, amount, methods_pending); + break; + case EngineID::MAXWELL_B: + maxwell_3d->CallMultiMethod(method, base_start, amount, methods_pending); + break; + case EngineID::KEPLER_COMPUTE_B: + kepler_compute->CallMultiMethod(method, base_start, amount, methods_pending); + break; + case EngineID::MAXWELL_DMA_COPY_A: + maxwell_dma->CallMultiMethod(method, base_start, amount, methods_pending); + break; + case EngineID::KEPLER_INLINE_TO_MEMORY_B: + kepler_memory->CallMultiMethod(method, base_start, amount, methods_pending); + break; + default: + UNIMPLEMENTED_MSG("Unimplemented engine"); + } +} + void GPU::ProcessBindMethod(const MethodCall& method_call) { // Bind the current subchannel to the desired engine id. LOG_DEBUG(HW_GPU, "Binding subchannel {} to engine {}", method_call.subchannel, diff --git a/src/video_core/gpu.h b/src/video_core/gpu.h index 5e3eb94e9..dd51c95b7 100644 --- a/src/video_core/gpu.h +++ b/src/video_core/gpu.h @@ -155,6 +155,10 @@ public: /// Calls a GPU method. void CallMethod(const MethodCall& method_call); + /// Calls a GPU multivalue method. + void CallMultiMethod(u32 method, u32 subchannel, const u32* base_start, u32 amount, + u32 methods_pending); + /// Flush all current written commands into the host GPU for execution. void FlushCommands(); /// Synchronizes CPU writes with Host GPU memory. @@ -309,8 +313,12 @@ private: /// Calls a GPU engine method. void CallEngineMethod(const MethodCall& method_call); + /// Calls a GPU engine multivalue method. + void CallEngineMultiMethod(u32 method, u32 subchannel, const u32* base_start, u32 amount, + u32 methods_pending); + /// Determines where the method should be executed. - bool ExecuteMethodOnEngine(const MethodCall& method_call); + bool ExecuteMethodOnEngine(u32 method); protected: std::unique_ptr<Tegra::DmaPusher> dma_pusher; diff --git a/src/video_core/renderer_opengl/gl_rasterizer.cpp b/src/video_core/renderer_opengl/gl_rasterizer.cpp index 4c16c89d2..725b4c32d 100644 --- a/src/video_core/renderer_opengl/gl_rasterizer.cpp +++ b/src/video_core/renderer_opengl/gl_rasterizer.cpp @@ -59,14 +59,12 @@ constexpr std::size_t NumSupportedVertexAttributes = 16; template <typename Engine, typename Entry> Tegra::Texture::FullTextureInfo GetTextureInfo(const Engine& engine, const Entry& entry, ShaderType shader_type, std::size_t index = 0) { - if (entry.IsBindless()) { - const Tegra::Texture::TextureHandle tex_handle = - engine.AccessConstBuffer32(shader_type, entry.GetBuffer(), entry.GetOffset()); + if (entry.is_bindless) { + const auto tex_handle = engine.AccessConstBuffer32(shader_type, entry.buffer, entry.offset); return engine.GetTextureInfo(tex_handle); } const auto& gpu_profile = engine.AccessGuestDriverProfile(); - const u32 offset = - entry.GetOffset() + static_cast<u32>(index * gpu_profile.GetTextureHandlerSize()); + const u32 offset = entry.offset + static_cast<u32>(index * gpu_profile.GetTextureHandlerSize()); if constexpr (std::is_same_v<Engine, Tegra::Engines::Maxwell3D>) { return engine.GetStageTexture(shader_type, offset); } else { @@ -186,8 +184,12 @@ void RasterizerOpenGL::SetupVertexBuffer() { const GPUVAddr start = vertex_array.StartAddress(); const GPUVAddr end = regs.vertex_array_limit[index].LimitAddress(); - ASSERT(end > start); - const u64 size = end - start + 1; + ASSERT(end >= start); + const u64 size = end - start; + if (size == 0) { + glBindVertexBuffer(static_cast<GLuint>(index), 0, 0, vertex_array.stride); + continue; + } const auto [vertex_buffer, vertex_buffer_offset] = buffer_cache.UploadMemory(start, size); glBindVertexBuffer(static_cast<GLuint>(index), vertex_buffer, vertex_buffer_offset, vertex_array.stride); @@ -311,8 +313,8 @@ std::size_t RasterizerOpenGL::CalculateVertexArraysSize() const { const GPUVAddr start = regs.vertex_array[index].StartAddress(); const GPUVAddr end = regs.vertex_array_limit[index].LimitAddress(); - ASSERT(end > start); - size += end - start + 1; + size += end - start; + ASSERT(end >= start); } return size; @@ -344,7 +346,7 @@ void RasterizerOpenGL::ConfigureFramebuffers() { texture_cache.GuardRenderTargets(true); - View depth_surface = texture_cache.GetDepthBufferSurface(); + View depth_surface = texture_cache.GetDepthBufferSurface(true); const auto& regs = gpu.regs; UNIMPLEMENTED_IF(regs.rt_separate_frag_data == 0); @@ -353,7 +355,7 @@ void RasterizerOpenGL::ConfigureFramebuffers() { FramebufferCacheKey key; const auto colors_count = static_cast<std::size_t>(regs.rt_control.count); for (std::size_t index = 0; index < colors_count; ++index) { - View color_surface{texture_cache.GetColorBufferSurface(index)}; + View color_surface{texture_cache.GetColorBufferSurface(index, true)}; if (!color_surface) { continue; } @@ -377,28 +379,52 @@ void RasterizerOpenGL::ConfigureFramebuffers() { glBindFramebuffer(GL_DRAW_FRAMEBUFFER, framebuffer_cache.GetFramebuffer(key)); } -void RasterizerOpenGL::ConfigureClearFramebuffer(bool using_color_fb, bool using_depth_fb, - bool using_stencil_fb) { +void RasterizerOpenGL::ConfigureClearFramebuffer(bool using_color, bool using_depth_stencil) { auto& gpu = system.GPU().Maxwell3D(); const auto& regs = gpu.regs; texture_cache.GuardRenderTargets(true); View color_surface; - if (using_color_fb) { + + if (using_color) { + // Determine if we have to preserve the contents. + // First we have to make sure all clear masks are enabled. + bool preserve_contents = !regs.clear_buffers.R || !regs.clear_buffers.G || + !regs.clear_buffers.B || !regs.clear_buffers.A; const std::size_t index = regs.clear_buffers.RT; - color_surface = texture_cache.GetColorBufferSurface(index); + if (regs.clear_flags.scissor) { + // Then we have to confirm scissor testing clears the whole image. + const auto& scissor = regs.scissor_test[0]; + preserve_contents |= scissor.min_x > 0; + preserve_contents |= scissor.min_y > 0; + preserve_contents |= scissor.max_x < regs.rt[index].width; + preserve_contents |= scissor.max_y < regs.rt[index].height; + } + + color_surface = texture_cache.GetColorBufferSurface(index, preserve_contents); texture_cache.MarkColorBufferInUse(index); } + View depth_surface; - if (using_depth_fb || using_stencil_fb) { - depth_surface = texture_cache.GetDepthBufferSurface(); + if (using_depth_stencil) { + bool preserve_contents = false; + if (regs.clear_flags.scissor) { + // For depth stencil clears we only have to confirm scissor test covers the whole image. + const auto& scissor = regs.scissor_test[0]; + preserve_contents |= scissor.min_x > 0; + preserve_contents |= scissor.min_y > 0; + preserve_contents |= scissor.max_x < regs.zeta_width; + preserve_contents |= scissor.max_y < regs.zeta_height; + } + + depth_surface = texture_cache.GetDepthBufferSurface(preserve_contents); texture_cache.MarkDepthBufferInUse(); } texture_cache.GuardRenderTargets(false); FramebufferCacheKey key; - key.colors[0] = color_surface; - key.zeta = depth_surface; + key.colors[0] = std::move(color_surface); + key.zeta = std::move(depth_surface); state_tracker.NotifyFramebuffer(); glBindFramebuffer(GL_DRAW_FRAMEBUFFER, framebuffer_cache.GetFramebuffer(key)); @@ -418,8 +444,7 @@ void RasterizerOpenGL::Clear() { if (regs.clear_buffers.R || regs.clear_buffers.G || regs.clear_buffers.B || regs.clear_buffers.A) { use_color = true; - } - if (use_color) { + state_tracker.NotifyColorMask0(); glColorMaski(0, regs.clear_buffers.R != 0, regs.clear_buffers.G != 0, regs.clear_buffers.B != 0, regs.clear_buffers.A != 0); @@ -457,7 +482,7 @@ void RasterizerOpenGL::Clear() { UNIMPLEMENTED_IF(regs.clear_flags.viewport); - ConfigureClearFramebuffer(use_color, use_depth, use_stencil); + ConfigureClearFramebuffer(use_color, use_depth || use_stencil); if (use_color) { glClearBufferfv(GL_COLOR, 0, regs.clear_color); @@ -829,9 +854,9 @@ void RasterizerOpenGL::SetupDrawGlobalMemory(std::size_t stage_index, const Shad u32 binding = device.GetBaseBindings(stage_index).shader_storage_buffer; for (const auto& entry : shader->GetEntries().global_memory_entries) { - const auto addr{cbufs.const_buffers[entry.GetCbufIndex()].address + entry.GetCbufOffset()}; - const auto gpu_addr{memory_manager.Read<u64>(addr)}; - const auto size{memory_manager.Read<u32>(addr + 8)}; + const GPUVAddr addr{cbufs.const_buffers[entry.cbuf_index].address + entry.cbuf_offset}; + const GPUVAddr gpu_addr{memory_manager.Read<u64>(addr)}; + const u32 size{memory_manager.Read<u32>(addr + 8)}; SetupGlobalMemory(binding++, entry, gpu_addr, size); } } @@ -843,7 +868,7 @@ void RasterizerOpenGL::SetupComputeGlobalMemory(const Shader& kernel) { u32 binding = 0; for (const auto& entry : kernel->GetEntries().global_memory_entries) { - const auto addr{cbufs[entry.GetCbufIndex()].Address() + entry.GetCbufOffset()}; + const auto addr{cbufs[entry.cbuf_index].Address() + entry.cbuf_offset}; const auto gpu_addr{memory_manager.Read<u64>(addr)}; const auto size{memory_manager.Read<u32>(addr + 8)}; SetupGlobalMemory(binding++, entry, gpu_addr, size); @@ -854,7 +879,7 @@ void RasterizerOpenGL::SetupGlobalMemory(u32 binding, const GlobalMemoryEntry& e GPUVAddr gpu_addr, std::size_t size) { const auto alignment{device.GetShaderStorageBufferAlignment()}; const auto [ssbo, buffer_offset] = - buffer_cache.UploadMemory(gpu_addr, size, alignment, entry.IsWritten()); + buffer_cache.UploadMemory(gpu_addr, size, alignment, entry.is_written); glBindBufferRange(GL_SHADER_STORAGE_BUFFER, binding, ssbo, buffer_offset, static_cast<GLsizeiptr>(size)); } @@ -865,7 +890,7 @@ void RasterizerOpenGL::SetupDrawTextures(std::size_t stage_index, const Shader& u32 binding = device.GetBaseBindings(stage_index).sampler; for (const auto& entry : shader->GetEntries().samplers) { const auto shader_type = static_cast<ShaderType>(stage_index); - for (std::size_t i = 0; i < entry.Size(); ++i) { + for (std::size_t i = 0; i < entry.size; ++i) { const auto texture = GetTextureInfo(maxwell3d, entry, shader_type, i); SetupTexture(binding++, texture, entry); } @@ -877,7 +902,7 @@ void RasterizerOpenGL::SetupComputeTextures(const Shader& kernel) { const auto& compute = system.GPU().KeplerCompute(); u32 binding = 0; for (const auto& entry : kernel->GetEntries().samplers) { - for (std::size_t i = 0; i < entry.Size(); ++i) { + for (std::size_t i = 0; i < entry.size; ++i) { const auto texture = GetTextureInfo(compute, entry, ShaderType::Compute, i); SetupTexture(binding++, texture, entry); } @@ -934,7 +959,7 @@ void RasterizerOpenGL::SetupImage(u32 binding, const Tegra::Texture::TICEntry& t if (!tic.IsBuffer()) { view->ApplySwizzle(tic.x_source, tic.y_source, tic.z_source, tic.w_source); } - if (entry.IsWritten()) { + if (entry.is_written) { view->MarkAsModified(texture_cache.Tick()); } glBindImageTexture(binding, view->GetTexture(), 0, GL_TRUE, 0, GL_READ_WRITE, @@ -995,11 +1020,7 @@ void RasterizerOpenGL::SyncDepthClamp() { } flags[Dirty::DepthClampEnabled] = false; - const auto& state = gpu.regs.view_volume_clip_control; - UNIMPLEMENTED_IF_MSG(state.depth_clamp_far != state.depth_clamp_near, - "Unimplemented depth clamp separation!"); - - oglEnable(GL_DEPTH_CLAMP, state.depth_clamp_far || state.depth_clamp_near); + oglEnable(GL_DEPTH_CLAMP, gpu.regs.view_volume_clip_control.depth_clamp_disabled == 0); } void RasterizerOpenGL::SyncClipEnabled(u32 clip_mask) { diff --git a/src/video_core/renderer_opengl/gl_rasterizer.h b/src/video_core/renderer_opengl/gl_rasterizer.h index ebd2173eb..87249fb6f 100644 --- a/src/video_core/renderer_opengl/gl_rasterizer.h +++ b/src/video_core/renderer_opengl/gl_rasterizer.h @@ -95,7 +95,8 @@ private: /// Configures the color and depth framebuffer states. void ConfigureFramebuffers(); - void ConfigureClearFramebuffer(bool using_color_fb, bool using_depth_fb, bool using_stencil_fb); + /// Configures the color and depth framebuffer for clearing. + void ConfigureClearFramebuffer(bool using_color, bool using_depth_stencil); /// Configures the current constbuffers to use for the draw command. void SetupDrawConstBuffers(std::size_t stage_index, const Shader& shader); diff --git a/src/video_core/renderer_opengl/gl_shader_cache.cpp b/src/video_core/renderer_opengl/gl_shader_cache.cpp index f63156b8d..9759a7078 100644 --- a/src/video_core/renderer_opengl/gl_shader_cache.cpp +++ b/src/video_core/renderer_opengl/gl_shader_cache.cpp @@ -10,8 +10,6 @@ #include <thread> #include <unordered_set> -#include <boost/functional/hash.hpp> - #include "common/alignment.h" #include "common/assert.h" #include "common/logging/log.h" @@ -28,76 +26,26 @@ #include "video_core/renderer_opengl/gl_shader_disk_cache.h" #include "video_core/renderer_opengl/gl_state_tracker.h" #include "video_core/renderer_opengl/utils.h" +#include "video_core/shader/memory_util.h" #include "video_core/shader/registry.h" #include "video_core/shader/shader_ir.h" namespace OpenGL { using Tegra::Engines::ShaderType; +using VideoCommon::Shader::GetShaderAddress; +using VideoCommon::Shader::GetShaderCode; +using VideoCommon::Shader::GetUniqueIdentifier; +using VideoCommon::Shader::KERNEL_MAIN_OFFSET; using VideoCommon::Shader::ProgramCode; using VideoCommon::Shader::Registry; using VideoCommon::Shader::ShaderIR; +using VideoCommon::Shader::STAGE_MAIN_OFFSET; namespace { -constexpr u32 STAGE_MAIN_OFFSET = 10; -constexpr u32 KERNEL_MAIN_OFFSET = 0; - constexpr VideoCommon::Shader::CompilerSettings COMPILER_SETTINGS{}; -/// 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) { - constexpr std::size_t start_offset = 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) { - ProgramCode code(VideoCommon::Shader::MAX_PROGRAM_LENGTH); - ASSERT_OR_EXECUTE(host_ptr != nullptr, { - std::fill(code.begin(), code.end(), 0); - return code; - }); - memory_manager.ReadBlockUnsafe(gpu_addr, code.data(), code.size() * sizeof(u64)); - code.resize(CalculateProgramSize(code)); - return code; -} - /// Gets the shader type from a Maxwell program type constexpr GLenum GetGLShaderType(ShaderType shader_type) { switch (shader_type) { @@ -114,17 +62,6 @@ constexpr GLenum GetGLShaderType(ShaderType shader_type) { } } -/// Hashes one (or two) program streams -u64 GetUniqueIdentifier(ShaderType shader_type, bool is_a, const ProgramCode& code, - const ProgramCode& code_b = {}) { - u64 unique_identifier = boost::hash_value(code); - if (is_a) { - // VertexA programs include two programs - boost::hash_combine(unique_identifier, boost::hash_value(code_b)); - } - return unique_identifier; -} - constexpr const char* GetShaderTypeName(ShaderType shader_type) { switch (shader_type) { case ShaderType::Vertex: @@ -456,11 +393,12 @@ Shader ShaderCacheOpenGL::GetStageProgram(Maxwell::ShaderProgram program) { const auto host_ptr{memory_manager.GetPointer(address)}; // No shader found - create a new one - ProgramCode code{GetShaderCode(memory_manager, address, host_ptr)}; + ProgramCode code{GetShaderCode(memory_manager, address, host_ptr, false)}; ProgramCode code_b; if (program == Maxwell::ShaderProgram::VertexA) { const GPUVAddr address_b{GetShaderAddress(system, Maxwell::ShaderProgram::VertexB)}; - code_b = GetShaderCode(memory_manager, address_b, memory_manager.GetPointer(address_b)); + const u8* host_ptr_b = memory_manager.GetPointer(address_b); + code_b = GetShaderCode(memory_manager, address_b, host_ptr_b, false); } const auto unique_identifier = GetUniqueIdentifier( @@ -498,7 +436,7 @@ Shader ShaderCacheOpenGL::GetComputeKernel(GPUVAddr code_addr) { const auto host_ptr{memory_manager.GetPointer(code_addr)}; // No kernel found, create a new one - auto code{GetShaderCode(memory_manager, code_addr, host_ptr)}; + auto code{GetShaderCode(memory_manager, code_addr, host_ptr, true)}; const auto unique_identifier{GetUniqueIdentifier(ShaderType::Compute, false, code)}; const ShaderParameters params{system, disk_cache, device, diff --git a/src/video_core/renderer_opengl/gl_shader_decompiler.cpp b/src/video_core/renderer_opengl/gl_shader_decompiler.cpp index 0cd3ad7e1..99fd4ae2c 100644 --- a/src/video_core/renderer_opengl/gl_shader_decompiler.cpp +++ b/src/video_core/renderer_opengl/gl_shader_decompiler.cpp @@ -870,13 +870,13 @@ private: for (const auto& sampler : ir.GetSamplers()) { const std::string name = GetSampler(sampler); const std::string description = fmt::format("layout (binding = {}) uniform", binding); - binding += sampler.IsIndexed() ? sampler.Size() : 1; + binding += sampler.is_indexed ? sampler.size : 1; std::string sampler_type = [&]() { - if (sampler.IsBuffer()) { + if (sampler.is_buffer) { return "samplerBuffer"; } - switch (sampler.GetType()) { + switch (sampler.type) { case Tegra::Shader::TextureType::Texture1D: return "sampler1D"; case Tegra::Shader::TextureType::Texture2D: @@ -890,17 +890,17 @@ private: return "sampler2D"; } }(); - if (sampler.IsArray()) { + if (sampler.is_array) { sampler_type += "Array"; } - if (sampler.IsShadow()) { + if (sampler.is_shadow) { sampler_type += "Shadow"; } - if (!sampler.IsIndexed()) { + if (!sampler.is_indexed) { code.AddLine("{} {} {};", description, sampler_type, name); } else { - code.AddLine("{} {} {}[{}];", description, sampler_type, name, sampler.Size()); + code.AddLine("{} {} {}[{}];", description, sampler_type, name, sampler.size); } } if (!ir.GetSamplers().empty()) { @@ -946,14 +946,14 @@ private: u32 binding = device.GetBaseBindings(stage).image; for (const auto& image : ir.GetImages()) { std::string qualifier = "coherent volatile"; - if (image.IsRead() && !image.IsWritten()) { + if (image.is_read && !image.is_written) { qualifier += " readonly"; - } else if (image.IsWritten() && !image.IsRead()) { + } else if (image.is_written && !image.is_read) { qualifier += " writeonly"; } - const char* format = image.IsAtomic() ? "r32ui, " : ""; - const char* type_declaration = GetImageTypeDeclaration(image.GetType()); + const char* format = image.is_atomic ? "r32ui, " : ""; + const char* type_declaration = GetImageTypeDeclaration(image.type); code.AddLine("layout ({}binding = {}) {} uniform uimage{} {};", format, binding++, qualifier, type_declaration, GetImage(image)); } @@ -1337,8 +1337,8 @@ private: ASSERT(meta); const std::size_t count = operation.GetOperandsCount(); - const bool has_array = meta->sampler.IsArray(); - const bool has_shadow = meta->sampler.IsShadow(); + const bool has_array = meta->sampler.is_array; + const bool has_shadow = meta->sampler.is_shadow; std::string expr = "texture" + function_suffix; if (!meta->aoffi.empty()) { @@ -1346,7 +1346,7 @@ private: } else if (!meta->ptp.empty()) { expr += "Offsets"; } - if (!meta->sampler.IsIndexed()) { + if (!meta->sampler.is_indexed) { expr += '(' + GetSampler(meta->sampler) + ", "; } else { expr += '(' + GetSampler(meta->sampler) + '[' + Visit(meta->index).AsUint() + "], "; @@ -1870,6 +1870,14 @@ private: return GenerateBinaryInfix(operation, ">=", Type::Bool, type, type); } + Expression LogicalAddCarry(Operation operation) { + const std::string carry = code.GenerateTemporary(); + code.AddLine("uint {};", carry); + code.AddLine("uaddCarry({}, {}, {});", VisitOperand(operation, 0).AsUint(), + VisitOperand(operation, 1).AsUint(), carry); + return {fmt::format("({} != 0)", carry), Type::Bool}; + } + Expression LogicalFIsNan(Operation operation) { return GenerateUnary(operation, "isnan", Type::Bool, Type::Float); } @@ -1974,7 +1982,7 @@ private: std::string expr = GenerateTexture( operation, "", {TextureOffset{}, TextureArgument{Type::Float, meta->bias}}); - if (meta->sampler.IsShadow()) { + if (meta->sampler.is_shadow) { expr = "vec4(" + expr + ')'; } return {expr + GetSwizzle(meta->element), Type::Float}; @@ -1986,7 +1994,7 @@ private: std::string expr = GenerateTexture( operation, "Lod", {TextureArgument{Type::Float, meta->lod}, TextureOffset{}}); - if (meta->sampler.IsShadow()) { + if (meta->sampler.is_shadow) { expr = "vec4(" + expr + ')'; } return {expr + GetSwizzle(meta->element), Type::Float}; @@ -1995,11 +2003,11 @@ private: Expression TextureGather(Operation operation) { const auto& meta = std::get<MetaTexture>(operation.GetMeta()); - const auto type = meta.sampler.IsShadow() ? Type::Float : Type::Int; - const bool separate_dc = meta.sampler.IsShadow(); + const auto type = meta.sampler.is_shadow ? Type::Float : Type::Int; + const bool separate_dc = meta.sampler.is_shadow; std::vector<TextureIR> ir; - if (meta.sampler.IsShadow()) { + if (meta.sampler.is_shadow) { ir = {TextureOffset{}}; } else { ir = {TextureOffset{}, TextureArgument{type, meta.component}}; @@ -2044,7 +2052,7 @@ private: constexpr std::array constructors = {"int", "ivec2", "ivec3", "ivec4"}; const auto meta = std::get_if<MetaTexture>(&operation.GetMeta()); ASSERT(meta); - UNIMPLEMENTED_IF(meta->sampler.IsArray()); + UNIMPLEMENTED_IF(meta->sampler.is_array); const std::size_t count = operation.GetOperandsCount(); std::string expr = "texelFetch("; @@ -2065,7 +2073,7 @@ private: } expr += ')'; - if (meta->lod && !meta->sampler.IsBuffer()) { + if (meta->lod && !meta->sampler.is_buffer) { expr += ", "; expr += Visit(meta->lod).AsInt(); } @@ -2076,12 +2084,10 @@ private: } Expression TextureGradient(Operation operation) { - const auto meta = std::get_if<MetaTexture>(&operation.GetMeta()); - ASSERT(meta); - + const auto& meta = std::get<MetaTexture>(operation.GetMeta()); std::string expr = GenerateTexture(operation, "Grad", {TextureDerivates{}, TextureOffset{}}); - return {std::move(expr) + GetSwizzle(meta->element), Type::Float}; + return {std::move(expr) + GetSwizzle(meta.element), Type::Float}; } Expression ImageLoad(Operation operation) { @@ -2441,6 +2447,8 @@ private: &GLSLDecompiler::LogicalNotEqual<Type::Uint>, &GLSLDecompiler::LogicalGreaterEqual<Type::Uint>, + &GLSLDecompiler::LogicalAddCarry, + &GLSLDecompiler::Logical2HLessThan<false>, &GLSLDecompiler::Logical2HEqual<false>, &GLSLDecompiler::Logical2HLessEqual<false>, @@ -2598,11 +2606,11 @@ private: } std::string GetSampler(const Sampler& sampler) const { - return AppendSuffix(static_cast<u32>(sampler.GetIndex()), "sampler"); + return AppendSuffix(sampler.index, "sampler"); } std::string GetImage(const Image& image) const { - return AppendSuffix(static_cast<u32>(image.GetIndex()), "image"); + return AppendSuffix(image.index, "image"); } std::string AppendSuffix(u32 index, std::string_view name) const { diff --git a/src/video_core/renderer_opengl/gl_shader_decompiler.h b/src/video_core/renderer_opengl/gl_shader_decompiler.h index e7dbd810c..e8a178764 100644 --- a/src/video_core/renderer_opengl/gl_shader_decompiler.h +++ b/src/video_core/renderer_opengl/gl_shader_decompiler.h @@ -33,36 +33,19 @@ public: } private: - u32 index{}; + u32 index = 0; }; -class GlobalMemoryEntry { -public: - explicit GlobalMemoryEntry(u32 cbuf_index, u32 cbuf_offset, bool is_read, bool is_written) +struct GlobalMemoryEntry { + constexpr explicit GlobalMemoryEntry(u32 cbuf_index, u32 cbuf_offset, bool is_read, + bool is_written) : cbuf_index{cbuf_index}, cbuf_offset{cbuf_offset}, is_read{is_read}, is_written{ is_written} {} - u32 GetCbufIndex() const { - return cbuf_index; - } - - u32 GetCbufOffset() const { - return cbuf_offset; - } - - bool IsRead() const { - return is_read; - } - - bool IsWritten() const { - return is_written; - } - -private: - u32 cbuf_index{}; - u32 cbuf_offset{}; - bool is_read{}; - bool is_written{}; + u32 cbuf_index = 0; + u32 cbuf_offset = 0; + bool is_read = false; + bool is_written = false; }; struct ShaderEntries { diff --git a/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp b/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp index be1c31978..648b1e71b 100644 --- a/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp +++ b/src/video_core/renderer_vulkan/fixed_pipeline_state.cpp @@ -81,7 +81,7 @@ void FixedPipelineState::Rasterizer::Fill(const Maxwell& regs) noexcept { primitive_restart_enable.Assign(regs.primitive_restart.enabled != 0 ? 1 : 0); cull_enable.Assign(regs.cull_test_enabled != 0 ? 1 : 0); depth_bias_enable.Assign(enabled_lut[POLYGON_OFFSET_ENABLE_LUT[topology_index]] != 0 ? 1 : 0); - depth_clamp_enable.Assign(clip.depth_clamp_near == 1 || clip.depth_clamp_far == 1 ? 1 : 0); + depth_clamp_disabled.Assign(regs.view_volume_clip_control.depth_clamp_disabled.Value()); ndc_minus_one_to_one.Assign(regs.depth_mode == Maxwell::DepthMode::MinusOneToOne ? 1 : 0); cull_face.Assign(PackCullFace(regs.cull_face)); front_face.Assign(packed_front_face); @@ -140,6 +140,12 @@ void FixedPipelineState::BlendingAttachment::Fill(const Maxwell& regs, std::size enable.Assign(1); } +void FixedPipelineState::Fill(const Maxwell& regs) { + rasterizer.Fill(regs); + depth_stencil.Fill(regs); + color_blending.Fill(regs); +} + std::size_t FixedPipelineState::Hash() const noexcept { const u64 hash = Common::CityHash64(reinterpret_cast<const char*>(this), sizeof *this); return static_cast<std::size_t>(hash); @@ -149,15 +155,6 @@ bool FixedPipelineState::operator==(const FixedPipelineState& rhs) const noexcep return std::memcmp(this, &rhs, sizeof *this) == 0; } -FixedPipelineState GetFixedPipelineState(const Maxwell& regs) { - FixedPipelineState fixed_state; - fixed_state.rasterizer.Fill(regs); - fixed_state.depth_stencil.Fill(regs); - fixed_state.color_blending.Fill(regs); - fixed_state.padding = {}; - return fixed_state; -} - u32 FixedPipelineState::PackComparisonOp(Maxwell::ComparisonOp op) noexcept { // OpenGL enums go from 0x200 to 0x207 and the others from 1 to 8 // If we substract 0x200 to OpenGL enums and 1 to the others we get a 0-7 range. diff --git a/src/video_core/renderer_vulkan/fixed_pipeline_state.h b/src/video_core/renderer_vulkan/fixed_pipeline_state.h index 9fe6bdbf9..8652067a7 100644 --- a/src/video_core/renderer_vulkan/fixed_pipeline_state.h +++ b/src/video_core/renderer_vulkan/fixed_pipeline_state.h @@ -17,7 +17,7 @@ namespace Vulkan { using Maxwell = Tegra::Engines::Maxwell3D::Regs; -struct alignas(32) FixedPipelineState { +struct FixedPipelineState { static u32 PackComparisonOp(Maxwell::ComparisonOp op) noexcept; static Maxwell::ComparisonOp UnpackComparisonOp(u32 packed) noexcept; @@ -129,7 +129,7 @@ struct alignas(32) FixedPipelineState { auto& binding = bindings[index]; binding.raw = 0; binding.enabled.Assign(enabled ? 1 : 0); - binding.stride.Assign(stride); + binding.stride.Assign(static_cast<u16>(stride)); binding_divisors[index] = divisor; } @@ -153,7 +153,7 @@ struct alignas(32) FixedPipelineState { BitField<4, 1, u32> primitive_restart_enable; BitField<5, 1, u32> cull_enable; BitField<6, 1, u32> depth_bias_enable; - BitField<7, 1, u32> depth_clamp_enable; + BitField<7, 1, u32> depth_clamp_disabled; BitField<8, 1, u32> ndc_minus_one_to_one; BitField<9, 2, u32> cull_face; BitField<11, 1, u32> front_face; @@ -237,7 +237,8 @@ struct alignas(32) FixedPipelineState { Rasterizer rasterizer; DepthStencil depth_stencil; ColorBlending color_blending; - std::array<u8, 20> padding; + + void Fill(const Maxwell& regs); std::size_t Hash() const noexcept; @@ -250,9 +251,6 @@ struct alignas(32) FixedPipelineState { static_assert(std::has_unique_object_representations_v<FixedPipelineState>); static_assert(std::is_trivially_copyable_v<FixedPipelineState>); static_assert(std::is_trivially_constructible_v<FixedPipelineState>); -static_assert(sizeof(FixedPipelineState) % 32 == 0, "Size is not aligned"); - -FixedPipelineState GetFixedPipelineState(const Maxwell& regs); } // namespace Vulkan diff --git a/src/video_core/renderer_vulkan/vk_device.h b/src/video_core/renderer_vulkan/vk_device.h index a4d841e26..c8640762d 100644 --- a/src/video_core/renderer_vulkan/vk_device.h +++ b/src/video_core/renderer_vulkan/vk_device.h @@ -82,11 +82,6 @@ public: return present_family; } - /// Returns true if the device is integrated with the host CPU. - bool IsIntegrated() const { - return properties.deviceType == VK_PHYSICAL_DEVICE_TYPE_INTEGRATED_GPU; - } - /// Returns the current Vulkan API version provided in Vulkan-formatted version numbers. u32 GetApiVersion() const { return properties.apiVersion; diff --git a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp index 8332b42aa..852a17a70 100644 --- a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp +++ b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp @@ -249,7 +249,7 @@ vk::Pipeline VKGraphicsPipeline::CreatePipeline(const RenderPassParams& renderpa rasterization_ci.sType = VK_STRUCTURE_TYPE_PIPELINE_RASTERIZATION_STATE_CREATE_INFO; rasterization_ci.pNext = nullptr; rasterization_ci.flags = 0; - rasterization_ci.depthClampEnable = rs.depth_clamp_enable; + rasterization_ci.depthClampEnable = rs.depth_clamp_disabled == 0 ? VK_TRUE : VK_FALSE; rasterization_ci.rasterizerDiscardEnable = VK_FALSE; rasterization_ci.polygonMode = VK_POLYGON_MODE_FILL; rasterization_ci.cullMode = @@ -288,7 +288,7 @@ vk::Pipeline VKGraphicsPipeline::CreatePipeline(const RenderPassParams& renderpa depth_stencil_ci.maxDepthBounds = 0.0f; std::array<VkPipelineColorBlendAttachmentState, Maxwell::NumRenderTargets> cb_attachments; - const std::size_t num_attachments = renderpass_params.color_attachments.size(); + const auto num_attachments = static_cast<std::size_t>(renderpass_params.num_color_attachments); for (std::size_t index = 0; index < num_attachments; ++index) { static constexpr std::array COMPONENT_TABLE = { VK_COLOR_COMPONENT_R_BIT, VK_COLOR_COMPONENT_G_BIT, VK_COLOR_COMPONENT_B_BIT, diff --git a/src/video_core/renderer_vulkan/vk_memory_manager.cpp b/src/video_core/renderer_vulkan/vk_memory_manager.cpp index 6a9e658bf..b4c650a63 100644 --- a/src/video_core/renderer_vulkan/vk_memory_manager.cpp +++ b/src/video_core/renderer_vulkan/vk_memory_manager.cpp @@ -118,8 +118,7 @@ private: }; VKMemoryManager::VKMemoryManager(const VKDevice& device) - : device{device}, properties{device.GetPhysical().GetMemoryProperties()}, - is_memory_unified{GetMemoryUnified(properties)} {} + : device{device}, properties{device.GetPhysical().GetMemoryProperties()} {} VKMemoryManager::~VKMemoryManager() = default; @@ -209,16 +208,6 @@ VKMemoryCommit VKMemoryManager::TryAllocCommit(const VkMemoryRequirements& requi return {}; } -bool VKMemoryManager::GetMemoryUnified(const VkPhysicalDeviceMemoryProperties& properties) { - for (u32 heap_index = 0; heap_index < properties.memoryHeapCount; ++heap_index) { - if (!(properties.memoryHeaps[heap_index].flags & VK_MEMORY_HEAP_DEVICE_LOCAL_BIT)) { - // Memory is considered unified when heaps are device local only. - return false; - } - } - return true; -} - VKMemoryCommitImpl::VKMemoryCommitImpl(const VKDevice& device, VKMemoryAllocation* allocation, const vk::DeviceMemory& memory, u64 begin, u64 end) : device{device}, memory{memory}, interval{begin, end}, allocation{allocation} {} diff --git a/src/video_core/renderer_vulkan/vk_memory_manager.h b/src/video_core/renderer_vulkan/vk_memory_manager.h index 5b6858e9b..1af88e3d4 100644 --- a/src/video_core/renderer_vulkan/vk_memory_manager.h +++ b/src/video_core/renderer_vulkan/vk_memory_manager.h @@ -40,11 +40,6 @@ public: /// Commits memory required by the image and binds it. VKMemoryCommit Commit(const vk::Image& image, bool host_visible); - /// Returns true if the memory allocations are done always in host visible and coherent memory. - bool IsMemoryUnified() const { - return is_memory_unified; - } - private: /// Allocates a chunk of memory. bool AllocMemory(VkMemoryPropertyFlags wanted_properties, u32 type_mask, u64 size); @@ -53,12 +48,8 @@ private: VKMemoryCommit TryAllocCommit(const VkMemoryRequirements& requirements, VkMemoryPropertyFlags wanted_properties); - /// Returns true if the device uses an unified memory model. - static bool GetMemoryUnified(const VkPhysicalDeviceMemoryProperties& properties); - - const VKDevice& device; ///< Device handler. - const VkPhysicalDeviceMemoryProperties properties; ///< Physical device properties. - const bool is_memory_unified; ///< True if memory model is unified. + const VKDevice& device; ///< Device handler. + const VkPhysicalDeviceMemoryProperties properties; ///< Physical device properties. std::vector<std::unique_ptr<VKMemoryAllocation>> allocations; ///< Current allocations. }; diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp index 91b1b16a5..8fbd63dbc 100644 --- a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp +++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp @@ -27,12 +27,18 @@ #include "video_core/renderer_vulkan/vk_update_descriptor.h" #include "video_core/renderer_vulkan/wrapper.h" #include "video_core/shader/compiler_settings.h" +#include "video_core/shader/memory_util.h" namespace Vulkan { MICROPROFILE_DECLARE(Vulkan_PipelineCache); using Tegra::Engines::ShaderType; +using VideoCommon::Shader::GetShaderAddress; +using VideoCommon::Shader::GetShaderCode; +using VideoCommon::Shader::KERNEL_MAIN_OFFSET; +using VideoCommon::Shader::ProgramCode; +using VideoCommon::Shader::STAGE_MAIN_OFFSET; namespace { @@ -45,60 +51,6 @@ constexpr VkDescriptorType STORAGE_IMAGE = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE; 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; } @@ -133,7 +85,7 @@ void AddBindings(std::vector<VkDescriptorSetLayoutBinding>& bindings, u32& bindi u32 count = 1; if constexpr (descriptor_type == VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER) { // Combined image samplers can be arrayed. - count = container[i].Size(); + count = container[i].size; } VkDescriptorSetLayoutBinding& entry = bindings.emplace_back(); entry.binding = binding++; @@ -161,6 +113,24 @@ u32 FillDescriptorLayout(const ShaderEntries& entries, } // Anonymous namespace +std::size_t GraphicsPipelineCacheKey::Hash() const noexcept { + const u64 hash = Common::CityHash64(reinterpret_cast<const char*>(this), sizeof *this); + return static_cast<std::size_t>(hash); +} + +bool GraphicsPipelineCacheKey::operator==(const GraphicsPipelineCacheKey& rhs) const noexcept { + return std::memcmp(&rhs, this, sizeof *this) == 0; +} + +std::size_t ComputePipelineCacheKey::Hash() const noexcept { + const u64 hash = Common::CityHash64(reinterpret_cast<const char*>(this), sizeof *this); + return static_cast<std::size_t>(hash); +} + +bool ComputePipelineCacheKey::operator==(const ComputePipelineCacheKey& rhs) const noexcept { + return std::memcmp(&rhs, this, sizeof *this) == 0; +} + CachedShader::CachedShader(Core::System& system, Tegra::Engines::ShaderType stage, GPUVAddr gpu_addr, VAddr cpu_addr, ProgramCode program_code, u32 main_offset) @@ -212,9 +182,9 @@ std::array<Shader, Maxwell::MaxShaderProgram> VKPipelineCache::GetShaders() { const auto host_ptr{memory_manager.GetPointer(program_addr)}; // No shader found - create a new one - constexpr u32 stage_offset = 10; + constexpr u32 stage_offset = STAGE_MAIN_OFFSET; const auto stage = static_cast<Tegra::Engines::ShaderType>(index == 0 ? 0 : index - 1); - auto code = GetShaderCode(memory_manager, program_addr, host_ptr, false); + ProgramCode code = GetShaderCode(memory_manager, program_addr, host_ptr, false); shader = std::make_shared<CachedShader>(system, stage, program_addr, *cpu_addr, std::move(code), stage_offset); @@ -270,11 +240,10 @@ VKComputePipeline& VKPipelineCache::GetComputePipeline(const ComputePipelineCach // No shader found - create a new one const auto host_ptr = memory_manager.GetPointer(program_addr); - auto code = GetShaderCode(memory_manager, program_addr, host_ptr, true); - constexpr u32 kernel_main_offset = 0; + ProgramCode code = GetShaderCode(memory_manager, program_addr, host_ptr, true); shader = std::make_shared<CachedShader>(system, Tegra::Engines::ShaderType::Compute, program_addr, *cpu_addr, std::move(code), - kernel_main_offset); + KERNEL_MAIN_OFFSET); if (cpu_addr) { Register(shader); } else { @@ -393,7 +362,7 @@ void AddEntry(std::vector<VkDescriptorUpdateTemplateEntry>& template_entries, u3 if constexpr (descriptor_type == COMBINED_IMAGE_SAMPLER) { for (u32 i = 0; i < count; ++i) { - const u32 num_samplers = container[i].Size(); + const u32 num_samplers = container[i].size; VkDescriptorUpdateTemplateEntry& entry = template_entries.emplace_back(); entry.dstBinding = binding; entry.dstArrayElement = 0; diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.h b/src/video_core/renderer_vulkan/vk_pipeline_cache.h index 602a0a340..ebddafb73 100644 --- a/src/video_core/renderer_vulkan/vk_pipeline_cache.h +++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.h @@ -7,7 +7,6 @@ #include <array> #include <cstddef> #include <memory> -#include <tuple> #include <type_traits> #include <unordered_map> #include <utility> @@ -25,6 +24,7 @@ #include "video_core/renderer_vulkan/vk_resource_manager.h" #include "video_core/renderer_vulkan/vk_shader_decompiler.h" #include "video_core/renderer_vulkan/wrapper.h" +#include "video_core/shader/memory_util.h" #include "video_core/shader/registry.h" #include "video_core/shader/shader_ir.h" #include "video_core/surface.h" @@ -47,46 +47,40 @@ 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::array<GPUVAddr, Maxwell::MaxShaderProgram> shaders; + u64 padding; // This is necessary for unique object representations - 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; - } + std::size_t Hash() const noexcept; - 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); + bool operator==(const GraphicsPipelineCacheKey& rhs) const noexcept; + + bool operator!=(const GraphicsPipelineCacheKey& rhs) const noexcept { + return !operator==(rhs); } }; +static_assert(std::has_unique_object_representations_v<GraphicsPipelineCacheKey>); +static_assert(std::is_trivially_copyable_v<GraphicsPipelineCacheKey>); +static_assert(std::is_trivially_constructible_v<GraphicsPipelineCacheKey>); 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); - } + GPUVAddr shader; + u32 shared_memory_size; + std::array<u32, 3> workgroup_size; + + std::size_t Hash() const noexcept; + + bool operator==(const ComputePipelineCacheKey& rhs) const noexcept; - 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); + bool operator!=(const ComputePipelineCacheKey& rhs) const noexcept { + return !operator==(rhs); } }; +static_assert(std::has_unique_object_representations_v<ComputePipelineCacheKey>); +static_assert(std::is_trivially_copyable_v<ComputePipelineCacheKey>); +static_assert(std::is_trivially_constructible_v<ComputePipelineCacheKey>); } // namespace Vulkan @@ -113,7 +107,8 @@ namespace Vulkan { class CachedShader final : public RasterizerCacheObject { public: explicit CachedShader(Core::System& system, Tegra::Engines::ShaderType stage, GPUVAddr gpu_addr, - VAddr cpu_addr, ProgramCode program_code, u32 main_offset); + VAddr cpu_addr, VideoCommon::Shader::ProgramCode program_code, + u32 main_offset); ~CachedShader(); GPUVAddr GetGpuAddr() const { @@ -145,7 +140,7 @@ private: Tegra::Engines::ShaderType stage); GPUVAddr gpu_addr{}; - ProgramCode program_code; + VideoCommon::Shader::ProgramCode program_code; VideoCommon::Shader::Registry registry; VideoCommon::Shader::ShaderIR shader_ir; ShaderEntries entries; diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.cpp b/src/video_core/renderer_vulkan/vk_rasterizer.cpp index 8a1f57891..127f77b2c 100644 --- a/src/video_core/renderer_vulkan/vk_rasterizer.cpp +++ b/src/video_core/renderer_vulkan/vk_rasterizer.cpp @@ -119,14 +119,13 @@ template <typename Engine, typename Entry> Tegra::Texture::FullTextureInfo GetTextureInfo(const Engine& engine, const Entry& entry, std::size_t stage, std::size_t index = 0) { const auto stage_type = static_cast<Tegra::Engines::ShaderType>(stage); - if (entry.IsBindless()) { - const Tegra::Texture::TextureHandle tex_handle = - engine.AccessConstBuffer32(stage_type, entry.GetBuffer(), entry.GetOffset()); + if (entry.is_bindless) { + const auto tex_handle = engine.AccessConstBuffer32(stage_type, entry.buffer, entry.offset); return engine.GetTextureInfo(tex_handle); } const auto& gpu_profile = engine.AccessGuestDriverProfile(); const u32 entry_offset = static_cast<u32>(index * gpu_profile.GetTextureHandlerSize()); - const u32 offset = entry.GetOffset() + entry_offset; + const u32 offset = entry.offset + entry_offset; if constexpr (std::is_same_v<Engine, Tegra::Engines::Maxwell3D>) { return engine.GetStageTexture(stage_type, offset); } else { @@ -316,7 +315,8 @@ void RasterizerVulkan::Draw(bool is_indexed, bool is_instanced) { query_cache.UpdateCounters(); const auto& gpu = system.GPU().Maxwell3D(); - GraphicsPipelineCacheKey key{GetFixedPipelineState(gpu.regs)}; + GraphicsPipelineCacheKey key; + key.fixed_state.Fill(gpu.regs); buffer_cache.Map(CalculateGraphicsStreamBufferSize(is_indexed)); @@ -334,10 +334,11 @@ void RasterizerVulkan::Draw(bool is_indexed, bool is_instanced) { buffer_cache.Unmap(); - const auto texceptions = UpdateAttachments(); + const Texceptions texceptions = UpdateAttachments(); SetupImageTransitions(texceptions, color_attachments, zeta_attachment); key.renderpass_params = GetRenderPassParams(texceptions); + key.padding = 0; auto& pipeline = pipeline_cache.GetGraphicsPipeline(key); scheduler.BindGraphicsPipeline(pipeline.GetHandle()); @@ -453,10 +454,12 @@ void RasterizerVulkan::DispatchCompute(GPUVAddr code_addr) { query_cache.UpdateCounters(); const auto& launch_desc = system.GPU().KeplerCompute().launch_description; - const ComputePipelineCacheKey key{ - code_addr, - launch_desc.shared_alloc, - {launch_desc.block_dim_x, launch_desc.block_dim_y, launch_desc.block_dim_z}}; + ComputePipelineCacheKey key; + key.shader = code_addr; + key.shared_memory_size = launch_desc.shared_alloc; + key.workgroup_size = {launch_desc.block_dim_x, launch_desc.block_dim_y, + launch_desc.block_dim_z}; + auto& pipeline = pipeline_cache.GetComputePipeline(key); // Compute dispatches can't be executed inside a renderpass @@ -652,7 +655,7 @@ RasterizerVulkan::Texceptions RasterizerVulkan::UpdateAttachments() { Texceptions texceptions; for (std::size_t rt = 0; rt < Maxwell::NumRenderTargets; ++rt) { if (update_rendertargets) { - color_attachments[rt] = texture_cache.GetColorBufferSurface(rt); + color_attachments[rt] = texture_cache.GetColorBufferSurface(rt, true); } if (color_attachments[rt] && WalkAttachmentOverlaps(*color_attachments[rt])) { texceptions[rt] = true; @@ -660,7 +663,7 @@ RasterizerVulkan::Texceptions RasterizerVulkan::UpdateAttachments() { } if (update_rendertargets) { - zeta_attachment = texture_cache.GetDepthBufferSurface(); + zeta_attachment = texture_cache.GetDepthBufferSurface(true); } if (zeta_attachment && WalkAttachmentOverlaps(*zeta_attachment)) { texceptions[ZETA_TEXCEPTION_INDEX] = true; @@ -688,7 +691,7 @@ std::tuple<VkFramebuffer, VkExtent2D> RasterizerVulkan::ConfigureFramebuffers( FramebufferCacheKey key{renderpass, std::numeric_limits<u32>::max(), std::numeric_limits<u32>::max(), std::numeric_limits<u32>::max()}; - const auto try_push = [&](const View& view) { + const auto try_push = [&key](const View& view) { if (!view) { return false; } @@ -699,7 +702,9 @@ std::tuple<VkFramebuffer, VkExtent2D> RasterizerVulkan::ConfigureFramebuffers( return true; }; - for (std::size_t index = 0; index < std::size(color_attachments); ++index) { + const auto& regs = system.GPU().Maxwell3D().regs; + const std::size_t num_attachments = static_cast<std::size_t>(regs.rt_control.count); + for (std::size_t index = 0; index < num_attachments; ++index) { if (try_push(color_attachments[index])) { texture_cache.MarkColorBufferInUse(index); } @@ -877,8 +882,12 @@ void RasterizerVulkan::SetupVertexArrays(FixedPipelineState::VertexInput& vertex const GPUVAddr start{vertex_array.StartAddress()}; const GPUVAddr end{regs.vertex_array_limit[index].LimitAddress()}; - ASSERT(end > start); - const std::size_t size{end - start + 1}; + ASSERT(end >= start); + const std::size_t size{end - start}; + if (size == 0) { + buffer_bindings.AddVertexBinding(DefaultBuffer(), 0); + continue; + } const auto [buffer, offset] = buffer_cache.UploadMemory(start, size); buffer_bindings.AddVertexBinding(buffer, offset); } @@ -886,6 +895,9 @@ void RasterizerVulkan::SetupVertexArrays(FixedPipelineState::VertexInput& vertex void RasterizerVulkan::SetupIndexBuffer(BufferBindings& buffer_bindings, DrawParameters& params, bool is_indexed) { + if (params.num_vertices == 0) { + return; + } const auto& regs = system.GPU().Maxwell3D().regs; switch (regs.draw.topology) { case Maxwell::PrimitiveTopology::Quads: { @@ -961,7 +973,7 @@ void RasterizerVulkan::SetupGraphicsTextures(const ShaderEntries& entries, std:: MICROPROFILE_SCOPE(Vulkan_Textures); const auto& gpu = system.GPU().Maxwell3D(); for (const auto& entry : entries.samplers) { - for (std::size_t i = 0; i < entry.Size(); ++i) { + for (std::size_t i = 0; i < entry.size; ++i) { const auto texture = GetTextureInfo(gpu, entry, stage, i); SetupTexture(texture, entry); } @@ -1013,7 +1025,7 @@ void RasterizerVulkan::SetupComputeTextures(const ShaderEntries& entries) { MICROPROFILE_SCOPE(Vulkan_Textures); const auto& gpu = system.GPU().KeplerCompute(); for (const auto& entry : entries.samplers) { - for (std::size_t i = 0; i < entry.Size(); ++i) { + for (std::size_t i = 0; i < entry.size; ++i) { const auto texture = GetTextureInfo(gpu, entry, ComputeShaderIndex, i); SetupTexture(texture, entry); } @@ -1033,8 +1045,7 @@ void RasterizerVulkan::SetupConstBuffer(const ConstBufferEntry& entry, const Tegra::Engines::ConstBufferInfo& buffer) { if (!buffer.enabled) { // Set values to zero to unbind buffers - update_descriptor_queue.AddBuffer(buffer_cache.GetEmptyBuffer(sizeof(float)), 0, - sizeof(float)); + update_descriptor_queue.AddBuffer(DefaultBuffer(), 0, DEFAULT_BUFFER_SIZE); return; } @@ -1057,7 +1068,9 @@ void RasterizerVulkan::SetupGlobalBuffer(const GlobalBufferEntry& entry, GPUVAdd if (size == 0) { // Sometimes global memory pointers don't have a proper size. Upload a dummy entry // because Vulkan doesn't like empty buffers. - constexpr std::size_t dummy_size = 4; + // Note: Do *not* use DefaultBuffer() here, storage buffers can be written breaking the + // default buffer. + static constexpr std::size_t dummy_size = 4; const auto buffer = buffer_cache.GetEmptyBuffer(dummy_size); update_descriptor_queue.AddBuffer(buffer, 0, dummy_size); return; @@ -1094,7 +1107,7 @@ void RasterizerVulkan::SetupTexture(const Tegra::Texture::FullTextureInfo& textu void RasterizerVulkan::SetupImage(const Tegra::Texture::TICEntry& tic, const ImageEntry& entry) { auto view = texture_cache.GetImageSurface(tic, entry); - if (entry.IsWritten()) { + if (entry.is_written) { view->MarkAsModified(texture_cache.Tick()); } @@ -1222,7 +1235,7 @@ std::size_t RasterizerVulkan::CalculateVertexArraysSize() const { const GPUVAddr end{regs.vertex_array_limit[index].LimitAddress()}; DEBUG_ASSERT(end >= start); - size += (end - start + 1) * regs.vertex_array[index].enable; + size += (end - start) * regs.vertex_array[index].enable; } return size; } @@ -1245,28 +1258,54 @@ std::size_t RasterizerVulkan::CalculateConstBufferSize( } RenderPassParams RasterizerVulkan::GetRenderPassParams(Texceptions texceptions) const { - using namespace VideoCore::Surface; - const auto& regs = system.GPU().Maxwell3D().regs; - RenderPassParams renderpass_params; + const std::size_t num_attachments = static_cast<std::size_t>(regs.rt_control.count); - for (std::size_t rt = 0; rt < static_cast<std::size_t>(regs.rt_control.count); ++rt) { + RenderPassParams params; + params.color_formats = {}; + std::size_t color_texceptions = 0; + + std::size_t index = 0; + for (std::size_t rt = 0; rt < num_attachments; ++rt) { const auto& rendertarget = regs.rt[rt]; if (rendertarget.Address() == 0 || rendertarget.format == Tegra::RenderTargetFormat::NONE) { continue; } - renderpass_params.color_attachments.push_back(RenderPassParams::ColorAttachment{ - static_cast<u32>(rt), PixelFormatFromRenderTargetFormat(rendertarget.format), - texceptions[rt]}); + params.color_formats[index] = static_cast<u8>(rendertarget.format); + color_texceptions |= (texceptions[rt] ? 1ULL : 0ULL) << index; + ++index; } + params.num_color_attachments = static_cast<u8>(index); + params.texceptions = static_cast<u8>(color_texceptions); + + params.zeta_format = regs.zeta_enable ? static_cast<u8>(regs.zeta.format) : 0; + params.zeta_texception = texceptions[ZETA_TEXCEPTION_INDEX]; + return params; +} - renderpass_params.has_zeta = regs.zeta_enable; - if (renderpass_params.has_zeta) { - renderpass_params.zeta_pixel_format = PixelFormatFromDepthFormat(regs.zeta.format); - renderpass_params.zeta_texception = texceptions[ZETA_TEXCEPTION_INDEX]; +VkBuffer RasterizerVulkan::DefaultBuffer() { + if (default_buffer) { + return *default_buffer; } - return renderpass_params; + VkBufferCreateInfo ci; + ci.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO; + ci.pNext = nullptr; + ci.flags = 0; + ci.size = DEFAULT_BUFFER_SIZE; + ci.usage = VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT | + VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT; + ci.sharingMode = VK_SHARING_MODE_EXCLUSIVE; + ci.queueFamilyIndexCount = 0; + ci.pQueueFamilyIndices = nullptr; + default_buffer = device.GetLogical().CreateBuffer(ci); + default_buffer_commit = memory_manager.Commit(default_buffer, false); + + scheduler.RequestOutsideRenderPassOperationContext(); + scheduler.Record([buffer = *default_buffer](vk::CommandBuffer cmdbuf) { + cmdbuf.FillBuffer(buffer, 0, DEFAULT_BUFFER_SIZE, 0); + }); + return *default_buffer; } } // namespace Vulkan diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.h b/src/video_core/renderer_vulkan/vk_rasterizer.h index 2fa46b0cc..d41a7929e 100644 --- a/src/video_core/renderer_vulkan/vk_rasterizer.h +++ b/src/video_core/renderer_vulkan/vk_rasterizer.h @@ -155,6 +155,7 @@ private: using Texceptions = std::bitset<Maxwell::NumRenderTargets + 1>; static constexpr std::size_t ZETA_TEXCEPTION_INDEX = 8; + static constexpr VkDeviceSize DEFAULT_BUFFER_SIZE = 4 * sizeof(float); void FlushWork(); @@ -247,6 +248,8 @@ private: RenderPassParams GetRenderPassParams(Texceptions texceptions) const; + VkBuffer DefaultBuffer(); + Core::System& system; Core::Frontend::EmuWindow& render_window; VKScreenInfo& screen_info; @@ -271,6 +274,9 @@ private: VKFenceManager fence_manager; VKQueryCache query_cache; + vk::Buffer default_buffer; + VKMemoryCommit default_buffer_commit; + std::array<View, Maxwell::NumRenderTargets> color_attachments; View zeta_attachment; diff --git a/src/video_core/renderer_vulkan/vk_renderpass_cache.cpp b/src/video_core/renderer_vulkan/vk_renderpass_cache.cpp index 4e5286a69..3f71d005e 100644 --- a/src/video_core/renderer_vulkan/vk_renderpass_cache.cpp +++ b/src/video_core/renderer_vulkan/vk_renderpass_cache.cpp @@ -2,9 +2,11 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. +#include <cstring> #include <memory> #include <vector> +#include "common/cityhash.h" #include "video_core/engines/maxwell_3d.h" #include "video_core/renderer_vulkan/maxwell_to_vk.h" #include "video_core/renderer_vulkan/vk_device.h" @@ -13,6 +15,15 @@ namespace Vulkan { +std::size_t RenderPassParams::Hash() const noexcept { + const u64 hash = Common::CityHash64(reinterpret_cast<const char*>(this), sizeof *this); + return static_cast<std::size_t>(hash); +} + +bool RenderPassParams::operator==(const RenderPassParams& rhs) const noexcept { + return std::memcmp(&rhs, this, sizeof *this) == 0; +} + VKRenderPassCache::VKRenderPassCache(const VKDevice& device) : device{device} {} VKRenderPassCache::~VKRenderPassCache() = default; @@ -27,20 +38,22 @@ VkRenderPass VKRenderPassCache::GetRenderPass(const RenderPassParams& params) { } vk::RenderPass VKRenderPassCache::CreateRenderPass(const RenderPassParams& params) const { + using namespace VideoCore::Surface; std::vector<VkAttachmentDescription> descriptors; std::vector<VkAttachmentReference> color_references; - for (std::size_t rt = 0; rt < params.color_attachments.size(); ++rt) { - const auto attachment = params.color_attachments[rt]; - const auto format = - MaxwellToVK::SurfaceFormat(device, FormatType::Optimal, attachment.pixel_format); + const std::size_t num_attachments = static_cast<std::size_t>(params.num_color_attachments); + for (std::size_t rt = 0; rt < num_attachments; ++rt) { + const auto guest_format = static_cast<Tegra::RenderTargetFormat>(params.color_formats[rt]); + const PixelFormat pixel_format = PixelFormatFromRenderTargetFormat(guest_format); + const auto format = MaxwellToVK::SurfaceFormat(device, FormatType::Optimal, pixel_format); ASSERT_MSG(format.attachable, "Trying to attach a non-attachable format with format={}", - static_cast<u32>(attachment.pixel_format)); + static_cast<int>(pixel_format)); - // TODO(Rodrigo): Add eMayAlias when it's needed. - const auto color_layout = attachment.is_texception - ? VK_IMAGE_LAYOUT_GENERAL - : VK_IMAGE_LAYOUT_COLOR_ATTACHMENT_OPTIMAL; + // TODO(Rodrigo): Add MAY_ALIAS_BIT when it's needed. + const VkImageLayout color_layout = ((params.texceptions >> rt) & 1) != 0 + ? VK_IMAGE_LAYOUT_GENERAL + : VK_IMAGE_LAYOUT_COLOR_ATTACHMENT_OPTIMAL; VkAttachmentDescription& descriptor = descriptors.emplace_back(); descriptor.flags = VK_ATTACHMENT_DESCRIPTION_MAY_ALIAS_BIT; descriptor.format = format.format; @@ -58,15 +71,17 @@ vk::RenderPass VKRenderPassCache::CreateRenderPass(const RenderPassParams& param } VkAttachmentReference zeta_attachment_ref; - if (params.has_zeta) { - const auto format = - MaxwellToVK::SurfaceFormat(device, FormatType::Optimal, params.zeta_pixel_format); + const bool has_zeta = params.zeta_format != 0; + if (has_zeta) { + const auto guest_format = static_cast<Tegra::DepthFormat>(params.zeta_format); + const PixelFormat pixel_format = PixelFormatFromDepthFormat(guest_format); + const auto format = MaxwellToVK::SurfaceFormat(device, FormatType::Optimal, pixel_format); ASSERT_MSG(format.attachable, "Trying to attach a non-attachable format with format={}", - static_cast<u32>(params.zeta_pixel_format)); + static_cast<int>(pixel_format)); - const auto zeta_layout = params.zeta_texception - ? VK_IMAGE_LAYOUT_GENERAL - : VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL; + const VkImageLayout zeta_layout = params.zeta_texception != 0 + ? VK_IMAGE_LAYOUT_GENERAL + : VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL; VkAttachmentDescription& descriptor = descriptors.emplace_back(); descriptor.flags = 0; descriptor.format = format.format; @@ -78,7 +93,7 @@ vk::RenderPass VKRenderPassCache::CreateRenderPass(const RenderPassParams& param descriptor.initialLayout = zeta_layout; descriptor.finalLayout = zeta_layout; - zeta_attachment_ref.attachment = static_cast<u32>(params.color_attachments.size()); + zeta_attachment_ref.attachment = static_cast<u32>(num_attachments); zeta_attachment_ref.layout = zeta_layout; } @@ -90,7 +105,7 @@ vk::RenderPass VKRenderPassCache::CreateRenderPass(const RenderPassParams& param subpass_description.colorAttachmentCount = static_cast<u32>(color_references.size()); subpass_description.pColorAttachments = color_references.data(); subpass_description.pResolveAttachments = nullptr; - subpass_description.pDepthStencilAttachment = params.has_zeta ? &zeta_attachment_ref : nullptr; + subpass_description.pDepthStencilAttachment = has_zeta ? &zeta_attachment_ref : nullptr; subpass_description.preserveAttachmentCount = 0; subpass_description.pPreserveAttachments = nullptr; @@ -101,7 +116,7 @@ vk::RenderPass VKRenderPassCache::CreateRenderPass(const RenderPassParams& param stage |= VK_PIPELINE_STAGE_COLOR_ATTACHMENT_OUTPUT_BIT; } - if (params.has_zeta) { + if (has_zeta) { access |= VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_READ_BIT | VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_WRITE_BIT; stage |= VK_PIPELINE_STAGE_LATE_FRAGMENT_TESTS_BIT; diff --git a/src/video_core/renderer_vulkan/vk_renderpass_cache.h b/src/video_core/renderer_vulkan/vk_renderpass_cache.h index 921b6efb5..8b0fec720 100644 --- a/src/video_core/renderer_vulkan/vk_renderpass_cache.h +++ b/src/video_core/renderer_vulkan/vk_renderpass_cache.h @@ -4,8 +4,7 @@ #pragma once -#include <memory> -#include <tuple> +#include <type_traits> #include <unordered_map> #include <boost/container/static_vector.hpp> @@ -19,51 +18,25 @@ namespace Vulkan { class VKDevice; -// TODO(Rodrigo): Optimize this structure for faster hashing - struct RenderPassParams { - struct ColorAttachment { - u32 index = 0; - VideoCore::Surface::PixelFormat pixel_format = VideoCore::Surface::PixelFormat::Invalid; - bool is_texception = false; - - std::size_t Hash() const noexcept { - return static_cast<std::size_t>(pixel_format) | - static_cast<std::size_t>(is_texception) << 6 | - static_cast<std::size_t>(index) << 7; - } - - bool operator==(const ColorAttachment& rhs) const noexcept { - return std::tie(index, pixel_format, is_texception) == - std::tie(rhs.index, rhs.pixel_format, rhs.is_texception); - } - }; - - boost::container::static_vector<ColorAttachment, - Tegra::Engines::Maxwell3D::Regs::NumRenderTargets> - color_attachments{}; - // TODO(Rodrigo): Unify has_zeta into zeta_pixel_format and zeta_component_type. - VideoCore::Surface::PixelFormat zeta_pixel_format = VideoCore::Surface::PixelFormat::Invalid; - bool has_zeta = false; - bool zeta_texception = false; - - std::size_t Hash() const noexcept { - std::size_t hash = 0; - for (const auto& rt : color_attachments) { - boost::hash_combine(hash, rt.Hash()); - } - boost::hash_combine(hash, zeta_pixel_format); - boost::hash_combine(hash, has_zeta); - boost::hash_combine(hash, zeta_texception); - return hash; - } + std::array<u8, Tegra::Engines::Maxwell3D::Regs::NumRenderTargets> color_formats; + u8 num_color_attachments; + u8 texceptions; + + u8 zeta_format; + u8 zeta_texception; + + std::size_t Hash() const noexcept; + + bool operator==(const RenderPassParams& rhs) const noexcept; - bool operator==(const RenderPassParams& rhs) const { - return std::tie(color_attachments, zeta_pixel_format, has_zeta, zeta_texception) == - std::tie(rhs.color_attachments, rhs.zeta_pixel_format, rhs.has_zeta, - rhs.zeta_texception); + bool operator!=(const RenderPassParams& rhs) const noexcept { + return !operator==(rhs); } }; +static_assert(std::has_unique_object_representations_v<RenderPassParams>); +static_assert(std::is_trivially_copyable_v<RenderPassParams>); +static_assert(std::is_trivially_constructible_v<RenderPassParams>); } // namespace Vulkan diff --git a/src/video_core/renderer_vulkan/vk_shader_decompiler.cpp b/src/video_core/renderer_vulkan/vk_shader_decompiler.cpp index aaa138f52..18678968c 100644 --- a/src/video_core/renderer_vulkan/vk_shader_decompiler.cpp +++ b/src/video_core/renderer_vulkan/vk_shader_decompiler.cpp @@ -103,8 +103,8 @@ struct GenericVaryingDescription { }; spv::Dim GetSamplerDim(const Sampler& sampler) { - ASSERT(!sampler.IsBuffer()); - switch (sampler.GetType()) { + ASSERT(!sampler.is_buffer); + switch (sampler.type) { case Tegra::Shader::TextureType::Texture1D: return spv::Dim::Dim1D; case Tegra::Shader::TextureType::Texture2D: @@ -114,13 +114,13 @@ spv::Dim GetSamplerDim(const Sampler& sampler) { case Tegra::Shader::TextureType::TextureCube: return spv::Dim::Cube; default: - UNIMPLEMENTED_MSG("Unimplemented sampler type={}", static_cast<u32>(sampler.GetType())); + UNIMPLEMENTED_MSG("Unimplemented sampler type={}", static_cast<int>(sampler.type)); return spv::Dim::Dim2D; } } std::pair<spv::Dim, bool> GetImageDim(const Image& image) { - switch (image.GetType()) { + switch (image.type) { case Tegra::Shader::ImageType::Texture1D: return {spv::Dim::Dim1D, false}; case Tegra::Shader::ImageType::TextureBuffer: @@ -134,7 +134,7 @@ std::pair<spv::Dim, bool> GetImageDim(const Image& image) { case Tegra::Shader::ImageType::Texture3D: return {spv::Dim::Dim3D, false}; default: - UNIMPLEMENTED_MSG("Unimplemented image type={}", static_cast<u32>(image.GetType())); + UNIMPLEMENTED_MSG("Unimplemented image type={}", static_cast<int>(image.type)); return {spv::Dim::Dim2D, false}; } } @@ -879,11 +879,11 @@ private: u32 DeclareTexelBuffers(u32 binding) { for (const auto& sampler : ir.GetSamplers()) { - if (!sampler.IsBuffer()) { + if (!sampler.is_buffer) { continue; } - ASSERT(!sampler.IsArray()); - ASSERT(!sampler.IsShadow()); + ASSERT(!sampler.is_array); + ASSERT(!sampler.is_shadow); constexpr auto dim = spv::Dim::Buffer; constexpr int depth = 0; @@ -894,23 +894,23 @@ private: const Id image_type = TypeImage(t_float, dim, depth, arrayed, ms, sampled, format); const Id pointer_type = TypePointer(spv::StorageClass::UniformConstant, image_type); const Id id = OpVariable(pointer_type, spv::StorageClass::UniformConstant); - AddGlobalVariable(Name(id, fmt::format("sampler_{}", sampler.GetIndex()))); + AddGlobalVariable(Name(id, fmt::format("sampler_{}", sampler.index))); Decorate(id, spv::Decoration::Binding, binding++); Decorate(id, spv::Decoration::DescriptorSet, DESCRIPTOR_SET); - texel_buffers.emplace(sampler.GetIndex(), TexelBuffer{image_type, id}); + texel_buffers.emplace(sampler.index, TexelBuffer{image_type, id}); } return binding; } u32 DeclareSamplers(u32 binding) { for (const auto& sampler : ir.GetSamplers()) { - if (sampler.IsBuffer()) { + if (sampler.is_buffer) { continue; } const auto dim = GetSamplerDim(sampler); - const int depth = sampler.IsShadow() ? 1 : 0; - const int arrayed = sampler.IsArray() ? 1 : 0; + const int depth = sampler.is_shadow ? 1 : 0; + const int arrayed = sampler.is_array ? 1 : 0; constexpr bool ms = false; constexpr int sampled = 1; constexpr auto format = spv::ImageFormat::Unknown; @@ -918,17 +918,17 @@ private: const Id sampler_type = TypeSampledImage(image_type); const Id sampler_pointer_type = TypePointer(spv::StorageClass::UniformConstant, sampler_type); - const Id type = sampler.IsIndexed() - ? TypeArray(sampler_type, Constant(t_uint, sampler.Size())) + const Id type = sampler.is_indexed + ? TypeArray(sampler_type, Constant(t_uint, sampler.size)) : sampler_type; const Id pointer_type = TypePointer(spv::StorageClass::UniformConstant, type); const Id id = OpVariable(pointer_type, spv::StorageClass::UniformConstant); - AddGlobalVariable(Name(id, fmt::format("sampler_{}", sampler.GetIndex()))); + AddGlobalVariable(Name(id, fmt::format("sampler_{}", sampler.index))); Decorate(id, spv::Decoration::Binding, binding++); Decorate(id, spv::Decoration::DescriptorSet, DESCRIPTOR_SET); - sampled_images.emplace(sampler.GetIndex(), SampledImage{image_type, sampler_type, - sampler_pointer_type, id}); + sampled_images.emplace( + sampler.index, SampledImage{image_type, sampler_type, sampler_pointer_type, id}); } return binding; } @@ -943,17 +943,17 @@ private: const Id image_type = TypeImage(t_uint, dim, depth, arrayed, ms, sampled, format, {}); const Id pointer_type = TypePointer(spv::StorageClass::UniformConstant, image_type); const Id id = OpVariable(pointer_type, spv::StorageClass::UniformConstant); - AddGlobalVariable(Name(id, fmt::format("image_{}", image.GetIndex()))); + AddGlobalVariable(Name(id, fmt::format("image_{}", image.index))); Decorate(id, spv::Decoration::Binding, binding++); Decorate(id, spv::Decoration::DescriptorSet, DESCRIPTOR_SET); - if (image.IsRead() && !image.IsWritten()) { + if (image.is_read && !image.is_written) { Decorate(id, spv::Decoration::NonWritable); - } else if (image.IsWritten() && !image.IsRead()) { + } else if (image.is_written && !image.is_read) { Decorate(id, spv::Decoration::NonReadable); } - images.emplace(static_cast<u32>(image.GetIndex()), StorageImage{image_type, id}); + images.emplace(image.index, StorageImage{image_type, id}); } return binding; } @@ -1584,6 +1584,15 @@ private: return {OpCompositeConstruct(t_half, low, high), Type::HalfFloat}; } + Expression LogicalAddCarry(Operation operation) { + const Id op_a = AsUint(Visit(operation[0])); + const Id op_b = AsUint(Visit(operation[1])); + + const Id result = OpIAddCarry(TypeStruct({t_uint, t_uint}), op_a, op_b); + const Id carry = OpCompositeExtract(t_uint, result, 1); + return {OpINotEqual(t_bool, carry, Constant(t_uint, 0)), Type::Bool}; + } + Expression LogicalAssign(Operation operation) { const Node& dest = operation[0]; const Node& src = operation[1]; @@ -1611,11 +1620,11 @@ private: Id GetTextureSampler(Operation operation) { const auto& meta = std::get<MetaTexture>(operation.GetMeta()); - ASSERT(!meta.sampler.IsBuffer()); + ASSERT(!meta.sampler.is_buffer); - const auto& entry = sampled_images.at(meta.sampler.GetIndex()); + const auto& entry = sampled_images.at(meta.sampler.index); Id sampler = entry.variable; - if (meta.sampler.IsIndexed()) { + if (meta.sampler.is_indexed) { const Id index = AsInt(Visit(meta.index)); sampler = OpAccessChain(entry.sampler_pointer_type, sampler, index); } @@ -1624,8 +1633,8 @@ private: Id GetTextureImage(Operation operation) { const auto& meta = std::get<MetaTexture>(operation.GetMeta()); - const u32 index = meta.sampler.GetIndex(); - if (meta.sampler.IsBuffer()) { + const u32 index = meta.sampler.index; + if (meta.sampler.is_buffer) { const auto& entry = texel_buffers.at(index); return OpLoad(entry.image_type, entry.image); } else { @@ -1636,7 +1645,7 @@ private: Id GetImage(Operation operation) { const auto& meta = std::get<MetaImage>(operation.GetMeta()); - const auto entry = images.at(meta.image.GetIndex()); + const auto entry = images.at(meta.image.index); return OpLoad(entry.image_type, entry.image); } @@ -1652,7 +1661,7 @@ private: } if (const auto meta = std::get_if<MetaTexture>(&operation.GetMeta())) { // Add array coordinate for textures - if (meta->sampler.IsArray()) { + if (meta->sampler.is_array) { Id array = AsInt(Visit(meta->array)); if (type == Type::Float) { array = OpConvertSToF(t_float, array); @@ -1758,7 +1767,7 @@ private: operands.push_back(GetOffsetCoordinates(operation)); } - if (meta.sampler.IsShadow()) { + if (meta.sampler.is_shadow) { const Id dref = AsFloat(Visit(meta.depth_compare)); return {OpImageSampleDrefExplicitLod(t_float, sampler, coords, dref, mask, operands), Type::Float}; @@ -1773,7 +1782,7 @@ private: const Id coords = GetCoordinates(operation, Type::Float); Id texture{}; - if (meta.sampler.IsShadow()) { + if (meta.sampler.is_shadow) { texture = OpImageDrefGather(t_float4, GetTextureSampler(operation), coords, AsFloat(Visit(meta.depth_compare))); } else { @@ -1800,8 +1809,8 @@ private: } const Id lod = AsUint(Visit(operation[0])); - const std::size_t coords_count = [&]() { - switch (const auto type = meta.sampler.GetType(); type) { + const std::size_t coords_count = [&meta] { + switch (const auto type = meta.sampler.type) { case Tegra::Shader::TextureType::Texture1D: return 1; case Tegra::Shader::TextureType::Texture2D: @@ -1810,7 +1819,7 @@ private: case Tegra::Shader::TextureType::Texture3D: return 3; default: - UNREACHABLE_MSG("Invalid texture type={}", static_cast<u32>(type)); + UNREACHABLE_MSG("Invalid texture type={}", static_cast<int>(type)); return 2; } }(); @@ -1853,7 +1862,7 @@ private: const Id image = GetTextureImage(operation); const Id coords = GetCoordinates(operation, Type::Int); Id fetch; - if (meta.lod && !meta.sampler.IsBuffer()) { + if (meta.lod && !meta.sampler.is_buffer) { fetch = OpImageFetch(t_float4, image, coords, spv::ImageOperandsMask::Lod, AsInt(Visit(meta.lod))); } else { @@ -2518,6 +2527,8 @@ private: &SPIRVDecompiler::Binary<&Module::OpINotEqual, Type::Bool, Type::Uint>, &SPIRVDecompiler::Binary<&Module::OpUGreaterThanEqual, Type::Bool, Type::Uint>, + &SPIRVDecompiler::LogicalAddCarry, + &SPIRVDecompiler::Binary<&Module::OpFOrdLessThan, Type::Bool2, Type::HalfFloat>, &SPIRVDecompiler::Binary<&Module::OpFOrdEqual, Type::Bool2, Type::HalfFloat>, &SPIRVDecompiler::Binary<&Module::OpFOrdLessThanEqual, Type::Bool2, Type::HalfFloat>, @@ -2969,7 +2980,7 @@ ShaderEntries GenerateShaderEntries(const VideoCommon::Shader::ShaderIR& ir) { entries.global_buffers.emplace_back(base.cbuf_index, base.cbuf_offset, usage.is_written); } for (const auto& sampler : ir.GetSamplers()) { - if (sampler.IsBuffer()) { + if (sampler.is_buffer) { entries.texel_buffers.emplace_back(sampler); } else { entries.samplers.emplace_back(sampler); diff --git a/src/video_core/renderer_vulkan/vk_staging_buffer_pool.cpp b/src/video_core/renderer_vulkan/vk_staging_buffer_pool.cpp index 94d954d7a..45c180221 100644 --- a/src/video_core/renderer_vulkan/vk_staging_buffer_pool.cpp +++ b/src/video_core/renderer_vulkan/vk_staging_buffer_pool.cpp @@ -39,8 +39,7 @@ VKStagingBufferPool::StagingBuffer& VKStagingBufferPool::StagingBuffer::operator VKStagingBufferPool::VKStagingBufferPool(const VKDevice& device, VKMemoryManager& memory_manager, VKScheduler& scheduler) - : device{device}, memory_manager{memory_manager}, scheduler{scheduler}, - is_device_integrated{device.IsIntegrated()} {} + : device{device}, memory_manager{memory_manager}, scheduler{scheduler} {} VKStagingBufferPool::~VKStagingBufferPool() = default; @@ -56,9 +55,7 @@ void VKStagingBufferPool::TickFrame() { current_delete_level = (current_delete_level + 1) % NumLevels; ReleaseCache(true); - if (!is_device_integrated) { - ReleaseCache(false); - } + ReleaseCache(false); } VKBuffer* VKStagingBufferPool::TryGetReservedBuffer(std::size_t size, bool host_visible) { @@ -81,7 +78,7 @@ VKBuffer& VKStagingBufferPool::CreateStagingBuffer(std::size_t size, bool host_v ci.size = 1ULL << log2; ci.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | - VK_BUFFER_USAGE_INDEX_BUFFER_BIT; + VK_BUFFER_USAGE_INDEX_BUFFER_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT; ci.sharingMode = VK_SHARING_MODE_EXCLUSIVE; ci.queueFamilyIndexCount = 0; ci.pQueueFamilyIndices = nullptr; @@ -95,7 +92,7 @@ VKBuffer& VKStagingBufferPool::CreateStagingBuffer(std::size_t size, bool host_v } VKStagingBufferPool::StagingBuffersCache& VKStagingBufferPool::GetCache(bool host_visible) { - return is_device_integrated || host_visible ? host_staging_buffers : device_staging_buffers; + return host_visible ? host_staging_buffers : device_staging_buffers; } void VKStagingBufferPool::ReleaseCache(bool host_visible) { diff --git a/src/video_core/renderer_vulkan/vk_staging_buffer_pool.h b/src/video_core/renderer_vulkan/vk_staging_buffer_pool.h index a0840ff8c..faf6418fd 100644 --- a/src/video_core/renderer_vulkan/vk_staging_buffer_pool.h +++ b/src/video_core/renderer_vulkan/vk_staging_buffer_pool.h @@ -71,7 +71,6 @@ private: const VKDevice& device; VKMemoryManager& memory_manager; VKScheduler& scheduler; - const bool is_device_integrated; StagingBuffersCache host_staging_buffers; StagingBuffersCache device_staging_buffers; diff --git a/src/video_core/renderer_vulkan/wrapper.cpp b/src/video_core/renderer_vulkan/wrapper.cpp index 539f3c974..7f5bc1404 100644 --- a/src/video_core/renderer_vulkan/wrapper.cpp +++ b/src/video_core/renderer_vulkan/wrapper.cpp @@ -2,6 +2,7 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. +#include <algorithm> #include <exception> #include <memory> #include <optional> @@ -16,6 +17,23 @@ namespace Vulkan::vk { namespace { +void SortPhysicalDevices(std::vector<VkPhysicalDevice>& devices, const InstanceDispatch& dld) { + std::stable_sort(devices.begin(), devices.end(), [&](auto lhs, auto rhs) { + // This will call Vulkan more than needed, but these calls are cheap. + const auto lhs_properties = vk::PhysicalDevice(lhs, dld).GetProperties(); + const auto rhs_properties = vk::PhysicalDevice(rhs, dld).GetProperties(); + + // Prefer discrete GPUs, Nvidia over AMD, AMD over Intel, Intel over the rest. + const bool preferred = + (lhs_properties.deviceType == VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU && + rhs_properties.deviceType != VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU) || + (lhs_properties.vendorID == 0x10DE && rhs_properties.vendorID != 0x10DE) || + (lhs_properties.vendorID == 0x1002 && rhs_properties.vendorID != 0x1002) || + (lhs_properties.vendorID == 0x8086 && rhs_properties.vendorID != 0x8086); + return !preferred; + }); +} + template <typename T> bool Proc(T& result, const InstanceDispatch& dld, const char* proc_name, VkInstance instance = nullptr) noexcept { @@ -389,7 +407,8 @@ std::optional<std::vector<VkPhysicalDevice>> Instance::EnumeratePhysicalDevices( if (dld->vkEnumeratePhysicalDevices(handle, &num, physical_devices.data()) != VK_SUCCESS) { return std::nullopt; } - return physical_devices; + SortPhysicalDevices(physical_devices, *dld); + return std::make_optional(std::move(physical_devices)); } DebugCallback Instance::TryCreateDebugCallback( diff --git a/src/video_core/shader/control_flow.cpp b/src/video_core/shader/control_flow.cpp index e00a3fb70..8d86020f6 100644 --- a/src/video_core/shader/control_flow.cpp +++ b/src/video_core/shader/control_flow.cpp @@ -13,6 +13,7 @@ #include "common/common_types.h" #include "video_core/shader/ast.h" #include "video_core/shader/control_flow.h" +#include "video_core/shader/memory_util.h" #include "video_core/shader/registry.h" #include "video_core/shader/shader_ir.h" @@ -115,17 +116,6 @@ Pred GetPredicate(u32 index, bool negated) { return static_cast<Pred>(static_cast<u64>(index) + (negated ? 8ULL : 0ULL)); } -/** - * Returns whether the instruction at the specified offset is a 'sched' instruction. - * Sched instructions always appear before a sequence of 3 instructions. - */ -constexpr bool IsSchedInstruction(u32 offset, u32 main_offset) { - constexpr u32 SchedPeriod = 4; - u32 absolute_offset = offset - main_offset; - - return (absolute_offset % SchedPeriod) == 0; -} - enum class ParseResult : u32 { ControlCaught, BlockEnd, diff --git a/src/video_core/shader/decode.cpp b/src/video_core/shader/decode.cpp index 87ac9ac6c..a75a5cc63 100644 --- a/src/video_core/shader/decode.cpp +++ b/src/video_core/shader/decode.cpp @@ -13,6 +13,7 @@ #include "video_core/engines/shader_bytecode.h" #include "video_core/engines/shader_header.h" #include "video_core/shader/control_flow.h" +#include "video_core/shader/memory_util.h" #include "video_core/shader/node_helper.h" #include "video_core/shader/shader_ir.h" @@ -23,17 +24,6 @@ using Tegra::Shader::OpCode; namespace { -/** - * Returns whether the instruction at the specified offset is a 'sched' instruction. - * Sched instructions always appear before a sequence of 3 instructions. - */ -constexpr bool IsSchedInstruction(u32 offset, u32 main_offset) { - constexpr u32 SchedPeriod = 4; - u32 absolute_offset = offset - main_offset; - - return (absolute_offset % SchedPeriod) == 0; -} - void DeduceTextureHandlerSize(VideoCore::GuestDriverProfile& gpu_driver, const std::list<Sampler>& used_samplers) { if (gpu_driver.IsTextureHandlerSizeKnown() || used_samplers.size() <= 1) { @@ -42,11 +32,11 @@ void DeduceTextureHandlerSize(VideoCore::GuestDriverProfile& gpu_driver, u32 count{}; std::vector<u32> bound_offsets; for (const auto& sampler : used_samplers) { - if (sampler.IsBindless()) { + if (sampler.is_bindless) { continue; } ++count; - bound_offsets.emplace_back(sampler.GetOffset()); + bound_offsets.emplace_back(sampler.offset); } if (count > 1) { gpu_driver.DeduceTextureHandlerSize(std::move(bound_offsets)); @@ -56,14 +46,14 @@ void DeduceTextureHandlerSize(VideoCore::GuestDriverProfile& gpu_driver, std::optional<u32> TryDeduceSamplerSize(const Sampler& sampler_to_deduce, VideoCore::GuestDriverProfile& gpu_driver, const std::list<Sampler>& used_samplers) { - const u32 base_offset = sampler_to_deduce.GetOffset(); + const u32 base_offset = sampler_to_deduce.offset; u32 max_offset{std::numeric_limits<u32>::max()}; for (const auto& sampler : used_samplers) { - if (sampler.IsBindless()) { + if (sampler.is_bindless) { continue; } - if (sampler.GetOffset() > base_offset) { - max_offset = std::min(sampler.GetOffset(), max_offset); + if (sampler.offset > base_offset) { + max_offset = std::min(sampler.offset, max_offset); } } if (max_offset == std::numeric_limits<u32>::max()) { @@ -363,14 +353,14 @@ void ShaderIR::PostDecode() { return; } for (auto& sampler : used_samplers) { - if (!sampler.IsIndexed()) { + if (!sampler.is_indexed) { continue; } if (const auto size = TryDeduceSamplerSize(sampler, gpu_driver, used_samplers)) { - sampler.SetSize(*size); + sampler.size = *size; } else { LOG_CRITICAL(HW_GPU, "Failed to deduce size of indexed sampler"); - sampler.SetSize(1); + sampler.size = 1; } } } diff --git a/src/video_core/shader/decode/arithmetic_half.cpp b/src/video_core/shader/decode/arithmetic_half.cpp index ee7d9a29d..a276aee44 100644 --- a/src/video_core/shader/decode/arithmetic_half.cpp +++ b/src/video_core/shader/decode/arithmetic_half.cpp @@ -19,22 +19,46 @@ u32 ShaderIR::DecodeArithmeticHalf(NodeBlock& bb, u32 pc) { const Instruction instr = {program_code[pc]}; const auto opcode = OpCode::Decode(instr); - if (opcode->get().GetId() == OpCode::Id::HADD2_C || - opcode->get().GetId() == OpCode::Id::HADD2_R) { + bool negate_a = false; + bool negate_b = false; + bool absolute_a = false; + bool absolute_b = false; + + switch (opcode->get().GetId()) { + case OpCode::Id::HADD2_R: if (instr.alu_half.ftz == 0) { LOG_DEBUG(HW_GPU, "{} without FTZ is not implemented", opcode->get().GetName()); } + negate_a = ((instr.value >> 43) & 1) != 0; + negate_b = ((instr.value >> 31) & 1) != 0; + absolute_a = ((instr.value >> 44) & 1) != 0; + absolute_b = ((instr.value >> 30) & 1) != 0; + break; + case OpCode::Id::HADD2_C: + if (instr.alu_half.ftz == 0) { + LOG_DEBUG(HW_GPU, "{} without FTZ is not implemented", opcode->get().GetName()); + } + negate_a = ((instr.value >> 43) & 1) != 0; + negate_b = ((instr.value >> 56) & 1) != 0; + absolute_a = ((instr.value >> 44) & 1) != 0; + absolute_b = ((instr.value >> 54) & 1) != 0; + break; + case OpCode::Id::HMUL2_R: + negate_a = ((instr.value >> 43) & 1) != 0; + absolute_a = ((instr.value >> 44) & 1) != 0; + absolute_b = ((instr.value >> 30) & 1) != 0; + break; + case OpCode::Id::HMUL2_C: + negate_b = ((instr.value >> 31) & 1) != 0; + absolute_a = ((instr.value >> 44) & 1) != 0; + absolute_b = ((instr.value >> 54) & 1) != 0; + break; } - const bool negate_a = - opcode->get().GetId() != OpCode::Id::HMUL2_R && instr.alu_half.negate_a != 0; - const bool negate_b = - opcode->get().GetId() != OpCode::Id::HMUL2_C && instr.alu_half.negate_b != 0; - Node op_a = UnpackHalfFloat(GetRegister(instr.gpr8), instr.alu_half.type_a); - op_a = GetOperandAbsNegHalf(op_a, instr.alu_half.abs_a, negate_a); + op_a = GetOperandAbsNegHalf(op_a, absolute_a, negate_a); - auto [type_b, op_b] = [&]() -> std::tuple<HalfType, Node> { + auto [type_b, op_b] = [this, instr, opcode]() -> std::pair<HalfType, Node> { switch (opcode->get().GetId()) { case OpCode::Id::HADD2_C: case OpCode::Id::HMUL2_C: @@ -48,17 +72,16 @@ u32 ShaderIR::DecodeArithmeticHalf(NodeBlock& bb, u32 pc) { } }(); op_b = UnpackHalfFloat(op_b, type_b); - // redeclaration to avoid a bug in clang with reusing local bindings in lambdas - Node op_b_alt = GetOperandAbsNegHalf(op_b, instr.alu_half.abs_b, negate_b); + op_b = GetOperandAbsNegHalf(op_b, absolute_b, negate_b); - Node value = [&]() { + Node value = [this, opcode, op_a, op_b = op_b] { switch (opcode->get().GetId()) { case OpCode::Id::HADD2_C: case OpCode::Id::HADD2_R: - return Operation(OperationCode::HAdd, PRECISE, op_a, op_b_alt); + return Operation(OperationCode::HAdd, PRECISE, op_a, op_b); case OpCode::Id::HMUL2_C: case OpCode::Id::HMUL2_R: - return Operation(OperationCode::HMul, PRECISE, op_a, op_b_alt); + return Operation(OperationCode::HMul, PRECISE, op_a, op_b); default: UNIMPLEMENTED_MSG("Unhandled half float instruction: {}", opcode->get().GetName()); return Immediate(0); diff --git a/src/video_core/shader/decode/arithmetic_integer.cpp b/src/video_core/shader/decode/arithmetic_integer.cpp index 0f4c3103a..a041519b7 100644 --- a/src/video_core/shader/decode/arithmetic_integer.cpp +++ b/src/video_core/shader/decode/arithmetic_integer.cpp @@ -35,15 +35,38 @@ u32 ShaderIR::DecodeArithmeticInteger(NodeBlock& bb, u32 pc) { case OpCode::Id::IADD_C: case OpCode::Id::IADD_R: case OpCode::Id::IADD_IMM: { - UNIMPLEMENTED_IF_MSG(instr.alu.saturate_d, "IADD saturation not implemented"); + UNIMPLEMENTED_IF_MSG(instr.alu.saturate_d, "IADD.SAT"); + UNIMPLEMENTED_IF_MSG(instr.iadd.x && instr.generates_cc, "IADD.X Rd.CC"); op_a = GetOperandAbsNegInteger(op_a, false, instr.alu_integer.negate_a, true); op_b = GetOperandAbsNegInteger(op_b, false, instr.alu_integer.negate_b, true); - const Node value = Operation(OperationCode::IAdd, PRECISE, op_a, op_b); + Node value = Operation(OperationCode::UAdd, op_a, op_b); - SetInternalFlagsFromInteger(bb, value, instr.generates_cc); - SetRegister(bb, instr.gpr0, value); + if (instr.iadd.x) { + Node carry = GetInternalFlag(InternalFlag::Carry); + Node x = Operation(OperationCode::Select, std::move(carry), Immediate(1), Immediate(0)); + value = Operation(OperationCode::UAdd, std::move(value), std::move(x)); + } + + if (instr.generates_cc) { + const Node i0 = Immediate(0); + + Node zero = Operation(OperationCode::LogicalIEqual, value, i0); + Node sign = Operation(OperationCode::LogicalILessThan, value, i0); + Node carry = Operation(OperationCode::LogicalAddCarry, op_a, op_b); + + Node pos_a = Operation(OperationCode::LogicalIGreaterThan, op_a, i0); + Node pos_b = Operation(OperationCode::LogicalIGreaterThan, op_b, i0); + Node pos = Operation(OperationCode::LogicalAnd, std::move(pos_a), std::move(pos_b)); + Node overflow = Operation(OperationCode::LogicalAnd, pos, sign); + + SetInternalFlag(bb, InternalFlag::Zero, std::move(zero)); + SetInternalFlag(bb, InternalFlag::Sign, std::move(sign)); + SetInternalFlag(bb, InternalFlag::Carry, std::move(carry)); + SetInternalFlag(bb, InternalFlag::Overflow, std::move(overflow)); + } + SetRegister(bb, instr.gpr0, std::move(value)); break; } case OpCode::Id::IADD3_C: @@ -249,8 +272,8 @@ u32 ShaderIR::DecodeArithmeticInteger(NodeBlock& bb, u32 pc) { } case OpCode::Id::LEA_IMM: { const bool neg = instr.lea.imm.neg != 0; - return {Immediate(static_cast<u32>(instr.lea.imm.entry_a)), - GetOperandAbsNegInteger(GetRegister(instr.gpr8), false, neg, true), + return {GetOperandAbsNegInteger(GetRegister(instr.gpr8), false, neg, true), + Immediate(static_cast<u32>(instr.lea.imm.entry_a)), Immediate(static_cast<u32>(instr.lea.imm.entry_b))}; } case OpCode::Id::LEA_RZ: { diff --git a/src/video_core/shader/decode/image.cpp b/src/video_core/shader/decode/image.cpp index 85ee9aa5e..60b6ad72a 100644 --- a/src/video_core/shader/decode/image.cpp +++ b/src/video_core/shader/decode/image.cpp @@ -485,11 +485,10 @@ u32 ShaderIR::DecodeImage(NodeBlock& bb, u32 pc) { Image& ShaderIR::GetImage(Tegra::Shader::Image image, Tegra::Shader::ImageType type) { const auto offset = static_cast<u32>(image.index.Value()); - const auto it = - std::find_if(std::begin(used_images), std::end(used_images), - [offset](const Image& entry) { return entry.GetOffset() == offset; }); + const auto it = std::find_if(std::begin(used_images), std::end(used_images), + [offset](const Image& entry) { return entry.offset == offset; }); if (it != std::end(used_images)) { - ASSERT(!it->IsBindless() && it->GetType() == it->GetType()); + ASSERT(!it->is_bindless && it->type == type); return *it; } @@ -505,13 +504,12 @@ Image& ShaderIR::GetBindlessImage(Tegra::Shader::Register reg, Tegra::Shader::Im const auto buffer = std::get<1>(result); const auto offset = std::get<2>(result); - const auto it = - std::find_if(std::begin(used_images), std::end(used_images), - [buffer = buffer, offset = offset](const Image& entry) { - return entry.GetBuffer() == buffer && entry.GetOffset() == offset; - }); + const auto it = std::find_if(std::begin(used_images), std::end(used_images), + [buffer, offset](const Image& entry) { + return entry.buffer == buffer && entry.offset == offset; + }); if (it != std::end(used_images)) { - ASSERT(it->IsBindless() && it->GetType() == it->GetType()); + ASSERT(it->is_bindless && it->type == type); return *it; } diff --git a/src/video_core/shader/decode/register_set_predicate.cpp b/src/video_core/shader/decode/register_set_predicate.cpp index 8d54cce34..6116c31aa 100644 --- a/src/video_core/shader/decode/register_set_predicate.cpp +++ b/src/video_core/shader/decode/register_set_predicate.cpp @@ -2,6 +2,8 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. +#include <utility> + #include "common/assert.h" #include "common/common_types.h" #include "video_core/engines/shader_bytecode.h" @@ -10,20 +12,20 @@ namespace VideoCommon::Shader { +using std::move; using Tegra::Shader::Instruction; using Tegra::Shader::OpCode; namespace { -constexpr u64 NUM_PROGRAMMABLE_PREDICATES = 7; -} +constexpr u64 NUM_CONDITION_CODES = 4; +constexpr u64 NUM_PREDICATES = 7; +} // namespace u32 ShaderIR::DecodeRegisterSetPredicate(NodeBlock& bb, u32 pc) { const Instruction instr = {program_code[pc]}; const auto opcode = OpCode::Decode(instr); - UNIMPLEMENTED_IF(instr.p2r_r2p.mode != Tegra::Shader::R2pMode::Pr); - - const Node apply_mask = [&] { + Node apply_mask = [this, opcode, instr] { switch (opcode->get().GetId()) { case OpCode::Id::R2P_IMM: case OpCode::Id::P2R_IMM: @@ -34,39 +36,43 @@ u32 ShaderIR::DecodeRegisterSetPredicate(NodeBlock& bb, u32 pc) { } }(); - const auto offset = static_cast<u32>(instr.p2r_r2p.byte) * 8; + const u32 offset = static_cast<u32>(instr.p2r_r2p.byte) * 8; + + const bool cc = instr.p2r_r2p.mode == Tegra::Shader::R2pMode::Cc; + const u64 num_entries = cc ? NUM_CONDITION_CODES : NUM_PREDICATES; + const auto get_entry = [this, cc](u64 entry) { + return cc ? GetInternalFlag(static_cast<InternalFlag>(entry)) : GetPredicate(entry); + }; switch (opcode->get().GetId()) { case OpCode::Id::R2P_IMM: { - const Node mask = GetRegister(instr.gpr8); + Node mask = GetRegister(instr.gpr8); - for (u64 pred = 0; pred < NUM_PROGRAMMABLE_PREDICATES; ++pred) { - const auto shift = static_cast<u32>(pred); + for (u64 entry = 0; entry < num_entries; ++entry) { + const u32 shift = static_cast<u32>(entry); - const Node apply_compare = BitfieldExtract(apply_mask, shift, 1); - const Node condition = - Operation(OperationCode::LogicalUNotEqual, apply_compare, Immediate(0)); + Node apply = BitfieldExtract(apply_mask, shift, 1); + Node condition = Operation(OperationCode::LogicalUNotEqual, apply, Immediate(0)); - const Node value_compare = BitfieldExtract(mask, offset + shift, 1); - const Node value = - Operation(OperationCode::LogicalUNotEqual, value_compare, Immediate(0)); + Node compare = BitfieldExtract(mask, offset + shift, 1); + Node value = Operation(OperationCode::LogicalUNotEqual, move(compare), Immediate(0)); - const Node code = Operation(OperationCode::LogicalAssign, GetPredicate(pred), value); - bb.push_back(Conditional(condition, {code})); + Node code = Operation(OperationCode::LogicalAssign, get_entry(entry), move(value)); + bb.push_back(Conditional(condition, {move(code)})); } break; } case OpCode::Id::P2R_IMM: { Node value = Immediate(0); - for (u64 pred = 0; pred < NUM_PROGRAMMABLE_PREDICATES; ++pred) { - Node bit = Operation(OperationCode::Select, GetPredicate(pred), Immediate(1U << pred), + for (u64 entry = 0; entry < num_entries; ++entry) { + Node bit = Operation(OperationCode::Select, get_entry(entry), Immediate(1U << entry), Immediate(0)); - value = Operation(OperationCode::UBitwiseOr, std::move(value), std::move(bit)); + value = Operation(OperationCode::UBitwiseOr, move(value), move(bit)); } - value = Operation(OperationCode::UBitwiseAnd, std::move(value), apply_mask); - value = BitfieldInsert(GetRegister(instr.gpr8), std::move(value), offset, 8); + value = Operation(OperationCode::UBitwiseAnd, move(value), apply_mask); + value = BitfieldInsert(GetRegister(instr.gpr8), move(value), offset, 8); - SetRegister(bb, instr.gpr0, std::move(value)); + SetRegister(bb, instr.gpr0, move(value)); break; } default: diff --git a/src/video_core/shader/decode/texture.cpp b/src/video_core/shader/decode/texture.cpp index e68f1d305..8f0bb996e 100644 --- a/src/video_core/shader/decode/texture.cpp +++ b/src/video_core/shader/decode/texture.cpp @@ -139,15 +139,15 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) { } const Node component = Immediate(static_cast<u32>(instr.tld4s.component)); - const SamplerInfo info{TextureType::Texture2D, false, is_depth_compare, false}; - const Sampler& sampler = *GetSampler(instr.sampler, info); + SamplerInfo info; + info.is_shadow = is_depth_compare; + const std::optional<Sampler> sampler = GetSampler(instr.sampler, info); Node4 values; for (u32 element = 0; element < values.size(); ++element) { - auto coords_copy = coords; - MetaTexture meta{sampler, {}, depth_compare, aoffi, {}, {}, - {}, {}, component, element, {}}; - values[element] = Operation(OperationCode::TextureGather, meta, std::move(coords_copy)); + MetaTexture meta{*sampler, {}, depth_compare, aoffi, {}, {}, + {}, {}, component, element, {}}; + values[element] = Operation(OperationCode::TextureGather, meta, coords); } if (instr.tld4s.fp16_flag) { @@ -165,18 +165,20 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) { "AOFFI is not implemented"); const bool is_array = instr.txd.is_array != 0; - u64 base_reg = instr.gpr8.Value(); const auto derivate_reg = instr.gpr20.Value(); const auto texture_type = instr.txd.texture_type.Value(); const auto coord_count = GetCoordCount(texture_type); - Node index_var{}; - const Sampler* sampler = - is_bindless - ? GetBindlessSampler(base_reg, index_var, {{texture_type, is_array, false, false}}) - : GetSampler(instr.sampler, {{texture_type, is_array, false, false}}); + u64 base_reg = instr.gpr8.Value(); + Node index_var; + SamplerInfo info; + info.type = texture_type; + info.is_array = is_array; + const std::optional<Sampler> sampler = is_bindless + ? GetBindlessSampler(base_reg, info, index_var) + : GetSampler(instr.sampler, info); Node4 values; - if (sampler == nullptr) { - std::generate(values.begin(), values.end(), [] { return Immediate(0); }); + if (!sampler) { + std::generate(values.begin(), values.end(), [this] { return Immediate(0); }); WriteTexInstructionFloat(bb, instr, values); break; } @@ -214,14 +216,12 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) { is_bindless = true; [[fallthrough]]; case OpCode::Id::TXQ: { - // TODO: The new commits on the texture refactor, change the way samplers work. - // Sadly, not all texture instructions specify the type of texture their sampler - // uses. This must be fixed at a later instance. - Node index_var{}; - const Sampler* sampler = - is_bindless ? GetBindlessSampler(instr.gpr8, index_var) : GetSampler(instr.sampler); - - if (sampler == nullptr) { + Node index_var; + const std::optional<Sampler> sampler = is_bindless + ? GetBindlessSampler(instr.gpr8, {}, index_var) + : GetSampler(instr.sampler, {}); + + if (!sampler) { u32 indexer = 0; for (u32 element = 0; element < 4; ++element) { if (!instr.txq.IsComponentEnabled(element)) { @@ -267,12 +267,17 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) { UNIMPLEMENTED_IF_MSG(instr.tmml.UsesMiscMode(Tegra::Shader::TextureMiscMode::NDV), "NDV is not implemented"); - auto texture_type = instr.tmml.texture_type.Value(); - Node index_var{}; - const Sampler* sampler = - is_bindless ? GetBindlessSampler(instr.gpr20, index_var) : GetSampler(instr.sampler); - - if (sampler == nullptr) { + const auto texture_type = instr.tmml.texture_type.Value(); + const bool is_array = instr.tmml.array != 0; + SamplerInfo info; + info.type = texture_type; + info.is_array = is_array; + Node index_var; + const std::optional<Sampler> sampler = + is_bindless ? GetBindlessSampler(instr.gpr20, info, index_var) + : GetSampler(instr.sampler, info); + + if (!sampler) { u32 indexer = 0; for (u32 element = 0; element < 2; ++element) { if (!instr.tmml.IsComponentEnabled(element)) { @@ -299,12 +304,11 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) { coords.push_back(GetRegister(instr.gpr8.Value() + 1)); break; default: - UNIMPLEMENTED_MSG("Unhandled texture type {}", static_cast<u32>(texture_type)); + UNIMPLEMENTED_MSG("Unhandled texture type {}", static_cast<int>(texture_type)); // Fallback to interpreting as a 2D texture for now coords.push_back(GetRegister(instr.gpr8.Value() + 0)); coords.push_back(GetRegister(instr.gpr8.Value() + 1)); - texture_type = TextureType::Texture2D; } u32 indexer = 0; for (u32 element = 0; element < 2; ++element) { @@ -353,98 +357,103 @@ u32 ShaderIR::DecodeTexture(NodeBlock& bb, u32 pc) { return pc; } -ShaderIR::SamplerInfo ShaderIR::GetSamplerInfo(std::optional<SamplerInfo> sampler_info, u32 offset, +ShaderIR::SamplerInfo ShaderIR::GetSamplerInfo(SamplerInfo info, u32 offset, std::optional<u32> buffer) { - if (sampler_info) { - return *sampler_info; + if (info.IsComplete()) { + return info; } const auto sampler = buffer ? registry.ObtainBindlessSampler(*buffer, offset) : registry.ObtainBoundSampler(offset); if (!sampler) { LOG_WARNING(HW_GPU, "Unknown sampler info"); - return SamplerInfo{TextureType::Texture2D, false, false, false}; - } - return SamplerInfo{sampler->texture_type, sampler->is_array != 0, sampler->is_shadow != 0, - sampler->is_buffer != 0}; + info.type = info.type.value_or(Tegra::Shader::TextureType::Texture2D); + info.is_array = info.is_array.value_or(false); + info.is_shadow = info.is_shadow.value_or(false); + info.is_buffer = info.is_buffer.value_or(false); + return info; + } + info.type = info.type.value_or(sampler->texture_type); + info.is_array = info.is_array.value_or(sampler->is_array != 0); + info.is_shadow = info.is_shadow.value_or(sampler->is_shadow != 0); + info.is_buffer = info.is_buffer.value_or(sampler->is_buffer != 0); + return info; } -const Sampler* ShaderIR::GetSampler(const Tegra::Shader::Sampler& sampler, - std::optional<SamplerInfo> sampler_info) { +std::optional<Sampler> ShaderIR::GetSampler(Tegra::Shader::Sampler sampler, + SamplerInfo sampler_info) { const auto offset = static_cast<u32>(sampler.index.Value()); const auto info = GetSamplerInfo(sampler_info, offset); // If this sampler has already been used, return the existing mapping. - const auto it = - std::find_if(used_samplers.begin(), used_samplers.end(), - [offset](const Sampler& entry) { return entry.GetOffset() == offset; }); + const auto it = std::find_if(used_samplers.begin(), used_samplers.end(), + [offset](const Sampler& entry) { return entry.offset == offset; }); if (it != used_samplers.end()) { - ASSERT(!it->IsBindless() && it->GetType() == info.type && it->IsArray() == info.is_array && - it->IsShadow() == info.is_shadow && it->IsBuffer() == info.is_buffer); - return &*it; + ASSERT(!it->is_bindless && it->type == info.type && it->is_array == info.is_array && + it->is_shadow == info.is_shadow && it->is_buffer == info.is_buffer); + return *it; } // Otherwise create a new mapping for this sampler const auto next_index = static_cast<u32>(used_samplers.size()); - return &used_samplers.emplace_back(next_index, offset, info.type, info.is_array, info.is_shadow, - info.is_buffer, false); + return used_samplers.emplace_back(next_index, offset, *info.type, *info.is_array, + *info.is_shadow, *info.is_buffer, false); } -const Sampler* ShaderIR::GetBindlessSampler(Tegra::Shader::Register reg, Node& index_var, - std::optional<SamplerInfo> sampler_info) { +std::optional<Sampler> ShaderIR::GetBindlessSampler(Tegra::Shader::Register reg, SamplerInfo info, + Node& index_var) { const Node sampler_register = GetRegister(reg); const auto [base_node, tracked_sampler_info] = TrackBindlessSampler(sampler_register, global_code, static_cast<s64>(global_code.size())); ASSERT(base_node != nullptr); if (base_node == nullptr) { - return nullptr; + return std::nullopt; } if (const auto bindless_sampler_info = std::get_if<BindlessSamplerNode>(&*tracked_sampler_info)) { const u32 buffer = bindless_sampler_info->GetIndex(); const u32 offset = bindless_sampler_info->GetOffset(); - const auto info = GetSamplerInfo(sampler_info, offset, buffer); + info = GetSamplerInfo(info, offset, buffer); // If this sampler has already been used, return the existing mapping. - const auto it = - std::find_if(used_samplers.begin(), used_samplers.end(), - [buffer = buffer, offset = offset](const Sampler& entry) { - return entry.GetBuffer() == buffer && entry.GetOffset() == offset; - }); + const auto it = std::find_if(used_samplers.begin(), used_samplers.end(), + [buffer = buffer, offset = offset](const Sampler& entry) { + return entry.buffer == buffer && entry.offset == offset; + }); if (it != used_samplers.end()) { - ASSERT(it->IsBindless() && it->GetType() == info.type && - it->IsArray() == info.is_array && it->IsShadow() == info.is_shadow); - return &*it; + ASSERT(it->is_bindless && it->type == info.type && it->is_array == info.is_array && + it->is_shadow == info.is_shadow); + return *it; } // Otherwise create a new mapping for this sampler const auto next_index = static_cast<u32>(used_samplers.size()); - return &used_samplers.emplace_back(next_index, offset, buffer, info.type, info.is_array, - info.is_shadow, info.is_buffer, false); - } else if (const auto array_sampler_info = - std::get_if<ArraySamplerNode>(&*tracked_sampler_info)) { + return used_samplers.emplace_back(next_index, offset, buffer, *info.type, *info.is_array, + *info.is_shadow, *info.is_buffer, false); + } + if (const auto array_sampler_info = std::get_if<ArraySamplerNode>(&*tracked_sampler_info)) { const u32 base_offset = array_sampler_info->GetBaseOffset() / 4; index_var = GetCustomVariable(array_sampler_info->GetIndexVar()); - const auto info = GetSamplerInfo(sampler_info, base_offset); + info = GetSamplerInfo(info, base_offset); // If this sampler has already been used, return the existing mapping. const auto it = std::find_if( used_samplers.begin(), used_samplers.end(), - [base_offset](const Sampler& entry) { return entry.GetOffset() == base_offset; }); + [base_offset](const Sampler& entry) { return entry.offset == base_offset; }); if (it != used_samplers.end()) { - ASSERT(!it->IsBindless() && it->GetType() == info.type && - it->IsArray() == info.is_array && it->IsShadow() == info.is_shadow && - it->IsBuffer() == info.is_buffer && it->IsIndexed()); - return &*it; + ASSERT(!it->is_bindless && it->type == info.type && it->is_array == info.is_array && + it->is_shadow == info.is_shadow && it->is_buffer == info.is_buffer && + it->is_indexed); + return *it; } uses_indexed_samplers = true; // Otherwise create a new mapping for this sampler const auto next_index = static_cast<u32>(used_samplers.size()); - return &used_samplers.emplace_back(next_index, base_offset, info.type, info.is_array, - info.is_shadow, info.is_buffer, true); + return used_samplers.emplace_back(next_index, base_offset, *info.type, *info.is_array, + *info.is_shadow, *info.is_buffer, true); } - return nullptr; + return std::nullopt; } void ShaderIR::WriteTexInstructionFloat(NodeBlock& bb, Instruction instr, const Node4& components) { @@ -529,10 +538,16 @@ Node4 ShaderIR::GetTextureCode(Instruction instr, TextureType texture_type, ASSERT_MSG(texture_type != TextureType::Texture3D || !is_array || !is_shadow, "Illegal texture type"); - const SamplerInfo info{texture_type, is_array, is_shadow, false}; + SamplerInfo info; + info.type = texture_type; + info.is_array = is_array; + info.is_shadow = is_shadow; + info.is_buffer = false; + Node index_var; - const Sampler* sampler = is_bindless ? GetBindlessSampler(*bindless_reg, index_var, info) - : GetSampler(instr.sampler, info); + const std::optional<Sampler> sampler = is_bindless + ? GetBindlessSampler(*bindless_reg, info, index_var) + : GetSampler(instr.sampler, info); if (!sampler) { return {Immediate(0), Immediate(0), Immediate(0), Immediate(0)}; } @@ -683,12 +698,17 @@ Node4 ShaderIR::GetTld4Code(Instruction instr, TextureType texture_type, bool de u64 parameter_register = instr.gpr20.Value(); - const SamplerInfo info{texture_type, is_array, depth_compare, false}; - Node index_var{}; - const Sampler* sampler = is_bindless ? GetBindlessSampler(parameter_register++, index_var, info) - : GetSampler(instr.sampler, info); + SamplerInfo info; + info.type = texture_type; + info.is_array = is_array; + info.is_shadow = depth_compare; + + Node index_var; + const std::optional<Sampler> sampler = + is_bindless ? GetBindlessSampler(parameter_register++, info, index_var) + : GetSampler(instr.sampler, info); Node4 values; - if (sampler == nullptr) { + if (!sampler) { for (u32 element = 0; element < values.size(); ++element) { values[element] = Immediate(0); } @@ -743,12 +763,12 @@ Node4 ShaderIR::GetTldCode(Tegra::Shader::Instruction instr) { // const Node aoffi_register{is_aoffi ? GetRegister(gpr20_cursor++) : nullptr}; // const Node multisample{is_multisample ? GetRegister(gpr20_cursor++) : nullptr}; - const auto& sampler = *GetSampler(instr.sampler); + const std::optional<Sampler> sampler = GetSampler(instr.sampler, {}); Node4 values; for (u32 element = 0; element < values.size(); ++element) { auto coords_copy = coords; - MetaTexture meta{sampler, array_register, {}, {}, {}, {}, {}, lod, {}, element, {}}; + MetaTexture meta{*sampler, array_register, {}, {}, {}, {}, {}, lod, {}, element, {}}; values[element] = Operation(OperationCode::TexelFetch, meta, std::move(coords_copy)); } @@ -756,7 +776,11 @@ Node4 ShaderIR::GetTldCode(Tegra::Shader::Instruction instr) { } Node4 ShaderIR::GetTldsCode(Instruction instr, TextureType texture_type, bool is_array) { - const Sampler& sampler = *GetSampler(instr.sampler); + SamplerInfo info; + info.type = texture_type; + info.is_array = is_array; + info.is_shadow = false; + const std::optional<Sampler> sampler = GetSampler(instr.sampler, info); const std::size_t type_coord_count = GetCoordCount(texture_type); const bool lod_enabled = instr.tlds.GetTextureProcessMode() == TextureProcessMode::LL; @@ -784,7 +808,7 @@ Node4 ShaderIR::GetTldsCode(Instruction instr, TextureType texture_type, bool is Node4 values; for (u32 element = 0; element < values.size(); ++element) { auto coords_copy = coords; - MetaTexture meta{sampler, array, {}, {}, {}, {}, {}, lod, {}, element, {}}; + MetaTexture meta{*sampler, array, {}, {}, {}, {}, {}, lod, {}, element, {}}; values[element] = Operation(OperationCode::TexelFetch, meta, std::move(coords_copy)); } return values; diff --git a/src/video_core/shader/memory_util.cpp b/src/video_core/shader/memory_util.cpp new file mode 100644 index 000000000..074f21691 --- /dev/null +++ b/src/video_core/shader/memory_util.cpp @@ -0,0 +1,77 @@ +// Copyright 2020 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include <algorithm> +#include <cstddef> + +#include <boost/container_hash/hash.hpp> + +#include "common/common_types.h" +#include "core/core.h" +#include "video_core/engines/maxwell_3d.h" +#include "video_core/memory_manager.h" +#include "video_core/shader/memory_util.h" +#include "video_core/shader/shader_ir.h" + +namespace VideoCommon::Shader { + +GPUVAddr GetShaderAddress(Core::System& system, + Tegra::Engines::Maxwell3D::Regs::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; +} + +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; +} + +std::size_t CalculateProgramSize(const ProgramCode& program, bool is_compute) { + // This is the encoded version of BRA that jumps to itself. All Nvidia + // shaders end with one. + static constexpr u64 SELF_JUMPING_BRANCH = 0xE2400FFFFF07000FULL; + static constexpr u64 MASK = 0xFFFFFFFFFF7FFFFFULL; + + const std::size_t start_offset = is_compute ? KERNEL_MAIN_OFFSET : STAGE_MAIN_OFFSET; + 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()); +} + +ProgramCode GetShaderCode(Tegra::MemoryManager& memory_manager, GPUVAddr gpu_addr, + const u8* host_ptr, bool is_compute) { + ProgramCode code(VideoCommon::Shader::MAX_PROGRAM_LENGTH); + ASSERT_OR_EXECUTE(host_ptr != nullptr, { return code; }); + memory_manager.ReadBlockUnsafe(gpu_addr, code.data(), code.size() * sizeof(u64)); + code.resize(CalculateProgramSize(code, is_compute)); + return code; +} + +u64 GetUniqueIdentifier(Tegra::Engines::ShaderType shader_type, bool is_a, const ProgramCode& code, + const ProgramCode& code_b) { + u64 unique_identifier = boost::hash_value(code); + if (is_a) { + // VertexA programs include two programs + boost::hash_combine(unique_identifier, boost::hash_value(code_b)); + } + return unique_identifier; +} + +} // namespace VideoCommon::Shader diff --git a/src/video_core/shader/memory_util.h b/src/video_core/shader/memory_util.h new file mode 100644 index 000000000..be90d24fd --- /dev/null +++ b/src/video_core/shader/memory_util.h @@ -0,0 +1,47 @@ +// Copyright 2020 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#pragma once + +#include <cstddef> +#include <vector> + +#include "common/common_types.h" +#include "video_core/engines/maxwell_3d.h" +#include "video_core/engines/shader_type.h" + +namespace Core { +class System; +} + +namespace Tegra { +class MemoryManager; +} + +namespace VideoCommon::Shader { + +using ProgramCode = std::vector<u64>; + +constexpr u32 STAGE_MAIN_OFFSET = 10; +constexpr u32 KERNEL_MAIN_OFFSET = 0; + +/// Gets the address for the specified shader stage program +GPUVAddr GetShaderAddress(Core::System& system, + Tegra::Engines::Maxwell3D::Regs::ShaderProgram program); + +/// Gets if the current instruction offset is a scheduler instruction +bool IsSchedInstruction(std::size_t offset, std::size_t main_offset); + +/// Calculates the size of a program stream +std::size_t CalculateProgramSize(const ProgramCode& program, bool is_compute); + +/// Gets the shader program code from memory for the specified address +ProgramCode GetShaderCode(Tegra::MemoryManager& memory_manager, GPUVAddr gpu_addr, + const u8* host_ptr, bool is_compute); + +/// Hashes one (or two) program streams +u64 GetUniqueIdentifier(Tegra::Engines::ShaderType shader_type, bool is_a, const ProgramCode& code, + const ProgramCode& code_b = {}); + +} // namespace VideoCommon::Shader diff --git a/src/video_core/shader/node.h b/src/video_core/shader/node.h index 3eee961f5..601c822d2 100644 --- a/src/video_core/shader/node.h +++ b/src/video_core/shader/node.h @@ -132,6 +132,8 @@ enum class OperationCode { LogicalUNotEqual, /// (uint a, uint b) -> bool LogicalUGreaterEqual, /// (uint a, uint b) -> bool + LogicalAddCarry, /// (uint a, uint b) -> bool + Logical2HLessThan, /// (MetaHalfArithmetic, f16vec2 a, f16vec2) -> bool2 Logical2HEqual, /// (MetaHalfArithmetic, f16vec2 a, f16vec2) -> bool2 Logical2HLessEqual, /// (MetaHalfArithmetic, f16vec2 a, f16vec2) -> bool2 @@ -265,76 +267,30 @@ class ArraySamplerNode; using TrackSamplerData = std::variant<BindlessSamplerNode, ArraySamplerNode>; using TrackSampler = std::shared_ptr<TrackSamplerData>; -class Sampler { -public: - /// This constructor is for bound samplers +struct Sampler { + /// Bound samplers constructor constexpr explicit Sampler(u32 index, u32 offset, Tegra::Shader::TextureType type, bool is_array, bool is_shadow, bool is_buffer, bool is_indexed) : index{index}, offset{offset}, type{type}, is_array{is_array}, is_shadow{is_shadow}, is_buffer{is_buffer}, is_indexed{is_indexed} {} - /// This constructor is for bindless samplers + /// Bindless samplers constructor constexpr explicit Sampler(u32 index, u32 offset, u32 buffer, Tegra::Shader::TextureType type, bool is_array, bool is_shadow, bool is_buffer, bool is_indexed) : index{index}, offset{offset}, buffer{buffer}, type{type}, is_array{is_array}, is_shadow{is_shadow}, is_buffer{is_buffer}, is_bindless{true}, is_indexed{is_indexed} {} - constexpr u32 GetIndex() const { - return index; - } - - constexpr u32 GetOffset() const { - return offset; - } - - constexpr u32 GetBuffer() const { - return buffer; - } - - constexpr Tegra::Shader::TextureType GetType() const { - return type; - } - - constexpr bool IsArray() const { - return is_array; - } - - constexpr bool IsShadow() const { - return is_shadow; - } - - constexpr bool IsBuffer() const { - return is_buffer; - } - - constexpr bool IsBindless() const { - return is_bindless; - } - - constexpr bool IsIndexed() const { - return is_indexed; - } - - constexpr u32 Size() const { - return size; - } - - constexpr void SetSize(u32 new_size) { - size = new_size; - } - -private: - u32 index{}; ///< Emulated index given for the this sampler. - u32 offset{}; ///< Offset in the const buffer from where the sampler is being read. - u32 buffer{}; ///< Buffer where the bindless sampler is being read (unused on bound samplers). - u32 size{1}; ///< Size of the sampler. + u32 index = 0; ///< Emulated index given for the this sampler. + u32 offset = 0; ///< Offset in the const buffer from where the sampler is being read. + u32 buffer = 0; ///< Buffer where the bindless sampler is being read (unused on bound samplers). + u32 size = 1; ///< Size of the sampler. Tegra::Shader::TextureType type{}; ///< The type used to sample this texture (Texture2D, etc) - bool is_array{}; ///< Whether the texture is being sampled as an array texture or not. - bool is_shadow{}; ///< Whether the texture is being sampled as a depth texture or not. - bool is_buffer{}; ///< Whether the texture is a texture buffer without sampler. - bool is_bindless{}; ///< Whether this sampler belongs to a bindless texture or not. - bool is_indexed{}; ///< Whether this sampler is an indexed array of textures. + bool is_array = false; ///< Whether the texture is being sampled as an array texture or not. + bool is_shadow = false; ///< Whether the texture is being sampled as a depth texture or not. + bool is_buffer = false; ///< Whether the texture is a texture buffer without sampler. + bool is_bindless = false; ///< Whether this sampler belongs to a bindless texture or not. + bool is_indexed = false; ///< Whether this sampler is an indexed array of textures. }; /// Represents a tracked bindless sampler into a direct const buffer @@ -379,13 +335,13 @@ private: u32 offset; }; -class Image final { +struct Image { public: - /// This constructor is for bound images + /// Bound images constructor constexpr explicit Image(u32 index, u32 offset, Tegra::Shader::ImageType type) : index{index}, offset{offset}, type{type} {} - /// This constructor is for bindless samplers + /// Bindless samplers constructor constexpr explicit Image(u32 index, u32 offset, u32 buffer, Tegra::Shader::ImageType type) : index{index}, offset{offset}, buffer{buffer}, type{type}, is_bindless{true} {} @@ -403,53 +359,20 @@ public: is_atomic = true; } - constexpr u32 GetIndex() const { - return index; - } - - constexpr u32 GetOffset() const { - return offset; - } - - constexpr u32 GetBuffer() const { - return buffer; - } - - constexpr Tegra::Shader::ImageType GetType() const { - return type; - } - - constexpr bool IsBindless() const { - return is_bindless; - } - - constexpr bool IsWritten() const { - return is_written; - } - - constexpr bool IsRead() const { - return is_read; - } - - constexpr bool IsAtomic() const { - return is_atomic; - } - -private: - u32 index{}; - u32 offset{}; - u32 buffer{}; + u32 index = 0; + u32 offset = 0; + u32 buffer = 0; Tegra::Shader::ImageType type{}; - bool is_bindless{}; - bool is_written{}; - bool is_read{}; - bool is_atomic{}; + bool is_bindless = false; + bool is_written = false; + bool is_read = false; + bool is_atomic = false; }; struct GlobalMemoryBase { - u32 cbuf_index{}; - u32 cbuf_offset{}; + u32 cbuf_index = 0; + u32 cbuf_offset = 0; bool operator<(const GlobalMemoryBase& rhs) const { return std::tie(cbuf_index, cbuf_offset) < std::tie(rhs.cbuf_index, rhs.cbuf_offset); @@ -463,7 +386,7 @@ struct MetaArithmetic { /// Parameters describing a texture sampler struct MetaTexture { - const Sampler& sampler; + Sampler sampler; Node array; Node depth_compare; std::vector<Node> aoffi; diff --git a/src/video_core/shader/shader_ir.h b/src/video_core/shader/shader_ir.h index c6e7bdf50..15ae152f2 100644 --- a/src/video_core/shader/shader_ir.h +++ b/src/video_core/shader/shader_ir.h @@ -18,6 +18,7 @@ #include "video_core/engines/shader_header.h" #include "video_core/shader/ast.h" #include "video_core/shader/compiler_settings.h" +#include "video_core/shader/memory_util.h" #include "video_core/shader/node.h" #include "video_core/shader/registry.h" @@ -25,16 +26,13 @@ namespace VideoCommon::Shader { struct ShaderBlock; -using ProgramCode = std::vector<u64>; - constexpr u32 MAX_PROGRAM_LENGTH = 0x1000; -class ConstBuffer { -public: - explicit ConstBuffer(u32 max_offset, bool is_indirect) +struct ConstBuffer { + constexpr explicit ConstBuffer(u32 max_offset, bool is_indirect) : max_offset{max_offset}, is_indirect{is_indirect} {} - ConstBuffer() = default; + constexpr ConstBuffer() = default; void MarkAsUsed(u64 offset) { max_offset = std::max(max_offset, static_cast<u32>(offset)); @@ -57,8 +55,8 @@ public: } private: - u32 max_offset{}; - bool is_indirect{}; + u32 max_offset = 0; + bool is_indirect = false; }; struct GlobalMemoryUsage { @@ -192,10 +190,14 @@ private: friend class ASTDecoder; struct SamplerInfo { - Tegra::Shader::TextureType type; - bool is_array; - bool is_shadow; - bool is_buffer; + std::optional<Tegra::Shader::TextureType> type; + std::optional<bool> is_array; + std::optional<bool> is_shadow; + std::optional<bool> is_buffer; + + constexpr bool IsComplete() const noexcept { + return type && is_array && is_shadow && is_buffer; + } }; void Decode(); @@ -328,16 +330,15 @@ private: OperationCode GetPredicateCombiner(Tegra::Shader::PredOperation operation); /// Queries the missing sampler info from the execution context. - SamplerInfo GetSamplerInfo(std::optional<SamplerInfo> sampler_info, u32 offset, + SamplerInfo GetSamplerInfo(SamplerInfo info, u32 offset, std::optional<u32> buffer = std::nullopt); - /// Accesses a texture sampler - const Sampler* GetSampler(const Tegra::Shader::Sampler& sampler, - std::optional<SamplerInfo> sampler_info = std::nullopt); + /// Accesses a texture sampler. + std::optional<Sampler> GetSampler(Tegra::Shader::Sampler sampler, SamplerInfo info); /// Accesses a texture sampler for a bindless texture. - const Sampler* GetBindlessSampler(Tegra::Shader::Register reg, Node& index_var, - std::optional<SamplerInfo> sampler_info = std::nullopt); + std::optional<Sampler> GetBindlessSampler(Tegra::Shader::Register reg, SamplerInfo info, + Node& index_var); /// Accesses an image. Image& GetImage(Tegra::Shader::Image image, Tegra::Shader::ImageType type); diff --git a/src/video_core/shader/track.cpp b/src/video_core/shader/track.cpp index 513e9bf49..eb97bfd41 100644 --- a/src/video_core/shader/track.cpp +++ b/src/video_core/shader/track.cpp @@ -153,21 +153,13 @@ std::tuple<Node, u32, u32> ShaderIR::TrackCbuf(Node tracked, const NodeBlock& co if (gpr->GetIndex() == Tegra::Shader::Register::ZeroIndex) { return {}; } - s64 current_cursor = cursor; - while (current_cursor > 0) { - // Reduce the cursor in one to avoid infinite loops when the instruction sets the same - // register that it uses as operand - const auto [source, new_cursor] = TrackRegister(gpr, code, current_cursor - 1); - current_cursor = new_cursor; - if (!source) { - continue; - } - const auto [base_address, index, offset] = TrackCbuf(source, code, current_cursor); - if (base_address != nullptr) { - return {base_address, index, offset}; - } + // Reduce the cursor in one to avoid infinite loops when the instruction sets the same + // register that it uses as operand + const auto [source, new_cursor] = TrackRegister(gpr, code, cursor - 1); + if (!source) { + return {}; } - return {}; + return TrackCbuf(source, code, new_cursor); } if (const auto operation = std::get_if<OperationNode>(&*tracked)) { for (std::size_t i = operation->GetOperandsCount(); i > 0; --i) { diff --git a/src/video_core/texture_cache/surface_params.cpp b/src/video_core/texture_cache/surface_params.cpp index 0de499946..884fabffe 100644 --- a/src/video_core/texture_cache/surface_params.cpp +++ b/src/video_core/texture_cache/surface_params.cpp @@ -81,7 +81,7 @@ SurfaceParams SurfaceParams::CreateForTexture(const FormatLookupTable& lookup_ta params.pixel_format = lookup_table.GetPixelFormat( tic.format, params.srgb_conversion, tic.r_type, tic.g_type, tic.b_type, tic.a_type); params.type = GetFormatType(params.pixel_format); - if (entry.IsShadow() && params.type == SurfaceType::ColorTexture) { + if (entry.is_shadow && params.type == SurfaceType::ColorTexture) { switch (params.pixel_format) { case PixelFormat::R16U: case PixelFormat::R16F: @@ -108,7 +108,7 @@ SurfaceParams SurfaceParams::CreateForTexture(const FormatLookupTable& lookup_ta params.emulated_levels = 1; params.is_layered = false; } else { - params.target = TextureTypeToSurfaceTarget(entry.GetType(), entry.IsArray()); + params.target = TextureTypeToSurfaceTarget(entry.type, entry.is_array); params.width = tic.Width(); params.height = tic.Height(); params.depth = tic.Depth(); @@ -138,7 +138,7 @@ SurfaceParams SurfaceParams::CreateForImage(const FormatLookupTable& lookup_tabl tic.format, params.srgb_conversion, tic.r_type, tic.g_type, tic.b_type, tic.a_type); params.type = GetFormatType(params.pixel_format); params.type = GetFormatType(params.pixel_format); - params.target = ImageTypeToSurfaceTarget(entry.GetType()); + params.target = ImageTypeToSurfaceTarget(entry.type); // TODO: on 1DBuffer we should use the tic info. if (tic.IsBuffer()) { params.target = SurfaceTarget::TextureBuffer; @@ -248,12 +248,12 @@ SurfaceParams SurfaceParams::CreateForFermiCopySurface( VideoCore::Surface::SurfaceTarget SurfaceParams::ExpectedTarget( const VideoCommon::Shader::Sampler& entry) { - return TextureTypeToSurfaceTarget(entry.GetType(), entry.IsArray()); + return TextureTypeToSurfaceTarget(entry.type, entry.is_array); } VideoCore::Surface::SurfaceTarget SurfaceParams::ExpectedTarget( const VideoCommon::Shader::Image& entry) { - return ImageTypeToSurfaceTarget(entry.GetType()); + return ImageTypeToSurfaceTarget(entry.type); } bool SurfaceParams::IsLayered() const { diff --git a/src/video_core/texture_cache/texture_cache.h b/src/video_core/texture_cache/texture_cache.h index cf6bd005a..d6efc34b2 100644 --- a/src/video_core/texture_cache/texture_cache.h +++ b/src/video_core/texture_cache/texture_cache.h @@ -143,7 +143,7 @@ public: } const auto params{SurfaceParams::CreateForTexture(format_lookup_table, tic, entry)}; - const auto [surface, view] = GetSurface(gpu_addr, *cpu_addr, params, false); + const auto [surface, view] = GetSurface(gpu_addr, *cpu_addr, params, true, false); if (guard_samplers) { sampled_textures.push_back(surface); } @@ -163,7 +163,7 @@ public: return GetNullSurface(SurfaceParams::ExpectedTarget(entry)); } const auto params{SurfaceParams::CreateForImage(format_lookup_table, tic, entry)}; - const auto [surface, view] = GetSurface(gpu_addr, *cpu_addr, params, false); + const auto [surface, view] = GetSurface(gpu_addr, *cpu_addr, params, true, false); if (guard_samplers) { sampled_textures.push_back(surface); } @@ -178,7 +178,7 @@ public: return any_rt; } - TView GetDepthBufferSurface() { + TView GetDepthBufferSurface(bool preserve_contents) { std::lock_guard lock{mutex}; auto& maxwell3d = system.GPU().Maxwell3D(); if (!maxwell3d.dirty.flags[VideoCommon::Dirty::ZetaBuffer]) { @@ -199,7 +199,7 @@ public: return {}; } const auto depth_params{SurfaceParams::CreateForDepthBuffer(system)}; - auto surface_view = GetSurface(gpu_addr, *cpu_addr, depth_params, true); + auto surface_view = GetSurface(gpu_addr, *cpu_addr, depth_params, preserve_contents, true); if (depth_buffer.target) depth_buffer.target->MarkAsRenderTarget(false, NO_RT); depth_buffer.target = surface_view.first; @@ -209,7 +209,7 @@ public: return surface_view.second; } - TView GetColorBufferSurface(std::size_t index) { + TView GetColorBufferSurface(std::size_t index, bool preserve_contents) { std::lock_guard lock{mutex}; ASSERT(index < Tegra::Engines::Maxwell3D::Regs::NumRenderTargets); auto& maxwell3d = system.GPU().Maxwell3D(); @@ -239,8 +239,9 @@ public: return {}; } - auto surface_view = GetSurface(gpu_addr, *cpu_addr, - SurfaceParams::CreateForFramebuffer(system, index), true); + auto surface_view = + GetSurface(gpu_addr, *cpu_addr, SurfaceParams::CreateForFramebuffer(system, index), + preserve_contents, true); if (render_targets[index].target) { auto& surface = render_targets[index].target; surface->MarkAsRenderTarget(false, NO_RT); @@ -300,9 +301,9 @@ public: const std::optional<VAddr> src_cpu_addr = system.GPU().MemoryManager().GpuToCpuAddress(src_gpu_addr); std::pair<TSurface, TView> dst_surface = - GetSurface(dst_gpu_addr, *dst_cpu_addr, dst_params, false); + GetSurface(dst_gpu_addr, *dst_cpu_addr, dst_params, true, false); std::pair<TSurface, TView> src_surface = - GetSurface(src_gpu_addr, *src_cpu_addr, src_params, false); + GetSurface(src_gpu_addr, *src_cpu_addr, src_params, true, false); ImageBlit(src_surface.second, dst_surface.second, copy_config); dst_surface.first->MarkAsModified(true, Tick()); } @@ -532,18 +533,22 @@ private: * @param overlaps The overlapping surfaces registered in the cache. * @param params The parameters for the new surface. * @param gpu_addr The starting address of the new surface. + * @param preserve_contents Indicates that the new surface should be loaded from memory or left + * blank. * @param untopological Indicates to the recycler that the texture has no way to match the * overlaps due to topological reasons. **/ std::pair<TSurface, TView> RecycleSurface(std::vector<TSurface>& overlaps, const SurfaceParams& params, const GPUVAddr gpu_addr, + const bool preserve_contents, const MatchTopologyResult untopological) { + const bool do_load = preserve_contents && Settings::IsGPULevelExtreme(); for (auto& surface : overlaps) { Unregister(surface); } switch (PickStrategy(overlaps, params, gpu_addr, untopological)) { case RecycleStrategy::Ignore: { - return InitializeSurface(gpu_addr, params, Settings::IsGPULevelExtreme()); + return InitializeSurface(gpu_addr, params, do_load); } case RecycleStrategy::Flush: { std::sort(overlaps.begin(), overlaps.end(), @@ -553,7 +558,7 @@ private: for (auto& surface : overlaps) { FlushSurface(surface); } - return InitializeSurface(gpu_addr, params); + return InitializeSurface(gpu_addr, params, preserve_contents); } case RecycleStrategy::BufferCopy: { auto new_surface = GetUncachedSurface(gpu_addr, params); @@ -562,7 +567,7 @@ private: } default: { UNIMPLEMENTED_MSG("Unimplemented Texture Cache Recycling Strategy!"); - return InitializeSurface(gpu_addr, params); + return InitializeSurface(gpu_addr, params, do_load); } } } @@ -700,11 +705,14 @@ private: * @param params The parameters on the new surface. * @param gpu_addr The starting address of the new surface. * @param cpu_addr The starting address of the new surface on physical memory. + * @param preserve_contents Indicates that the new surface should be loaded from memory or + * left blank. */ std::optional<std::pair<TSurface, TView>> Manage3DSurfaces(std::vector<TSurface>& overlaps, const SurfaceParams& params, const GPUVAddr gpu_addr, - const VAddr cpu_addr) { + const VAddr cpu_addr, + bool preserve_contents) { if (params.target == SurfaceTarget::Texture3D) { bool failed = false; if (params.num_levels > 1) { @@ -754,7 +762,7 @@ private: return std::nullopt; } Unregister(surface); - return InitializeSurface(gpu_addr, params); + return InitializeSurface(gpu_addr, params, preserve_contents); } return std::nullopt; } @@ -765,7 +773,7 @@ private: return {{surface, surface->GetMainView()}}; } } - return InitializeSurface(gpu_addr, params); + return InitializeSurface(gpu_addr, params, preserve_contents); } } @@ -788,10 +796,13 @@ private: * * @param gpu_addr The starting address of the candidate surface. * @param params The parameters on the candidate surface. + * @param preserve_contents Indicates that the new surface should be loaded from memory or + * left blank. * @param is_render Whether or not the surface is a render target. **/ std::pair<TSurface, TView> GetSurface(const GPUVAddr gpu_addr, const VAddr cpu_addr, - const SurfaceParams& params, bool is_render) { + const SurfaceParams& params, bool preserve_contents, + bool is_render) { // Step 1 // Check Level 1 Cache for a fast structural match. If candidate surface // matches at certain level we are pretty much done. @@ -800,7 +811,8 @@ private: const auto topological_result = current_surface->MatchesTopology(params); if (topological_result != MatchTopologyResult::FullMatch) { std::vector<TSurface> overlaps{current_surface}; - return RecycleSurface(overlaps, params, gpu_addr, topological_result); + return RecycleSurface(overlaps, params, gpu_addr, preserve_contents, + topological_result); } const auto struct_result = current_surface->MatchesStructure(params); @@ -825,7 +837,7 @@ private: // If none are found, we are done. we just load the surface and create it. if (overlaps.empty()) { - return InitializeSurface(gpu_addr, params); + return InitializeSurface(gpu_addr, params, preserve_contents); } // Step 3 @@ -835,13 +847,15 @@ private: for (const auto& surface : overlaps) { const auto topological_result = surface->MatchesTopology(params); if (topological_result != MatchTopologyResult::FullMatch) { - return RecycleSurface(overlaps, params, gpu_addr, topological_result); + return RecycleSurface(overlaps, params, gpu_addr, preserve_contents, + topological_result); } } // Check if it's a 3D texture if (params.block_depth > 0) { - auto surface = Manage3DSurfaces(overlaps, params, gpu_addr, cpu_addr); + auto surface = + Manage3DSurfaces(overlaps, params, gpu_addr, cpu_addr, preserve_contents); if (surface) { return *surface; } @@ -861,7 +875,8 @@ private: return *view; } } - return RecycleSurface(overlaps, params, gpu_addr, MatchTopologyResult::FullMatch); + return RecycleSurface(overlaps, params, gpu_addr, preserve_contents, + MatchTopologyResult::FullMatch); } // Now we check if the candidate is a mipmap/layer of the overlap std::optional<TView> view = @@ -885,7 +900,7 @@ private: pair.first->EmplaceView(params, gpu_addr, candidate_size); if (mirage_view) return {pair.first, *mirage_view}; - return RecycleSurface(overlaps, params, gpu_addr, + return RecycleSurface(overlaps, params, gpu_addr, preserve_contents, MatchTopologyResult::FullMatch); } return {current_surface, *view}; @@ -901,7 +916,8 @@ private: } } // We failed all the tests, recycle the overlaps into a new texture. - return RecycleSurface(overlaps, params, gpu_addr, MatchTopologyResult::FullMatch); + return RecycleSurface(overlaps, params, gpu_addr, preserve_contents, + MatchTopologyResult::FullMatch); } /** @@ -1059,10 +1075,10 @@ private: } std::pair<TSurface, TView> InitializeSurface(GPUVAddr gpu_addr, const SurfaceParams& params, - bool do_load = true) { + bool preserve_contents) { auto new_surface{GetUncachedSurface(gpu_addr, params)}; Register(new_surface); - if (do_load) { + if (preserve_contents) { LoadSurface(new_surface); } return {new_surface, new_surface->GetMainView()}; @@ -1156,7 +1172,7 @@ private: /// Returns true the shader sampler entry is compatible with the TIC texture type. static bool IsTypeCompatible(Tegra::Texture::TextureType tic_type, const VideoCommon::Shader::Sampler& entry) { - const auto shader_type = entry.GetType(); + const auto shader_type = entry.type; switch (tic_type) { case Tegra::Texture::TextureType::Texture1D: case Tegra::Texture::TextureType::Texture1DArray: @@ -1177,7 +1193,7 @@ private: if (shader_type == Tegra::Shader::TextureType::TextureCube) { return true; } - return shader_type == Tegra::Shader::TextureType::Texture2D && entry.IsArray(); + return shader_type == Tegra::Shader::TextureType::Texture2D && entry.is_array; } UNREACHABLE(); return true; |