emu-mirror/yuzu
1
0
Fork 0
mirror of https://github.com/yuzu-mirror/yuzu.git synced 2024-11-18 00:47:33 +01:00
Code Issues Packages Projects Releases Wiki Activity

shader_environment: Fix local memory size calculations

Browse Source
This commit is contained in:
ReinUsesLisp 2021-06-10 02:27:00 -03:00 committed by ameerj
parent 60a96c49e5
commit 5befc0bf87
1 changed files with 3 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
src/video_core/shader_environment.cpp
View File

@ -69,7 +69,7 @@ u32 GenericEnvironment::TextureBoundBuffer() const {
} }
u32 GenericEnvironment::LocalMemorySize() const { u32 GenericEnvironment::LocalMemorySize() const {
return local_memory_size + sph.common3.shader_local_memory_crs_size; return local_memory_size;
} }
u32 GenericEnvironment::SharedMemorySize() const { u32 GenericEnvironment::SharedMemorySize() const {
@ -233,7 +233,7 @@ GraphicsEnvironment::GraphicsEnvironment(Tegra::Engines::Maxwell3D& maxwell3d_,
} }
const u64 local_size{sph.LocalMemorySize()}; const u64 local_size{sph.LocalMemorySize()};
ASSERT(local_size <= std::numeric_limits<u32>::max()); ASSERT(local_size <= std::numeric_limits<u32>::max());
local_memory_size = static_cast<u32>(local_size); local_memory_size = static_cast<u32>(local_size) + sph.common3.shader_local_memory_crs_size;
texture_bound = maxwell3d->regs.tex_cb_index; texture_bound = maxwell3d->regs.tex_cb_index;
} }
@ -261,7 +261,7 @@ ComputeEnvironment::ComputeEnvironment(Tegra::Engines::KeplerCompute& kepler_com
&kepler_compute_} { &kepler_compute_} {
const auto& qmd{kepler_compute->launch_description}; const auto& qmd{kepler_compute->launch_description};
stage = Shader::Stage::Compute; stage = Shader::Stage::Compute;
local_memory_size = qmd.local_pos_alloc; local_memory_size = qmd.local_pos_alloc + qmd.local_crs_alloc;
texture_bound = kepler_compute->regs.tex_cb_index; texture_bound = kepler_compute->regs.tex_cb_index;
shared_memory_size = qmd.shared_alloc; shared_memory_size = qmd.shared_alloc;
workgroup_size = {qmd.block_dim_x, qmd.block_dim_y, qmd.block_dim_z}; workgroup_size = {qmd.block_dim_x, qmd.block_dim_y, qmd.block_dim_z};