recompiler: proper VS inputs initialization

This commit is contained in:
psucien 2024-07-12 18:12:06 +02:00
parent a75851f7e2
commit 1b94f07a6a
6 changed files with 32 additions and 8 deletions

View File

@ -131,6 +131,13 @@ Id EmitReadConstBufferU32(EmitContext& ctx, u32 handle, Id index) {
return ctx.OpBitcast(ctx.U32[1], EmitReadConstBuffer(ctx, handle, index));
}
Id EmitReadStepRate(EmitContext& ctx, int rate_idx) {
return ctx.OpLoad(
ctx.U32[1], ctx.OpAccessChain(ctx.TypePointer(spv::StorageClass::PushConstant, ctx.U32[1]),
ctx.instance_step_rates,
rate_idx == 0 ? ctx.u32_zero_value : ctx.u32_one_value));
}
Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr, u32 comp) {
if (IR::IsParam(attr)) {
const u32 index{u32(attr) - u32(IR::Attribute::Param0)};
@ -149,11 +156,7 @@ Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr, u32 comp) {
return ctx.OpLoad(param.component_type, param.id);
}
} else {
const auto rate_idx = param.id.value == 0 ? ctx.u32_zero_value : ctx.u32_one_value;
const auto step_rate = ctx.OpLoad(
ctx.U32[1],
ctx.OpAccessChain(ctx.TypePointer(spv::StorageClass::PushConstant, ctx.U32[1]),
ctx.instance_step_rates, rate_idx));
const auto step_rate = EmitReadStepRate(ctx, param.id.value);
const auto offset = ctx.OpIAdd(
ctx.U32[1],
ctx.OpIMul(
@ -182,6 +185,12 @@ Id EmitGetAttributeU32(EmitContext& ctx, IR::Attribute attr, u32 comp) {
switch (attr) {
case IR::Attribute::VertexId:
return ctx.OpLoad(ctx.U32[1], ctx.vertex_index);
case IR::Attribute::InstanceId:
return ctx.OpLoad(ctx.U32[1], ctx.instance_id);
case IR::Attribute::InstanceId0:
return EmitReadStepRate(ctx, 0);
case IR::Attribute::InstanceId1:
return EmitReadStepRate(ctx, 1);
case IR::Attribute::WorkgroupId:
return ctx.OpCompositeExtract(ctx.U32[1], ctx.OpLoad(ctx.U32[3], ctx.workgroup_id), comp);
case IR::Attribute::LocalInvocationId:

View File

@ -35,10 +35,20 @@ void Translator::EmitPrologue() {
IR::VectorReg dst_vreg = IR::VectorReg::V0;
switch (info.stage) {
case Stage::Vertex:
// https://github.com/chaotic-cx/mesa-mirror/blob/72326e15/src/amd/vulkan/radv_shader_args.c#L146C1-L146C23
// v0: vertex ID, always present
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::VertexId));
// v1: instance ID, step rate 0
if (info.num_input_vgprs > 0) {
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::InstanceId0));
}
// v2: instance ID, step rate 1
if (info.num_input_vgprs > 1) {
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::InstanceId1));
}
// v3: instance ID, plain
if (info.num_input_vgprs > 2) {
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::InstanceId));
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::PrimitiveId));
}
break;
case Stage::Fragment:
// https://github.com/chaotic-cx/mesa-mirror/blob/72326e15/src/amd/vulkan/radv_shader_args.c#L258

View File

@ -72,6 +72,8 @@ enum class Attribute : u64 {
LocalInvocationId = 75,
LocalInvocationIndex = 76,
FragCoord = 77,
InstanceId0 = 78, // step rate 0
InstanceId1 = 79, // step rate 1
Max,
};

View File

@ -163,6 +163,7 @@ struct Info {
std::array<u32, 3> workgroup_size{};
u32 num_user_data;
u32 num_input_vgprs;
std::span<const u32> user_data;
Stage stage;

View File

@ -80,6 +80,7 @@ struct Liverpool {
union {
BitField<0, 6, u64> num_vgprs;
BitField<6, 4, u64> num_sgprs;
BitField<24, 2, u64> vgpr_comp_cnt; // SPI provided per-thread inputs
BitField<33, 5, u64> num_user_regs;
} settings;
UserData user_data;

View File

@ -72,6 +72,7 @@ Shader::Info MakeShaderInfo(Shader::Stage stage, std::span<const u32, 16> user_d
switch (stage) {
case Shader::Stage::Vertex: {
info.num_user_data = regs.vs_program.settings.num_user_regs;
info.num_input_vgprs = regs.vs_program.settings.vgpr_comp_cnt;
BuildVsOutputs(info, regs.vs_output_control);
break;
}