diff options
author | ameerj <52414509+ameerj@users.noreply.github.com> | 2021-05-10 04:01:03 +0200 |
---|---|---|
committer | ameerj <52414509+ameerj@users.noreply.github.com> | 2021-07-23 03:51:31 +0200 |
commit | 80813b1d144a7f0f11047e7348620b720def93a9 (patch) | |
tree | 705aafbb73d550a9ee8f3b5fd24fae9c4beb46d3 /src/shader_recompiler/backend/glasm/emit_glasm.cpp | |
parent | glasm: Add conversion instructions to GLASM (diff) | |
download | yuzu-80813b1d144a7f0f11047e7348620b720def93a9.tar yuzu-80813b1d144a7f0f11047e7348620b720def93a9.tar.gz yuzu-80813b1d144a7f0f11047e7348620b720def93a9.tar.bz2 yuzu-80813b1d144a7f0f11047e7348620b720def93a9.tar.lz yuzu-80813b1d144a7f0f11047e7348620b720def93a9.tar.xz yuzu-80813b1d144a7f0f11047e7348620b720def93a9.tar.zst yuzu-80813b1d144a7f0f11047e7348620b720def93a9.zip |
Diffstat (limited to '')
-rw-r--r-- | src/shader_recompiler/backend/glasm/emit_glasm.cpp | 13 |
1 files changed, 13 insertions, 0 deletions
diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp index 0e4b189c9..e6e065e7f 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp @@ -149,6 +149,18 @@ void EmitInst(EmitContext& ctx, IR::Inst* inst) { } throw LogicError("Invalid opcode {}", inst->GetOpcode()); } + +void SetupOptions(std::string& header, Info info) { + if (info.uses_int64_bit_atomics) { + header += "OPTION NV_shader_atomic_int64;"; + } + if (info.uses_atomic_f32_add) { + header += "OPTION NV_shader_atomic_float;"; + } + if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) { + header += "OPTION NV_shader_atomic_fp16_vector;"; + } +} } // Anonymous namespace std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) { @@ -160,6 +172,7 @@ std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) { } std::string header = "!!NVcp5.0\n" "OPTION NV_internal;"; + SetupOptions(header, program.info); switch (program.stage) { case Stage::Compute: header += fmt::format("GROUP_SIZE {} {} {};", program.workgroup_size[0], |