code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include <string_view>

#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
#include "shader_recompiler/frontend/ir/value.h"
#include "shader_recompiler/profile.h"

namespace Shader::Backend::GLSL {
namespace {
constexpr char cas_loop[]{"for(;;){{uint old_value={};uint "
                          "cas_result=atomicCompSwap({},old_value,bitfieldInsert({},{},{},{}));"
                          "if(cas_result==old_value){{break;}}}}"};

void SsboWriteCas(EmitContext& ctx, const IR::Value& binding, std::string_view offset_var,
                  std::string_view value, std::string_view bit_offset, u32 num_bits) {
    const auto ssbo{fmt::format("{}_ssbo{}[{}>>2]", ctx.stage_name, binding.U32(), offset_var)};
    ctx.Add(cas_loop, ssbo, ssbo, ssbo, value, bit_offset, num_bits);
}
} // Anonymous namespace

void EmitLoadGlobalU8(EmitContext&) {
    NotImplemented();
}

void EmitLoadGlobalS8(EmitContext&) {
    NotImplemented();
}

void EmitLoadGlobalU16(EmitContext&) {
    NotImplemented();
}

void EmitLoadGlobalS16(EmitContext&) {
    NotImplemented();
}

void EmitLoadGlobal32(EmitContext& ctx, IR::Inst& inst, std::string_view address) {
    if (ctx.profile.support_int64) {
        return ctx.AddU32("{}=LoadGlobal32({});", inst, address);
    }
    LOG_WARNING(Shader_GLSL, "Int64 not supported, ignoring memory operation");
    ctx.AddU32("{}=0u;", inst);
}

void EmitLoadGlobal64(EmitContext& ctx, IR::Inst& inst, std::string_view address) {
    if (ctx.profile.support_int64) {
        return ctx.AddU32x2("{}=LoadGlobal64({});", inst, address);
    }
    LOG_WARNING(Shader_GLSL, "Int64 not supported, ignoring memory operation");
    ctx.AddU32x2("{}=uvec2(0);", inst);
}

void EmitLoadGlobal128(EmitContext& ctx, IR::Inst& inst, std::string_view address) {
    if (ctx.profile.support_int64) {
        return ctx.AddU32x4("{}=LoadGlobal128({});", inst, address);
    }
    LOG_WARNING(Shader_GLSL, "Int64 not supported, ignoring memory operation");
    ctx.AddU32x4("{}=uvec4(0);", inst);
}

void EmitWriteGlobalU8(EmitContext&) {
    NotImplemented();
}

void EmitWriteGlobalS8(EmitContext&) {
    NotImplemented();
}

void EmitWriteGlobalU16(EmitContext&) {
    NotImplemented();
}

void EmitWriteGlobalS16(EmitContext&) {
    NotImplemented();
}

void EmitWriteGlobal32(EmitContext& ctx, std::string_view address, std::string_view value) {
    if (ctx.profile.support_int64) {
        return ctx.Add("WriteGlobal32({},{});", address, value);
    }
    LOG_WARNING(Shader_GLSL, "Int64 not supported, ignoring memory operation");
}

void EmitWriteGlobal64(EmitContext& ctx, std::string_view address, std::string_view value) {
    if (ctx.profile.support_int64) {
        return ctx.Add("WriteGlobal64({},{});", address, value);
    }
    LOG_WARNING(Shader_GLSL, "Int64 not supported, ignoring memory operation");
}

void EmitWriteGlobal128(EmitContext& ctx, std::string_view address, std::string_view value) {
    if (ctx.profile.support_int64) {
        return ctx.Add("WriteGlobal128({},{});", address, value);
    }
    LOG_WARNING(Shader_GLSL, "Int64 not supported, ignoring memory operation");
}

void EmitLoadStorageU8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
                       const IR::Value& offset) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    ctx.AddU32("{}=bitfieldExtract({}_ssbo{}[{}>>2],int({}%4)*8,8);", inst, ctx.stage_name,
               binding.U32(), offset_var, offset_var);
}

void EmitLoadStorageS8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
                       const IR::Value& offset) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    ctx.AddU32("{}=bitfieldExtract(int({}_ssbo{}[{}>>2]),int({}%4)*8,8);", inst, ctx.stage_name,
               binding.U32(), offset_var, offset_var);
}

void EmitLoadStorageU16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
                        const IR::Value& offset) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    ctx.AddU32("{}=bitfieldExtract({}_ssbo{}[{}>>2],int(({}>>1)%2)*16,16);", inst, ctx.stage_name,
               binding.U32(), offset_var, offset_var);
}

void EmitLoadStorageS16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
                        const IR::Value& offset) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    ctx.AddU32("{}=bitfieldExtract(int({}_ssbo{}[{}>>2]),int(({}>>1)%2)*16,16);", inst,
               ctx.stage_name, binding.U32(), offset_var, offset_var);
}

void EmitLoadStorage32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
                       const IR::Value& offset) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    ctx.AddU32("{}={}_ssbo{}[{}>>2];", inst, ctx.stage_name, binding.U32(), offset_var);
}

void EmitLoadStorage64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
                       const IR::Value& offset) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    ctx.AddU32x2("{}=uvec2({}_ssbo{}[{}>>2],{}_ssbo{}[({}+4)>>2]);", inst, ctx.stage_name,
                 binding.U32(), offset_var, ctx.stage_name, binding.U32(), offset_var);
}

void EmitLoadStorage128(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
                        const IR::Value& offset) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    ctx.AddU32x4("{}=uvec4({}_ssbo{}[{}>>2],{}_ssbo{}[({}+4)>>2],{}_ssbo{}[({}+8)>>2],{}_ssbo{}[({}"
                 "+12)>>2]);",
                 inst, ctx.stage_name, binding.U32(), offset_var, ctx.stage_name, binding.U32(),
                 offset_var, ctx.stage_name, binding.U32(), offset_var, ctx.stage_name,
                 binding.U32(), offset_var);
}

void EmitWriteStorageU8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
                        std::string_view value) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    const auto bit_offset{fmt::format("int({}%4)*8", offset_var)};
    SsboWriteCas(ctx, binding, offset_var, value, bit_offset, 8);
}

void EmitWriteStorageS8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
                        std::string_view value) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    const auto bit_offset{fmt::format("int({}%4)*8", offset_var)};
    SsboWriteCas(ctx, binding, offset_var, value, bit_offset, 8);
}

void EmitWriteStorageU16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
                         std::string_view value) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    const auto bit_offset{fmt::format("int(({}>>1)%2)*16", offset_var)};
    SsboWriteCas(ctx, binding, offset_var, value, bit_offset, 16);
}

void EmitWriteStorageS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
                         std::string_view value) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    const auto bit_offset{fmt::format("int(({}>>1)%2)*16", offset_var)};
    SsboWriteCas(ctx, binding, offset_var, value, bit_offset, 16);
}

void EmitWriteStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
                        std::string_view value) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    ctx.Add("{}_ssbo{}[{}>>2]={};", ctx.stage_name, binding.U32(), offset_var, value);
}

void EmitWriteStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
                        std::string_view value) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    ctx.Add("{}_ssbo{}[{}>>2]={}.x;", ctx.stage_name, binding.U32(), offset_var, value);
    ctx.Add("{}_ssbo{}[({}+4)>>2]={}.y;", ctx.stage_name, binding.U32(), offset_var, value);
}

void EmitWriteStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
                         std::string_view value) {
    const auto offset_var{ctx.var_alloc.Consume(offset)};
    ctx.Add("{}_ssbo{}[{}>>2]={}.x;", ctx.stage_name, binding.U32(), offset_var, value);
    ctx.Add("{}_ssbo{}[({}+4)>>2]={}.y;", ctx.stage_name, binding.U32(), offset_var, value);
    ctx.Add("{}_ssbo{}[({}+8)>>2]={}.z;", ctx.stage_name, binding.U32(), offset_var, value);
    ctx.Add("{}_ssbo{}[({}+12)>>2]={}.w;", ctx.stage_name, binding.U32(), offset_var, value);
}
} // namespace Shader::Backend::GLSL


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include <string_view>

#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
#include "shader_recompiler/frontend/ir/value.h"

