glsl: Cleanup and address feedback
This commit is contained in:
parent
5355568a2d
commit
34fdb6471d
|
@ -77,7 +77,6 @@ std::string_view SamplerType(TextureType type, bool is_depth) {
|
||||||
case TextureType::ColorArrayCube:
|
case TextureType::ColorArrayCube:
|
||||||
return "samplerCubeArrayShadow";
|
return "samplerCubeArrayShadow";
|
||||||
default:
|
default:
|
||||||
fmt::print("Texture type: {}", type);
|
|
||||||
throw NotImplementedException("Texture type: {}", type);
|
throw NotImplementedException("Texture type: {}", type);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
@ -191,29 +190,27 @@ void SetupOutPerVertex(EmitContext& ctx, std::string& header) {
|
||||||
if (!StoresPerVertexAttributes(ctx.stage)) {
|
if (!StoresPerVertexAttributes(ctx.stage)) {
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
header += "out gl_PerVertex{";
|
header += "out gl_PerVertex{vec4 gl_Position;";
|
||||||
header += "vec4 gl_Position;";
|
|
||||||
if (ctx.info.stores_point_size) {
|
if (ctx.info.stores_point_size) {
|
||||||
header += "float gl_PointSize;";
|
header += "float gl_PointSize;";
|
||||||
}
|
}
|
||||||
if (ctx.info.stores_clip_distance) {
|
if (ctx.info.stores_clip_distance) {
|
||||||
header += "float gl_ClipDistance[];";
|
header += "float gl_ClipDistance[];";
|
||||||
}
|
}
|
||||||
if (ctx.info.stores_viewport_index && ctx.supports_viewport_layer &&
|
if (ctx.info.stores_viewport_index && ctx.profile.support_gl_vertex_viewport_layer &&
|
||||||
ctx.stage != Stage::Geometry) {
|
ctx.stage != Stage::Geometry) {
|
||||||
header += "int gl_ViewportIndex;";
|
header += "int gl_ViewportIndex;";
|
||||||
}
|
}
|
||||||
header += "};\n";
|
header += "};";
|
||||||
if (ctx.info.stores_viewport_index && ctx.stage == Stage::Geometry) {
|
if (ctx.info.stores_viewport_index && ctx.stage == Stage::Geometry) {
|
||||||
header += "out int gl_ViewportIndex;";
|
header += "out int gl_ViewportIndex;";
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
} // namespace
|
} // Anonymous namespace
|
||||||
|
|
||||||
EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_,
|
EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_,
|
||||||
const RuntimeInfo& runtime_info_)
|
const RuntimeInfo& runtime_info_)
|
||||||
: info{program.info}, profile{profile_}, runtime_info{runtime_info_} {
|
: info{program.info}, profile{profile_}, runtime_info{runtime_info_} {
|
||||||
supports_viewport_layer = profile.support_gl_vertex_viewport_layer;
|
|
||||||
SetupExtensions(header);
|
SetupExtensions(header);
|
||||||
stage = program.stage;
|
stage = program.stage;
|
||||||
switch (program.stage) {
|
switch (program.stage) {
|
||||||
|
@ -222,18 +219,18 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile
|
||||||
stage_name = "vs";
|
stage_name = "vs";
|
||||||
break;
|
break;
|
||||||
case Stage::TessellationControl:
|
case Stage::TessellationControl:
|
||||||
stage_name = "tsc";
|
stage_name = "tcs";
|
||||||
header += fmt::format("layout(vertices={})out;\n", program.invocations);
|
header += fmt::format("layout(vertices={})out;", program.invocations);
|
||||||
break;
|
break;
|
||||||
case Stage::TessellationEval:
|
case Stage::TessellationEval:
|
||||||
stage_name = "tse";
|
stage_name = "tes";
|
||||||
header += fmt::format("layout({},{},{})in;\n", GetTessMode(runtime_info.tess_primitive),
|
header += fmt::format("layout({},{},{})in;", GetTessMode(runtime_info.tess_primitive),
|
||||||
GetTessSpacing(runtime_info.tess_spacing),
|
GetTessSpacing(runtime_info.tess_spacing),
|
||||||
runtime_info.tess_clockwise ? "cw" : "ccw");
|
runtime_info.tess_clockwise ? "cw" : "ccw");
|
||||||
break;
|
break;
|
||||||
case Stage::Geometry:
|
case Stage::Geometry:
|
||||||
stage_name = "gs";
|
stage_name = "gs";
|
||||||
header += fmt::format("layout({})in;layout({},max_vertices={})out;\n",
|
header += fmt::format("layout({})in;layout({},max_vertices={})out;",
|
||||||
InputPrimitive(runtime_info.input_topology),
|
InputPrimitive(runtime_info.input_topology),
|
||||||
OutputPrimitive(program.output_topology), program.output_vertices);
|
OutputPrimitive(program.output_topology), program.output_vertices);
|
||||||
break;
|
break;
|
||||||
|
@ -242,7 +239,7 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile
|
||||||
break;
|
break;
|
||||||
case Stage::Compute:
|
case Stage::Compute:
|
||||||
stage_name = "cs";
|
stage_name = "cs";
|
||||||
header += fmt::format("layout(local_size_x={},local_size_y={},local_size_z={}) in;\n",
|
header += fmt::format("layout(local_size_x={},local_size_y={},local_size_z={}) in;",
|
||||||
program.workgroup_size[0], program.workgroup_size[1],
|
program.workgroup_size[0], program.workgroup_size[1],
|
||||||
program.workgroup_size[2]);
|
program.workgroup_size[2]);
|
||||||
break;
|
break;
|
||||||
|
@ -260,11 +257,8 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile
|
||||||
if (!info.uses_patches[index]) {
|
if (!info.uses_patches[index]) {
|
||||||
continue;
|
continue;
|
||||||
}
|
}
|
||||||
if (stage == Stage::TessellationControl) {
|
const auto qualifier{stage == Stage::TessellationControl ? "out" : "in"};
|
||||||
header += fmt::format("layout(location={})patch out vec4 patch{};", index, index);
|
header += fmt::format("layout(location={})patch {} vec4 patch{};", index, qualifier, index);
|
||||||
} else {
|
|
||||||
header += fmt::format("layout(location={})patch in vec4 patch{};", index, index);
|
|
||||||
}
|
|
||||||
}
|
}
|
||||||
for (size_t index = 0; index < info.stores_frag_color.size(); ++index) {
|
for (size_t index = 0; index < info.stores_frag_color.size(); ++index) {
|
||||||
if (!info.stores_frag_color[index]) {
|
if (!info.stores_frag_color[index]) {
|
||||||
|
@ -278,18 +272,18 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile
|
||||||
DefineGenericOutput(index, program.invocations);
|
DefineGenericOutput(index, program.invocations);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
header += "\n";
|
|
||||||
DefineConstantBuffers(bindings);
|
DefineConstantBuffers(bindings);
|
||||||
DefineStorageBuffers(bindings);
|
DefineStorageBuffers(bindings);
|
||||||
SetupImages(bindings);
|
SetupImages(bindings);
|
||||||
|
SetupTextures(bindings);
|
||||||
DefineHelperFunctions();
|
DefineHelperFunctions();
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitContext::SetupExtensions(std::string&) {
|
void EmitContext::SetupExtensions(std::string&) {
|
||||||
// TODO: track this usage
|
// TODO: track this usage
|
||||||
header += "#extension GL_ARB_sparse_texture2 : enable\n";
|
header += "#extension GL_ARB_sparse_texture2 : enable\n"
|
||||||
header += "#extension GL_EXT_texture_shadow_lod : enable\n";
|
"#extension GL_EXT_texture_shadow_lod : enable\n"
|
||||||
header += "#extension GL_EXT_shader_image_load_formatted : enable\n";
|
"#extension GL_EXT_shader_image_load_formatted : enable\n";
|
||||||
if (info.uses_int64) {
|
if (info.uses_int64) {
|
||||||
header += "#extension GL_ARB_gpu_shader_int64 : enable\n";
|
header += "#extension GL_ARB_gpu_shader_int64 : enable\n";
|
||||||
}
|
}
|
||||||
|
@ -312,13 +306,14 @@ void EmitContext::SetupExtensions(std::string&) {
|
||||||
}
|
}
|
||||||
if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask || info.uses_subgroup_vote ||
|
if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask || info.uses_subgroup_vote ||
|
||||||
info.uses_subgroup_shuffles || info.uses_fswzadd) {
|
info.uses_subgroup_shuffles || info.uses_fswzadd) {
|
||||||
header += "#extension GL_ARB_shader_ballot : enable\n";
|
header += "#extension GL_ARB_shader_ballot : enable\n"
|
||||||
header += "#extension GL_ARB_shader_group_vote : enable\n";
|
"#extension GL_ARB_shader_group_vote : enable\n";
|
||||||
if (!info.uses_int64) {
|
if (!info.uses_int64) {
|
||||||
header += "#extension GL_ARB_gpu_shader_int64 : enable\n";
|
header += "#extension GL_ARB_gpu_shader_int64 : enable\n";
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
if (info.stores_viewport_index && supports_viewport_layer && stage != Stage::Geometry) {
|
if (info.stores_viewport_index && profile.support_gl_vertex_viewport_layer &&
|
||||||
|
stage != Stage::Geometry) {
|
||||||
header += "#extension GL_ARB_shader_viewport_layer_array : enable\n";
|
header += "#extension GL_ARB_shader_viewport_layer_array : enable\n";
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
@ -386,46 +381,45 @@ void EmitContext::DefineGenericOutput(size_t index, u32 invocations) {
|
||||||
std::fill_n(output_generics[index].begin() + element, num_components, element_info);
|
std::fill_n(output_generics[index].begin() + element, num_components, element_info);
|
||||||
element += num_components;
|
element += num_components;
|
||||||
}
|
}
|
||||||
header += "\n";
|
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitContext::DefineHelperFunctions() {
|
void EmitContext::DefineHelperFunctions() {
|
||||||
header += "\n#define ftoi floatBitsToInt\n#define ftou floatBitsToUint\n"
|
header += "\n#define ftoi floatBitsToInt\n#define ftou floatBitsToUint\n"
|
||||||
"#define itof intBitsToFloat\n#define utof uintBitsToFloat\n";
|
"#define itof intBitsToFloat\n#define utof uintBitsToFloat\n";
|
||||||
if (info.uses_global_increment || info.uses_shared_increment) {
|
if (info.uses_global_increment || info.uses_shared_increment) {
|
||||||
header += "uint CasIncrement(uint op_a,uint op_b){return(op_a>=op_b)?0u:(op_a+1u);}\n";
|
header += "uint CasIncrement(uint op_a,uint op_b){return op_a>=op_b?0u:(op_a+1u);}";
|
||||||
}
|
}
|
||||||
if (info.uses_global_decrement || info.uses_shared_decrement) {
|
if (info.uses_global_decrement || info.uses_shared_decrement) {
|
||||||
header += "uint CasDecrement(uint op_a,uint "
|
header += "uint CasDecrement(uint op_a,uint "
|
||||||
"op_b){return(op_a==0||op_a>op_b)?op_b:(op_a-1u);}\n";
|
"op_b){return op_a==0||op_a>op_b?op_b:(op_a-1u);}";
|
||||||
}
|
}
|
||||||
if (info.uses_atomic_f32_add) {
|
if (info.uses_atomic_f32_add) {
|
||||||
header += "uint CasFloatAdd(uint op_a,float op_b){return "
|
header += "uint CasFloatAdd(uint op_a,float op_b){return "
|
||||||
"ftou(utof(op_a)+op_b);}\n";
|
"ftou(utof(op_a)+op_b);}";
|
||||||
}
|
}
|
||||||
if (info.uses_atomic_f32x2_add) {
|
if (info.uses_atomic_f32x2_add) {
|
||||||
header += "uint CasFloatAdd32x2(uint op_a,vec2 op_b){return "
|
header += "uint CasFloatAdd32x2(uint op_a,vec2 op_b){return "
|
||||||
"packHalf2x16(unpackHalf2x16(op_a)+op_b);}\n";
|
"packHalf2x16(unpackHalf2x16(op_a)+op_b);}";
|
||||||
}
|
}
|
||||||
if (info.uses_atomic_f32x2_min) {
|
if (info.uses_atomic_f32x2_min) {
|
||||||
header += "uint CasFloatMin32x2(uint op_a,vec2 op_b){return "
|
header += "uint CasFloatMin32x2(uint op_a,vec2 op_b){return "
|
||||||
"packHalf2x16(min(unpackHalf2x16(op_a),op_b));}\n";
|
"packHalf2x16(min(unpackHalf2x16(op_a),op_b));}";
|
||||||
}
|
}
|
||||||
if (info.uses_atomic_f32x2_max) {
|
if (info.uses_atomic_f32x2_max) {
|
||||||
header += "uint CasFloatMax32x2(uint op_a,vec2 op_b){return "
|
header += "uint CasFloatMax32x2(uint op_a,vec2 op_b){return "
|
||||||
"packHalf2x16(max(unpackHalf2x16(op_a),op_b));}\n";
|
"packHalf2x16(max(unpackHalf2x16(op_a),op_b));}";
|
||||||
}
|
}
|
||||||
if (info.uses_atomic_f16x2_add) {
|
if (info.uses_atomic_f16x2_add) {
|
||||||
header += "uint CasFloatAdd16x2(uint op_a,f16vec2 op_b){return "
|
header += "uint CasFloatAdd16x2(uint op_a,f16vec2 op_b){return "
|
||||||
"packFloat2x16(unpackFloat2x16(op_a)+op_b);}\n";
|
"packFloat2x16(unpackFloat2x16(op_a)+op_b);}";
|
||||||
}
|
}
|
||||||
if (info.uses_atomic_f16x2_min) {
|
if (info.uses_atomic_f16x2_min) {
|
||||||
header += "uint CasFloatMin16x2(uint op_a,f16vec2 op_b){return "
|
header += "uint CasFloatMin16x2(uint op_a,f16vec2 op_b){return "
|
||||||
"packFloat2x16(min(unpackFloat2x16(op_a),op_b));}\n";
|
"packFloat2x16(min(unpackFloat2x16(op_a),op_b));}";
|
||||||
}
|
}
|
||||||
if (info.uses_atomic_f16x2_max) {
|
if (info.uses_atomic_f16x2_max) {
|
||||||
header += "uint CasFloatMax16x2(uint op_a,f16vec2 op_b){return "
|
header += "uint CasFloatMax16x2(uint op_a,f16vec2 op_b){return "
|
||||||
"packFloat2x16(max(unpackFloat2x16(op_a),op_b));}\n";
|
"packFloat2x16(max(unpackFloat2x16(op_a),op_b));}";
|
||||||
}
|
}
|
||||||
if (info.uses_atomic_s32_min) {
|
if (info.uses_atomic_s32_min) {
|
||||||
header += "uint CasMinS32(uint op_a,uint op_b){return uint(min(int(op_a),int(op_b)));}";
|
header += "uint CasMinS32(uint op_a,uint op_b){return uint(min(int(op_a),int(op_b)));}";
|
||||||
|
@ -534,6 +528,9 @@ void EmitContext::SetupImages(Bindings& bindings) {
|
||||||
}
|
}
|
||||||
bindings.image += desc.count;
|
bindings.image += desc.count;
|
||||||
}
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
void EmitContext::SetupTextures(Bindings& bindings) {
|
||||||
texture_buffer_bindings.reserve(info.texture_buffer_descriptors.size());
|
texture_buffer_bindings.reserve(info.texture_buffer_descriptors.size());
|
||||||
for (const auto& desc : info.texture_buffer_descriptors) {
|
for (const auto& desc : info.texture_buffer_descriptors) {
|
||||||
texture_buffer_bindings.push_back(bindings.texture);
|
texture_buffer_bindings.push_back(bindings.texture);
|
||||||
|
|
|
@ -31,7 +31,7 @@ struct Program;
|
||||||
namespace Shader::Backend::GLSL {
|
namespace Shader::Backend::GLSL {
|
||||||
|
|
||||||
struct GenericElementInfo {
|
struct GenericElementInfo {
|
||||||
std::string name{};
|
std::string name;
|
||||||
u32 first_element{};
|
u32 first_element{};
|
||||||
u32 num_components{};
|
u32 num_components{};
|
||||||
};
|
};
|
||||||
|
@ -159,7 +159,6 @@ public:
|
||||||
|
|
||||||
bool uses_y_direction{};
|
bool uses_y_direction{};
|
||||||
bool uses_cc_carry{};
|
bool uses_cc_carry{};
|
||||||
bool supports_viewport_layer{};
|
|
||||||
|
|
||||||
private:
|
private:
|
||||||
void SetupExtensions(std::string& header);
|
void SetupExtensions(std::string& header);
|
||||||
|
@ -169,6 +168,7 @@ private:
|
||||||
void DefineHelperFunctions();
|
void DefineHelperFunctions();
|
||||||
std::string DefineGlobalMemoryFunctions();
|
std::string DefineGlobalMemoryFunctions();
|
||||||
void SetupImages(Bindings& bindings);
|
void SetupImages(Bindings& bindings);
|
||||||
|
void SetupTextures(Bindings& bindings);
|
||||||
};
|
};
|
||||||
|
|
||||||
} // namespace Shader::Backend::GLSL
|
} // namespace Shader::Backend::GLSL
|
||||||
|
|
|
@ -83,7 +83,6 @@ void Invoke(EmitContext& ctx, IR::Inst* inst) {
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitInst(EmitContext& ctx, IR::Inst* inst) {
|
void EmitInst(EmitContext& ctx, IR::Inst* inst) {
|
||||||
// ctx.Add("/* $ {} $ */", inst->GetOpcode());
|
|
||||||
switch (inst->GetOpcode()) {
|
switch (inst->GetOpcode()) {
|
||||||
#define OPCODE(name, result_type, ...) \
|
#define OPCODE(name, result_type, ...) \
|
||||||
case IR::Opcode::name: \
|
case IR::Opcode::name: \
|
||||||
|
@ -156,12 +155,10 @@ void EmitCode(EmitContext& ctx, const IR::Program& program) {
|
||||||
ctx.Add("for(;;){{");
|
ctx.Add("for(;;){{");
|
||||||
break;
|
break;
|
||||||
case IR::AbstractSyntaxNode::Type::Repeat:
|
case IR::AbstractSyntaxNode::Type::Repeat:
|
||||||
ctx.Add("if({}){{", ctx.var_alloc.Consume(node.data.repeat.cond));
|
ctx.Add("if({}){{continue;}}else{{break;}}}}",
|
||||||
ctx.Add("continue;\n}}else{{");
|
ctx.var_alloc.Consume(node.data.repeat.cond));
|
||||||
ctx.Add("break;\n}}\n}}");
|
|
||||||
break;
|
break;
|
||||||
default:
|
default:
|
||||||
fmt::print("{}", node.type);
|
|
||||||
throw NotImplementedException("AbstractSyntaxNode::Type {}", node.type);
|
throw NotImplementedException("AbstractSyntaxNode::Type {}", node.type);
|
||||||
break;
|
break;
|
||||||
}
|
}
|
||||||
|
@ -200,7 +197,7 @@ std::string EmitGLSL(const Profile& profile, const RuntimeInfo& runtime_info, IR
|
||||||
EmitContext ctx{program, bindings, profile, runtime_info};
|
EmitContext ctx{program, bindings, profile, runtime_info};
|
||||||
Precolor(program);
|
Precolor(program);
|
||||||
EmitCode(ctx, program);
|
EmitCode(ctx, program);
|
||||||
const std::string version{fmt::format("#version 460{}\n", GlslVersionSpecifier(ctx))};
|
const std::string version{fmt::format("#version 450{}\n", GlslVersionSpecifier(ctx))};
|
||||||
ctx.header.insert(0, version);
|
ctx.header.insert(0, version);
|
||||||
if (program.local_memory_size > 0) {
|
if (program.local_memory_size > 0) {
|
||||||
ctx.header += fmt::format("uint lmem[{}];", program.local_memory_size / 4);
|
ctx.header += fmt::format("uint lmem[{}];", program.local_memory_size / 4);
|
||||||
|
@ -225,10 +222,8 @@ std::string EmitGLSL(const Profile& profile, const RuntimeInfo& runtime_info, IR
|
||||||
if (program.info.uses_subgroup_shuffles) {
|
if (program.info.uses_subgroup_shuffles) {
|
||||||
ctx.header += "bool shfl_in_bounds;";
|
ctx.header += "bool shfl_in_bounds;";
|
||||||
}
|
}
|
||||||
ctx.header += "\n";
|
|
||||||
ctx.code.insert(0, ctx.header);
|
ctx.code.insert(0, ctx.header);
|
||||||
ctx.code += "}";
|
ctx.code += '}';
|
||||||
// fmt::print("\n{}\n", ctx.code);
|
|
||||||
return ctx.code;
|
return ctx.code;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
@ -11,7 +11,7 @@
|
||||||
|
|
||||||
namespace Shader::Backend::GLSL {
|
namespace Shader::Backend::GLSL {
|
||||||
namespace {
|
namespace {
|
||||||
static constexpr std::string_view cas_loop{R"(for (;;){{
|
constexpr const char cas_loop[]{R"(for (;;){{
|
||||||
uint old_value={};
|
uint old_value={};
|
||||||
{}=atomicCompSwap({},old_value,{}({},{}));
|
{}=atomicCompSwap({},old_value,{}({},{}));
|
||||||
if ({}==old_value){{break;}}
|
if ({}==old_value){{break;}}
|
||||||
|
@ -21,7 +21,7 @@ void SharedCasFunction(EmitContext& ctx, IR::Inst& inst, std::string_view offset
|
||||||
std::string_view value, std::string_view function) {
|
std::string_view value, std::string_view function) {
|
||||||
const auto ret{ctx.var_alloc.Define(inst, GlslVarType::U32)};
|
const auto ret{ctx.var_alloc.Define(inst, GlslVarType::U32)};
|
||||||
const std::string smem{fmt::format("smem[{}>>2]", offset)};
|
const std::string smem{fmt::format("smem[{}>>2]", offset)};
|
||||||
ctx.Add(cas_loop.data(), smem, ret, smem, function, smem, value, ret);
|
ctx.Add(cas_loop, smem, ret, smem, function, smem, value, ret);
|
||||||
}
|
}
|
||||||
|
|
||||||
void SsboCasFunction(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
void SsboCasFunction(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||||
|
@ -29,7 +29,7 @@ void SsboCasFunction(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||||
const auto ret{ctx.var_alloc.Define(inst, GlslVarType::U32)};
|
const auto ret{ctx.var_alloc.Define(inst, GlslVarType::U32)};
|
||||||
const std::string ssbo{fmt::format("{}_ssbo{}[{}>>2]", ctx.stage_name, binding.U32(),
|
const std::string ssbo{fmt::format("{}_ssbo{}[{}>>2]", ctx.stage_name, binding.U32(),
|
||||||
ctx.var_alloc.Consume(offset))};
|
ctx.var_alloc.Consume(offset))};
|
||||||
ctx.Add(cas_loop.data(), ssbo, ret, ssbo, function, ssbo, value, ret);
|
ctx.Add(cas_loop, ssbo, ret, ssbo, function, ssbo, value, ret);
|
||||||
}
|
}
|
||||||
|
|
||||||
void SsboCasFunctionF32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
void SsboCasFunctionF32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||||
|
@ -38,10 +38,10 @@ void SsboCasFunctionF32(EmitContext& ctx, IR::Inst& inst, const IR::Value& bindi
|
||||||
const std::string ssbo{fmt::format("{}_ssbo{}[{}>>2]", ctx.stage_name, binding.U32(),
|
const std::string ssbo{fmt::format("{}_ssbo{}[{}>>2]", ctx.stage_name, binding.U32(),
|
||||||
ctx.var_alloc.Consume(offset))};
|
ctx.var_alloc.Consume(offset))};
|
||||||
const auto ret{ctx.var_alloc.Define(inst, GlslVarType::U32)};
|
const auto ret{ctx.var_alloc.Define(inst, GlslVarType::U32)};
|
||||||
ctx.Add(cas_loop.data(), ssbo, ret, ssbo, function, ssbo, value, ret);
|
ctx.Add(cas_loop, ssbo, ret, ssbo, function, ssbo, value, ret);
|
||||||
ctx.AddF32("{}=utof({});", inst, ret);
|
ctx.AddF32("{}=utof({});", inst, ret);
|
||||||
}
|
}
|
||||||
} // namespace
|
} // Anonymous namespace
|
||||||
|
|
||||||
void EmitSharedAtomicIAdd32(EmitContext& ctx, IR::Inst& inst, std::string_view pointer_offset,
|
void EmitSharedAtomicIAdd32(EmitContext& ctx, IR::Inst& inst, std::string_view pointer_offset,
|
||||||
std::string_view value) {
|
std::string_view value) {
|
||||||
|
|
|
@ -10,7 +10,7 @@
|
||||||
|
|
||||||
namespace Shader::Backend::GLSL {
|
namespace Shader::Backend::GLSL {
|
||||||
namespace {
|
namespace {
|
||||||
static void Alias(IR::Inst& inst, const IR::Value& value) {
|
void Alias(IR::Inst& inst, const IR::Value& value) {
|
||||||
if (value.IsImmediate()) {
|
if (value.IsImmediate()) {
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
|
|
|
@ -10,13 +10,14 @@
|
||||||
|
|
||||||
namespace Shader::Backend::GLSL {
|
namespace Shader::Backend::GLSL {
|
||||||
namespace {
|
namespace {
|
||||||
static constexpr std::string_view SWIZZLE{"xyzw"};
|
constexpr std::string_view SWIZZLE{"xyzw"};
|
||||||
void CompositeInsert(EmitContext& ctx, std::string_view result, std::string_view composite,
|
void CompositeInsert(EmitContext& ctx, std::string_view result, std::string_view composite,
|
||||||
std::string_view object, u32 index) {
|
std::string_view object, u32 index) {
|
||||||
ctx.Add("{}={};", result, composite);
|
ctx.Add("{}={};", result, composite);
|
||||||
ctx.Add("{}.{}={};", result, SWIZZLE[index], object);
|
ctx.Add("{}.{}={};", result, SWIZZLE[index], object);
|
||||||
}
|
}
|
||||||
} // namespace
|
} // Anonymous namespace
|
||||||
|
|
||||||
void EmitCompositeConstructU32x2(EmitContext& ctx, IR::Inst& inst, std::string_view e1,
|
void EmitCompositeConstructU32x2(EmitContext& ctx, IR::Inst& inst, std::string_view e1,
|
||||||
std::string_view e2) {
|
std::string_view e2) {
|
||||||
ctx.AddU32x2("{}=uvec2({},{});", inst, e1, e2);
|
ctx.AddU32x2("{}=uvec2({},{});", inst, e1, e2);
|
||||||
|
|
|
@ -7,6 +7,7 @@
|
||||||
#include "shader_recompiler/backend/glsl/emit_context.h"
|
#include "shader_recompiler/backend/glsl/emit_context.h"
|
||||||
#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
|
#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
|
||||||
#include "shader_recompiler/frontend/ir/value.h"
|
#include "shader_recompiler/frontend/ir/value.h"
|
||||||
|
#include "shader_recompiler/profile.h"
|
||||||
|
|
||||||
namespace Shader::Backend::GLSL {
|
namespace Shader::Backend::GLSL {
|
||||||
namespace {
|
namespace {
|
||||||
|
@ -39,11 +40,10 @@ std::string OutputVertexIndex(EmitContext& ctx, std::string_view vertex) {
|
||||||
return "";
|
return "";
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
} // namespace
|
} // Anonymous namespace
|
||||||
|
|
||||||
void EmitGetCbufU8([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
void EmitGetCbufU8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||||
[[maybe_unused]] const IR::Value& binding,
|
const IR::Value& offset) {
|
||||||
[[maybe_unused]] const IR::Value& offset) {
|
|
||||||
if (offset.IsImmediate()) {
|
if (offset.IsImmediate()) {
|
||||||
ctx.AddU32("{}=bitfieldExtract(ftou({}_cbuf{}[{}].{}),int({}),8);", inst, ctx.stage_name,
|
ctx.AddU32("{}=bitfieldExtract(ftou({}_cbuf{}[{}].{}),int({}),8);", inst, ctx.stage_name,
|
||||||
binding.U32(), offset.U32() / 16, OffsetSwizzle(offset.U32()),
|
binding.U32(), offset.U32() / 16, OffsetSwizzle(offset.U32()),
|
||||||
|
@ -55,9 +55,8 @@ void EmitGetCbufU8([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst&
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitGetCbufS8([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
void EmitGetCbufS8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||||
[[maybe_unused]] const IR::Value& binding,
|
const IR::Value& offset) {
|
||||||
[[maybe_unused]] const IR::Value& offset) {
|
|
||||||
if (offset.IsImmediate()) {
|
if (offset.IsImmediate()) {
|
||||||
ctx.AddU32("{}=bitfieldExtract(ftoi({}_cbuf{}[{}].{}),int({}),8);", inst, ctx.stage_name,
|
ctx.AddU32("{}=bitfieldExtract(ftoi({}_cbuf{}[{}].{}),int({}),8);", inst, ctx.stage_name,
|
||||||
binding.U32(), offset.U32() / 16, OffsetSwizzle(offset.U32()),
|
binding.U32(), offset.U32() / 16, OffsetSwizzle(offset.U32()),
|
||||||
|
@ -69,9 +68,8 @@ void EmitGetCbufS8([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst&
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitGetCbufU16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
void EmitGetCbufU16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||||
[[maybe_unused]] const IR::Value& binding,
|
const IR::Value& offset) {
|
||||||
[[maybe_unused]] const IR::Value& offset) {
|
|
||||||
if (offset.IsImmediate()) {
|
if (offset.IsImmediate()) {
|
||||||
ctx.AddU32("{}=bitfieldExtract(ftou({}_cbuf{}[{}].{}),int({}),16);", inst, ctx.stage_name,
|
ctx.AddU32("{}=bitfieldExtract(ftou({}_cbuf{}[{}].{}),int({}),16);", inst, ctx.stage_name,
|
||||||
binding.U32(), offset.U32() / 16, OffsetSwizzle(offset.U32()),
|
binding.U32(), offset.U32() / 16, OffsetSwizzle(offset.U32()),
|
||||||
|
@ -84,9 +82,8 @@ void EmitGetCbufU16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitGetCbufS16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
void EmitGetCbufS16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||||
[[maybe_unused]] const IR::Value& binding,
|
const IR::Value& offset) {
|
||||||
[[maybe_unused]] const IR::Value& offset) {
|
|
||||||
if (offset.IsImmediate()) {
|
if (offset.IsImmediate()) {
|
||||||
ctx.AddU32("{}=bitfieldExtract(ftoi({}_cbuf{}[{}].{}),int({}),16);", inst, ctx.stage_name,
|
ctx.AddU32("{}=bitfieldExtract(ftoi({}_cbuf{}[{}].{}),int({}),16);", inst, ctx.stage_name,
|
||||||
binding.U32(), offset.U32() / 16, OffsetSwizzle(offset.U32()),
|
binding.U32(), offset.U32() / 16, OffsetSwizzle(offset.U32()),
|
||||||
|
@ -196,7 +193,7 @@ void EmitGetAttribute(EmitContext& ctx, IR::Inst& inst, IR::Attribute attr,
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, std::string_view value,
|
void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, std::string_view value,
|
||||||
[[maybe_unused]] std::string_view vertex) {
|
std::string_view vertex) {
|
||||||
if (IR::IsGeneric(attr)) {
|
if (IR::IsGeneric(attr)) {
|
||||||
const u32 index{IR::GenericAttributeIndex(attr)};
|
const u32 index{IR::GenericAttributeIndex(attr)};
|
||||||
const u32 element{IR::GenericAttributeElement(attr)};
|
const u32 element{IR::GenericAttributeElement(attr)};
|
||||||
|
@ -223,7 +220,7 @@ void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, std::string_view val
|
||||||
ctx.Add("gl_Position.{}={};", swizzle, value);
|
ctx.Add("gl_Position.{}={};", swizzle, value);
|
||||||
break;
|
break;
|
||||||
case IR::Attribute::ViewportIndex:
|
case IR::Attribute::ViewportIndex:
|
||||||
if (ctx.stage != Stage::Geometry && !ctx.supports_viewport_layer) {
|
if (ctx.stage != Stage::Geometry && !ctx.profile.support_gl_vertex_viewport_layer) {
|
||||||
// LOG_WARNING(..., "Shader stores viewport index but device does not support viewport
|
// LOG_WARNING(..., "Shader stores viewport index but device does not support viewport
|
||||||
// layer extension");
|
// layer extension");
|
||||||
break;
|
break;
|
||||||
|
@ -247,8 +244,7 @@ void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, std::string_view val
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitGetPatch([[maybe_unused]] EmitContext& ctx, IR::Inst& inst,
|
void EmitGetPatch(EmitContext& ctx, IR::Inst& inst, IR::Patch patch) {
|
||||||
[[maybe_unused]] IR::Patch patch) {
|
|
||||||
if (!IR::IsGeneric(patch)) {
|
if (!IR::IsGeneric(patch)) {
|
||||||
throw NotImplementedException("Non-generic patch load");
|
throw NotImplementedException("Non-generic patch load");
|
||||||
}
|
}
|
||||||
|
|
|
@ -25,7 +25,7 @@ void Compare(EmitContext& ctx, IR::Inst& inst, std::string_view lhs, std::string
|
||||||
bool Precise(IR::Inst& inst) {
|
bool Precise(IR::Inst& inst) {
|
||||||
return {inst.Flags<IR::FpControl>().no_contraction};
|
return {inst.Flags<IR::FpControl>().no_contraction};
|
||||||
}
|
}
|
||||||
} // namespace
|
} // Anonymous namespace
|
||||||
|
|
||||||
void EmitFPAbs16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
void EmitFPAbs16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
||||||
[[maybe_unused]] std::string_view value) {
|
[[maybe_unused]] std::string_view value) {
|
||||||
|
|
|
@ -102,7 +102,7 @@ IR::Inst* PrepareSparse(IR::Inst& inst) {
|
||||||
}
|
}
|
||||||
return sparse_inst;
|
return sparse_inst;
|
||||||
}
|
}
|
||||||
} // namespace
|
} // Anonymous namespace
|
||||||
|
|
||||||
void EmitImageSampleImplicitLod([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
void EmitImageSampleImplicitLod([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
||||||
[[maybe_unused]] const IR::Value& index,
|
[[maybe_unused]] const IR::Value& index,
|
||||||
|
|
|
@ -25,9 +25,7 @@ void EmitPhi(EmitContext& ctx, IR::Inst& phi) {
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitVoid(EmitContext& ctx) {
|
void EmitVoid(EmitContext& ctx) {}
|
||||||
// NotImplemented();
|
|
||||||
}
|
|
||||||
|
|
||||||
void EmitReference(EmitContext& ctx, const IR::Value& value) {
|
void EmitReference(EmitContext& ctx, const IR::Value& value) {
|
||||||
ctx.var_alloc.Consume(value);
|
ctx.var_alloc.Consume(value);
|
||||||
|
@ -94,13 +92,9 @@ void EmitDeviceMemoryBarrier(EmitContext& ctx) {
|
||||||
NotImplemented();
|
NotImplemented();
|
||||||
}
|
}
|
||||||
|
|
||||||
void EmitPrologue(EmitContext& ctx) {
|
void EmitPrologue(EmitContext& ctx) {}
|
||||||
// NotImplemented();
|
|
||||||
}
|
|
||||||
|
|
||||||
void EmitEpilogue(EmitContext& ctx) {
|
void EmitEpilogue(EmitContext& ctx) {}
|
||||||
// NotImplemented();
|
|
||||||
}
|
|
||||||
|
|
||||||
void EmitEmitVertex(EmitContext& ctx, const IR::Value& stream) {
|
void EmitEmitVertex(EmitContext& ctx, const IR::Value& stream) {
|
||||||
ctx.Add("EmitStreamVertex(int({}));", ctx.var_alloc.Consume(stream));
|
ctx.Add("EmitStreamVertex(int({}));", ctx.var_alloc.Consume(stream));
|
||||||
|
|
Reference in New Issue