spirv: Handle tgid enable bits

This commit is contained in:
IndecisiveTurtle 2024-08-15 20:25:18 +03:00
parent 1ca8a5c3c9
commit e47a61dec9
9 changed files with 19 additions and 15 deletions

View File

@ -10,7 +10,6 @@
#include <arpa/inet.h> #include <arpa/inet.h>
#endif #endif
#include <thread>
#include "common/assert.h" #include "common/assert.h"
#include "common/logging/log.h" #include "common/logging/log.h"
#include "core/libraries/error_codes.h" #include "core/libraries/error_codes.h"
@ -60,7 +59,6 @@ int PS4_SYSV_ABI sce_net_in6addr_nodelocal_allnodes() {
} }
OrbisNetId PS4_SYSV_ABI sceNetAccept(OrbisNetId s, OrbisNetSockaddr* addr, u32* paddrlen) { OrbisNetId PS4_SYSV_ABI sceNetAccept(OrbisNetId s, OrbisNetSockaddr* addr, u32* paddrlen) {
std::this_thread::sleep_for(std::chrono::seconds(5));
LOG_ERROR(Lib_Net, "(STUBBED) called"); LOG_ERROR(Lib_Net, "(STUBBED) called");
return ORBIS_OK; return ORBIS_OK;
} }

View File

@ -146,9 +146,6 @@ void CFG::EmitDivergenceLabels() {
// Sort labels to make sure block insertion is correct. // Sort labels to make sure block insertion is correct.
std::ranges::sort(labels); std::ranges::sort(labels);
for (const auto label : labels) {
LOG_INFO(Render_Vulkan, "Emitting label {:#x}", label);
}
} }
void CFG::EmitBlocks() { void CFG::EmitBlocks() {
@ -165,7 +162,7 @@ void CFG::EmitBlocks() {
const Label end = *next_it; const Label end = *next_it;
const size_t end_index = GetIndex(end) - 1; const size_t end_index = GetIndex(end) - 1;
const auto& end_inst = inst_list[end_index]; const auto& end_inst = inst_list[end_index];
LOG_INFO(Render_Vulkan, "Emitting block {:#x}-{:#x}", start, end);
// Insert block between the labels using the last instruction // Insert block between the labels using the last instruction
// as an indicator for branching type. // as an indicator for branching type.
Block* block = block_pool.Create(); Block* block = block_pool.Create();

View File

@ -64,9 +64,15 @@ void Translator::EmitPrologue() {
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::LocalInvocationId, 1)); ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::LocalInvocationId, 1));
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::LocalInvocationId, 2)); ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::LocalInvocationId, 2));
if (info.tgid_enable[0]) {
ir.SetScalarReg(dst_sreg++, ir.GetAttributeU32(IR::Attribute::WorkgroupId, 0)); ir.SetScalarReg(dst_sreg++, ir.GetAttributeU32(IR::Attribute::WorkgroupId, 0));
}
if (info.tgid_enable[1]) {
ir.SetScalarReg(dst_sreg++, ir.GetAttributeU32(IR::Attribute::WorkgroupId, 1)); ir.SetScalarReg(dst_sreg++, ir.GetAttributeU32(IR::Attribute::WorkgroupId, 1));
}
if (info.tgid_enable[2]) {
ir.SetScalarReg(dst_sreg++, ir.GetAttributeU32(IR::Attribute::WorkgroupId, 2)); ir.SetScalarReg(dst_sreg++, ir.GetAttributeU32(IR::Attribute::WorkgroupId, 2));
}
break; break;
default: default:
throw NotImplementedException("Unknown shader stage"); throw NotImplementedException("Unknown shader stage");

View File

@ -180,6 +180,7 @@ struct Info {
SamplerResourceList samplers; SamplerResourceList samplers;
std::array<u32, 3> workgroup_size{}; std::array<u32, 3> workgroup_size{};
std::array<bool, 3> tgid_enable;
u32 num_user_data; u32 num_user_data;
u32 num_input_vgprs; u32 num_input_vgprs;

View File

@ -180,10 +180,6 @@ Liverpool::Task Liverpool::ProcessGraphics(std::span<const u32> dcb, std::span<c
UNREACHABLE_MSG("Invalid PM4 type {}", type); UNREACHABLE_MSG("Invalid PM4 type {}", type);
} }
if (VAddr(dcb.data()) == 0x3c24df50) {
printf("bad\n");
}
const u32 count = header->type3.NumWords(); const u32 count = header->type3.NumWords();
const PM4ItOpcode opcode = header->type3.opcode; const PM4ItOpcode opcode = header->type3.opcode;
switch (opcode) { switch (opcode) {

View File

@ -130,6 +130,7 @@ struct Liverpool {
BitField<0, 6, u64> num_vgprs; BitField<0, 6, u64> num_vgprs;
BitField<6, 4, u64> num_sgprs; BitField<6, 4, u64> num_sgprs;
BitField<33, 5, u64> num_user_regs; BitField<33, 5, u64> num_user_regs;
BitField<39, 3, u64> tgid_enable;
BitField<47, 9, u64> lds_dwords; BitField<47, 9, u64> lds_dwords;
} settings; } settings;
INSERT_PADDING_WORDS(1); INSERT_PADDING_WORDS(1);
@ -148,6 +149,10 @@ struct Liverpool {
return settings.lds_dwords.Value() * 128 * 4; return settings.lds_dwords.Value() * 128 * 4;
} }
bool IsTgidEnabled(u32 i) const noexcept {
return (settings.tgid_enable.Value() >> i) & 1;
}
std::span<const u32> Code() const { std::span<const u32> Code() const {
const u32* code = Address<u32*>(); const u32* code = Address<u32*>();
BinaryInfo bininfo; BinaryInfo bininfo;

View File

@ -1,6 +1,6 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project // SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later // SPDX-License-Identifier: GPL-2.0-or-later
#pragma clang optimize off
#include <algorithm> #include <algorithm>
#include "common/alignment.h" #include "common/alignment.h"
#include "common/scope_exit.h" #include "common/scope_exit.h"

View File

@ -3,7 +3,6 @@
#pragma once #pragma once
#include <array>
#include <mutex> #include <mutex>
#include <boost/container/small_vector.hpp> #include <boost/container/small_vector.hpp>
#include <boost/icl/interval_map.hpp> #include <boost/icl/interval_map.hpp>

View File

@ -93,6 +93,8 @@ Shader::Info MakeShaderInfo(Shader::Stage stage, std::span<const u32, 16> user_d
info.num_user_data = cs_pgm.settings.num_user_regs; info.num_user_data = cs_pgm.settings.num_user_regs;
info.workgroup_size = {cs_pgm.num_thread_x.full, cs_pgm.num_thread_y.full, info.workgroup_size = {cs_pgm.num_thread_x.full, cs_pgm.num_thread_y.full,
cs_pgm.num_thread_z.full}; cs_pgm.num_thread_z.full};
info.tgid_enable = {cs_pgm.IsTgidEnabled(0), cs_pgm.IsTgidEnabled(1),
cs_pgm.IsTgidEnabled(2)};
info.shared_memory_size = cs_pgm.SharedMemSize(); info.shared_memory_size = cs_pgm.SharedMemSize();
break; break;
} }