diff --git a/.github/workflows/source-regressions.yml b/.github/workflows/source-regressions.yml index ae298b434..4958df690 100644 --- a/.github/workflows/source-regressions.yml +++ b/.github/workflows/source-regressions.yml @@ -150,6 +150,21 @@ jobs: for module in "${{ runner.temp }}"/global-load-16/modules/*.spv; do spirv-val --target-env vulkan1.3 --uniform-buffer-standard-layout "$module" done + - name: Configure the authored direct surface atomic fixture + run: >- + CC=clang CXX=clang++ cmake -S tests/native_suatom -B ${{ runner.temp }}/native-suatom + -DREPO_ROOT=${{ github.workspace }} + -DSPIRV_HEADERS_ROOT=${{ github.workspace }}/_deps/spirv-headers + - name: Build the direct surface atomic fixture + run: cmake --build ${{ runner.temp }}/native-suatom --parallel 2 + - name: Translate authored direct surface atomic programs + run: ${{ runner.temp }}/native-suatom/native_suatom ${{ runner.temp }}/native-suatom/modules + - name: Validate every direct surface atomic module + shell: bash + run: | + for module in "${{ runner.temp }}"/native-suatom/modules/*.spv; do + spirv-val --target-env vulkan1.3 "$module" + done actionlint: name: Workflow lint (actionlint) diff --git a/gpu/vk/eden/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/gpu/vk/eden/shader_recompiler/frontend/ampere/lower_to_maxwell.h index a26518b3e..3374edef8 100644 --- a/gpu/vk/eden/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/gpu/vk/eden/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -64,6 +64,7 @@ inline u32 RenameableFields(u32 opcode) { switch (opcode) { case 0x381: case 0x981: return 1 | 2; // LDG Rd, [Ra] (0x381: Ra a register pair) case 0x986: return 2 | 4; // STG [Ra], Rb + case 0x593: return 1 | 2 | 4; // SUATOM: scalar destination, coordinate and addend default: break; } if (extended_renaming) { @@ -94,6 +95,7 @@ inline u32 RenameableFields(u32 opcode) { // 32-bit operand taken for a quad reached R239 or the scratch from R236 or R244 and failed whole mesh programs (OoT's // foliage). inline u32 FieldSpan(u32 opcode, u32 field) { + if (opcode == 0x593) return field < 3 ? 1 : 0; if (!extended_renaming) return 0; const u32 family = opcode & 0x1ff, form = (opcode >> 9) & 7; if (opcode == 0x805) return field == 0 ? 2 : 0; @@ -1444,6 +1446,19 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 if (strong) emit(0xef98000000000000ULL | predicate | (u64{1} << 8)); break; } + case 0x593: { // SUATOM.D.1D.ADD.U32.STRONG.GPU.IGN, direct descriptor, no fault predicate. + const u32 binding = insn.Bits(54, 5), offset = insn.Bits(40, 14); + if (insn.words[2] != 0x001ee000 || insn.Bits(59, 5) || binding >= 18 || offset > 0x1fff || + (insn.words[3] & 0xc00001ffU)) + return fail(index, "native direct surface atomic mode"); + if (!bind_texture_buffer(binding)) + return fail(index, "native surface descriptor buffer conflict"); + // Strong(GPU) is the atomic's coherence/scope, not acquire/release. The existing image + // atomic backend uses Device/Relaxed; explicit native MEMBARs are lowered separately. + emit(base(0xea00000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(32, 8)} << 20) | (u64{offset} << 36)); + break; + } case 0x59f: { // SURED.D.2D..STRONG.GPU.IGN U32, direct descriptor (title 1's race compute). const u32 binding = insn.Bits(54, 5), offset = insn.Bits(40, 14); const u32 data = insn.Bits(32, 8), coordinate = insn.Bits(24, 8); diff --git a/tests/native_suatom.cpp b/tests/native_suatom.cpp new file mode 100644 index 000000000..ab26ad9c5 --- /dev/null +++ b/tests/native_suatom.cpp @@ -0,0 +1,251 @@ +// Independently authored encoding fields; no shader dump or game program is used. +#include +#include +#include +#include +#include +#include +#include "shader_recompiler/frontend/ampere/lower_to_maxwell.h" + +using Shader::Ampere::Instruction; +using Shader::Ampere::LoweredProgram; +#define CHECK(x) do { if (!(x)) { std::fprintf(stderr, "FAIL:%d %s\n", __LINE__, #x); std::exit(1); } } while (0) + +static void field(Instruction& i, unsigned at, unsigned count, uint32_t value) { + CHECK(count <= 32 && (count == 32 || value < (uint64_t{1} << count))); + for (unsigned bit = 0; bit < count; ++bit) { + auto& word = i.words[(at + bit) / 32]; + const uint32_t mask = uint32_t{1} << ((at + bit) % 32); + word = (word & ~mask) | ((value >> bit & 1) ? mask : 0); + } +} +static Instruction op(uint32_t code, uint32_t pred = 7) { + Instruction i{}; + field(i, 0, 12, code); field(i, 12, 4, pred); + return i; +} +static Instruction suatom(uint32_t d = 13, uint32_t a = 9, uint32_t b = 17, + uint32_t bank = 3, uint32_t offset = 0x24, uint32_t pred = 7) { + auto i = op(0x593, pred); + field(i, 16, 8, d); field(i, 24, 8, a); field(i, 32, 8, b); + field(i, 40, 14, offset); field(i, 54, 5, bank); + field(i, 77, 4, 7); // Strong(GPU), not acquire/release + field(i, 81, 3, 7); // no fault-predicate destination + field(i, 84, 3, 1); // normal eviction; U32 ADD, 1D, sample addressing, IGN are zero + return i; +} +static Instruction exit_op() { + auto i = op(0x94d); field(i, 87, 3, 7); return i; +} +static LoweredProgram lower(Instruction i) { + return Shader::Ampere::LowerToMaxwell(std::vector{i, exit_op()}); +} +static uint64_t address(unsigned index) { return 0x50 + index * 8 + (index + 1) / 3 * 8; } + +static void lowering_checks() { + const auto r = lower(suatom()); + CHECK(r); + CHECK(r.required_texture_buffer == 3 && r.instructions.size() == 2); + const uint64_t w = r.instructions[0]; + CHECK(w >> 55 == 0xea00000000000000ULL >> 55); + CHECK((w & 255) == 13 && (w >> 8 & 255) == 9 && (w >> 20 & 255) == 17); + CHECK((w >> 36 & 0x1fff) == 0x24 && (w >> 54 & 1) == 0); + CHECK((w >> 29 & 15) == 0 && (w >> 33 & 7) == 0 && (w >> 49 & 31) == 0); + for (uint32_t p : {2U, 10U, 7U}) { + const auto predicated = lower(suatom(13, 9, 17, 3, 0x24, p)); + CHECK(predicated && (predicated.instructions[0] >> 16 & 15) == p); + } + CHECK(lower(suatom(255, 255, 255, 17, 0x1fff))); + for (auto [at, count, value] : std::vector>{ + {0,12,0x793}, {0,12,0xf93}, {54,5,18}, {40,14,0x2000}, + {59,2,1}, {61,3,1}, {64,8,1}, {72,1,1}, {73,3,1}, {76,1,1}, + {77,4,0}, {77,4,5}, {77,4,10}, {81,3,2}, {84,3,0}, {87,4,1}, + {91,1,1}, {95,1,1}, {96,1,1}, {104,1,1}, {126,1,1}, {127,1,1}}) { + auto bad = suatom(); field(bad, at, count, value); + const auto refused = lower(bad); + CHECK(!refused && refused.instructions.empty() && refused.error_offset == 0); + } + auto scheduled = suatom(); field(scheduled, 105, 21, 0x155555); + CHECK(lower(scheduled).instructions == r.instructions); + const auto same = Shader::Ampere::LowerToMaxwell(std::vector{suatom(), suatom(7), exit_op()}); + CHECK(same && same.required_texture_buffer == 3); + const auto conflict = Shader::Ampere::LowerToMaxwell(std::vector{suatom(), suatom(7,9,17,2), exit_op()}); + CHECK(!conflict && conflict.error == "native surface descriptor buffer conflict" && conflict.error_offset == 16); + // The addend alone is high: it must participate in REGMAP discovery, and has scalar width. + const auto high = lower(suatom(13, 9, 250)); + CHECK(high && high.register_map_words > 0); + CHECK(Shader::Ampere::RenameableFields(0x593) == 7); + CHECK(Shader::Ampere::FieldSpan(0x593, 0) == 1 && Shader::Ampere::FieldSpan(0x593, 1) == 1 && + Shader::Ampere::FieldSpan(0x593, 2) == 1 && Shader::Ampere::FieldSpan(0x593, 3) == 0); + // A conditional BRA crosses an instruction expanded by high-register mapping. + auto branch = op(0x947, 2); field(branch, 32, 32, 16); field(branch, 87, 3, 7); + const auto relocated = Shader::Ampere::LowerToMaxwell(std::vector{branch, suatom(250,250,250), exit_op()}); + CHECK(relocated && relocated.register_map_words > 0); + const auto delta = static_cast((relocated.instructions[0] >> 20 & 0xffffff) << 8) >> 8; + CHECK(address(0) + 8 + delta == address(relocated.source_instruction_offsets[2])); + // Explicit MEMBAR.SC.GPU remains a separate instruction; SUATOM adds no fences. + auto barrier = op(0x992); field(barrier, 76, 3, 2); + const auto barriers = Shader::Ampere::LowerToMaxwell(std::vector{barrier, suatom(), barrier, exit_op()}); + CHECK(barriers && barriers.instructions.size() == 4); + CHECK(barriers.instructions[0] >> 48 == 0xef98 && barriers.instructions[2] == barriers.instructions[0]); + std::puts("PASS: direct SUATOM fields, predicates, scalar REGMAP, relocation, descriptor bounds/conflicts and mode refusals"); +} + +#ifndef SUATOM_LOWERING_ONLY +#include +#include +#include +#include +#include +#include "shader_recompiler/backend/spirv/emit_spirv.h" +#include "shader_recompiler/frontend/maxwell/translate_program.h" +#include "shader_recompiler/host_translate_info.h" + +struct Environment final : Shader::Environment { + std::vector code; + uint32_t bank, offset; + Environment(const LoweredProgram& lowered, uint32_t o) : bank(lowered.required_texture_buffer), offset(o) { + stage = Shader::Stage::Compute; + code.resize(0x50 / 8); + for (auto word : lowered.instructions) { + if (code.size() % 4 == 0) code.push_back(0); + code.push_back(word); + } + } + uint64_t ReadInstruction(uint32_t at) override { CHECK(at / 8 < code.size()); return code[at/8]; } + uint32_t ReadCbufValue(uint32_t b, uint32_t o) override { + CHECK(b == bank && o == offset); return 0x1234; + } + Shader::TextureType ReadTextureType(uint32_t h) override { CHECK(h == 0x1234); return Shader::TextureType::Color1D; } + Shader::TexturePixelFormat ReadTexturePixelFormat(uint32_t) override { return Shader::TexturePixelFormat::R32_UINT; } + bool IsTexturePixelFormatInteger(uint32_t h) override { CHECK(h == 0x1234); return true; } + uint32_t ReadViewportTransformState() override { return 0; } + uint32_t TextureBoundBuffer() const override { return bank; } + uint32_t LocalMemorySize() const override { return 0; } + uint32_t SharedMemorySize() const override { return 4; } + std::array WorkgroupSize() const override { return {32,1,1}; } + bool HasHLEMacroState() const override { return false; } + std::optional GetReplaceConstBuffer(uint32_t,uint32_t) override { return std::nullopt; } + void Dump(uint64_t,uint64_t) override {} +}; +static Instruction move(uint32_t reg, uint32_t value) { + auto i = op(0x802); field(i,16,8,reg); field(i,32,32,value); field(i,72,4,15); return i; +} +static void module_check(const std::filesystem::path& directory, const char* name, + uint32_t d, uint32_t a, uint32_t b, uint32_t pred = 7, + bool branch = false, bool explicit_barriers = false) { + std::vector source; + if (d == 239) { + auto uniform = op(0x882); field(uniform,32,32,42); source.push_back(uniform); + } + // Initialize old destination, coordinate and addend, respecting input aliases. + source.push_back(move(d, 99)); source.push_back(move(a, 3)); + if (b != a) source.push_back(move(b, 11)); + if (pred != 7 || branch) { + auto tid = op(0x919); field(tid,16,8,1); field(tid,72,8,33); source.push_back(tid); + auto compare = op(0x20c); field(compare,24,8,1); field(compare,32,8,255); + field(compare,76,3,5); field(compare,81,3,2); field(compare,84,3,7); field(compare,87,3,7); + source.push_back(compare); // P2 = threadIdx.x != 0, so its value is not constant + } + if (branch) { + auto jump = op(0x947,2); field(jump,32,32,16); field(jump,87,3,7); source.push_back(jump); + } + auto barrier = op(0x992); field(barrier,76,3,2); + if (explicit_barriers) source.push_back(barrier); + source.push_back(suatom(d,a,b,3,0x24,pred)); + if (explicit_barriers) source.push_back(barrier); + // A side effect consumes the old value returned by the actual SUATOM translation. + auto store = op(0x388); field(store,24,8,255); field(store,32,8,d); field(store,73,3,4); + source.push_back(store); source.push_back(exit_op()); + const auto lowered = Shader::Ampere::LowerToMaxwell(source); CHECK(lowered); + if (d == 239) CHECK(lowered.register_map_words > 0); + Environment env(lowered,0x90); + Shader::ObjectPool insts; + Shader::ObjectPool blocks; + Shader::ObjectPool flow_blocks; + Shader::Maxwell::Flow::CFG cfg{env,flow_blocks,0x50}; + Shader::HostTranslateInfo host{}; host.support_int64 = host.support_float64 = host.support_float16 = true; + host.support_conditional_barrier = true; + auto program = Shader::Maxwell::TranslateProgram(insts,blocks,env,cfg,host); + CHECK(program.info.image_descriptors.size() == 1); + const auto& descriptor = program.info.image_descriptors[0]; + CHECK(descriptor.cbuf_index == 3 && descriptor.cbuf_offset == 0x90 && descriptor.is_read && descriptor.is_written); + CHECK(descriptor.format == Shader::ImageFormat::R32_UINT && descriptor.type == Shader::TextureType::Color1D); + Shader::Profile profile{}; profile.supported_spirv = 0x00010600; + profile.support_int64 = true; profile.support_descriptor_aliasing = true; + const auto spirv = Shader::Backend::SPIRV::EmitSPIRV(profile,program); + std::map constants; + std::map> phis; + std::vector stored; + uint32_t atomic = 0, atomic_count = 0, barriers = 0, selects = 0; + bool texel_pointer = false; + for (size_t at = 5; at < spirv.size();) { + const uint32_t count = spirv[at] >> 16; CHECK(count && at + count <= spirv.size()); + const auto opcode = static_cast(spirv[at] & 0xffff); + const auto* w = &spirv[at+1]; + if (opcode == spv::Op::OpConstant && count == 4) constants[w[1]] = w[2]; + if (opcode == spv::Op::OpImageTexelPointer) { + CHECK(constants.at(w[3]) == 3); texel_pointer = true; + } + if (opcode == spv::Op::OpAtomicIAdd) { + CHECK(count == 7 && constants.at(w[3]) == uint32_t(spv::ScopeDevice) && constants.at(w[4]) == 0); + CHECK(constants.at(w[5]) == (a == b ? 3U : 11U)); + atomic = w[1]; ++atomic_count; + } + if (opcode == spv::Op::OpMemoryBarrier) { + CHECK(constants.at(w[0]) == uint32_t(spv::ScopeDevice)); + CHECK(constants.at(w[1]) & uint32_t(spv::MemorySemanticsImageMemoryMask)); + CHECK(constants.at(w[1]) & uint32_t(spv::MemorySemanticsAcquireReleaseMask)); + ++barriers; + } + if (opcode == spv::Op::OpSelectionMerge) ++selects; + if (opcode == spv::Op::OpPhi) { + for (uint32_t n = 2; n + 1 < count - 1; n += 2) phis[w[1]].push_back(w[n]); + } + if (opcode == spv::Op::OpStore) stored.push_back(w[1]); + at += count; + } + CHECK(atomic_count == 1 && texel_pointer); + CHECK(barriers == (explicit_barriers ? 2U : 0U)); + CHECK((pred == 7 && !branch) || selects > 0); + std::set values{atomic}; + bool changed; + do { + changed = false; + for (const auto& [result,inputs] : phis) + for (auto input : inputs) if (values.contains(input)) changed |= values.insert(result).second; + } while (changed); + bool consumed = false; + for (auto value : stored) consumed |= values.contains(value); + CHECK(consumed == (d != 255)); // discarding the result must still retain the atomic side effect + std::filesystem::create_directories(directory); + std::ofstream out(directory / (std::string(name) + ".spv"),std::ios::binary); + out.write(reinterpret_cast(spirv.data()),static_cast(spirv.size()*4)); CHECK(out.good()); + std::printf("PASS: %s actual lowering/Maxwell/IR/SPIR-V, descriptor and atomic result contract\n",name); +} +#endif + +int main(int argc, char** argv) { + if (argc == 2 && std::string(argv[1]) == "--expect-original-rejection") { + const auto r = lower(suatom()); + CHECK(!r && r.error == "unsupported native opcode" && r.error_offset == 0); + std::puts("PASS: original production lowering rejects the authored direct SUATOM"); return 0; + } + lowering_checks(); +#ifndef SUATOM_LOWERING_ONLY + CHECK(argc == 2); + const std::filesystem::path directory{argv[1]}; + module_check(directory,"separate",13,9,17); + module_check(directory,"dest-coordinate-alias",9,9,17); + module_check(directory,"dest-addend-alias",17,9,17); + module_check(directory,"all-alias",9,9,9); + module_check(directory,"high-addend",13,9,250); + module_check(directory,"high-all-alias",250,250,250); + module_check(directory,"reserved239-alias",239,239,239); + module_check(directory,"predicate",13,9,17,2); + module_check(directory,"inverted-predicate",13,9,17,10); + module_check(directory,"relocated-branch",250,250,250,7,true); + module_check(directory,"explicit-barriers",13,9,17,7,false,true); + module_check(directory,"discard-old-value",255,9,17); +#endif +} diff --git a/tests/native_suatom/CMakeLists.txt b/tests/native_suatom/CMakeLists.txt new file mode 100644 index 000000000..eb01434dd --- /dev/null +++ b/tests/native_suatom/CMakeLists.txt @@ -0,0 +1,26 @@ +cmake_minimum_required(VERSION 3.20) +project(native_suatom LANGUAGES CXX) +set(CMAKE_CXX_STANDARD 20) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +if(NOT DEFINED REPO_ROOT OR NOT DEFINED SPIRV_HEADERS_ROOT) + message(FATAL_ERROR "REPO_ROOT and SPIRV_HEADERS_ROOT are required") +endif() +set(SHADER "${REPO_ROOT}/gpu/vk/eden/shader_recompiler") +file(GLOB_RECURSE SOURCES CONFIGURE_DEPENDS + "${SHADER}/backend/spirv/*.cpp" + "${SHADER}/frontend/ir/*.cpp" + "${SHADER}/frontend/maxwell/*.cpp" + "${SHADER}/ir_opt/*.cpp" + "${REPO_ROOT}/gpu/vk/sirit/src/*.cpp") +add_executable(native_suatom "${REPO_ROOT}/tests/native_suatom.cpp" ${SOURCES}) +target_include_directories(native_suatom PRIVATE + "${REPO_ROOT}/gpu/vk/compat" "${REPO_ROOT}/gpu/compat" + "${REPO_ROOT}/gpu/vk/eden" "${REPO_ROOT}/gpu/third_party/suyu" + "${REPO_ROOT}/gpu/vk/sirit/include" "${REPO_ROOT}/gpu/vk/sirit/src" + "${SPIRV_HEADERS_ROOT}/include") +target_compile_definitions(native_suatom PRIVATE NDEBUG MKWVK_INCLUDES=2 VK_NO_PROTOTYPES) +if(MSVC) + target_compile_options(native_suatom PRIVATE /W0 /O1 /EHsc) +else() + target_compile_options(native_suatom PRIVATE -w -O1) +endif() diff --git a/tests/native_suatom/README.md b/tests/native_suatom/README.md new file mode 100644 index 000000000..1278c1657 --- /dev/null +++ b/tests/native_suatom/README.md @@ -0,0 +1,34 @@ +# Direct surface atomic regression + +The fields in `../native_suatom.cpp` are independently authored. No game shader, +dump, cache or private SDK instruction is used. The small Python check compiles +the production lowering header; the CMake target additionally runs the same +instructions through the actual Maxwell CFG/translator, optimization passes and +SPIR-V emitter. Every emitted module is validated in the existing Linux SPIR-V +job. No GPU or game executes. + +The accepted form is direct-descriptor `0x593`, 1D U32 ADD, sample addressing, +IGN, Strong(GPU), normal eviction, and no fault-predicate destination. Other +forms and modifiers are refused. Registers and descriptor fields are authored +separately, with alias/high-register, predication, branch relocation, offset/bank +boundary and conflict cases. The full path checks the descriptor bank and byte +offset, atomic coordinate/addend, and a store consuming the returned old value. + +Strong(GPU) is not an acquire/release qualifier. Mesa emits image atomics with +`MemOrder::Strong(MemScope::GPU)` and handles explicit barriers separately. The +existing image atomic backend emits a Device-scoped relaxed atomic; no fences +are added by this lowering. A case with explicit native MEMBARs verifies those +barriers survive independently. + +Encoding references: + +- [Independent decoder](https://github.com/seedds/nxnative/blob/7f50222003e4270e3655d4f798e85744d390edf0/dm/third_party/ryujinx/src/Ryujinx.Graphics.Shader/Instructions/InstEmitAmpereMem.cs#L529): direct descriptor bank/word offset and SUATOM operands. +- [Mesa encoder](https://chromium.googlesource.com/external/gitlab.freedesktop.org/mesa/mesa/+/b1fe47e94409d81924c8663f58fac97e66ab4527/src/nouveau/compiler/nak/sm70_encode.rs): shared atomic/dimension/type/order fields. Mesa encodes its register-handle form, not this direct-descriptor opcode. +- [Mesa NIR translation](https://chromium.googlesource.com/external/gitlab.freedesktop.org/mesa/mesa/+/b1fe47e94409d81924c8663f58fac97e66ab4527/src/nouveau/compiler/nak/from_nir.rs): Strong(GPU) image atomics and separate barriers. +- [Public SM87 encoding tables](https://mbuchel.github.io/sass-encoding-lab/): `0x593` and IGN. The opcode table SHA-256 is `559048b6a840e5548ba5fa577e31230f0a217f370d50c7c378e12997d96dd747`. + +Configure with Clang C++20 and `-DREPO_ROOT=` plus +`-DSPIRV_HEADERS_ROOT=`, then build `native_suatom`. Run it with an +authored output directory and validate its `.spv` files using `spirv-val +--target-env vulkan1.3`. These checks establish translation and module validity, +not driver execution, general surface-atomic coverage or a visible game fix. diff --git a/tests/run_all.py b/tests/run_all.py index e853b0b28..398da8843 100644 --- a/tests/run_all.py +++ b/tests/run_all.py @@ -38,6 +38,7 @@ "tests/test_cpp_preflight.py": ["--cxx", "g++"], "tests/test_invalidate_vcore.py": ["--cxx", "g++"], "tests/test_lower_atomg_cas.py": ["--cxx", "g++"], + "tests/test_native_suatom.py": ["--cxx", "clang++"], "nx2/test_cbuf_table_extent.py": ["--cxx", "g++"], "nx2/test_dump_ssbo_read.py": ["--cxx", "g++"], "nx2/test_mesh_dynamic_state.py": ["--cxx", "g++"], diff --git a/tests/test_native_suatom.py b/tests/test_native_suatom.py new file mode 100644 index 000000000..72d380e01 --- /dev/null +++ b/tests/test_native_suatom.py @@ -0,0 +1,34 @@ +"""Authored direct SUATOM fields exercise the real production lowering; no GPU or shader dump. + +The same C++ cases also run through Maxwell/IR/SPIR-V in tests/native_suatom/CMakeLists.txt. +""" +import argparse +import os +from pathlib import Path +import subprocess +import tempfile + +ROOT = Path(__file__).resolve().parents[1] + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--cxx", default="clang++") + args = parser.parse_args() + flags = (subprocess.CREATE_NO_WINDOW | subprocess.BELOW_NORMAL_PRIORITY_CLASS) if os.name == "nt" else 0 + def run(command): + result = subprocess.run(command, stdout=subprocess.PIPE, stderr=subprocess.STDOUT, + text=True, creationflags=flags) + print(result.stdout, end="", flush=True) + result.check_returncode() + + with tempfile.TemporaryDirectory(prefix="native-suatom-") as tmp: + exe = Path(tmp) / "fixture.exe" + run([args.cxx, "-std=c++20", "-O1", "-DSUATOM_LOWERING_ONLY", + "-I" + str(ROOT / "gpu/vk/compat"), "-I" + str(ROOT / "gpu/vk/eden"), + str(ROOT / "tests/native_suatom.cpp"), "-o", str(exe)]) + run([str(exe)]) + + +if __name__ == "__main__": + main()