namespace Shader::Backend::GLSL {

void EmitGetRegister(EmitContext& ctx) {
    NotImplemented();
}

void EmitSetRegister(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetPred(EmitContext& ctx) {
    NotImplemented();
}

void EmitSetPred(EmitContext& ctx) {
    NotImplemented();
}

void EmitSetGotoVariable(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetGotoVariable(EmitContext& ctx) {
    NotImplemented();
}

void EmitSetIndirectBranchVariable(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetIndirectBranchVariable(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetZFlag(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetSFlag(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetCFlag(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetOFlag(EmitContext& ctx) {
    NotImplemented();
}

void EmitSetZFlag(EmitContext& ctx) {
    NotImplemented();
}

void EmitSetSFlag(EmitContext& ctx) {
    NotImplemented();
}

void EmitSetCFlag(EmitContext& ctx) {
    NotImplemented();
}

void EmitSetOFlag(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetZeroFromOp(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetSignFromOp(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetCarryFromOp(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetOverflowFromOp(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetSparseFromOp(EmitContext& ctx) {
    NotImplemented();
}

void EmitGetInBoundsFromOp(EmitContext& ctx) {
    NotImplemented();
}

} // namespace Shader::Backend::GLSL

repo_name: litucks/torzu

license: mit


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include <string_view>

#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
#include "shader_recompiler/frontend/ir/value.h"

namespace Shader::Backend::GLSL {
void EmitSelectU1(EmitContext& ctx, IR::Inst& inst, std::string_view cond,
                  std::string_view true_value, std::string_view false_value) {
    ctx.AddU1("{}={}?{}:{};", inst, cond, true_value, false_value);
}

void EmitSelectU8([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view cond,
                  [[maybe_unused]] std::string_view true_value,
                  [[maybe_unused]] std::string_view false_value) {
    NotImplemented();
}

void EmitSelectU16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view cond,
                   [[maybe_unused]] std::string_view true_value,
                   [[maybe_unused]] std::string_view false_value) {
    NotImplemented();
}

void EmitSelectU32(EmitContext& ctx, IR::Inst& inst, std::string_view cond,
                   std::string_view true_value, std::string_view false_value) {
    ctx.AddU32("{}={}?{}:{};", inst, cond, true_value, false_value);
}

void EmitSelectU64(EmitContext& ctx, IR::Inst& inst, std::string_view cond,
                   std::string_view true_value, std::string_view false_value) {
    ctx.AddU64("{}={}?{}:{};", inst, cond, true_value, false_value);
}

void EmitSelectF16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view cond,
                   [[maybe_unused]] std::string_view true_value,
                   [[maybe_unused]] std::string_view false_value) {
    NotImplemented();
}

void EmitSelectF32(EmitContext& ctx, IR::Inst& inst, std::string_view cond,
                   std::string_view true_value, std::string_view false_value) {
    ctx.AddF32("{}={}?{}:{};", inst, cond, true_value, false_value);
}

void EmitSelectF64(EmitContext& ctx, IR::Inst& inst, std::string_view cond,
                   std::string_view true_value, std::string_view false_value) {
    ctx.AddF64("{}={}?{}:{};", inst, cond, true_value, false_value);
}

} // namespace Shader::Backend::GLSL


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
#include "shader_recompiler/frontend/ir/value.h"

namespace Shader::Backend::GLSL {
namespace {
constexpr char cas_loop[]{"for(;;){{uint old_value={};uint "
                          "cas_result=atomicCompSwap({},old_value,bitfieldInsert({},{},{},{}));"
                          "if(cas_result==old_value){{break;}}}}"};

void SharedWriteCas(EmitContext& ctx, std::string_view offset, std::string_view value,
                    std::string_view bit_offset, u32 num_bits) {
    const auto smem{fmt::format("smem[{}>>2]", offset)};
    ctx.Add(cas_loop, smem, smem, smem, value, bit_offset, num_bits);
}
} // Anonymous namespace

void EmitLoadSharedU8(EmitContext& ctx, IR::Inst& inst, std::string_view offset) {
    ctx.AddU32("{}=bitfieldExtract(smem[{}>>2],int({}%4)*8,8);", inst, offset, offset);
}

void EmitLoadSharedS8(EmitContext& ctx, IR::Inst& inst, std::string_view offset) {
    ctx.AddU32("{}=bitfieldExtract(int(smem[{}>>2]),int({}%4)*8,8);", inst, offset, offset);
}

void EmitLoadSharedU16(EmitContext& ctx, IR::Inst& inst, std::string_view offset) {
    ctx.AddU32("{}=bitfieldExtract(smem[{}>>2],int(({}>>1)%2)*16,16);", inst, offset, offset);
}

void EmitLoadSharedS16(EmitContext& ctx, IR::Inst& inst, std::string_view offset) {
    ctx.AddU32("{}=bitfieldExtract(int(smem[{}>>2]),int(({}>>1)%2)*16,16);", inst, offset, offset);
}

void EmitLoadSharedU32(EmitContext& ctx, IR::Inst& inst, std::string_view offset) {
    ctx.AddU32("{}=smem[{}>>2];", inst, offset);
}

void EmitLoadSharedU64(EmitContext& ctx, IR::Inst& inst, std::string_view offset) {
    ctx.AddU32x2("{}=uvec2(smem[{}>>2],smem[({}+4)>>2]);", inst, offset, offset);
}

void EmitLoadSharedU128(EmitContext& ctx, IR::Inst& inst, std::string_view offset) {
    ctx.AddU32x4("{}=uvec4(smem[{}>>2],smem[({}+4)>>2],smem[({}+8)>>2],smem[({}+12)>>2]);", inst,
                 offset, offset, offset, offset);
}

void EmitWriteSharedU8(EmitContext& ctx, std::string_view offset, std::string_view value) {
    const auto bit_offset{fmt::format("int({}%4)*8", offset)};
    SharedWriteCas(ctx, offset, value, bit_offset, 8);
}

void EmitWriteSharedU16(EmitContext& ctx, std::string_view offset, std::string_view value) {
    const auto bit_offset{fmt::format("int(({}>>1)%2)*16", offset)};
    SharedWriteCas(ctx, offset, value, bit_offset, 16);
}

void EmitWriteSharedU32(EmitContext& ctx, std::string_view offset, std::string_view value) {
    ctx.Add("smem[{}>>2]={};", offset, value);
}

void EmitWriteSharedU64(EmitContext& ctx, std::string_view offset, std::string_view value) {
    ctx.Add("smem[{}>>2]={}.x;", offset, value);
    ctx.Add("smem[({}+4)>>2]={}.y;", offset, value);
}

void EmitWriteSharedU128(EmitContext& ctx, std::string_view offset, std::string_view value) {
    ctx.Add("smem[{}>>2]={}.x;", offset, value);
    ctx.Add("smem[({}+4)>>2]={}.y;", offset, value);
    ctx.Add("smem[({}+8)>>2]={}.z;", offset, value);
    ctx.Add("smem[({}+12)>>2]={}.w;", offset, value);
}

} // namespace Shader::Backend::GLSL


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
#include "shader_recompiler/frontend/ir/program.h"
#include "shader_recompiler/frontend/ir/value.h"
#include "shader_recompiler/profile.h"

namespace Shader::Backend::GLSL {
namespace {
std::string_view OutputVertexIndex(EmitContext& ctx) {
    return ctx.stage == Stage::TessellationControl ? "[gl_InvocationID]" : "";
}

void InitializeOutputVaryings(EmitContext& ctx) {
    if (ctx.uses_geometry_passthrough) {
        return;
    }
    if (ctx.stage == Stage::VertexB || ctx.stage == Stage::Geometry) {
        ctx.Add("gl_Position=vec4(0,0,0,1);");
    }
    for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
        if (!ctx.info.stores.Generic(index)) {
            continue;
        }
        const auto& info_array{ctx.output_generics.at(index)};
        const auto output_decorator{OutputVertexIndex(ctx)};
        size_t element{};
        while (element < info_array.size()) {
            const auto& info{info_array.at(element)};
            const auto varying_name{fmt::format("{}{}", info.name, output_decorator)};
            switch (info.num_components) {
            case 1: {
                const char value{element == 3 ? '1' : '0'};
                ctx.Add("{}={}.f;", varying_name, value);
                break;
            }
            case 2:
            case 3:
                if (element + info.num_components < 4) {
                    ctx.Add("{}=vec{}(0);", varying_name, info.num_components);
                } else {
                    // last element is the w component, must be initialized to 1
                    const auto zeros{info.num_components == 3 ? "0,0," : "0,"};
                    ctx.Add("{}=vec{}({}1);", varying_name, info.num_components, zeros);
                }
                break;
            case 4:
                ctx.Add("{}=vec4(0,0,0,1);", varying_name);
                break;
            default:
                break;
            }
            element += info.num_components;
        }
    }
}
} // Anonymous namespace

void EmitPhi(EmitContext& ctx, IR::Inst& phi) {
    const size_t num_args{phi.NumArgs()};
    for (size_t i = 0; i < num_args; ++i) {
        ctx.var_alloc.Consume(phi.Arg(i));
    }
    if (!phi.Definition<Id>().is_valid) {
        // The phi node wasn't forward defined
        ctx.var_alloc.PhiDefine(phi, phi.Type());
    }
}

void EmitVoid(EmitContext&) {}

void EmitReference(EmitContext& ctx, const IR::Value& value) {
    ctx.var_alloc.Consume(value);
}

void EmitPhiMove(EmitContext& ctx, const IR::Value& phi_value, const IR::Value& value) {
    IR::Inst& phi{*phi_value.InstRecursive()};
    const auto phi_type{phi.Type()};
    if (!phi.Definition<Id>().is_valid) {
        // The phi node wasn't forward defined
        ctx.var_alloc.PhiDefine(phi, phi_type);
    }
    const auto phi_reg{ctx.var_alloc.Consume(IR::Value{&phi})};
    const auto val_reg{ctx.var_alloc.Consume(value)};
    if (phi_reg == val_reg) {
        return;
    }
    const bool needs_workaround{ctx.profile.has_gl_bool_ref_bug && phi_type == IR::Type::U1};
    const auto suffix{needs_workaround ? "?true:false" : ""};
    ctx.Add("{}={}{};", phi_reg, val_reg, suffix);
}

void EmitPrologue(EmitContext& ctx) {
    InitializeOutputVaryings(ctx);
}

void EmitEpilogue(EmitContext&) {}

void EmitEmitVertex(EmitContext& ctx, const IR::Value& stream) {
    ctx.Add("EmitStreamVertex(int({}));", ctx.var_alloc.Consume(stream));
    InitializeOutputVaryings(ctx);
}

void EmitEndPrimitive(EmitContext& ctx, const IR::Value& stream) {
    ctx.Add("EndStreamPrimitive(int({}));", ctx.var_alloc.Consume(stream));
}

} // namespace Shader::Backend::GLSL


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
#include "shader_recompiler/backend/glsl/glsl_emit_context.h"

namespace Shader::Backend::GLSL {

void EmitUndefU1(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU1("{}=false;", inst);
}

void EmitUndefU8(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU32("{}=0u;", inst);
}

void EmitUndefU16(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU32("{}=0u;", inst);
}

void EmitUndefU32(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU32("{}=0u;", inst);
}

void EmitUndefU64(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU64("{}=0u;", inst);
}

} // namespace Shader::Backend::GLSL


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include <string_view>

#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
#include "shader_recompiler/frontend/ir/value.h"
#include "shader_recompiler/profile.h"

namespace Shader::Backend::GLSL {
namespace {
constexpr char THREAD_ID[]{"gl_SubGroupInvocationARB"};

void SetInBoundsFlag(EmitContext& ctx, IR::Inst& inst) {
    IR::Inst* const in_bounds{inst.GetAssociatedPseudoOperation(IR::Opcode::GetInBoundsFromOp)};
    if (!in_bounds) {
        return;
    }
    ctx.AddU1("{}=shfl_in_bounds;", *in_bounds);
    in_bounds->Invalidate();
}

std::string ComputeMinThreadId(std::string_view thread_id, std::string_view segmentation_mask) {
    return fmt::format("({}&{})", thread_id, segmentation_mask);
}

std::string ComputeMaxThreadId(std::string_view min_thread_id, std::string_view clamp,
                               std::string_view not_seg_mask) {
    return fmt::format("({})|({}&{})", min_thread_id, clamp, not_seg_mask);
}

std::string GetMaxThreadId(std::string_view thread_id, std::string_view clamp,
                           std::string_view segmentation_mask) {
    const auto not_seg_mask{fmt::format("(~{})", segmentation_mask)};
    const auto min_thread_id{ComputeMinThreadId(thread_id, segmentation_mask)};
    return ComputeMaxThreadId(min_thread_id, clamp, not_seg_mask);
}

void UseShuffleNv(EmitContext& ctx, IR::Inst& inst, std::string_view shfl_op,
                  std::string_view value, std::string_view index,
                  [[maybe_unused]] std::string_view clamp, std::string_view segmentation_mask) {
    const auto width{fmt::format("32u>>(bitCount({}&31u))", segmentation_mask)};
    ctx.AddU32("{}={}({},{},{},shfl_in_bounds);", inst, shfl_op, value, index, width);
    SetInBoundsFlag(ctx, inst);
}

std::string_view BallotIndex(EmitContext& ctx) {
    if (!ctx.profile.warp_size_potentially_larger_than_guest) {
        return ".x";
    }
    return "[gl_SubGroupInvocationARB>>5]";
}

std::string GetMask(EmitContext& ctx, std::string_view mask) {
    const auto ballot_index{BallotIndex(ctx)};
    return fmt::format("uint(uvec2({}){})", mask, ballot_index);
}
} // Anonymous namespace

void EmitLaneId(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU32("{}={}&31u;", inst, THREAD_ID);
}

void EmitVoteAll(EmitContext& ctx, IR::Inst& inst, std::string_view pred) {
    if (!ctx.profile.warp_size_potentially_larger_than_guest) {
        ctx.AddU1("{}=allInvocationsEqualARB({});", inst, pred);
        return;
    }
    const auto ballot_index{BallotIndex(ctx)};
    const auto active_mask{fmt::format("uvec2(ballotARB(true)){}", ballot_index)};
    const auto ballot{fmt::format("uvec2(ballotARB({})){}", pred, ballot_index)};
    ctx.AddU1("{}=({}&{})=={};", inst, ballot, active_mask, active_mask);
}

void EmitVoteAny(EmitContext& ctx, IR::Inst& inst, std::string_view pred) {
    if (!ctx.profile.warp_size_potentially_larger_than_guest) {
        ctx.AddU1("{}=anyInvocationARB({});", inst, pred);
        return;
    }
    const auto ballot_index{BallotIndex(ctx)};
    const auto active_mask{fmt::format("uvec2(ballotARB(true)){}", ballot_index)};
    const auto ballot{fmt::format("uvec2(ballotARB({})){}", pred, ballot_index)};
    ctx.AddU1("{}=({}&{})!=0u;", inst, ballot, active_mask, active_mask);
}

void EmitVoteEqual(EmitContext& ctx, IR::Inst& inst, std::string_view pred) {
    if (!ctx.profile.warp_size_potentially_larger_than_guest) {
        ctx.AddU1("{}=allInvocationsEqualARB({});", inst, pred);
        return;
    }
    const auto ballot_index{BallotIndex(ctx)};
    const auto active_mask{fmt::format("uvec2(ballotARB(true)){}", ballot_index)};
    const auto ballot{fmt::format("uvec2(ballotARB({})){}", pred, ballot_index)};
    const auto value{fmt::format("({}^{})", ballot, active_mask)};
    ctx.AddU1("{}=({}==0)||({}=={});", inst, value, value, active_mask);
}

void EmitSubgroupBallot(EmitContext& ctx, IR::Inst& inst, std::string_view pred) {
    const auto ballot_index{BallotIndex(ctx)};
    ctx.AddU32("{}=uvec2(ballotARB({})){};", inst, pred, ballot_index);
}

void EmitSubgroupEqMask(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU32("{}={};", inst, GetMask(ctx, "gl_SubGroupEqMaskARB"));
}

void EmitSubgroupLtMask(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU32("{}={};", inst, GetMask(ctx, "gl_SubGroupLtMaskARB"));
}

void EmitSubgroupLeMask(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU32("{}={};", inst, GetMask(ctx, "gl_SubGroupLeMaskARB"));
}

void EmitSubgroupGtMask(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU32("{}={};", inst, GetMask(ctx, "gl_SubGroupGtMaskARB"));
}

void EmitSubgroupGeMask(EmitContext& ctx, IR::Inst& inst) {
    ctx.AddU32("{}={};", inst, GetMask(ctx, "gl_SubGroupGeMaskARB"));
}

void EmitShuffleIndex(EmitContext& ctx, IR::Inst& inst, std::string_view value,
                      std::string_view index, std::string_view clamp, std::string_view seg_mask) {
    if (ctx.profile.support_gl_warp_intrinsics) {
        UseShuffleNv(ctx, inst, "shuffleNV", value, index, clamp, seg_mask);
        return;
    }
    const bool big_warp{ctx.profile.warp_size_potentially_larger_than_guest};
    const auto is_upper_partition{"int(gl_SubGroupInvocationARB)>=32"};
    const auto upper_index{fmt::format("{}?{}+32:{}", is_upper_partition, index, index)};
    const auto upper_clamp{fmt::format("{}?{}+32:{}", is_upper_partition, clamp, clamp)};

    const auto not_seg_mask{fmt::format("(~{})", seg_mask)};
    const auto min_thread_id{ComputeMinThreadId(THREAD_ID, seg_mask)};
    const auto max_thread_id{
        ComputeMaxThreadId(min_thread_id, big_warp ? upper_clamp : clamp, not_seg_mask)};

    const auto lhs{fmt::format("({}&{})", big_warp ? upper_index : index, not_seg_mask)};
    const auto src_thread_id{fmt::format("({})|({})", lhs, min_thread_id)};
    ctx.Add("shfl_in_bounds=int({})<=int({});", src_thread_id, max_thread_id);
    SetInBoundsFlag(ctx, inst);
    ctx.Add("shfl_result=readInvocationARB({},{});", value, src_thread_id);
    ctx.AddU32("{}=shfl_in_bounds?shfl_result:{};", inst, value);
}

void EmitShuffleUp(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view index,
                   std::string_view clamp, std::string_view seg_mask) {
    if (ctx.profile.support_gl_warp_intrinsics) {
        UseShuffleNv(ctx, inst, "shuffleUpNV", value, index, clamp, seg_mask);
        return;
    }
    const bool big_warp{ctx.profile.warp_size_potentially_larger_than_guest};
    const auto is_upper_partition{"int(gl_SubGroupInvocationARB)>=32"};
    const auto upper_clamp{fmt::format("{}?{}+32:{}", is_upper_partition, clamp, clamp)};

    const auto max_thread_id{GetMaxThreadId(THREAD_ID, big_warp ? upper_clamp : clamp, seg_mask)};
    const auto src_thread_id{fmt::format("({}-{})", THREAD_ID, index)};
    ctx.Add("shfl_in_bounds=int({})>=int({});", src_thread_id, max_thread_id);
    SetInBoundsFlag(ctx, inst);
    ctx.Add("shfl_result=readInvocationARB({},{});", value, src_thread_id);
    ctx.AddU32("{}=shfl_in_bounds?shfl_result:{};", inst, value);
}

void EmitShuffleDown(EmitContext& ctx, IR::Inst& inst, std::string_view value,
                     std::string_view index, std::string_view clamp, std::string_view seg_mask) {
    if (ctx.profile.support_gl_warp_intrinsics) {
        UseShuffleNv(ctx, inst, "shuffleDownNV", value, index, clamp, seg_mask);
        return;
    }
    const bool big_warp{ctx.profile.warp_size_potentially_larger_than_guest};
    const auto is_upper_partition{"int(gl_SubGroupInvocationARB)>=32"};
    const auto upper_clamp{fmt::format("{}?{}+32:{}", is_upper_partition, clamp, clamp)};

    const auto max_thread_id{GetMaxThreadId(THREAD_ID, big_warp ? upper_clamp : clamp, seg_mask)};
    const auto src_thread_id{fmt::format("({}+{})", THREAD_ID, index)};
    ctx.Add("shfl_in_bounds=int({})<=int({});", src_thread_id, max_thread_id);
    SetInBoundsFlag(ctx, inst);
    ctx.Add("shfl_result=readInvocationARB({},{});", value, src_thread_id);
    ctx.AddU32("{}=shfl_in_bounds?shfl_result:{};", inst, value);
}

void EmitShuffleButterfly(EmitContext& ctx, IR::Inst& inst, std::string_view value,
                          std::string_view index, std::string_view clamp,
                          std::string_view seg_mask) {
    if (ctx.profile.support_gl_warp_intrinsics) {
        UseShuffleNv(ctx, inst, "shuffleXorNV", value, index, clamp, seg_mask);
        return;
    }
    const bool big_warp{ctx.profile.warp_size_potentially_larger_than_guest};
    const auto is_upper_partition{"int(gl_SubGroupInvocationARB)>=32"};
    const auto upper_clamp{fmt::format("{}?{}+32:{}", is_upper_partition, clamp, clamp)};

    const auto max_thread_id{GetMaxThreadId(THREAD_ID, big_warp ? upper_clamp : clamp, seg_mask)};
    const auto src_thread_id{fmt::format("({}^{})", THREAD_ID, index)};
    ctx.Add("shfl_in_bounds=int({})<=int({});", src_thread_id, max_thread_id);
    SetInBoundsFlag(ctx, inst);
    ctx.Add("shfl_result=readInvocationARB({},{});", value, src_thread_id);
    ctx.AddU32("{}=shfl_in_bounds?shfl_result:{};", inst, value);
}

void EmitFSwizzleAdd(EmitContext& ctx, IR::Inst& inst, std::string_view op_a, std::string_view op_b,
                     std::string_view swizzle) {
    const auto mask{fmt::format("({}>>((gl_SubGroupInvocationARB&3)<<1))&3", swizzle)};
    const std::string modifier_a = fmt::format("FSWZ_A[{}]", mask);
    const std::string modifier_b = fmt::format("FSWZ_B[{}]", mask);
    ctx.AddF32("{}=({}*{})+({}*{});", inst, op_a, modifier_a, op_b, modifier_b);
}

void EmitDPdxFine(EmitContext& ctx, IR::Inst& inst, std::string_view op_a) {
    if (ctx.profile.support_gl_derivative_control) {
        ctx.AddF32("{}=dFdxFine({});", inst, op_a);
    } else {
        LOG_WARNING(Shader_GLSL, "Device does not support dFdxFine, fallback to dFdx");
        ctx.AddF32("{}=dFdx({});", inst, op_a);
    }
}

void EmitDPdyFine(EmitContext& ctx, IR::Inst& inst, std::string_view op_a) {
    if (ctx.profile.support_gl_derivative_control) {
        ctx.AddF32("{}=dFdyFine({});", inst, op_a);
    } else {
        LOG_WARNING(Shader_GLSL, "Device does not support dFdyFine, fallback to dFdy");
        ctx.AddF32("{}=dFdy({});", inst, op_a);
    }
}

void EmitDPdxCoarse(EmitContext& ctx, IR::Inst& inst, std::string_view op_a) {
    if (ctx.profile.support_gl_derivative_control) {
        ctx.AddF32("{}=dFdxCoarse({});", inst, op_a);
    } else {
        LOG_WARNING(Shader_GLSL, "Device does not support dFdxCoarse, fallback to dFdx");
        ctx.AddF32("{}=dFdx({});", inst, op_a);
    }
}

void EmitDPdyCoarse(EmitContext& ctx, IR::Inst& inst, std::string_view op_a) {
    if (ctx.profile.support_gl_derivative_control) {
        ctx.AddF32("{}=dFdyCoarse({});", inst, op_a);
    } else {
        LOG_WARNING(Shader_GLSL, "Device does not support dFdyCoarse, fallback to dFdy");
        ctx.AddF32("{}=dFdy({});", inst, op_a);
    }
}
} // namespace Shader::Backend::GLSL


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "common/div_ceil.h"
#include "shader_recompiler/backend/bindings.h"
#include "shader_recompiler/backend/glsl/glsl_emit_context.h"
#include "shader_recompiler/frontend/ir/program.h"
#include "shader_recompiler/profile.h"
#include "shader_recompiler/runtime_info.h"

namespace Shader::Backend::GLSL {
namespace {
u32 CbufIndex(size_t offset) {
    return (offset / 4) % 4;
}

char Swizzle(size_t offset) {
    return "xyzw"[CbufIndex(offset)];
}

std::string_view InterpDecorator(Interpolation interp) {
    switch (interp) {
    case Interpolation::Smooth:
        return "";
    case Interpolation::Flat:
        return "flat ";
    case Interpolation::NoPerspective:
        return "noperspective ";
    }
    throw InvalidArgument("Invalid interpolation {}", interp);
}

std::string_view InputArrayDecorator(Stage stage) {
    switch (stage) {
    case Stage::Geometry:
    case Stage::TessellationControl:
    case Stage::TessellationEval:
        return "[]";
    default:
        return "";
    }
}

bool StoresPerVertexAttributes(Stage stage) {
    switch (stage) {
    case Stage::VertexA:
    case Stage::VertexB:
    case Stage::Geometry:
    case Stage::TessellationEval:
        return true;
    default:
        return false;
    }
}

std::string OutputDecorator(Stage stage, u32 size) {
    switch (stage) {
    case Stage::TessellationControl:
        return fmt::format("[{}]", size);
    default:
        return "";
    }
}

std::string_view DepthSamplerType(TextureType type) {
    switch (type) {
    case TextureType::Color1D:
        return "sampler1DShadow";
    case TextureType::ColorArray1D:
        return "sampler1DArrayShadow";
    case TextureType::Color2D:
        return "sampler2DShadow";
    case TextureType::ColorArray2D:
        return "sampler2DArrayShadow";
    case TextureType::ColorCube:
        return "samplerCubeShadow";
    case TextureType::ColorArrayCube:
        return "samplerCubeArrayShadow";
    default:
        throw NotImplementedException("Texture type: {}", type);
    }
}

std::string_view ColorSamplerType(TextureType type, bool is_multisample = false) {
    if (is_multisample) {
        ASSERT(type == TextureType::Color2D || type == TextureType::ColorArray2D);
    }
    switch (type) {
    case TextureType::Color1D:
        return "sampler1D";
    case TextureType::ColorArray1D:
        return "sampler1DArray";
    case TextureType::Color2D:
    case TextureType::Color2DRect:
        return is_multisample ? "sampler2DMS" : "sampler2D";
    case TextureType::ColorArray2D:
        return is_multisample ? "sampler2DMSArray" : "sampler2DArray";
    case TextureType::Color3D:
        return "sampler3D";
    case TextureType::ColorCube:
        return "samplerCube";
    case TextureType::ColorArrayCube:
        return "samplerCubeArray";
    case TextureType::Buffer:
        return "samplerBuffer";
    default:
        throw NotImplementedException("Texture type: {}", type);
    }
}

std::string_view ImageType(TextureType type) {
    switch (type) {
    case TextureType::Color1D:
        return "uimage1D";
    case TextureType::ColorArray1D:
        return "uimage1DArray";
    case TextureType::Color2D:
        return "uimage2D";
    case TextureType::ColorArray2D:
        return "uimage2DArray";
    case TextureType::Color3D:
        return "uimage3D";
    case TextureType::ColorCube:
        return "uimageCube";
    case TextureType::ColorArrayCube:
        return "uimageCubeArray";
    case TextureType::Buffer:
        return "uimageBuffer";
    default:
        throw NotImplementedException("Image type: {}", type);
    }
}

std::string_view ImageFormatString(ImageFormat format) {
    switch (format) {
    case ImageFormat::Typeless:
        return "";
    case ImageFormat::R8_UINT:
        return ",r8ui";
    case ImageFormat::R8_SINT:
        return ",r8i";
    case ImageFormat::R16_UINT:
        return ",r16ui";
    case ImageFormat::R16_SINT:
        return ",r16i";
    case ImageFormat::R32_UINT:
        return ",r32ui";
    case ImageFormat::R32G32_UINT:
        return ",rg32ui";
    case ImageFormat::R32G32B32A32_UINT:
        return ",rgba32ui";
    default:
        throw NotImplementedException("Image format: {}", format);
    }
}

std::string_view ImageAccessQualifier(bool is_written, bool is_read) {
    if (is_written && !is_read) {
        return "writeonly ";
    }
    if (is_read && !is_written) {
        return "readonly ";
    }
    return "";
}

std::string_view GetTessMode(TessPrimitive primitive) {
    switch (primitive) {
    case TessPrimitive::Triangles:
        return "triangles";
    case TessPrimitive::Quads:
        return "quads";
    case TessPrimitive::Isolines:
        return "isolines";
    }
    throw InvalidArgument("Invalid tessellation primitive {}", primitive);
}

std::string_view GetTessSpacing(TessSpacing spacing) {
    switch (spacing) {
    case TessSpacing::Equal:
        return "equal_spacing";
    case TessSpacing::FractionalOdd:
        return "fractional_odd_spacing";
    case TessSpacing::FractionalEven:
        return "fractional_even_spacing";
    }
    throw InvalidArgument("Invalid tessellation spacing {}", spacing);
}

std::string_view InputPrimitive(InputTopology topology) {
    switch (topology) {
    case InputTopology::Points:
        return "points";
    case InputTopology::Lines:
        return "lines";
    case InputTopology::LinesAdjacency:
        return "lines_adjacency";
    case InputTopology::Triangles:
        return "triangles";
    case InputTopology::TrianglesAdjacency:
        return "triangles_adjacency";
    }
    throw InvalidArgument("Invalid input topology {}", topology);
}

std::string_view OutputPrimitive(OutputTopology topology) {
    switch (topology) {
    case OutputTopology::PointList:
        return "points";
    case OutputTopology::LineStrip:
        return "line_strip";
    case OutputTopology::TriangleStrip:
        return "triangle_strip";
    }
    throw InvalidArgument("Invalid output topology {}", topology);
}

void SetupOutPerVertex(EmitContext& ctx, std::string& header) {
    if (!StoresPerVertexAttributes(ctx.stage)) {
        return;
    }
    if (ctx.uses_geometry_passthrough) {
        return;
    }
    header += "out gl_PerVertex{vec4 gl_Position;";
    if (ctx.info.stores[IR::Attribute::PointSize]) {
        header += "float gl_PointSize;";
    }
    if (ctx.info.stores.ClipDistances()) {
        header += "float gl_ClipDistance[];";
    }
    if (ctx.info.stores[IR::Attribute::ViewportIndex] &&
        ctx.profile.support_viewport_index_layer_non_geometry && ctx.stage != Stage::Geometry) {
        header += "int gl_ViewportIndex;";
    }
    header += "};";
    if (ctx.info.stores[IR::Attribute::ViewportIndex] && ctx.stage == Stage::Geometry) {
        header += "out int gl_ViewportIndex;";
    }
}

void SetupInPerVertex(EmitContext& ctx, std::string& header) {
    // Currently only required for TessellationControl to adhere to
    // ARB_separate_shader_objects requirements
    if (ctx.stage != Stage::TessellationControl) {
        return;
    }
    const bool loads_position{ctx.info.loads.AnyComponent(IR::Attribute::PositionX)};
    const bool loads_point_size{ctx.info.loads[IR::Attribute::PointSize]};
    const bool loads_clip_distance{ctx.info.loads.ClipDistances()};
    const bool loads_per_vertex{loads_position || loads_point_size || loads_clip_distance};
    if (!loads_per_vertex) {
        return;
    }
    header += "in gl_PerVertex{";
    if (loads_position) {
        header += "vec4 gl_Position;";
    }
    if (loads_point_size) {
        header += "float gl_PointSize;";
    }
    if (loads_clip_distance) {
        header += "float gl_ClipDistance[];";
    }
    header += "}gl_in[gl_MaxPatchVertices];";
}
} // Anonymous namespace

EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_,
                         const RuntimeInfo& runtime_info_)
    : info{program.info}, profile{profile_}, runtime_info{runtime_info_}, stage{program.stage},
      uses_geometry_passthrough{program.is_geometry_passthrough &&
                                profile.support_geometry_shader_passthrough} {
    if (profile.need_fastmath_off) {
        header += "#pragma optionNV(fastmath off)\n";
    }
    SetupExtensions();
    switch (program.stage) {
    case Stage::VertexA:
    case Stage::VertexB:
        stage_name = "vs";
        break;
    case Stage::TessellationControl:
        stage_name = "tcs";
        header += fmt::format("layout(vertices={})out;", program.invocations);
        break;
    case Stage::TessellationEval:
        stage_name = "tes";
        header += fmt::format("layout({},{},{})in;", GetTessMode(runtime_info.tess_primitive),
                              GetTessSpacing(runtime_info.tess_spacing),
                              runtime_info.tess_clockwise ? "cw" : "ccw");
        break;
    case Stage::Geometry:
        stage_name = "gs";
        header += fmt::format("layout({})in;", InputPrimitive(runtime_info.input_topology));
        if (uses_geometry_passthrough) {
            header += "layout(passthrough)in gl_PerVertex{vec4 gl_Position;};";
            break;
        } else if (program.is_geometry_passthrough &&
                   !profile.support_geometry_shader_passthrough) {
            LOG_WARNING(Shader_GLSL, "Passthrough geometry program used but not supported");
        }
        header += fmt::format(
            "layout({},max_vertices={})out;in gl_PerVertex{{vec4 gl_Position;}}gl_in[];",
            OutputPrimitive(program.output_topology), program.output_vertices);
        break;
    case Stage::Fragment:
        stage_name = "fs";
        position_name = "gl_FragCoord";
        if (runtime_info.force_early_z) {
            header += "layout(early_fragment_tests)in;";
        }
        break;
    case Stage::Compute:
        stage_name = "cs";
        const u32 local_x{std::max(program.workgroup_size[0], 1u)};
        const u32 local_y{std::max(program.workgroup_size[1], 1u)};
        const u32 local_z{std::max(program.workgroup_size[2], 1u)};
        header += fmt::format("layout(local_size_x={},local_size_y={},local_size_z={}) in;",
                              local_x, local_y, local_z);
        break;
    }
    SetupOutPerVertex(*this, header);
    SetupInPerVertex(*this, header);

    for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
        if (!info.loads.Generic(index) || !runtime_info.previous_stage_stores.Generic(index)) {
            continue;
        }
        const auto qualifier{uses_geometry_passthrough ? "passthrough"
                                                       : fmt::format("location={}", index)};
        header += fmt::format("layout({}){}in vec4 in_attr{}{};", qualifier,
                              InterpDecorator(info.interpolation[index]), index,
                              InputArrayDecorator(stage));
    }
    for (size_t index = 0; index < info.uses_patches.size(); ++index) {
        if (!info.uses_patches[index]) {
            continue;
        }
        const auto qualifier{stage == Stage::TessellationControl ? "out" : "in"};
        header += fmt::format("layout(location={})patch {} vec4 patch{};", index, qualifier, index);
    }
    if (stage == Stage::Fragment) {
        for (size_t index = 0; index < info.stores_frag_color.size(); ++index) {
            if (!info.stores_frag_color[index] && !profile.need_declared_frag_colors) {
                continue;
            }
            header += fmt::format("layout(location={})out vec4 frag_color{};", index, index);
        }
    }
    for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
        if (info.stores.Generic(index)) {
            DefineGenericOutput(index, program.invocations);
        }
    }
    if (info.uses_rescaling_uniform) {
        header += "layout(location=0) uniform vec4 scaling;";
    }
    if (info.uses_render_area) {
        header += "layout(location=1) uniform vec4 render_area;";
    }
    DefineConstantBuffers(bindings);
    DefineConstantBufferIndirect();
    DefineStorageBuffers(bindings);
    SetupImages(bindings);
    SetupTextures(bindings);
    DefineHelperFunctions();
    DefineConstants();
}

void EmitContext::SetupExtensions() {
    header += "#extension GL_ARB_separate_shader_objects : enable\n";
    if (info.uses_shadow_lod && profile.support_gl_texture_shadow_lod) {
        header += "#extension GL_EXT_texture_shadow_lod : enable\n";
    }
    if (info.uses_int64 && profile.support_int64) {
        header += "#extension GL_ARB_gpu_shader_int64 : enable\n";
    }
    if (info.uses_int64_bit_atomics) {
        header += "#extension GL_NV_shader_atomic_int64 : enable\n";
    }
    if (info.uses_atomic_f32_add) {
        header += "#extension GL_NV_shader_atomic_float : enable\n";
    }
    if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) {
        header += "#extension GL_NV_shader_atomic_fp16_vector : enable\n";
    }
    if (info.uses_fp16) {
        if (profile.support_gl_nv_gpu_shader_5) {
            header += "#extension GL_NV_gpu_shader5 : enable\n";
        }
        if (profile.support_gl_amd_gpu_shader_half_float) {
            header += "#extension GL_AMD_gpu_shader_half_float : enable\n";
        }
    }
    if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask || info.uses_subgroup_vote ||
        info.uses_subgroup_shuffles || info.uses_fswzadd) {
        header += "#extension GL_ARB_shader_ballot : enable\n"
                  "#extension GL_ARB_shader_group_vote : enable\n";
        if (!info.uses_int64 && profile.support_int64) {
            header += "#extension GL_ARB_gpu_shader_int64 : enable\n";
        }
        if (profile.support_gl_warp_intrinsics) {
            header += "#extension GL_NV_shader_thread_shuffle : enable\n";
        }
    }
    if ((info.stores[IR::Attribute::ViewportIndex] || info.stores[IR::Attribute::Layer]) &&
        profile.support_viewport_index_layer_non_geometry && stage != Stage::Geometry) {
        header += "#extension GL_ARB_shader_viewport_layer_array : enable\n";
    }
    if (info.uses_sparse_residency && profile.support_gl_sparse_textures) {
        header += "#extension GL_ARB_sparse_texture2 : enable\n";
    }
    if (info.stores[IR::Attribute::ViewportMask] && profile.support_viewport_mask) {
        header += "#extension GL_NV_viewport_array2 : enable\n";
    }
    if (info.uses_typeless_image_reads) {
        header += "#extension GL_EXT_shader_image_load_formatted : enable\n";
    }
    if (info.uses_derivatives && profile.support_gl_derivative_control) {
        header += "#extension GL_ARB_derivative_control : enable\n";
    }
    if (uses_geometry_passthrough) {
        header += "#extension GL_NV_geometry_shader_passthrough : enable\n";
    }
}

void EmitContext::DefineConstantBuffers(Bindings& bindings) {
    if (info.constant_buffer_descriptors.empty()) {
        return;
    }
    for (const auto& desc : info.constant_buffer_descriptors) {
        const auto cbuf_type{profile.has_gl_cbuf_ftou_bug ? "uvec4" : "vec4"};
        const u32 cbuf_used_size{Common::DivCeil(info.constant_buffer_used_sizes[desc.index], 16U)};
        const u32 cbuf_binding_size{info.uses_global_memory ? 0x1000U : cbuf_used_size};
        header += fmt::format("layout(std140,binding={}) uniform {}_cbuf_{}{{{} {}_cbuf{}[{}];}};",
                              bindings.uniform_buffer, stage_name, desc.index, cbuf_type,
                              stage_name, desc.index, cbuf_binding_size);
        bindings.uniform_buffer += desc.count;
    }
}

void EmitContext::DefineConstantBufferIndirect() {
    if (!info.uses_cbuf_indirect) {
        return;
    }

    header += profile.has_gl_cbuf_ftou_bug ? "uvec4 " : "vec4 ";
    header += "GetCbufIndirect(uint binding, uint offset){"
              "switch(binding){"
              "default:";

    for (const auto& desc : info.constant_buffer_descriptors) {
        header +=
            fmt::format("case {}:return {}_cbuf{}[offset];", desc.index, stage_name, desc.index);
    }

    header += "}}";
}

void EmitContext::DefineStorageBuffers(Bindings& bindings) {
    if (info.storage_buffers_descriptors.empty()) {
        return;
    }
    u32 index{};
    for (const auto& desc : info.storage_buffers_descriptors) {
        header += fmt::format("layout(std430,binding={}) buffer {}_ssbo_{}{{uint {}_ssbo{}[];}};",
                              bindings.storage_buffer, stage_name, bindings.storage_buffer,
                              stage_name, index);
        bindings.storage_buffer += desc.count;
        index += desc.count;
    }
}

void EmitContext::DefineGenericOutput(size_t index, u32 invocations) {
    static constexpr std::string_view swizzle{"xyzw"};
    const size_t base_index{static_cast<size_t>(IR::Attribute::Generic0X) + index * 4};
    u32 element{0};
    while (element < 4) {
        std::string definition{fmt::format("layout(location={}", index)};
        const u32 remainder{4 - element};
        const TransformFeedbackVarying* xfb_varying{};
        const size_t xfb_varying_index{base_index + element};
        if (xfb_varying_index < runtime_info.xfb_count) {
            xfb_varying = &runtime_info.xfb_varyings[xfb_varying_index];
            xfb_varying = xfb_varying->components > 0 ? xfb_varying : nullptr;
        }
        const u32 num_components{xfb_varying ? xfb_varying->components : remainder};
        if (element > 0) {
            definition += fmt::format(",component={}", element);
        }
        if (xfb_varying) {
            definition +=
                fmt::format(",xfb_buffer={},xfb_stride={},xfb_offset={}", xfb_varying->buffer,
                            xfb_varying->stride, xfb_varying->offset);
        }
        std::string name{fmt::format("out_attr{}", index)};
        if (num_components < 4 || element > 0) {
            name += fmt::format("_{}", swizzle.substr(element, num_components));
        }
        const auto type{num_components == 1 ? "float" : fmt::format("vec{}", num_components)};
        definition += fmt::format(")out {} {}{};", type, name, OutputDecorator(stage, invocations));
        header += definition;

        const GenericElementInfo element_info{
            .name = name,
            .first_element = element,
            .num_components = num_components,
        };
        std::fill_n(output_generics[index].begin() + element, num_components, element_info);
        element += num_components;
    }
}

void EmitContext::DefineHelperFunctions() {
    header += "\n#define ftoi floatBitsToInt\n#define ftou floatBitsToUint\n"
              "#define itof intBitsToFloat\n#define utof uintBitsToFloat\n";
    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);}";
    }
    if (info.uses_global_decrement || info.uses_shared_decrement) {
        header += "uint CasDecrement(uint op_a,uint op_b){"
                  "return op_a==0||op_a>op_b?op_b:(op_a-1u);}";
    }
    if (info.uses_atomic_f32_add) {
        header += "uint CasFloatAdd(uint op_a,float op_b){"
                  "return ftou(utof(op_a)+op_b);}";
    }
    if (info.uses_atomic_f32x2_add) {
        header += "uint CasFloatAdd32x2(uint op_a,vec2 op_b){"
                  "return packHalf2x16(unpackHalf2x16(op_a)+op_b);}";
    }
    if (info.uses_atomic_f32x2_min) {
        header += "uint CasFloatMin32x2(uint op_a,vec2 op_b){return "
                  "packHalf2x16(min(unpackHalf2x16(op_a),op_b));}";
    }
    if (info.uses_atomic_f32x2_max) {
        header += "uint CasFloatMax32x2(uint op_a,vec2 op_b){return "
                  "packHalf2x16(max(unpackHalf2x16(op_a),op_b));}";
    }
    if (info.uses_atomic_f16x2_add) {
        header += "uint CasFloatAdd16x2(uint op_a,f16vec2 op_b){return "
                  "packFloat2x16(unpackFloat2x16(op_a)+op_b);}";
    }
    if (info.uses_atomic_f16x2_min) {
        header += "uint CasFloatMin16x2(uint op_a,f16vec2 op_b){return "
                  "packFloat2x16(min(unpackFloat2x16(op_a),op_b));}";
    }
    if (info.uses_atomic_f16x2_max) {
        header += "uint CasFloatMax16x2(uint op_a,f16vec2 op_b){return "
                  "packFloat2x16(max(unpackFloat2x16(op_a),op_b));}";
    }
    if (info.uses_atomic_s32_min) {
        header += "uint CasMinS32(uint op_a,uint op_b){return uint(min(int(op_a),int(op_b)));}";
    }
    if (info.uses_atomic_s32_max) {
        header += "uint CasMaxS32(uint op_a,uint op_b){return uint(max(int(op_a),int(op_b)));}";
    }
    if (info.uses_global_memory && profile.support_int64) {
        header += DefineGlobalMemoryFunctions();
    }
    if (info.loads_indexed_attributes) {
        const bool is_array{stage == Stage::Geometry};
        const auto vertex_arg{is_array ? ",uint vertex" : ""};
        std::string func{
            fmt::format("float IndexedAttrLoad(int offset{}){{int base_index=offset>>2;uint "
                        "masked_index=uint(base_index)&3u;switch(base_index>>2){{",
                        vertex_arg)};
        if (info.loads.AnyComponent(IR::Attribute::PositionX)) {
            const auto position_idx{is_array ? "gl_in[vertex]." : ""};
            func += fmt::format("case {}:return {}{}[masked_index];",
                                static_cast<u32>(IR::Attribute::PositionX) >> 2, position_idx,
                                position_name);
        }
        const u32 base_attribute_value = static_cast<u32>(IR::Attribute::Generic0X) >> 2;
        for (u32 index = 0; index < IR::NUM_GENERICS; ++index) {
            if (!info.loads.Generic(index)) {
                continue;
            }
            const auto vertex_idx{is_array ? "[vertex]" : ""};
            func += fmt::format("case {}:return in_attr{}{}[masked_index];",
                                base_attribute_value + index, index, vertex_idx);
        }
        func += "default: return 0.0;}}";
        header += func;
    }
    if (info.stores_indexed_attributes) {
        // TODO
    }
}

std::string EmitContext::DefineGlobalMemoryFunctions() {
    const auto define_body{[&](std::string& func, size_t index, std::string_view return_statement) {
        const auto& ssbo{info.storage_buffers_descriptors[index]};
        const u32 size_cbuf_offset{ssbo.cbuf_offset + 8};
        const auto ssbo_addr{fmt::format("ssbo_addr{}", index)};
        const auto cbuf{fmt::format("{}_cbuf{}", stage_name, ssbo.cbuf_index)};
        std::array<std::string, 2> addr_xy;
        std::array<std::string, 2> size_xy;
        for (size_t i = 0; i < addr_xy.size(); ++i) {
            const auto addr_loc{ssbo.cbuf_offset + 4 * i};
            const auto size_loc{size_cbuf_offset + 4 * i};
            addr_xy[i] = fmt::format("ftou({}[{}].{})", cbuf, addr_loc / 16, Swizzle(addr_loc));
            size_xy[i] = fmt::format("ftou({}[{}].{})", cbuf, size_loc / 16, Swizzle(size_loc));
        }
        const u32 ssbo_align_mask{~(static_cast<u32>(profile.min_ssbo_alignment) - 1U)};
        const auto aligned_low_addr{fmt::format("{}&{}", addr_xy[0], ssbo_align_mask)};
        const auto aligned_addr{fmt::format("uvec2({},{})", aligned_low_addr, addr_xy[1])};
        const auto addr_pack{fmt::format("packUint2x32({})", aligned_addr)};
        const auto addr_statement{fmt::format("uint64_t {}={};", ssbo_addr, addr_pack)};
        func += addr_statement;

        const auto size_vec{fmt::format("uvec2({},{})", size_xy[0], size_xy[1])};
        const auto comp_lhs{fmt::format("(addr>={})", ssbo_addr)};
        const auto comp_rhs{fmt::format("(addr<({}+uint64_t({})))", ssbo_addr, size_vec)};
        const auto comparison{fmt::format("if({}&&{}){{", comp_lhs, comp_rhs)};
        func += comparison;

        const auto ssbo_name{fmt::format("{}_ssbo{}", stage_name, index)};
        func += fmt::format(fmt::runtime(return_statement), ssbo_name, ssbo_addr);
    }};
    std::string write_func{"void WriteGlobal32(uint64_t addr,uint data){"};
    std::string write_func_64{"void WriteGlobal64(uint64_t addr,uvec2 data){"};
    std::string write_func_128{"void WriteGlobal128(uint64_t addr,uvec4 data){"};
    std::string load_func{"uint LoadGlobal32(uint64_t addr){"};
    std::string load_func_64{"uvec2 LoadGlobal64(uint64_t addr){"};
    std::string load_func_128{"uvec4 LoadGlobal128(uint64_t addr){"};
    const size_t num_buffers{info.storage_buffers_descriptors.size()};
    for (size_t index = 0; index < num_buffers; ++index) {
        if (!info.nvn_buffer_used[index]) {
            continue;
        }
        define_body(write_func, index, "{0}[uint(addr-{1})>>2]=data;return;}}");
        define_body(write_func_64, index,
                    "{0}[uint(addr-{1})>>2]=data.x;{0}[uint(addr-{1}+4)>>2]=data.y;return;}}");
        define_body(write_func_128, index,
                    "{0}[uint(addr-{1})>>2]=data.x;{0}[uint(addr-{1}+4)>>2]=data.y;{0}[uint("
                    "addr-{1}+8)>>2]=data.z;{0}[uint(addr-{1}+12)>>2]=data.w;return;}}");
        define_body(load_func, index, "return {0}[uint(addr-{1})>>2];}}");
        define_body(load_func_64, index,
                    "return uvec2({0}[uint(addr-{1})>>2],{0}[uint(addr-{1}+4)>>2]);}}");
        define_body(load_func_128, index,
                    "return uvec4({0}[uint(addr-{1})>>2],{0}[uint(addr-{1}+4)>>2],{0}["
                    "uint(addr-{1}+8)>>2],{0}[uint(addr-{1}+12)>>2]);}}");
    }
    write_func += '}';
    write_func_64 += '}';
    write_func_128 += '}';
    load_func += "return 0u;}";
    load_func_64 += "return uvec2(0);}";
    load_func_128 += "return uvec4(0);}";
    return write_func + write_func_64 + write_func_128 + load_func + load_func_64 + load_func_128;
}

void EmitContext::SetupImages(Bindings& bindings) {
    image_buffers.reserve(info.image_buffer_descriptors.size());
    for (const auto& desc : info.image_buffer_descriptors) {
        image_buffers.push_back({bindings.image, desc.count});
        const auto format{ImageFormatString(desc.format)};
        const auto qualifier{ImageAccessQualifier(desc.is_written, desc.is_read)};
        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
        header += fmt::format("layout(binding={}{}) uniform {}uimageBuffer img{}{};",
                              bindings.image, format, qualifier, bindings.image, array_decorator);
        bindings.image += desc.count;
    }
    images.reserve(info.image_descriptors.size());
    for (const auto& desc : info.image_descriptors) {
        images.push_back({bindings.image, desc.count});
        const auto format{ImageFormatString(desc.format)};
        const auto image_type{ImageType(desc.type)};
        const auto qualifier{ImageAccessQualifier(desc.is_written, desc.is_read)};
        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
        header += fmt::format("layout(binding={}{})uniform {}{} img{}{};", bindings.image, format,
                              qualifier, image_type, bindings.image, array_decorator);
        bindings.image += desc.count;
    }
}

void EmitContext::SetupTextures(Bindings& bindings) {
    texture_buffers.reserve(info.texture_buffer_descriptors.size());
    for (const auto& desc : info.texture_buffer_descriptors) {
        texture_buffers.push_back({bindings.texture, desc.count});
        const auto sampler_type{ColorSamplerType(TextureType::Buffer)};
        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
        header += fmt::format("layout(binding={}) uniform {} tex{}{};", bindings.texture,
                              sampler_type, bindings.texture, array_decorator);
        bindings.texture += desc.count;
    }
    textures.reserve(info.texture_descriptors.size());
    for (const auto& desc : info.texture_descriptors) {
        textures.push_back({bindings.texture, desc.count});
        const auto sampler_type{desc.is_depth ? DepthSamplerType(desc.type)
                                              : ColorSamplerType(desc.type, desc.is_multisample)};
        const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""};
        header += fmt::format("layout(binding={}) uniform {} tex{}{};", bindings.texture,
                              sampler_type, bindings.texture, array_decorator);
        bindings.texture += desc.count;
    }
}

void EmitContext::DefineConstants() {
    if (info.uses_fswzadd) {
        header += "const float FSWZ_A[]=float[4](-1.f,1.f,-1.f,0.f);"
                  "const float FSWZ_B[]=float[4](-1.f,-1.f,1.f,-1.f);";
    }
}

} // namespace Shader::Backend::GLSL


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#pragma once

#include <string>
#include <utility>
#include <vector>

#include <fmt/format.h>

#include "shader_recompiler/backend/glsl/var_alloc.h"
#include "shader_recompiler/stage.h"

namespace Shader {
struct Info;
struct Profile;
struct RuntimeInfo;
} // namespace Shader

namespace Shader::Backend {
struct Bindings;
}

namespace Shader::IR {
class Inst;
struct Program;
} // namespace Shader::IR

namespace Shader::Backend::GLSL {

struct GenericElementInfo {
    std::string name;
    u32 first_element{};
    u32 num_components{};
};

struct TextureImageDefinition {
    u32 binding;
    u32 count;
};

class EmitContext {
public:
    explicit EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_,
                         const RuntimeInfo& runtime_info_);

    template <GlslVarType type, typename... Args>
    void Add(const char* format_str, IR::Inst& inst, Args&&... args) {
        const auto var_def{var_alloc.AddDefine(inst, type)};
        if (var_def.empty()) {
            // skip assignment.
            code += fmt::format(fmt::runtime(format_str + 3), std::forward<Args>(args)...);
        } else {
            code += fmt::format(fmt::runtime(format_str), var_def, std::forward<Args>(args)...);
        }
        // TODO: Remove this
        code += '\n';
    }

    template <typename... Args>
    void AddU1(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::U1>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddF16x2(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::F16x2>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddU32(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::U32>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddF32(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::F32>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddU64(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::U64>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddF64(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::F64>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddU32x2(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::U32x2>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddF32x2(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::F32x2>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddU32x3(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::U32x3>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddF32x3(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::F32x3>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddU32x4(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::U32x4>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddF32x4(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::F32x4>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddPrecF32(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::PrecF32>(format_str, inst, args...);
    }

    template <typename... Args>
    void AddPrecF64(const char* format_str, IR::Inst& inst, Args&&... args) {
        Add<GlslVarType::PrecF64>(format_str, inst, args...);
    }

    template <typename... Args>
    void Add(const char* format_str, Args&&... args) {
        code += fmt::format(fmt::runtime(format_str), std::forward<Args>(args)...);
        // TODO: Remove this
        code += '\n';
    }

    std::string header;
    std::string code;
    VarAlloc var_alloc;
    const Info& info;
    const Profile& profile;
    const RuntimeInfo& runtime_info;

    Stage stage{};
    std::string_view stage_name = "invalid";
    std::string_view position_name = "gl_Position";

    std::vector<TextureImageDefinition> texture_buffers;
    std::vector<TextureImageDefinition> image_buffers;
    std::vector<TextureImageDefinition> textures;
    std::vector<TextureImageDefinition> images;
    std::array<std::array<GenericElementInfo, 4>, 32> output_generics{};

    u32 num_safety_loop_vars{};

    bool uses_y_direction{};
    bool uses_cc_carry{};
    bool uses_geometry_passthrough{};

private:
    void SetupExtensions();
    void DefineConstantBuffers(Bindings& bindings);
    void DefineConstantBufferIndirect();
    void DefineStorageBuffers(Bindings& bindings);
    void DefineGenericOutput(size_t index, u32 invocations);
    void DefineHelperFunctions();
    void DefineConstants();
    std::string DefineGlobalMemoryFunctions();
    void SetupImages(Bindings& bindings);
    void SetupTextures(Bindings& bindings);
};

} // namespace Shader::Backend::GLSL


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include <string>
#include <string_view>

#include <fmt/format.h>

#include "shader_recompiler/backend/glsl/var_alloc.h"
#include "shader_recompiler/exception.h"
#include "shader_recompiler/frontend/ir/value.h"

namespace Shader::Backend::GLSL {
namespace {
std::string TypePrefix(GlslVarType type) {
    switch (type) {
    case GlslVarType::U1:
        return "b_";
    case GlslVarType::F16x2:
        return "f16x2_";
    case GlslVarType::U32:
        return "u_";
    case GlslVarType::F32:
        return "f_";
    case GlslVarType::U64:
        return "u64_";
    case GlslVarType::F64:
        return "d_";
    case GlslVarType::U32x2:
        return "u2_";
    case GlslVarType::F32x2:
        return "f2_";
    case GlslVarType::U32x3:
        return "u3_";
    case GlslVarType::F32x3:
        return "f3_";
    case GlslVarType::U32x4:
        return "u4_";
    case GlslVarType::F32x4:
        return "f4_";
    case GlslVarType::PrecF32:
        return "pf_";
    case GlslVarType::PrecF64:
        return "pd_";
    case GlslVarType::Void:
        return "";
    default:
        throw NotImplementedException("Type {}", type);
    }
}

std::string FormatFloat(std::string_view value, IR::Type type) {
    // TODO: Confirm FP64 nan/inf
    if (type == IR::Type::F32) {
        if (value == "nan") {
            return "utof(0x7fc00000)";
        }
        if (value == "inf") {
            return "utof(0x7f800000)";
        }
        if (value == "-inf") {
            return "utof(0xff800000)";
        }
    }
    if (value.find_first_of('e') != std::string_view::npos) {
        // scientific notation
        const auto cast{type == IR::Type::F32 ? "float" : "double"};
        return fmt::format("{}({})", cast, value);
    }
    const bool needs_dot{value.find_first_of('.') == std::string_view::npos};
    const bool needs_suffix{!value.ends_with('f')};
    const auto suffix{type == IR::Type::F32 ? "f" : "lf"};
    return fmt::format("{}{}{}", value, needs_dot ? "." : "", needs_suffix ? suffix : "");
}

std::string MakeImm(const IR::Value& value) {
    switch (value.Type()) {
    case IR::Type::U1:
        return fmt::format("{}", value.U1() ? "true" : "false");
    case IR::Type::U32:
        return fmt::format("{}u", value.U32());
    case IR::Type::F32:
        return FormatFloat(fmt::format("{}", value.F32()), IR::Type::F32);
    case IR::Type::U64:
        return fmt::format("{}ul", value.U64());
    case IR::Type::F64:
        return FormatFloat(fmt::format("{}", value.F64()), IR::Type::F64);
    case IR::Type::Void:
        return "";
    default:
        throw NotImplementedException("Immediate type {}", value.Type());
    }
}
} // Anonymous namespace

std::string VarAlloc::Representation(u32 index, GlslVarType type) const {
    const auto prefix{TypePrefix(type)};
    return fmt::format("{}{}", prefix, index);
}

std::string VarAlloc::Representation(Id id) const {
    return Representation(id.index, id.type);
}

std::string VarAlloc::Define(IR::Inst& inst, GlslVarType type) {
    if (inst.HasUses()) {
        inst.SetDefinition<Id>(Alloc(type));
        return Representation(inst.Definition<Id>());
    } else {
        Id id{};
        id.type.Assign(type);
        GetUseTracker(type).uses_temp = true;
        inst.SetDefinition<Id>(id);
        return 't' + Representation(inst.Definition<Id>());
    }
}

std::string VarAlloc::Define(IR::Inst& inst, IR::Type type) {
    return Define(inst, RegType(type));
}

std::string VarAlloc::PhiDefine(IR::Inst& inst, IR::Type type) {
    return AddDefine(inst, RegType(type));
}

std::string VarAlloc::AddDefine(IR::Inst& inst, GlslVarType type) {
    if (inst.HasUses()) {
        inst.SetDefinition<Id>(Alloc(type));
        return Representation(inst.Definition<Id>());
    } else {
        return "";
    }
}

std::string VarAlloc::Consume(const IR::Value& value) {
    return value.IsImmediate() ? MakeImm(value) : ConsumeInst(*value.InstRecursive());
}

std::string VarAlloc::ConsumeInst(IR::Inst& inst) {
    inst.DestructiveRemoveUsage();
    if (!inst.HasUses()) {
        Free(inst.Definition<Id>());
    }
    return Representation(inst.Definition<Id>());
}

std::string VarAlloc::GetGlslType(IR::Type type) const {
    return GetGlslType(RegType(type));
}

Id VarAlloc::Alloc(GlslVarType type) {
    auto& use_tracker{GetUseTracker(type)};
    const auto num_vars{use_tracker.var_use.size()};
    for (size_t var = 0; var < num_vars; ++var) {
        if (use_tracker.var_use[var]) {
            continue;
        }
        use_tracker.num_used = std::max(use_tracker.num_used, var + 1);
        use_tracker.var_use[var] = true;
        Id ret{};
        ret.is_valid.Assign(1);
        ret.type.Assign(type);
        ret.index.Assign(static_cast<u32>(var));
        return ret;
    }
    // Allocate a new variable
    use_tracker.var_use.push_back(true);
    Id ret{};
    ret.is_valid.Assign(1);
    ret.type.Assign(type);
    ret.index.Assign(static_cast<u32>(use_tracker.num_used));
    ++use_tracker.num_used;
    return ret;
}

void VarAlloc::Free(Id id) {
    if (id.is_valid == 0) {
        throw LogicError("Freeing invalid variable");
    }
    auto& use_tracker{GetUseTracker(id.type)};
    use_tracker.var_use[id.index] = false;
}

GlslVarType VarAlloc::RegType(IR::Type type) const {
    switch (type) {
    case IR::Type::U1:
        return GlslVarType::U1;
    case IR::Type::U32:
        return GlslVarType::U32;
    case IR::Type::F32:
        return GlslVarType::F32;
    case IR::Type::U64:
        return GlslVarType::U64;
    case IR::Type::F64:
        return GlslVarType::F64;
    default:
        throw NotImplementedException("IR type {}", type);
    }
}

std::string VarAlloc::GetGlslType(GlslVarType type) const {
    switch (type) {
    case GlslVarType::U1:
        return "bool";
    case GlslVarType::F16x2:
        return "f16vec2";
    case GlslVarType::U32:
        return "uint";
    case GlslVarType::F32:
    case GlslVarType::PrecF32:
        return "float";
    case GlslVarType::U64:
        return "uint64_t";
    case GlslVarType::F64:
    case GlslVarType::PrecF64:
        return "double";
    case GlslVarType::U32x2:
        return "uvec2";
    case GlslVarType::F32x2:
        return "vec2";
    case GlslVarType::U32x3:
        return "uvec3";
    case GlslVarType::F32x3:
        return "vec3";
    case GlslVarType::U32x4:
        return "uvec4";
    case GlslVarType::F32x4:
        return "vec4";
    case GlslVarType::Void:
        return "";
    default:
        throw NotImplementedException("Type {}", type);
    }
}

VarAlloc::UseTracker& VarAlloc::GetUseTracker(GlslVarType type) {
    switch (type) {
    case GlslVarType::U1:
        return var_bool;
    case GlslVarType::F16x2:
        return var_f16x2;
    case GlslVarType::U32:
        return var_u32;
    case GlslVarType::F32:
        return var_f32;
    case GlslVarType::U64:
        return var_u64;
    case GlslVarType::F64:
        return var_f64;
    case GlslVarType::U32x2:
        return var_u32x2;
    case GlslVarType::F32x2:
        return var_f32x2;
    case GlslVarType::U32x3:
        return var_u32x3;
    case GlslVarType::F32x3:
        return var_f32x3;
    case GlslVarType::U32x4:
        return var_u32x4;
    case GlslVarType::F32x4:
        return var_f32x4;
    case GlslVarType::PrecF32:
        return var_precf32;
    case GlslVarType::PrecF64:
        return var_precf64;
    default:
        throw NotImplementedException("Type {}", type);
    }
}

const VarAlloc::UseTracker& VarAlloc::GetUseTracker(GlslVarType type) const {
    switch (type) {
    case GlslVarType::U1:
        return var_bool;
    case GlslVarType::F16x2:
        return var_f16x2;
    case GlslVarType::U32:
        return var_u32;
    case GlslVarType::F32:
        return var_f32;
    case GlslVarType::U64:
        return var_u64;
    case GlslVarType::F64:
        return var_f64;
    case GlslVarType::U32x2:
        return var_u32x2;
    case GlslVarType::F32x2:
        return var_f32x2;
    case GlslVarType::U32x3:
        return var_u32x3;
    case GlslVarType::F32x3:
        return var_f32x3;
    case GlslVarType::U32x4:
        return var_u32x4;
    case GlslVarType::F32x4:
        return var_f32x4;
    case GlslVarType::PrecF32:
        return var_precf32;
    case GlslVarType::PrecF64:
        return var_precf64;
    default:
        throw NotImplementedException("Type {}", type);
    }
}

} // namespace Shader::Backend::GLSL


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#pragma once

#include <bitset>
#include <string>
#include <vector>

#include "common/bit_field.h"
#include "common/common_types.h"

namespace Shader::IR {
class Inst;
class Value;
enum class Type;
} // namespace Shader::IR

namespace Shader::Backend::GLSL {
enum class GlslVarType : u32 {
    U1,
    F16x2,
    U32,
    F32,
    U64,
    F64,
    U32x2,
    F32x2,
    U32x3,
    F32x3,
    U32x4,
    F32x4,
    PrecF32,
    PrecF64,
    Void,
};

struct Id {
    union {
        u32 raw;
        BitField<0, 1, u32> is_valid;
        BitField<1, 4, GlslVarType> type;
        BitField<6, 26, u32> index;
    };

    bool operator==(Id rhs) const noexcept {
        return raw == rhs.raw;
    }
    bool operator!=(Id rhs) const noexcept {
        return !operator==(rhs);
    }
};
static_assert(sizeof(Id) == sizeof(u32));

class VarAlloc {
public:
    struct UseTracker {
        bool uses_temp{};
        size_t num_used{};
        std::vector<bool> var_use;
    };

    /// Used for explicit usages of variables, may revert to temporaries
    std::string Define(IR::Inst& inst, GlslVarType type);
    std::string Define(IR::Inst& inst, IR::Type type);

    /// Used to assign variables used by the IR. May return a blank string if
    /// the instruction's result is unused in the IR.
    std::string AddDefine(IR::Inst& inst, GlslVarType type);
    std::string PhiDefine(IR::Inst& inst, IR::Type type);

    std::string Consume(const IR::Value& value);
    std::string ConsumeInst(IR::Inst& inst);

    std::string GetGlslType(GlslVarType type) const;
    std::string GetGlslType(IR::Type type) const;

    const UseTracker& GetUseTracker(GlslVarType type) const;
    std::string Representation(u32 index, GlslVarType type) const;

private:
    GlslVarType RegType(IR::Type type) const;
    Id Alloc(GlslVarType type);
    void Free(Id id);
    UseTracker& GetUseTracker(GlslVarType type);
    std::string Representation(Id id) const;

    UseTracker var_bool{};
    UseTracker var_f16x2{};
    UseTracker var_u32{};
    UseTracker var_u32x2{};
    UseTracker var_u32x3{};
    UseTracker var_u32x4{};
    UseTracker var_f32{};
    UseTracker var_f32x2{};
    UseTracker var_f32x3{};
    UseTracker var_f32x4{};
    UseTracker var_u64{};
    UseTracker var_f64{};
    UseTracker var_precf32{};
    UseTracker var_precf64{};
};

} // namespace Shader::Backend::GLSL


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include <span>
#include <tuple>
#include <type_traits>
#include <utility>
#include <vector>

#include "common/settings.h"
#include "shader_recompiler/backend/spirv/emit_spirv.h"
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
#include "shader_recompiler/frontend/ir/basic_block.h"
#include "shader_recompiler/frontend/ir/program.h"

namespace Shader::Backend::SPIRV {
namespace {
template <class Func>
struct FuncTraits {};

template <class ReturnType_, class... Args>
struct FuncTraits<ReturnType_ (*)(Args...)> {
    using ReturnType = ReturnType_;

    static constexpr size_t NUM_ARGS = sizeof...(Args);

    template <size_t I>
    using ArgType = std::tuple_element_t<I, std::tuple<Args...>>;
};

#ifdef _MSC_VER
#pragma warning(push)
#pragma warning(disable : 4702) // Ignore unreachable code warning
#endif

template <auto func, typename... Args>
void SetDefinition(EmitContext& ctx, IR::Inst* inst, Args... args) {
    inst->SetDefinition<Id>(func(ctx, std::forward<Args>(args)...));
}

#ifdef _MSC_VER
#pragma warning(pop)
#endif

template <typename ArgType>
ArgType Arg(EmitContext& ctx, const IR::Value& arg) {
    if constexpr (std::is_same_v<ArgType, Id>) {
        return ctx.Def(arg);
    } else if constexpr (std::is_same_v<ArgType, const IR::Value&>) {
        return arg;
    } else if constexpr (std::is_same_v<ArgType, u32>) {
        return arg.U32();
    } else if constexpr (std::is_same_v<ArgType, IR::Attribute>) {
        return arg.Attribute();
    } else if constexpr (std::is_same_v<ArgType, IR::Patch>) {
        return arg.Patch();
    } else if constexpr (std::is_same_v<ArgType, IR::Reg>) {
        return arg.Reg();
    }
}

template <auto func, bool is_first_arg_inst, size_t... I>
void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) {
    using Traits = FuncTraits<decltype(func)>;
    if constexpr (std::is_same_v<typename Traits::ReturnType, Id>) {
        if constexpr (is_first_arg_inst) {
            SetDefinition<func>(
                ctx, inst, inst,
                Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...);
        } else {
            SetDefinition<func>(
                ctx, inst, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I))...);
        }
    } else {
        if constexpr (is_first_arg_inst) {
            func(ctx, inst, Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...);
        } else {
            func(ctx, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I))...);
        }
    }
}

template <auto func>
void Invoke(EmitContext& ctx, IR::Inst* inst) {
    using Traits = FuncTraits<decltype(func)>;
    static_assert(Traits::NUM_ARGS >= 1, "Insufficient arguments");
    if constexpr (Traits::NUM_ARGS == 1) {
        Invoke<func, false>(ctx, inst, std::make_index_sequence<0>{});
    } else {
        using FirstArgType = typename Traits::template ArgType<1>;
        static constexpr bool is_first_arg_inst = std::is_same_v<FirstArgType, IR::Inst*>;
        using Indices = std::make_index_sequence<Traits::NUM_ARGS - (is_first_arg_inst ? 2 : 1)>;
        Invoke<func, is_first_arg_inst>(ctx, inst, Indices{});
    }
}

void EmitInst(EmitContext& ctx, IR::Inst* inst) {
    switch (inst->GetOpcode()) {
#define OPCODE(name, result_type, ...)                                                             \
    case IR::Opcode::name:                                                                         \
        return Invoke<&Emit##name>(ctx, inst);
#include "shader_recompiler/frontend/ir/opcodes.inc"
#undef OPCODE
    }
    throw LogicError("Invalid opcode {}", inst->GetOpcode());
}

Id TypeId(const EmitContext& ctx, IR::Type type) {
    switch (type) {
    case IR::Type::U1:
        return ctx.U1;
    case IR::Type::U32:
        return ctx.U32[1];
    default:
        throw NotImplementedException("Phi node type {}", type);
    }
}

void Traverse(EmitContext& ctx, IR::Program& program) {
    IR::Block* current_block{};
    for (const IR::AbstractSyntaxNode& node : program.syntax_list) {
        switch (node.type) {
        case IR::AbstractSyntaxNode::Type::Block: {
            const Id label{node.data.block->Definition<Id>()};
            if (current_block) {
                ctx.OpBranch(label);
            }
            current_block = node.data.block;
            ctx.AddLabel(label);
            for (IR::Inst& inst : node.data.block->Instructions()) {
                EmitInst(ctx, &inst);
            }
            break;
        }
        case IR::AbstractSyntaxNode::Type::If: {
            const Id if_label{node.data.if_node.body->Definition<Id>()};
            const Id endif_label{node.data.if_node.merge->Definition<Id>()};
            ctx.OpSelectionMerge(endif_label, spv::SelectionControlMask::MaskNone);
            ctx.OpBranchConditional(ctx.Def(node.data.if_node.cond), if_label, endif_label);
            break;
        }
        case IR::AbstractSyntaxNode::Type::Loop: {
            const Id body_label{node.data.loop.body->Definition<Id>()};
            const Id continue_label{node.data.loop.continue_block->Definition<Id>()};
            const Id endloop_label{node.data.loop.merge->Definition<Id>()};

            ctx.OpLoopMerge(endloop_label, continue_label, spv::LoopControlMask::MaskNone);
            ctx.OpBranch(body_label);
            break;
        }
        case IR::AbstractSyntaxNode::Type::Break: {
            const Id break_label{node.data.break_node.merge->Definition<Id>()};
            const Id skip_label{node.data.break_node.skip->Definition<Id>()};
            ctx.OpBranchConditional(ctx.Def(node.data.break_node.cond), break_label, skip_label);
            break;
        }
        case IR::AbstractSyntaxNode::Type::EndIf:
            if (current_block) {
                ctx.OpBranch(node.data.end_if.merge->Definition<Id>());
            }
            break;
        case IR::AbstractSyntaxNode::Type::Repeat: {
            Id cond{ctx.Def(node.data.repeat.cond)};
            if (!Settings::values.disable_shader_loop_safety_checks) {
                const Id pointer_type{ctx.TypePointer(spv::StorageClass::Private, ctx.U32[1])};
                const Id safety_counter{ctx.AddGlobalVariable(
                    pointer_type, spv::StorageClass::Private, ctx.Const(0x2000u))};
                if (ctx.profile.supported_spirv >= 0x00010400) {
                    ctx.interfaces.push_back(safety_counter);
                }
                const Id old_counter{ctx.OpLoad(ctx.U32[1], safety_counter)};
                const Id new_counter{ctx.OpISub(ctx.U32[1], old_counter, ctx.Const(1u))};
                ctx.OpStore(safety_counter, new_counter);

                const Id safety_cond{
                    ctx.OpSGreaterThanEqual(ctx.U1, new_counter, ctx.u32_zero_value)};
                cond = ctx.OpLogicalAnd(ctx.U1, cond, safety_cond);
            }
            const Id loop_header_label{node.data.repeat.loop_header->Definition<Id>()};
            const Id merge_label{node.data.repeat.merge->Definition<Id>()};
            ctx.OpBranchConditional(cond, loop_header_label, merge_label);
            break;
        }
        case IR::AbstractSyntaxNode::Type::Return:
            ctx.OpReturn();
            break;
        case IR::AbstractSyntaxNode::Type::Unreachable:
            ctx.OpUnreachable();
            break;
        }
        if (node.type != IR::AbstractSyntaxNode::Type::Block) {
            current_block = nullptr;
        }
    }
}

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) {
        block->SetDefinition(ctx.OpLabel());
    }
    Traverse(ctx, program);
    ctx.OpFunctionEnd();
    return main;
}

spv::ExecutionMode ExecutionMode(TessPrimitive primitive) {
    switch (primitive) {
    case TessPrimitive::Isolines:
        return spv::ExecutionMode::Isolines;
    case TessPrimitive::Triangles:
        return spv::ExecutionMode::Triangles;
    case TessPrimitive::Quads:
        return spv::ExecutionMode::Quads;
    }
    throw InvalidArgument("Tessellation primitive {}", primitive);
}

spv::ExecutionMode ExecutionMode(TessSpacing spacing) {
    switch (spacing) {
    case TessSpacing::Equal:
        return spv::ExecutionMode::SpacingEqual;
    case TessSpacing::FractionalOdd:
        return spv::ExecutionMode::SpacingFractionalOdd;
    case TessSpacing::FractionalEven:
        return spv::ExecutionMode::SpacingFractionalEven;
    }
    throw InvalidArgument("Tessellation spacing {}", spacing);
}

void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
    const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size());
    spv::ExecutionModel execution_model{};
    switch (program.stage) {
    case Stage::Compute: {
        const std::array<u32, 3> workgroup_size{program.workgroup_size};
        execution_model = spv::ExecutionModel::GLCompute;
        ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, workgroup_size[0],
                             workgroup_size[1], workgroup_size[2]);
        break;
    }
    case Stage::VertexB:
        execution_model = spv::ExecutionModel::Vertex;
        break;
    case Stage::TessellationControl:
        execution_model = spv::ExecutionModel::TessellationControl;
        ctx.AddCapability(spv::Capability::Tessellation);
        ctx.AddExecutionMode(main, spv::ExecutionMode::OutputVertices, program.invocations);
        break;
    case Stage::TessellationEval:
        execution_model = spv::ExecutionModel::TessellationEvaluation;
        ctx.AddCapability(spv::Capability::Tessellation);
        ctx.AddExecutionMode(main, ExecutionMode(ctx.runtime_info.tess_primitive));
        ctx.AddExecutionMode(main, ExecutionMode(ctx.runtime_info.tess_spacing));
        ctx.AddExecutionMode(main, ctx.runtime_info.tess_clockwise
                                       ? spv::ExecutionMode::VertexOrderCw
                                       : spv::ExecutionMode::VertexOrderCcw);
        break;
    case Stage::Geometry:
        execution_model = spv::ExecutionModel::Geometry;
        ctx.AddCapability(spv::Capability::Geometry);
        if (ctx.profile.support_geometry_streams) {
            ctx.AddCapability(spv::Capability::GeometryStreams);
        }
        switch (ctx.runtime_info.input_topology) {
        case InputTopology::Points:
            ctx.AddExecutionMode(main, spv::ExecutionMode::InputPoints);
            break;
        case InputTopology::Lines:
            ctx.AddExecutionMode(main, spv::ExecutionMode::InputLines);
            break;
        case InputTopology::LinesAdjacency:
            ctx.AddExecutionMode(main, spv::ExecutionMode::InputLinesAdjacency);
            break;
        case InputTopology::Triangles:
            ctx.AddExecutionMode(main, spv::ExecutionMode::Triangles);
            break;
        case InputTopology::TrianglesAdjacency:
            ctx.AddExecutionMode(main, spv::ExecutionMode::InputTrianglesAdjacency);
            break;
        }
        switch (program.output_topology) {
        case OutputTopology::PointList:
            ctx.AddExecutionMode(main, spv::ExecutionMode::OutputPoints);
            break;
        case OutputTopology::LineStrip:
            ctx.AddExecutionMode(main, spv::ExecutionMode::OutputLineStrip);
            break;
        case OutputTopology::TriangleStrip:
            ctx.AddExecutionMode(main, spv::ExecutionMode::OutputTriangleStrip);
            break;
        }
        if (program.info.stores[IR::Attribute::PointSize]) {
            ctx.AddCapability(spv::Capability::GeometryPointSize);
        }
        ctx.AddExecutionMode(main, spv::ExecutionMode::OutputVertices, program.output_vertices);
        ctx.AddExecutionMode(main, spv::ExecutionMode::Invocations, program.invocations);
        if (program.is_geometry_passthrough) {
            if (ctx.profile.support_geometry_shader_passthrough) {
                ctx.AddExtension("SPV_NV_geometry_shader_passthrough");
                ctx.AddCapability(spv::Capability::GeometryShaderPassthroughNV);
            } else {
                LOG_WARNING(Shader_SPIRV, "Geometry shader passthrough used with no support");
            }
        }
        break;
    case Stage::Fragment:
        execution_model = spv::ExecutionModel::Fragment;
        if (ctx.profile.lower_left_origin_mode) {
            ctx.AddExecutionMode(main, spv::ExecutionMode::OriginLowerLeft);
        } else {
            ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft);
        }
        if (program.info.stores_frag_depth) {
            ctx.AddExecutionMode(main, spv::ExecutionMode::DepthReplacing);
        }
        if (ctx.runtime_info.force_early_z) {
            ctx.AddExecutionMode(main, spv::ExecutionMode::EarlyFragmentTests);
        }
        break;
    default:
        throw NotImplementedException("Stage {}", program.stage);
    }
    ctx.AddEntryPoint(execution_model, main, "main", interfaces);
}

void SetupDenormControl(const Profile& profile, const IR::Program& program, EmitContext& ctx,
                        Id main_func) {
    const Info& info{program.info};
    if (info.uses_fp32_denorms_flush && info.uses_fp32_denorms_preserve) {
        LOG_DEBUG(Shader_SPIRV, "Fp32 denorm flush and preserve on the same shader");
    } else if (info.uses_fp32_denorms_flush) {
        if (profile.support_fp32_denorm_flush) {
            ctx.AddCapability(spv::Capability::DenormFlushToZero);
            ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormFlushToZero, 32U);
        } else {
            // Drivers will most likely flush denorms by default, no need to warn
        }
    } else if (info.uses_fp32_denorms_preserve) {
        if (profile.support_fp32_denorm_preserve) {
            ctx.AddCapability(spv::Capability::DenormPreserve);
            ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 32U);
        } else {
            LOG_DEBUG(Shader_SPIRV, "Fp32 denorm preserve used in shader without host support");
        }
    }
    if (!profile.support_separate_denorm_behavior || profile.has_broken_fp16_float_controls) {
        // No separate denorm behavior
        return;
    }
    if (info.uses_fp16_denorms_flush && info.uses_fp16_denorms_preserve) {
        LOG_DEBUG(Shader_SPIRV, "Fp16 denorm flush and preserve on the same shader");
    } else if (info.uses_fp16_denorms_flush) {
        if (profile.support_fp16_denorm_flush) {
            ctx.AddCapability(spv::Capability::DenormFlushToZero);
            ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormFlushToZero, 16U);
        } else {
            // Same as fp32, no need to warn as most drivers will flush by default
        }
    } else if (info.uses_fp16_denorms_preserve) {
        if (profile.support_fp16_denorm_preserve) {
            ctx.AddCapability(spv::Capability::DenormPreserve);
            ctx.AddExecutionMode(main_func, spv::ExecutionMode::DenormPreserve, 16U);
        } else {
            LOG_DEBUG(Shader_SPIRV, "Fp16 denorm preserve used in shader without host support");
        }
    }
}

void SetupSignedNanCapabilities(const Profile& profile, const IR::Program& program,
                                EmitContext& ctx, Id main_func) {
    if (profile.has_broken_fp16_float_controls && program.info.uses_fp16) {
        return;
    }
    if (program.info.uses_fp16 && profile.support_fp16_signed_zero_nan_preserve) {
        ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve);
        ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 16U);
    }
    if (profile.support_fp32_signed_zero_nan_preserve) {
        ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve);
        ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 32U);
    }
    if (program.info.uses_fp64 && profile.support_fp64_signed_zero_nan_preserve) {
        ctx.AddCapability(spv::Capability::SignedZeroInfNanPreserve);
        ctx.AddExecutionMode(main_func, spv::ExecutionMode::SignedZeroInfNanPreserve, 64U);
    }
}

void SetupTransformFeedbackCapabilities(EmitContext& ctx, Id main_func) {
    if (ctx.runtime_info.xfb_count == 0) {
        return;
    }
    ctx.AddCapability(spv::Capability::TransformFeedback);
    ctx.AddExecutionMode(main_func, spv::ExecutionMode::Xfb);
}

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 && profile.support_demote_to_helper_invocation) {
        if (profile.supported_spirv < 0x00010600) {
            ctx.AddExtension("SPV_EXT_demote_to_helper_invocation");
        }
        ctx.AddCapability(spv::Capability::DemoteToHelperInvocation);
    }
    if (info.stores[IR::Attribute::ViewportIndex] && profile.support_multi_viewport) {
        ctx.AddCapability(spv::Capability::MultiViewport);
    }
    if (info.stores[IR::Attribute::ViewportMask] && profile.support_viewport_mask) {
        ctx.AddExtension("SPV_NV_viewport_array2");
        ctx.AddCapability(spv::Capability::ShaderViewportMaskNV);
    }
    if (info.stores[IR::Attribute::Layer] || info.stores[IR::Attribute::ViewportIndex]) {
        if (profile.support_viewport_index_layer_non_geometry && ctx.stage != Stage::Geometry) {
            ctx.AddExtension("SPV_EXT_shader_viewport_index_layer");
            ctx.AddCapability(spv::Capability::ShaderViewportIndexLayerEXT);
        }
    }
    if (!profile.support_vertex_instance_id &&
        (info.loads[IR::Attribute::InstanceId] || info.loads[IR::Attribute::VertexId])) {
        ctx.AddExtension("SPV_KHR_shader_draw_parameters");
        ctx.AddCapability(spv::Capability::DrawParameters);
    }
    if ((info.uses_subgroup_vote || info.uses_subgroup_invocation_id ||
         info.uses_subgroup_shuffles) &&
        profile.support_vote) {
        ctx.AddCapability(spv::Capability::GroupNonUniformBallot);
        ctx.AddCapability(spv::Capability::GroupNonUniformShuffle);
        if (!profile.warp_size_potentially_larger_than_guest) {
            // vote ops are only used when not taking the long path
            ctx.AddCapability(spv::Capability::GroupNonUniformVote);
        }
    }
    if (info.uses_int64_bit_atomics && profile.support_int64_atomics) {
        ctx.AddCapability(spv::Capability::Int64Atomics);
    }
    if (info.uses_typeless_image_reads && profile.support_typeless_image_loads) {
        ctx.AddCapability(spv::Capability::StorageImageReadWithoutFormat);
    }
    if (info.uses_typeless_image_writes) {
        ctx.AddCapability(spv::Capability::StorageImageWriteWithoutFormat);
    }
    if (info.uses_image_buffers) {
        ctx.AddCapability(spv::Capability::ImageBuffer);
    }
    if (info.uses_sample_id) {
        ctx.AddCapability(spv::Capability::SampleRateShading);
    }
    if (info.uses_derivatives) {
        ctx.AddCapability(spv::Capability::DerivativeControl);
    }
    // TODO: Track this usage
    ctx.AddCapability(spv::Capability::ImageGatherExtended);
    ctx.AddCapability(spv::Capability::ImageQuery);
    ctx.AddCapability(spv::Capability::SampledBuffer);
}

void PatchPhiNodes(IR::Program& program, EmitContext& ctx) {
    auto inst{program.blocks.front()->begin()};
    size_t block_index{0};
    ctx.PatchDeferredPhi([&](size_t phi_arg) {
        if (phi_arg == 0) {
            ++inst;
            if (inst == program.blocks[block_index]->end() ||
                inst->GetOpcode() != IR::Opcode::Phi) {
                do {
                    ++block_index;
                    inst = program.blocks[block_index]->begin();
                } while (inst->GetOpcode() != IR::Opcode::Phi);
            }
        }
        return ctx.Def(inst->Arg(phi_arg));
    });
}
} // Anonymous namespace

std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info,
                           IR::Program& program, Bindings& bindings) {
    EmitContext ctx{profile, runtime_info, program, bindings};
    const Id main{DefineMain(ctx, program)};
    DefineEntryPoint(program, ctx, main);
    if (profile.support_float_controls) {
        ctx.AddExtension("SPV_KHR_float_controls");
        SetupDenormControl(profile, program, ctx, main);
        SetupSignedNanCapabilities(profile, program, ctx, main);
    }
    SetupCapabilities(profile, program.info, ctx);
    SetupTransformFeedbackCapabilities(ctx, main);
    PatchPhiNodes(program, ctx);
    return ctx.Assemble();
}

Id EmitPhi(EmitContext& ctx, IR::Inst* inst) {
    const size_t num_args{inst->NumArgs()};
    boost::container::small_vector<Id, 32> blocks;
    blocks.reserve(num_args);
    for (size_t index = 0; index < num_args; ++index) {
        blocks.push_back(inst->PhiBlock(index)->Definition<Id>());
    }
    // The type of a phi instruction is stored in its flags
    const Id result_type{TypeId(ctx, inst->Flags<IR::Type>())};
    return ctx.DeferredOpPhi(result_type, std::span(blocks.data(), blocks.size()));
}

void EmitVoid(EmitContext&) {}

Id EmitIdentity(EmitContext& ctx, const IR::Value& value) {
    const Id id{ctx.Def(value)};
    if (!Sirit::ValidId(id)) {
        throw NotImplementedException("Forward identity declaration");
    }
    return id;
}

Id EmitConditionRef(EmitContext& ctx, const IR::Value& value) {
    const Id id{ctx.Def(value)};
    if (!Sirit::ValidId(id)) {
        throw NotImplementedException("Forward identity declaration");
    }
    return id;
}

void EmitReference(EmitContext&) {}

void EmitPhiMove(EmitContext&) {
    throw LogicError("Unreachable instruction");
}

void EmitGetZeroFromOp(EmitContext&) {
    throw LogicError("Unreachable instruction");
}

void EmitGetSignFromOp(EmitContext&) {
    throw LogicError("Unreachable instruction");
}

void EmitGetCarryFromOp(EmitContext&) {
    throw LogicError("Unreachable instruction");
}

void EmitGetOverflowFromOp(EmitContext&) {
    throw LogicError("Unreachable instruction");
}

void EmitGetSparseFromOp(EmitContext&) {
    throw LogicError("Unreachable instruction");
}

void EmitGetInBoundsFromOp(EmitContext&) {
    throw LogicError("Unreachable instruction");
}

} // namespace Shader::Backend::SPIRV