spirv: Implement VertexId and InstanceId, refactor code
This commit is contained in:
parent
e802512d8e
commit
76c8a962ac
|
@ -48,6 +48,25 @@ Id ImageType(EmitContext& ctx, const TextureDescriptor& desc) {
|
||||||
}
|
}
|
||||||
throw InvalidArgument("Invalid texture type {}", desc.type);
|
throw InvalidArgument("Invalid texture type {}", desc.type);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
Id DefineVariable(EmitContext& ctx, Id type, std::optional<spv::BuiltIn> builtin,
|
||||||
|
spv::StorageClass storage_class) {
|
||||||
|
const Id pointer_type{ctx.TypePointer(storage_class, type)};
|
||||||
|
const Id id{ctx.AddGlobalVariable(pointer_type, storage_class)};
|
||||||
|
if (builtin) {
|
||||||
|
ctx.Decorate(id, spv::Decoration::BuiltIn, *builtin);
|
||||||
|
}
|
||||||
|
ctx.interfaces.push_back(id);
|
||||||
|
return id;
|
||||||
|
}
|
||||||
|
|
||||||
|
Id DefineInput(EmitContext& ctx, Id type, std::optional<spv::BuiltIn> builtin = std::nullopt) {
|
||||||
|
return DefineVariable(ctx, type, builtin, spv::StorageClass::Input);
|
||||||
|
}
|
||||||
|
|
||||||
|
Id DefineOutput(EmitContext& ctx, Id type, std::optional<spv::BuiltIn> builtin = std::nullopt) {
|
||||||
|
return DefineVariable(ctx, type, builtin, spv::StorageClass::Output);
|
||||||
|
}
|
||||||
} // Anonymous namespace
|
} // Anonymous namespace
|
||||||
|
|
||||||
void VectorTypes::Define(Sirit::Module& sirit_ctx, Id base_type, std::string_view name) {
|
void VectorTypes::Define(Sirit::Module& sirit_ctx, Id base_type, std::string_view name) {
|
||||||
|
@ -144,59 +163,8 @@ void EmitContext::DefineCommonConstants() {
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitContext::DefineInterfaces(const Info& info, Stage stage) {
|
void EmitContext::DefineInterfaces(const Info& info, Stage stage) {
|
||||||
const auto define{
|
DefineInputs(info, stage);
|
||||||
[this](Id type, std::optional<spv::BuiltIn> builtin, spv::StorageClass storage_class) {
|
DefineOutputs(info, stage);
|
||||||
const Id pointer_type{TypePointer(storage_class, type)};
|
|
||||||
const Id id{AddGlobalVariable(pointer_type, storage_class)};
|
|
||||||
if (builtin) {
|
|
||||||
Decorate(id, spv::Decoration::BuiltIn, *builtin);
|
|
||||||
}
|
|
||||||
interfaces.push_back(id);
|
|
||||||
return id;
|
|
||||||
}};
|
|
||||||
using namespace std::placeholders;
|
|
||||||
const auto define_input{std::bind(define, _1, _2, spv::StorageClass::Input)};
|
|
||||||
const auto define_output{std::bind(define, _1, _2, spv::StorageClass::Output)};
|
|
||||||
|
|
||||||
if (info.uses_workgroup_id) {
|
|
||||||
workgroup_id = define_input(U32[3], spv::BuiltIn::WorkgroupId);
|
|
||||||
}
|
|
||||||
if (info.uses_local_invocation_id) {
|
|
||||||
local_invocation_id = define_input(U32[3], spv::BuiltIn::LocalInvocationId);
|
|
||||||
}
|
|
||||||
if (info.loads_position) {
|
|
||||||
const bool is_fragment{stage != Stage::Fragment};
|
|
||||||
const spv::BuiltIn built_in{is_fragment ? spv::BuiltIn::Position : spv::BuiltIn::FragCoord};
|
|
||||||
input_position = define_input(F32[4], built_in);
|
|
||||||
}
|
|
||||||
for (size_t i = 0; i < info.loads_generics.size(); ++i) {
|
|
||||||
if (info.loads_generics[i]) {
|
|
||||||
// FIXME: Declare size from input
|
|
||||||
input_generics[i] = define_input(F32[4], std::nullopt);
|
|
||||||
Decorate(input_generics[i], spv::Decoration::Location, static_cast<u32>(i));
|
|
||||||
Name(input_generics[i], fmt::format("in_attr{}", i));
|
|
||||||
}
|
|
||||||
}
|
|
||||||
if (info.stores_position) {
|
|
||||||
output_position = define_output(F32[4], spv::BuiltIn::Position);
|
|
||||||
}
|
|
||||||
for (size_t i = 0; i < info.stores_generics.size(); ++i) {
|
|
||||||
if (info.stores_generics[i]) {
|
|
||||||
output_generics[i] = define_output(F32[4], std::nullopt);
|
|
||||||
Decorate(output_generics[i], spv::Decoration::Location, static_cast<u32>(i));
|
|
||||||
Name(output_generics[i], fmt::format("out_attr{}", i));
|
|
||||||
}
|
|
||||||
}
|
|
||||||
if (stage == Stage::Fragment) {
|
|
||||||
for (size_t i = 0; i < 8; ++i) {
|
|
||||||
if (!info.stores_frag_color[i]) {
|
|
||||||
continue;
|
|
||||||
}
|
|
||||||
frag_color[i] = define_output(F32[4], std::nullopt);
|
|
||||||
Decorate(frag_color[i], spv::Decoration::Location, static_cast<u32>(i));
|
|
||||||
Name(frag_color[i], fmt::format("frag_color{}", i));
|
|
||||||
}
|
|
||||||
}
|
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitContext::DefineConstantBuffers(const Info& info, u32& binding) {
|
void EmitContext::DefineConstantBuffers(const Info& info, u32& binding) {
|
||||||
|
@ -225,33 +193,6 @@ void EmitContext::DefineConstantBuffers(const Info& info, u32& binding) {
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitContext::DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type,
|
|
||||||
u32 binding, Id type, char type_char, u32 element_size) {
|
|
||||||
const Id array_type{TypeArray(type, Constant(U32[1], 65536U / element_size))};
|
|
||||||
Decorate(array_type, spv::Decoration::ArrayStride, element_size);
|
|
||||||
|
|
||||||
const Id struct_type{TypeStruct(array_type)};
|
|
||||||
Name(struct_type, fmt::format("cbuf_block_{}{}", type_char, element_size * CHAR_BIT));
|
|
||||||
Decorate(struct_type, spv::Decoration::Block);
|
|
||||||
MemberName(struct_type, 0, "data");
|
|
||||||
MemberDecorate(struct_type, 0, spv::Decoration::Offset, 0U);
|
|
||||||
|
|
||||||
const Id struct_pointer_type{TypePointer(spv::StorageClass::Uniform, struct_type)};
|
|
||||||
const Id uniform_type{TypePointer(spv::StorageClass::Uniform, type)};
|
|
||||||
uniform_types.*member_type = uniform_type;
|
|
||||||
|
|
||||||
for (const ConstantBufferDescriptor& desc : info.constant_buffer_descriptors) {
|
|
||||||
const Id id{AddGlobalVariable(struct_pointer_type, spv::StorageClass::Uniform)};
|
|
||||||
Decorate(id, spv::Decoration::Binding, binding);
|
|
||||||
Decorate(id, spv::Decoration::DescriptorSet, 0U);
|
|
||||||
Name(id, fmt::format("c{}", desc.index));
|
|
||||||
for (size_t i = 0; i < desc.count; ++i) {
|
|
||||||
cbufs[desc.index + i].*member_type = id;
|
|
||||||
}
|
|
||||||
binding += desc.count;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
void EmitContext::DefineStorageBuffers(const Info& info, u32& binding) {
|
void EmitContext::DefineStorageBuffers(const Info& info, u32& binding) {
|
||||||
if (info.storage_buffers_descriptors.empty()) {
|
if (info.storage_buffers_descriptors.empty()) {
|
||||||
return;
|
return;
|
||||||
|
@ -311,4 +252,94 @@ void EmitContext::DefineLabels(IR::Program& program) {
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
void EmitContext::DefineInputs(const Info& info, Stage stage) {
|
||||||
|
if (info.uses_workgroup_id) {
|
||||||
|
workgroup_id = DefineInput(*this, U32[3], spv::BuiltIn::WorkgroupId);
|
||||||
|
}
|
||||||
|
if (info.uses_local_invocation_id) {
|
||||||
|
local_invocation_id = DefineInput(*this, U32[3], spv::BuiltIn::LocalInvocationId);
|
||||||
|
}
|
||||||
|
if (info.loads_position) {
|
||||||
|
const bool is_fragment{stage != Stage::Fragment};
|
||||||
|
const spv::BuiltIn built_in{is_fragment ? spv::BuiltIn::Position : spv::BuiltIn::FragCoord};
|
||||||
|
input_position = DefineInput(*this, F32[4], built_in);
|
||||||
|
}
|
||||||
|
if (info.loads_instance_id) {
|
||||||
|
if (profile.support_vertex_instance_id) {
|
||||||
|
instance_id = DefineInput(*this, U32[1], spv::BuiltIn::InstanceId);
|
||||||
|
} else {
|
||||||
|
instance_index = DefineInput(*this, U32[1], spv::BuiltIn::InstanceIndex);
|
||||||
|
base_instance = DefineInput(*this, U32[1], spv::BuiltIn::BaseInstance);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
if (info.loads_vertex_id) {
|
||||||
|
if (profile.support_vertex_instance_id) {
|
||||||
|
vertex_id = DefineInput(*this, U32[1], spv::BuiltIn::VertexId);
|
||||||
|
} else {
|
||||||
|
vertex_index = DefineInput(*this, U32[1], spv::BuiltIn::VertexIndex);
|
||||||
|
base_vertex = DefineInput(*this, U32[1], spv::BuiltIn::BaseVertex);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
for (size_t index = 0; index < info.loads_generics.size(); ++index) {
|
||||||
|
if (!info.loads_generics[index]) {
|
||||||
|
continue;
|
||||||
|
}
|
||||||
|
// FIXME: Declare size from input
|
||||||
|
const Id id{DefineInput(*this, F32[4])};
|
||||||
|
Decorate(id, spv::Decoration::Location, static_cast<u32>(index));
|
||||||
|
Name(id, fmt::format("in_attr{}", index));
|
||||||
|
input_generics[index] = id;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
void EmitContext::DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type,
|
||||||
|
u32 binding, Id type, char type_char, u32 element_size) {
|
||||||
|
const Id array_type{TypeArray(type, Constant(U32[1], 65536U / element_size))};
|
||||||
|
Decorate(array_type, spv::Decoration::ArrayStride, element_size);
|
||||||
|
|
||||||
|
const Id struct_type{TypeStruct(array_type)};
|
||||||
|
Name(struct_type, fmt::format("cbuf_block_{}{}", type_char, element_size * CHAR_BIT));
|
||||||
|
Decorate(struct_type, spv::Decoration::Block);
|
||||||
|
MemberName(struct_type, 0, "data");
|
||||||
|
MemberDecorate(struct_type, 0, spv::Decoration::Offset, 0U);
|
||||||
|
|
||||||
|
const Id struct_pointer_type{TypePointer(spv::StorageClass::Uniform, struct_type)};
|
||||||
|
const Id uniform_type{TypePointer(spv::StorageClass::Uniform, type)};
|
||||||
|
uniform_types.*member_type = uniform_type;
|
||||||
|
|
||||||
|
for (const ConstantBufferDescriptor& desc : info.constant_buffer_descriptors) {
|
||||||
|
const Id id{AddGlobalVariable(struct_pointer_type, spv::StorageClass::Uniform)};
|
||||||
|
Decorate(id, spv::Decoration::Binding, binding);
|
||||||
|
Decorate(id, spv::Decoration::DescriptorSet, 0U);
|
||||||
|
Name(id, fmt::format("c{}", desc.index));
|
||||||
|
for (size_t i = 0; i < desc.count; ++i) {
|
||||||
|
cbufs[desc.index + i].*member_type = id;
|
||||||
|
}
|
||||||
|
binding += desc.count;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
void EmitContext::DefineOutputs(const Info& info, Stage stage) {
|
||||||
|
if (info.stores_position) {
|
||||||
|
output_position = DefineOutput(*this, F32[4], spv::BuiltIn::Position);
|
||||||
|
}
|
||||||
|
for (size_t i = 0; i < info.stores_generics.size(); ++i) {
|
||||||
|
if (info.stores_generics[i]) {
|
||||||
|
output_generics[i] = DefineOutput(*this, F32[4]);
|
||||||
|
Decorate(output_generics[i], spv::Decoration::Location, static_cast<u32>(i));
|
||||||
|
Name(output_generics[i], fmt::format("out_attr{}", i));
|
||||||
|
}
|
||||||
|
}
|
||||||
|
if (stage == Stage::Fragment) {
|
||||||
|
for (size_t i = 0; i < 8; ++i) {
|
||||||
|
if (!info.stores_frag_color[i]) {
|
||||||
|
continue;
|
||||||
|
}
|
||||||
|
frag_color[i] = DefineOutput(*this, F32[4]);
|
||||||
|
Decorate(frag_color[i], spv::Decoration::Location, static_cast<u32>(i));
|
||||||
|
Name(frag_color[i], fmt::format("frag_color{}", i));
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
} // namespace Shader::Backend::SPIRV
|
} // namespace Shader::Backend::SPIRV
|
||||||
|
|
|
@ -82,6 +82,12 @@ public:
|
||||||
|
|
||||||
Id workgroup_id{};
|
Id workgroup_id{};
|
||||||
Id local_invocation_id{};
|
Id local_invocation_id{};
|
||||||
|
Id instance_id{};
|
||||||
|
Id instance_index{};
|
||||||
|
Id base_instance{};
|
||||||
|
Id vertex_id{};
|
||||||
|
Id vertex_index{};
|
||||||
|
Id base_vertex{};
|
||||||
|
|
||||||
Id input_position{};
|
Id input_position{};
|
||||||
std::array<Id, 32> input_generics{};
|
std::array<Id, 32> input_generics{};
|
||||||
|
@ -99,11 +105,15 @@ private:
|
||||||
void DefineCommonConstants();
|
void DefineCommonConstants();
|
||||||
void DefineInterfaces(const Info& info, Stage stage);
|
void DefineInterfaces(const Info& info, Stage stage);
|
||||||
void DefineConstantBuffers(const Info& info, u32& binding);
|
void DefineConstantBuffers(const Info& info, u32& binding);
|
||||||
void DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type, u32 binding,
|
|
||||||
Id type, char type_char, u32 element_size);
|
|
||||||
void DefineStorageBuffers(const Info& info, u32& binding);
|
void DefineStorageBuffers(const Info& info, u32& binding);
|
||||||
void DefineTextures(const Info& info, u32& binding);
|
void DefineTextures(const Info& info, u32& binding);
|
||||||
void DefineLabels(IR::Program& program);
|
void DefineLabels(IR::Program& program);
|
||||||
|
|
||||||
|
void DefineConstantBuffers(const Info& info, Id UniformDefinitions::*member_type, u32 binding,
|
||||||
|
Id type, char type_char, u32 element_size);
|
||||||
|
|
||||||
|
void DefineInputs(const Info& info, Stage stage);
|
||||||
|
void DefineOutputs(const Info& info, Stage stage);
|
||||||
};
|
};
|
||||||
|
|
||||||
} // namespace Shader::Backend::SPIRV
|
} // namespace Shader::Backend::SPIRV
|
||||||
|
|
|
@ -113,6 +113,43 @@ Id TypeId(const EmitContext& ctx, IR::Type type) {
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
Id DefineMain(EmitContext& ctx, IR::Program& program) {
|
||||||
|
const Id void_function{ctx.TypeFunction(ctx.void_id)};
|
||||||
|
const Id main{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)};
|
||||||
|
for (IR::Block* const block : program.blocks) {
|
||||||
|
ctx.AddLabel(block->Definition<Id>());
|
||||||
|
for (IR::Inst& inst : block->Instructions()) {
|
||||||
|
EmitInst(ctx, &inst);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
ctx.OpFunctionEnd();
|
||||||
|
return main;
|
||||||
|
}
|
||||||
|
|
||||||
|
void DefineEntryPoint(Environment& env, EmitContext& ctx, Id main) {
|
||||||
|
const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size());
|
||||||
|
spv::ExecutionModel execution_model{};
|
||||||
|
switch (env.ShaderStage()) {
|
||||||
|
case Shader::Stage::Compute: {
|
||||||
|
const std::array<u32, 3> workgroup_size{env.WorkgroupSize()};
|
||||||
|
execution_model = spv::ExecutionModel::GLCompute;
|
||||||
|
ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, workgroup_size[0],
|
||||||
|
workgroup_size[1], workgroup_size[2]);
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
case Shader::Stage::VertexB:
|
||||||
|
execution_model = spv::ExecutionModel::Vertex;
|
||||||
|
break;
|
||||||
|
case Shader::Stage::Fragment:
|
||||||
|
execution_model = spv::ExecutionModel::Fragment;
|
||||||
|
ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft);
|
||||||
|
break;
|
||||||
|
default:
|
||||||
|
throw NotImplementedException("Stage {}", env.ShaderStage());
|
||||||
|
}
|
||||||
|
ctx.AddEntryPoint(execution_model, main, "main", interfaces);
|
||||||
|
}
|
||||||
|
|
||||||
void SetupDenormControl(const Profile& profile, const IR::Program& program, EmitContext& ctx,
|
void SetupDenormControl(const Profile& profile, const IR::Program& program, EmitContext& ctx,
|
||||||
Id main_func) {
|
Id main_func) {
|
||||||
if (!profile.support_float_controls) {
|
if (!profile.support_float_controls) {
|
||||||
|
@ -173,6 +210,25 @@ void SetupDenormControl(const Profile& profile, const IR::Program& program, Emit
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ctx) {
|
||||||
|
if (info.uses_sampled_1d) {
|
||||||
|
ctx.AddCapability(spv::Capability::Sampled1D);
|
||||||
|
}
|
||||||
|
if (info.uses_sparse_residency) {
|
||||||
|
ctx.AddCapability(spv::Capability::SparseResidency);
|
||||||
|
}
|
||||||
|
if (info.uses_demote_to_helper_invocation) {
|
||||||
|
ctx.AddExtension("SPV_EXT_demote_to_helper_invocation");
|
||||||
|
ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT);
|
||||||
|
}
|
||||||
|
if (!profile.support_vertex_instance_id && (info.loads_instance_id || info.loads_vertex_id)) {
|
||||||
|
ctx.AddExtension("SPV_KHR_shader_draw_parameters");
|
||||||
|
ctx.AddCapability(spv::Capability::DrawParameters);
|
||||||
|
}
|
||||||
|
// TODO: Track this usage
|
||||||
|
ctx.AddCapability(spv::Capability::ImageGatherExtended);
|
||||||
|
}
|
||||||
|
|
||||||
Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) {
|
Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) {
|
||||||
// Phi nodes can have forward declarations, if an argument is not defined provide a forward
|
// Phi nodes can have forward declarations, if an argument is not defined provide a forward
|
||||||
// declaration of it. Invoke will take care of giving it the right definition when it's
|
// declaration of it. Invoke will take care of giving it the right definition when it's
|
||||||
|
@ -202,53 +258,10 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) {
|
||||||
std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program,
|
std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program,
|
||||||
u32& binding) {
|
u32& binding) {
|
||||||
EmitContext ctx{profile, program, binding};
|
EmitContext ctx{profile, program, binding};
|
||||||
const Id void_function{ctx.TypeFunction(ctx.void_id)};
|
const Id main{DefineMain(ctx, program)};
|
||||||
const Id func{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)};
|
DefineEntryPoint(env, ctx, main);
|
||||||
for (IR::Block* const block : program.blocks) {
|
SetupDenormControl(profile, program, ctx, main);
|
||||||
ctx.AddLabel(block->Definition<Id>());
|
SetupCapabilities(profile, program.info, ctx);
|
||||||
for (IR::Inst& inst : block->Instructions()) {
|
|
||||||
EmitInst(ctx, &inst);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
ctx.OpFunctionEnd();
|
|
||||||
|
|
||||||
const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size());
|
|
||||||
spv::ExecutionModel execution_model{};
|
|
||||||
switch (env.ShaderStage()) {
|
|
||||||
case Shader::Stage::Compute: {
|
|
||||||
const std::array<u32, 3> workgroup_size{env.WorkgroupSize()};
|
|
||||||
execution_model = spv::ExecutionModel::GLCompute;
|
|
||||||
ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0],
|
|
||||||
workgroup_size[1], workgroup_size[2]);
|
|
||||||
break;
|
|
||||||
}
|
|
||||||
case Shader::Stage::VertexB:
|
|
||||||
execution_model = spv::ExecutionModel::Vertex;
|
|
||||||
break;
|
|
||||||
case Shader::Stage::Fragment:
|
|
||||||
execution_model = spv::ExecutionModel::Fragment;
|
|
||||||
ctx.AddExecutionMode(func, spv::ExecutionMode::OriginUpperLeft);
|
|
||||||
break;
|
|
||||||
default:
|
|
||||||
throw NotImplementedException("Stage {}", env.ShaderStage());
|
|
||||||
}
|
|
||||||
ctx.AddEntryPoint(execution_model, func, "main", interfaces);
|
|
||||||
|
|
||||||
SetupDenormControl(profile, program, ctx, func);
|
|
||||||
const Info& info{program.info};
|
|
||||||
if (info.uses_sampled_1d) {
|
|
||||||
ctx.AddCapability(spv::Capability::Sampled1D);
|
|
||||||
}
|
|
||||||
if (info.uses_sparse_residency) {
|
|
||||||
ctx.AddCapability(spv::Capability::SparseResidency);
|
|
||||||
}
|
|
||||||
if (info.uses_demote_to_helper_invocation) {
|
|
||||||
ctx.AddExtension("SPV_EXT_demote_to_helper_invocation");
|
|
||||||
ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT);
|
|
||||||
}
|
|
||||||
// TODO: Track this usage
|
|
||||||
ctx.AddCapability(spv::Capability::ImageGatherExtended);
|
|
||||||
|
|
||||||
return ctx.Assemble();
|
return ctx.Assemble();
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
@ -81,8 +81,8 @@ void EmitLoadStorageS8(EmitContext& ctx);
|
||||||
void EmitLoadStorageU16(EmitContext& ctx);
|
void EmitLoadStorageU16(EmitContext& ctx);
|
||||||
void EmitLoadStorageS16(EmitContext& ctx);
|
void EmitLoadStorageS16(EmitContext& ctx);
|
||||||
Id EmitLoadStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
Id EmitLoadStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||||
void EmitLoadStorage64(EmitContext& ctx);
|
Id EmitLoadStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||||
void EmitLoadStorage128(EmitContext& ctx);
|
Id EmitLoadStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||||
void EmitWriteStorageU8(EmitContext& ctx);
|
void EmitWriteStorageU8(EmitContext& ctx);
|
||||||
void EmitWriteStorageS8(EmitContext& ctx);
|
void EmitWriteStorageS8(EmitContext& ctx);
|
||||||
void EmitWriteStorageU16(EmitContext& ctx);
|
void EmitWriteStorageU16(EmitContext& ctx);
|
||||||
|
|
|
@ -19,6 +19,10 @@ Id InputAttrPointer(EmitContext& ctx, IR::Attribute attr) {
|
||||||
case IR::Attribute::PositionZ:
|
case IR::Attribute::PositionZ:
|
||||||
case IR::Attribute::PositionW:
|
case IR::Attribute::PositionW:
|
||||||
return ctx.OpAccessChain(ctx.input_f32, ctx.input_position, element_id());
|
return ctx.OpAccessChain(ctx.input_f32, ctx.input_position, element_id());
|
||||||
|
case IR::Attribute::InstanceId:
|
||||||
|
return ctx.OpLoad(ctx.U32[1], ctx.instance_id);
|
||||||
|
case IR::Attribute::VertexId:
|
||||||
|
return ctx.OpLoad(ctx.U32[1], ctx.vertex_id);
|
||||||
default:
|
default:
|
||||||
throw NotImplementedException("Read attribute {}", attr);
|
throw NotImplementedException("Read attribute {}", attr);
|
||||||
}
|
}
|
||||||
|
@ -125,6 +129,18 @@ Id EmitGetCbufU64(EmitContext& ctx, const IR::Value& binding, const IR::Value& o
|
||||||
}
|
}
|
||||||
|
|
||||||
Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr) {
|
Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr) {
|
||||||
|
if (!ctx.profile.support_vertex_instance_id) {
|
||||||
|
switch (attr) {
|
||||||
|
case IR::Attribute::InstanceId:
|
||||||
|
return ctx.OpISub(ctx.U32[1], ctx.OpLoad(ctx.U32[1], ctx.instance_index),
|
||||||
|
ctx.OpLoad(ctx.U32[1], ctx.base_instance));
|
||||||
|
case IR::Attribute::VertexId:
|
||||||
|
return ctx.OpISub(ctx.U32[1], ctx.OpLoad(ctx.U32[1], ctx.vertex_index),
|
||||||
|
ctx.OpLoad(ctx.U32[1], ctx.base_vertex));
|
||||||
|
default:
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
}
|
||||||
return ctx.OpLoad(ctx.F32[1], InputAttrPointer(ctx, attr));
|
return ctx.OpLoad(ctx.F32[1], InputAttrPointer(ctx, attr));
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
@ -7,8 +7,8 @@
|
||||||
#include "shader_recompiler/backend/spirv/emit_spirv.h"
|
#include "shader_recompiler/backend/spirv/emit_spirv.h"
|
||||||
|
|
||||||
namespace Shader::Backend::SPIRV {
|
namespace Shader::Backend::SPIRV {
|
||||||
|
namespace {
|
||||||
static Id StorageIndex(EmitContext& ctx, const IR::Value& offset, size_t element_size) {
|
Id StorageIndex(EmitContext& ctx, const IR::Value& offset, size_t element_size) {
|
||||||
if (offset.IsImmediate()) {
|
if (offset.IsImmediate()) {
|
||||||
const u32 imm_offset{static_cast<u32>(offset.U32() / element_size)};
|
const u32 imm_offset{static_cast<u32>(offset.U32() / element_size)};
|
||||||
return ctx.Constant(ctx.U32[1], imm_offset);
|
return ctx.Constant(ctx.U32[1], imm_offset);
|
||||||
|
@ -22,6 +22,32 @@ static Id StorageIndex(EmitContext& ctx, const IR::Value& offset, size_t element
|
||||||
return ctx.OpShiftRightLogical(ctx.U32[1], index, shift_id);
|
return ctx.OpShiftRightLogical(ctx.U32[1], index, shift_id);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
Id EmitLoadStorage(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||||
|
u32 num_components) {
|
||||||
|
// TODO: Support reinterpreting bindings, guaranteed to be aligned
|
||||||
|
if (!binding.IsImmediate()) {
|
||||||
|
throw NotImplementedException("Dynamic storage buffer indexing");
|
||||||
|
}
|
||||||
|
const Id ssbo{ctx.ssbos[binding.U32()]};
|
||||||
|
const Id base_index{StorageIndex(ctx, offset, sizeof(u32))};
|
||||||
|
std::array<Id, 4> components;
|
||||||
|
for (u32 element = 0; element < num_components; ++element) {
|
||||||
|
Id index{base_index};
|
||||||
|
if (element > 0) {
|
||||||
|
index = ctx.OpIAdd(ctx.U32[1], base_index, ctx.Constant(ctx.U32[1], element));
|
||||||
|
}
|
||||||
|
const Id pointer{ctx.OpAccessChain(ctx.storage_u32, ssbo, ctx.u32_zero_value, index)};
|
||||||
|
components[element] = ctx.OpLoad(ctx.U32[1], pointer);
|
||||||
|
}
|
||||||
|
if (num_components == 1) {
|
||||||
|
return components[0];
|
||||||
|
} else {
|
||||||
|
const std::span components_span(components.data(), num_components);
|
||||||
|
return ctx.OpCompositeConstruct(ctx.U32[num_components], components_span);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
} // Anonymous namespace
|
||||||
|
|
||||||
void EmitLoadGlobalU8(EmitContext&) {
|
void EmitLoadGlobalU8(EmitContext&) {
|
||||||
throw NotImplementedException("SPIR-V Instruction");
|
throw NotImplementedException("SPIR-V Instruction");
|
||||||
}
|
}
|
||||||
|
@ -95,21 +121,15 @@ void EmitLoadStorageS16(EmitContext&) {
|
||||||
}
|
}
|
||||||
|
|
||||||
Id EmitLoadStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) {
|
Id EmitLoadStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) {
|
||||||
if (!binding.IsImmediate()) {
|
return EmitLoadStorage(ctx, binding, offset, 1);
|
||||||
throw NotImplementedException("Dynamic storage buffer indexing");
|
|
||||||
}
|
|
||||||
const Id ssbo{ctx.ssbos[binding.U32()]};
|
|
||||||
const Id index{StorageIndex(ctx, offset, sizeof(u32))};
|
|
||||||
const Id pointer{ctx.OpAccessChain(ctx.storage_u32, ssbo, ctx.u32_zero_value, index)};
|
|
||||||
return ctx.OpLoad(ctx.U32[1], pointer);
|
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitLoadStorage64(EmitContext&) {
|
Id EmitLoadStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) {
|
||||||
throw NotImplementedException("SPIR-V Instruction");
|
return EmitLoadStorage(ctx, binding, offset, 2);
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitLoadStorage128(EmitContext&) {
|
Id EmitLoadStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) {
|
||||||
throw NotImplementedException("SPIR-V Instruction");
|
return EmitLoadStorage(ctx, binding, offset, 4);
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitWriteStorageU8(EmitContext&) {
|
void EmitWriteStorageU8(EmitContext&) {
|
||||||
|
|
|
@ -38,6 +38,12 @@ void GetAttribute(Info& info, IR::Attribute attribute) {
|
||||||
case IR::Attribute::PositionW:
|
case IR::Attribute::PositionW:
|
||||||
info.loads_position = true;
|
info.loads_position = true;
|
||||||
break;
|
break;
|
||||||
|
case IR::Attribute::InstanceId:
|
||||||
|
info.loads_instance_id = true;
|
||||||
|
break;
|
||||||
|
case IR::Attribute::VertexId:
|
||||||
|
info.loads_vertex_id = true;
|
||||||
|
break;
|
||||||
default:
|
default:
|
||||||
throw NotImplementedException("Get attribute {}", attribute);
|
throw NotImplementedException("Get attribute {}", attribute);
|
||||||
}
|
}
|
||||||
|
|
|
@ -8,6 +8,7 @@ namespace Shader {
|
||||||
|
|
||||||
struct Profile {
|
struct Profile {
|
||||||
bool unified_descriptor_binding{};
|
bool unified_descriptor_binding{};
|
||||||
|
bool support_vertex_instance_id{};
|
||||||
bool support_float_controls{};
|
bool support_float_controls{};
|
||||||
bool support_separate_denorm_behavior{};
|
bool support_separate_denorm_behavior{};
|
||||||
bool support_separate_rounding_mode{};
|
bool support_separate_rounding_mode{};
|
||||||
|
|
|
@ -59,6 +59,8 @@ struct Info {
|
||||||
|
|
||||||
std::array<bool, 32> loads_generics{};
|
std::array<bool, 32> loads_generics{};
|
||||||
bool loads_position{};
|
bool loads_position{};
|
||||||
|
bool loads_instance_id{};
|
||||||
|
bool loads_vertex_id{};
|
||||||
|
|
||||||
std::array<bool, 8> stores_frag_color{};
|
std::array<bool, 8> stores_frag_color{};
|
||||||
bool stores_frag_depth{};
|
bool stores_frag_depth{};
|
||||||
|
|
|
@ -230,6 +230,7 @@ PipelineCache::PipelineCache(RasterizerVulkan& rasterizer_, Tegra::GPU& gpu_,
|
||||||
const VkDriverIdKHR driver_id{device.GetDriverID()};
|
const VkDriverIdKHR driver_id{device.GetDriverID()};
|
||||||
profile = Shader::Profile{
|
profile = Shader::Profile{
|
||||||
.unified_descriptor_binding = true,
|
.unified_descriptor_binding = true,
|
||||||
|
.support_vertex_instance_id = false,
|
||||||
.support_float_controls = true,
|
.support_float_controls = true,
|
||||||
.support_separate_denorm_behavior = float_control.denormBehaviorIndependence ==
|
.support_separate_denorm_behavior = float_control.denormBehaviorIndependence ==
|
||||||
VK_SHADER_FLOAT_CONTROLS_INDEPENDENCE_ALL_KHR,
|
VK_SHADER_FLOAT_CONTROLS_INDEPENDENCE_ALL_KHR,
|
||||||
|
|
Reference in New Issue