shader: Add partial rasterizer integration
This commit is contained in:
ReinUsesLisp
ameerj
parent
72990df7ba
commit
260743f371
54 changed files with 1929 additions and 568 deletions
|
|
@ -62,18 +62,15 @@ void VectorTypes::Define(Sirit::Module& sirit_ctx, Id base_type, std::string_vie
|
|||
}
|
||||
}
|
||||
|
||||
EmitContext::EmitContext(const Profile& profile_, IR::Program& program)
|
||||
EmitContext::EmitContext(const Profile& profile_, IR::Program& program, u32& binding)
|
||||
: Sirit::Module(0x00010000), profile{profile_} {
|
||||
AddCapability(spv::Capability::Shader);
|
||||
DefineCommonTypes(program.info);
|
||||
DefineCommonConstants();
|
||||
DefineSpecialVariables(program.info);
|
||||
|
||||
u32 binding{};
|
||||
DefineInterfaces(program.info, program.stage);
|
||||
DefineConstantBuffers(program.info, binding);
|
||||
DefineStorageBuffers(program.info, binding);
|
||||
DefineTextures(program.info, binding);
|
||||
|
||||
DefineLabels(program);
|
||||
}
|
||||
|
||||
|
|
@ -96,6 +93,8 @@ Id EmitContext::Def(const IR::Value& value) {
|
|||
return Constant(F32[1], value.F32());
|
||||
case IR::Type::F64:
|
||||
return Constant(F64[1], value.F64());
|
||||
case IR::Type::Label:
|
||||
return value.Label()->Definition<Id>();
|
||||
default:
|
||||
throw NotImplementedException("Immediate type {}", value.Type());
|
||||
}
|
||||
|
|
@ -109,6 +108,9 @@ void EmitContext::DefineCommonTypes(const Info& info) {
|
|||
F32.Define(*this, TypeFloat(32), "f32");
|
||||
U32.Define(*this, TypeInt(32, false), "u32");
|
||||
|
||||
input_f32 = Name(TypePointer(spv::StorageClass::Input, F32[1]), "input_f32");
|
||||
output_f32 = Name(TypePointer(spv::StorageClass::Output, F32[1]), "output_f32");
|
||||
|
||||
if (info.uses_int8) {
|
||||
AddCapability(spv::Capability::Int8);
|
||||
U8 = Name(TypeInt(8, false), "u8");
|
||||
|
|
@ -139,15 +141,20 @@ void EmitContext::DefineCommonConstants() {
|
|||
u32_zero_value = Constant(U32[1], 0U);
|
||||
}
|
||||
|
||||
void EmitContext::DefineSpecialVariables(const Info& info) {
|
||||
const auto define{[this](Id type, spv::BuiltIn builtin, spv::StorageClass storage_class) {
|
||||
const Id pointer_type{TypePointer(storage_class, type)};
|
||||
const Id id{AddGlobalVariable(pointer_type, spv::StorageClass::Input)};
|
||||
Decorate(id, spv::Decoration::BuiltIn, builtin);
|
||||
return id;
|
||||
}};
|
||||
void EmitContext::DefineInterfaces(const Info& info, Stage stage) {
|
||||
const auto define{
|
||||
[this](Id type, std::optional<spv::BuiltIn> builtin, spv::StorageClass storage_class) {
|
||||
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);
|
||||
|
|
@ -155,6 +162,39 @@ void EmitContext::DefineSpecialVariables(const Info& info) {
|
|||
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) {
|
||||
|
|
|
|||
|
|
@ -46,7 +46,7 @@ struct UniformDefinitions {
|
|||
|
||||
class EmitContext final : public Sirit::Module {
|
||||
public:
|
||||
explicit EmitContext(const Profile& profile, IR::Program& program);
|
||||
explicit EmitContext(const Profile& profile, IR::Program& program, u32& binding);
|
||||
~EmitContext();
|
||||
|
||||
[[nodiscard]] Id Def(const IR::Value& value);
|
||||
|
|
@ -71,6 +71,9 @@ public:
|
|||
|
||||
UniformDefinitions uniform_types;
|
||||
|
||||
Id input_f32{};
|
||||
Id output_f32{};
|
||||
|
||||
Id storage_u32{};
|
||||
|
||||
std::array<UniformDefinitions, Info::MAX_CBUFS> cbufs{};
|
||||
|
|
@ -80,10 +83,21 @@ public:
|
|||
Id workgroup_id{};
|
||||
Id local_invocation_id{};
|
||||
|
||||
Id input_position{};
|
||||
std::array<Id, 32> input_generics{};
|
||||
|
||||
Id output_position{};
|
||||
std::array<Id, 32> output_generics{};
|
||||
|
||||
std::array<Id, 8> frag_color{};
|
||||
Id frag_depth {};
|
||||
|
||||
std::vector<Id> interfaces;
|
||||
|
||||
private:
|
||||
void DefineCommonTypes(const Info& info);
|
||||
void DefineCommonConstants();
|
||||
void DefineSpecialVariables(const Info& info);
|
||||
void DefineInterfaces(const Info& info, Stage stage);
|
||||
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);
|
||||
|
|
|
|||
|
|
@ -54,6 +54,8 @@ ArgType Arg(EmitContext& ctx, const IR::Value& arg) {
|
|||
return arg.U32();
|
||||
} else if constexpr (std::is_same_v<ArgType, IR::Block*>) {
|
||||
return arg.Label();
|
||||
} else if constexpr (std::is_same_v<ArgType, IR::Attribute>) {
|
||||
return arg.Attribute();
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -197,8 +199,9 @@ Id PhiArgDef(EmitContext& ctx, IR::Inst* inst, size_t index) {
|
|||
}
|
||||
} // Anonymous namespace
|
||||
|
||||
std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program) {
|
||||
EmitContext ctx{profile, program};
|
||||
std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program& program,
|
||||
u32& binding) {
|
||||
EmitContext ctx{profile, program, binding};
|
||||
const Id void_function{ctx.TypeFunction(ctx.void_id)};
|
||||
const Id func{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)};
|
||||
for (IR::Block* const block : program.blocks) {
|
||||
|
|
@ -208,28 +211,41 @@ std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env, IR::Program
|
|||
}
|
||||
}
|
||||
ctx.OpFunctionEnd();
|
||||
boost::container::small_vector<Id, 32> interfaces;
|
||||
const Info& info{program.info};
|
||||
if (info.uses_workgroup_id) {
|
||||
interfaces.push_back(ctx.workgroup_id);
|
||||
}
|
||||
if (info.uses_local_invocation_id) {
|
||||
interfaces.push_back(ctx.local_invocation_id);
|
||||
}
|
||||
const std::span interfaces_span(interfaces.data(), interfaces.size());
|
||||
ctx.AddEntryPoint(spv::ExecutionModel::GLCompute, func, "main", interfaces_span);
|
||||
|
||||
const std::array<u32, 3> workgroup_size{env.WorkgroupSize()};
|
||||
ctx.AddExecutionMode(func, spv::ExecutionMode::LocalSize, workgroup_size[0], workgroup_size[1],
|
||||
workgroup_size[2]);
|
||||
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);
|
||||
|
||||
|
|
|
|||
|
|
@ -16,18 +16,18 @@
|
|||
namespace Shader::Backend::SPIRV {
|
||||
|
||||
[[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, Environment& env,
|
||||
IR::Program& program);
|
||||
IR::Program& program, u32& binding);
|
||||
|
||||
// Microinstruction emitters
|
||||
Id EmitPhi(EmitContext& ctx, IR::Inst* inst);
|
||||
void EmitVoid(EmitContext& ctx);
|
||||
Id EmitIdentity(EmitContext& ctx, const IR::Value& value);
|
||||
void EmitBranch(EmitContext& ctx, IR::Block* label);
|
||||
void EmitBranchConditional(EmitContext& ctx, Id condition, IR::Block* true_label,
|
||||
IR::Block* false_label);
|
||||
void EmitLoopMerge(EmitContext& ctx, IR::Block* merge_label, IR::Block* continue_label);
|
||||
void EmitSelectionMerge(EmitContext& ctx, IR::Block* merge_label);
|
||||
void EmitBranch(EmitContext& ctx, Id label);
|
||||
void EmitBranchConditional(EmitContext& ctx, Id condition, Id true_label, Id false_label);
|
||||
void EmitLoopMerge(EmitContext& ctx, Id merge_label, Id continue_label);
|
||||
void EmitSelectionMerge(EmitContext& ctx, Id merge_label);
|
||||
void EmitReturn(EmitContext& ctx);
|
||||
void EmitDemoteToHelperInvocation(EmitContext& ctx, Id continue_label);
|
||||
void EmitGetRegister(EmitContext& ctx);
|
||||
void EmitSetRegister(EmitContext& ctx);
|
||||
void EmitGetPred(EmitContext& ctx);
|
||||
|
|
@ -41,10 +41,12 @@ Id EmitGetCbufS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& o
|
|||
Id EmitGetCbufU32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||
Id EmitGetCbufF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||
Id EmitGetCbufU64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||
void EmitGetAttribute(EmitContext& ctx);
|
||||
void EmitSetAttribute(EmitContext& ctx);
|
||||
Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr);
|
||||
void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, Id value);
|
||||
void EmitGetAttributeIndexed(EmitContext& ctx);
|
||||
void EmitSetAttributeIndexed(EmitContext& ctx);
|
||||
void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, Id value);
|
||||
void EmitSetFragDepth(EmitContext& ctx, Id value);
|
||||
void EmitGetZFlag(EmitContext& ctx);
|
||||
void EmitGetSFlag(EmitContext& ctx);
|
||||
void EmitGetCFlag(EmitContext& ctx);
|
||||
|
|
|
|||
|
|
@ -5,6 +5,43 @@
|
|||
#include "shader_recompiler/backend/spirv/emit_spirv.h"
|
||||
|
||||
namespace Shader::Backend::SPIRV {
|
||||
namespace {
|
||||
Id InputAttrPointer(EmitContext& ctx, IR::Attribute attr) {
|
||||
const u32 element{static_cast<u32>(attr) % 4};
|
||||
const auto element_id{[&] { return ctx.Constant(ctx.U32[1], element); }};
|
||||
if (IR::IsGeneric(attr)) {
|
||||
const u32 index{IR::GenericAttributeIndex(attr)};
|
||||
return ctx.OpAccessChain(ctx.input_f32, ctx.input_generics.at(index), element_id());
|
||||
}
|
||||
switch (attr) {
|
||||
case IR::Attribute::PositionX:
|
||||
case IR::Attribute::PositionY:
|
||||
case IR::Attribute::PositionZ:
|
||||
case IR::Attribute::PositionW:
|
||||
return ctx.OpAccessChain(ctx.input_f32, ctx.input_position, element_id());
|
||||
default:
|
||||
throw NotImplementedException("Read attribute {}", attr);
|
||||
}
|
||||
}
|
||||
|
||||
Id OutputAttrPointer(EmitContext& ctx, IR::Attribute attr) {
|
||||
const u32 element{static_cast<u32>(attr) % 4};
|
||||
const auto element_id{[&] { return ctx.Constant(ctx.U32[1], element); }};
|
||||
if (IR::IsGeneric(attr)) {
|
||||
const u32 index{IR::GenericAttributeIndex(attr)};
|
||||
return ctx.OpAccessChain(ctx.output_f32, ctx.output_generics.at(index), element_id());
|
||||
}
|
||||
switch (attr) {
|
||||
case IR::Attribute::PositionX:
|
||||
case IR::Attribute::PositionY:
|
||||
case IR::Attribute::PositionZ:
|
||||
case IR::Attribute::PositionW:
|
||||
return ctx.OpAccessChain(ctx.output_f32, ctx.output_position, element_id());
|
||||
default:
|
||||
throw NotImplementedException("Read attribute {}", attr);
|
||||
}
|
||||
}
|
||||
} // Anonymous namespace
|
||||
|
||||
void EmitGetRegister(EmitContext&) {
|
||||
throw NotImplementedException("SPIR-V Instruction");
|
||||
|
|
@ -87,12 +124,12 @@ Id EmitGetCbufU64(EmitContext& ctx, const IR::Value& binding, const IR::Value& o
|
|||
return GetCbuf(ctx, ctx.U64, &UniformDefinitions::U64, sizeof(u64), binding, offset);
|
||||
}
|
||||
|
||||
void EmitGetAttribute(EmitContext&) {
|
||||
throw NotImplementedException("SPIR-V Instruction");
|
||||
Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr) {
|
||||
return ctx.OpLoad(ctx.F32[1], InputAttrPointer(ctx, attr));
|
||||
}
|
||||
|
||||
void EmitSetAttribute(EmitContext&) {
|
||||
throw NotImplementedException("SPIR-V Instruction");
|
||||
void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, Id value) {
|
||||
ctx.OpStore(OutputAttrPointer(ctx, attr), value);
|
||||
}
|
||||
|
||||
void EmitGetAttributeIndexed(EmitContext&) {
|
||||
|
|
@ -103,6 +140,16 @@ void EmitSetAttributeIndexed(EmitContext&) {
|
|||
throw NotImplementedException("SPIR-V Instruction");
|
||||
}
|
||||
|
||||
void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, Id value) {
|
||||
const Id component_id{ctx.Constant(ctx.U32[1], component)};
|
||||
const Id pointer{ctx.OpAccessChain(ctx.output_f32, ctx.frag_color.at(index), component_id)};
|
||||
ctx.OpStore(pointer, value);
|
||||
}
|
||||
|
||||
void EmitSetFragDepth(EmitContext& ctx, Id value) {
|
||||
ctx.OpStore(ctx.frag_depth, value);
|
||||
}
|
||||
|
||||
void EmitGetZFlag(EmitContext&) {
|
||||
throw NotImplementedException("SPIR-V Instruction");
|
||||
}
|
||||
|
|
|
|||
|
|
@ -6,26 +6,29 @@
|
|||
|
||||
namespace Shader::Backend::SPIRV {
|
||||
|
||||
void EmitBranch(EmitContext& ctx, IR::Block* label) {
|
||||
ctx.OpBranch(label->Definition<Id>());
|
||||
void EmitBranch(EmitContext& ctx, Id label) {
|
||||
ctx.OpBranch(label);
|
||||
}
|
||||
|
||||
void EmitBranchConditional(EmitContext& ctx, Id condition, IR::Block* true_label,
|
||||
IR::Block* false_label) {
|
||||
ctx.OpBranchConditional(condition, true_label->Definition<Id>(), false_label->Definition<Id>());
|
||||
void EmitBranchConditional(EmitContext& ctx, Id condition, Id true_label, Id false_label) {
|
||||
ctx.OpBranchConditional(condition, true_label, false_label);
|
||||
}
|
||||
|
||||
void EmitLoopMerge(EmitContext& ctx, IR::Block* merge_label, IR::Block* continue_label) {
|
||||
ctx.OpLoopMerge(merge_label->Definition<Id>(), continue_label->Definition<Id>(),
|
||||
spv::LoopControlMask::MaskNone);
|
||||
void EmitLoopMerge(EmitContext& ctx, Id merge_label, Id continue_label) {
|
||||
ctx.OpLoopMerge(merge_label, continue_label, spv::LoopControlMask::MaskNone);
|
||||
}
|
||||
|
||||
void EmitSelectionMerge(EmitContext& ctx, IR::Block* merge_label) {
|
||||
ctx.OpSelectionMerge(merge_label->Definition<Id>(), spv::SelectionControlMask::MaskNone);
|
||||
void EmitSelectionMerge(EmitContext& ctx, Id merge_label) {
|
||||
ctx.OpSelectionMerge(merge_label, spv::SelectionControlMask::MaskNone);
|
||||
}
|
||||
|
||||
void EmitReturn(EmitContext& ctx) {
|
||||
ctx.OpReturn();
|
||||
}
|
||||
|
||||
void EmitDemoteToHelperInvocation(EmitContext& ctx, Id continue_label) {
|
||||
ctx.OpDemoteToHelperInvocationEXT();
|
||||
ctx.OpBranch(continue_label);
|
||||
}
|
||||
|
||||
} // namespace Shader::Backend::SPIRV
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue