glsl: Wip storage atomic ops
This commit is contained in:
parent
df793fc049
commit
3d9ecbe998
|
@ -19,8 +19,10 @@ EmitContext::EmitContext(IR::Program& program, [[maybe_unused]] Bindings& bindin
|
|||
program.workgroup_size[2]);
|
||||
}
|
||||
code += header;
|
||||
|
||||
DefineConstantBuffers();
|
||||
DefineStorageBuffers();
|
||||
DefineHelperFunctions();
|
||||
code += "void main(){\n";
|
||||
}
|
||||
|
||||
|
@ -28,6 +30,15 @@ void EmitContext::SetupExtensions(std::string& header) {
|
|||
if (info.uses_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 NV_shader_atomic_fp16_vector : enable\n";
|
||||
}
|
||||
}
|
||||
|
||||
void EmitContext::DefineConstantBuffers() {
|
||||
|
@ -48,18 +59,39 @@ void EmitContext::DefineStorageBuffers() {
|
|||
}
|
||||
u32 binding{};
|
||||
for (const auto& desc : info.storage_buffers_descriptors) {
|
||||
if (True(info.used_storage_buffer_types & IR::Type::U32) ||
|
||||
True(info.used_storage_buffer_types & IR::Type::F32)) {
|
||||
if (info.uses_s32_atomics) {
|
||||
Add("layout(std430,binding={}) buffer ssbo_{}_s32{{int ssbo{}_s32[];}};", binding,
|
||||
binding, desc.cbuf_index, desc.count);
|
||||
}
|
||||
if (True(info.used_storage_buffer_types & IR::Type::U32)) {
|
||||
Add("layout(std430,binding={}) buffer ssbo_{}_u32{{uint ssbo{}_u32[];}};", binding,
|
||||
binding, desc.cbuf_index, desc.count);
|
||||
}
|
||||
if (True(info.used_storage_buffer_types & IR::Type::U32x2) ||
|
||||
True(info.used_storage_buffer_types & IR::Type::F32x2)) {
|
||||
Add("layout(std430,binding={}) buffer ssbo_{}_u64{{uvec2 ssbo{}_u64[];}};", binding,
|
||||
if (True(info.used_storage_buffer_types & IR::Type::F32)) {
|
||||
Add("layout(std430,binding={}) buffer ssbo_{}_f32{{float ssbo{}_f32[];}};", binding,
|
||||
binding, desc.cbuf_index, desc.count);
|
||||
}
|
||||
if (True(info.used_storage_buffer_types & IR::Type::U32x2)) {
|
||||
Add("layout(std430,binding={}) buffer ssbo_{}_u32x2{{uvec2 ssbo{}_u32x2[];}};", binding,
|
||||
binding, desc.cbuf_index, desc.count);
|
||||
}
|
||||
if (True(info.used_storage_buffer_types & IR::Type::U64) ||
|
||||
True(info.used_storage_buffer_types & IR::Type::F64)) {
|
||||
Add("layout(std430,binding={}) buffer ssbo_{}_u64{{uint64_t ssbo{}_u64[];}};", binding,
|
||||
binding, desc.cbuf_index, desc.count);
|
||||
}
|
||||
++binding;
|
||||
}
|
||||
}
|
||||
|
||||
void EmitContext::DefineHelperFunctions() {
|
||||
if (info.uses_global_increment) {
|
||||
code += "uint CasIncrement(uint op_a,uint op_b){return(op_a>=op_b)?0u:(op_a+1u);}\n";
|
||||
}
|
||||
if (info.uses_global_decrement) {
|
||||
code +=
|
||||
"uint CasDecrement(uint op_a,uint op_b){return(op_a==0||op_a>op_b)?op_b:(op_a-1u);}\n";
|
||||
}
|
||||
}
|
||||
|
||||
} // namespace Shader::Backend::GLSL
|
||||
|
|
|
@ -31,13 +31,6 @@ class EmitContext {
|
|||
public:
|
||||
explicit EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_);
|
||||
|
||||
// template <typename... Args>
|
||||
// void Add(const char* format_str, IR::Inst& inst, Args&&... args) {
|
||||
// code += fmt::format(format_str, reg_alloc.Define(inst), std::forward<Args>(args)...);
|
||||
// // TODO: Remove this
|
||||
// code += '\n';
|
||||
// }
|
||||
|
||||
template <Type type, typename... Args>
|
||||
void Add(const char* format_str, IR::Inst& inst, Args&&... args) {
|
||||
code += fmt::format(format_str, reg_alloc.Define(inst, type), std::forward<Args>(args)...);
|
||||
|
@ -106,6 +99,7 @@ private:
|
|||
void SetupExtensions(std::string& header);
|
||||
void DefineConstantBuffers();
|
||||
void DefineStorageBuffers();
|
||||
void DefineHelperFunctions();
|
||||
};
|
||||
|
||||
} // namespace Shader::Backend::GLSL
|
||||
|
|
|
@ -0,0 +1,301 @@
|
|||
|
||||
// Copyright 2021 yuzu Emulator Project
|
||||
// Licensed under GPLv2 or any later version
|
||||
// Refer to the license.txt file included.
|
||||
|
||||
#include <string_view>
|
||||
|
||||
#include "shader_recompiler/backend/glsl/emit_context.h"
|
||||
#include "shader_recompiler/backend/glsl/emit_glsl_instructions.h"
|
||||
#include "shader_recompiler/frontend/ir/value.h"
|
||||
#include "shader_recompiler/profile.h"
|
||||
|
||||
namespace Shader::Backend::GLSL {
|
||||
namespace {
|
||||
static constexpr std::string_view cas_loop{R"(
|
||||
{} {};
|
||||
for (;;){{
|
||||
{} old_value={};
|
||||
{} = atomicCompSwap({},old_value,{}({},{}));
|
||||
if ({}==old_value){{break;}}
|
||||
}})"};
|
||||
|
||||
void CasFunction(EmitContext& ctx, IR::Inst& inst, std::string_view ssbo, std::string_view value,
|
||||
std::string_view type, std::string_view function) {
|
||||
const auto ret{ctx.reg_alloc.Define(inst)};
|
||||
ctx.Add(cas_loop.data(), type, ret, type, ssbo, ret, ssbo, function, ssbo, value, ret);
|
||||
}
|
||||
} // namespace
|
||||
|
||||
void EmitStorageAtomicIAdd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU32("{}=atomicAdd(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicSMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddS32("{}=atomicMin(ssbo{}_s32[{}],int({}));", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicUMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU32("{}=atomicMin(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicSMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddS32("{}=atomicMax(ssbo{}_s32[{}],int({}));", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicUMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU32("{}=atomicMax(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicInc32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
[[maybe_unused]] const IR::Value& offset, std::string_view value) {
|
||||
// const auto ret{ctx.reg_alloc.Define(inst)};
|
||||
// const auto type{"uint"};
|
||||
// ctx.Add(cas_loop.data(), type, ret, type, ssbo, ret, ssbo, "CasIncrement", ssbo, value, ret);
|
||||
const std::string ssbo{fmt::format("ssbo{}_u32[{}]", binding.U32(), offset.U32())};
|
||||
CasFunction(ctx, inst, ssbo, value, "uint", "CasIncrement");
|
||||
}
|
||||
|
||||
void EmitStorageAtomicDec32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
const std::string ssbo{fmt::format("ssbo{}_u32[{}]", binding.U32(), offset.U32())};
|
||||
CasFunction(ctx, inst, ssbo, value, "uint", "CasDecrement");
|
||||
}
|
||||
|
||||
void EmitStorageAtomicAnd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU32("{}=atomicAnd(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicOr32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU32("{}=atomicOr(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicXor32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU32("{}=atomicXor(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicExchange32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU32("{}=atomicExchange(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicIAdd64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
// ctx.AddU64("{}=atomicAdd(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
ctx.AddU64("{}=ssbo{}_u64[{}];", inst, binding.U32(), offset.U32());
|
||||
ctx.Add("ssbo{}_u64[{}]+={};", binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicSMin64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddS64("{}=atomicMin(int64_t(ssbo{}_u64[{}]),int64_t({}));", inst, binding.U32(),
|
||||
offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicUMin64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU64("{}=atomicMin(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicSMax64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddS64("{}=atomicMax(int64_t(ssbo{}_u64[{}]),int64_t({}));", inst, binding.U32(),
|
||||
offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicUMax64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU64("{}=atomicMax(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicAnd64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU64("{}=atomicAnd(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicOr64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU64("{}=atomicOr(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicXor64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU64("{}=atomicXor(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicExchange64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddU64("{}=atomicExchange(ssbo{}_u64[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicAddF32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
ctx.AddF32("{}=atomicAdd(ssbo{}_u32[{}],{});", inst, binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitStorageAtomicAddF16x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
||||
[[maybe_unused]] const IR::Value& binding,
|
||||
[[maybe_unused]] const IR::Value& offset,
|
||||
[[maybe_unused]] std::string_view value) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitStorageAtomicAddF32x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
||||
[[maybe_unused]] const IR::Value& binding,
|
||||
[[maybe_unused]] const IR::Value& offset,
|
||||
[[maybe_unused]] std::string_view value) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitStorageAtomicMinF16x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
||||
[[maybe_unused]] const IR::Value& binding,
|
||||
[[maybe_unused]] const IR::Value& offset,
|
||||
[[maybe_unused]] std::string_view value) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitStorageAtomicMinF32x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
||||
[[maybe_unused]] const IR::Value& binding,
|
||||
[[maybe_unused]] const IR::Value& offset,
|
||||
[[maybe_unused]] std::string_view value) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitStorageAtomicMaxF16x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
||||
[[maybe_unused]] const IR::Value& binding,
|
||||
[[maybe_unused]] const IR::Value& offset,
|
||||
[[maybe_unused]] std::string_view value) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitStorageAtomicMaxF32x2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst,
|
||||
[[maybe_unused]] const IR::Value& binding,
|
||||
[[maybe_unused]] const IR::Value& offset,
|
||||
[[maybe_unused]] std::string_view value) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicIAdd32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicSMin32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicUMin32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicSMax32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicUMax32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicInc32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicDec32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicAnd32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicOr32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicXor32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicExchange32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicIAdd64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicSMin64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicUMin64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicSMax64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicUMax64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicInc64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicDec64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicAnd64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicOr64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicXor64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicExchange64(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicAddF32(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicAddF16x2(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicAddF32x2(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicMinF16x2(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicMinF32x2(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicMaxF16x2(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicMaxF32x2(EmitContext&) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
} // namespace Shader::Backend::GLSL
|
|
@ -15,10 +15,7 @@ class Inst;
|
|||
class Value;
|
||||
} // namespace Shader::IR
|
||||
|
||||
#pragma optimize("", off)
|
||||
|
||||
namespace Shader::Backend::GLSL {
|
||||
|
||||
class EmitContext;
|
||||
|
||||
inline void EmitSetLoopSafetyVariable(EmitContext&) {}
|
||||
|
@ -114,7 +111,8 @@ void EmitLoadStorageU8(EmitContext& ctx, const IR::Value& binding, const IR::Val
|
|||
void EmitLoadStorageS8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||
void EmitLoadStorageU16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||
void EmitLoadStorageS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||
void EmitLoadStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||
void EmitLoadStorage32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset);
|
||||
void EmitLoadStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||
void EmitLoadStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset);
|
||||
void EmitWriteStorageU8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
|
@ -431,60 +429,60 @@ void EmitSharedAtomicExchange32(EmitContext& ctx, std::string_view pointer_offse
|
|||
std::string_view value);
|
||||
void EmitSharedAtomicExchange64(EmitContext& ctx, std::string_view pointer_offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicIAdd32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicSMin32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicUMin32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicSMax32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicUMax32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicInc32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicDec32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicAnd32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicOr32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicXor32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicExchange32(EmitContext& ctx, const IR::Value& binding,
|
||||
void EmitStorageAtomicIAdd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicSMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicUMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicSMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicUMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicInc32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicDec32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicAnd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicOr32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicXor32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicExchange32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicIAdd64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicSMin64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicUMin64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicSMax64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicUMax64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicAnd64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicOr64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicXor64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicExchange64(EmitContext& ctx, const IR::Value& binding,
|
||||
void EmitStorageAtomicIAdd64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicSMin64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicUMin64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicSMax64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicUMax64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicAnd64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicOr64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicXor64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicExchange64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicAddF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicAddF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicAddF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicMinF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicMinF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicMaxF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicMaxF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value);
|
||||
void EmitStorageAtomicAddF32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicAddF16x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicAddF32x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicMinF16x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicMinF32x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicMaxF16x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitStorageAtomicMaxF32x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value);
|
||||
void EmitGlobalAtomicIAdd32(EmitContext& ctx);
|
||||
void EmitGlobalAtomicSMin32(EmitContext& ctx);
|
||||
void EmitGlobalAtomicUMin32(EmitContext& ctx);
|
||||
|
|
|
@ -32,9 +32,9 @@ void EmitLoadStorageS16([[maybe_unused]] EmitContext& ctx,
|
|||
throw NotImplementedException("GLSL Instrucion");
|
||||
}
|
||||
|
||||
void EmitLoadStorage32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] const IR::Value& binding,
|
||||
[[maybe_unused]] const IR::Value& offset) {
|
||||
throw NotImplementedException("GLSL Instrucion");
|
||||
void EmitLoadStorage32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding,
|
||||
const IR::Value& offset) {
|
||||
ctx.AddU32("{}=ssbo{}_u32[{}];", inst, binding.U32(), offset.U32());
|
||||
}
|
||||
|
||||
void EmitLoadStorage64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] const IR::Value& binding,
|
||||
|
@ -83,7 +83,7 @@ void EmitWriteStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Va
|
|||
|
||||
void EmitWriteStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
ctx.Add("ssbo{}_u64[{}]={};", binding.U32(), offset.U32(), value);
|
||||
ctx.Add("ssbo{}_u32x2[{}]={};", binding.U32(), offset.U32(), value);
|
||||
}
|
||||
|
||||
void EmitWriteStorage128([[maybe_unused]] EmitContext& ctx,
|
||||
|
|
|
@ -20,7 +20,7 @@ static void NotImplemented() {
|
|||
}
|
||||
|
||||
void EmitPhi(EmitContext& ctx, IR::Inst& inst) {
|
||||
NotImplemented();
|
||||
// NotImplemented();
|
||||
}
|
||||
|
||||
void EmitVoid(EmitContext& ctx) {
|
||||
|
@ -439,257 +439,6 @@ void EmitSharedAtomicExchange64(EmitContext& ctx, std::string_view pointer_offse
|
|||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicIAdd32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicSMin32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicUMin32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicSMax32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicUMax32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicInc32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicDec32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicAnd32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicOr32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicXor32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicExchange32(EmitContext& ctx, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicIAdd64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicSMin64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicUMin64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicSMax64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicUMax64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicAnd64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicOr64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicXor64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicExchange64(EmitContext& ctx, const IR::Value& binding,
|
||||
const IR::Value& offset, std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicAddF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicAddF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicAddF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicMinF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicMinF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicMaxF16x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitStorageAtomicMaxF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset,
|
||||
std::string_view value) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicIAdd32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicSMin32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicUMin32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicSMax32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicUMax32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicInc32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicDec32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicAnd32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicOr32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicXor32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicExchange32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicIAdd64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicSMin64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicUMin64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicSMax64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicUMax64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicInc64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicDec64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicAnd64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicOr64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicXor64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicExchange64(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicAddF32(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicAddF16x2(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicAddF32x2(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicMinF16x2(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicMinF32x2(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicMaxF16x2(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitGlobalAtomicMaxF32x2(EmitContext& ctx) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
||||
void EmitBindlessImageSampleImplicitLod(EmitContext&) {
|
||||
NotImplemented();
|
||||
}
|
||||
|
|
|
@ -61,6 +61,12 @@ std::string MakeImm(const IR::Value& value) {
|
|||
}
|
||||
} // Anonymous namespace
|
||||
|
||||
std::string RegAlloc::Define(IR::Inst& inst) {
|
||||
const Id id{Alloc()};
|
||||
inst.SetDefinition<Id>(id);
|
||||
return Representation(id);
|
||||
}
|
||||
|
||||
std::string RegAlloc::Define(IR::Inst& inst, Type type) {
|
||||
const Id id{Alloc()};
|
||||
const auto type_str{GetType(type, id.index)};
|
||||
|
|
|
@ -48,7 +48,8 @@ static_assert(sizeof(Id) == sizeof(u32));
|
|||
|
||||
class RegAlloc {
|
||||
public:
|
||||
std::string Define(IR::Inst& inst, Type type = Type::Void);
|
||||
std::string Define(IR::Inst& inst);
|
||||
std::string Define(IR::Inst& inst, Type type);
|
||||
|
||||
std::string Consume(const IR::Value& value);
|
||||
|
||||
|
|
|
@ -687,9 +687,7 @@ void VisitUsages(Info& info, IR::Inst& inst) {
|
|||
case IR::Opcode::LoadStorage32:
|
||||
case IR::Opcode::WriteStorage32:
|
||||
case IR::Opcode::StorageAtomicIAdd32:
|
||||
case IR::Opcode::StorageAtomicSMin32:
|
||||
case IR::Opcode::StorageAtomicUMin32:
|
||||
case IR::Opcode::StorageAtomicSMax32:
|
||||
case IR::Opcode::StorageAtomicUMax32:
|
||||
case IR::Opcode::StorageAtomicAnd32:
|
||||
case IR::Opcode::StorageAtomicOr32:
|
||||
|
@ -714,6 +712,13 @@ void VisitUsages(Info& info, IR::Inst& inst) {
|
|||
case IR::Opcode::SharedAtomicExchange64:
|
||||
info.uses_int64_bit_atomics = true;
|
||||
break;
|
||||
case IR::Opcode::GlobalAtomicSMin32:
|
||||
case IR::Opcode::GlobalAtomicSMax32:
|
||||
case IR::Opcode::StorageAtomicSMin32:
|
||||
case IR::Opcode::StorageAtomicSMax32:
|
||||
info.used_storage_buffer_types |= IR::Type::U32;
|
||||
info.uses_s32_atomics = true;
|
||||
break;
|
||||
case IR::Opcode::GlobalAtomicInc32:
|
||||
case IR::Opcode::StorageAtomicInc32:
|
||||
info.used_storage_buffer_types |= IR::Type::U32;
|
||||
|
|
|
@ -190,6 +190,7 @@ struct Info {
|
|||
bool uses_atomic_f32x2_min{};
|
||||
bool uses_atomic_f32x2_max{};
|
||||
bool uses_int64_bit_atomics{};
|
||||
bool uses_s32_atomics{};
|
||||
bool uses_global_memory{};
|
||||
bool uses_atomic_image_u32{};
|
||||
|
||||
|
|
Loading…
Reference in New Issue