From 15652afe14c87641bbeccc02b3946c28c40ad06e Mon Sep 17 00:00:00 2001 From: monsterlag <64719697+monsterlag@users.noreply.github.com> Date: Sun, 8 Nov 2026 17:55:49 -0500 Subject: [PATCH] fix: lower direct 1D surface atomic add Preserve the old U32 image value, descriptor location and scalar register aliases for the supported native SUATOM form. Keep Strong(GPU) atomic coherence on the existing Device/Relaxed path and leave explicit MEMBAR instructions separate. Exercise authored instructions through production lowering, Maxwell CFG translation and SPIR-V emission, including strict mode rejection and high-register mapping. Assisted-by: OpenAI Codex Co-Authored-By: OpenAI Codex --- .github/workflows/source-regressions.yml | 15 ++ .../frontend/ampere/lower_to_maxwell.h | 15 ++ tests/native_suatom.cpp | 251 ++++++++++++++++++ tests/native_suatom/CMakeLists.txt | 26 ++ tests/native_suatom/README.md | 34 +++ tests/run_all.py | 1 + tests/test_native_suatom.py | 34 +++ 7 files changed, 376 insertions(+) create mode 100644 tests/native_suatom.cpp create mode 100644 tests/native_suatom/CMakeLists.txt create mode 100644 tests/native_suatom/README.md create mode 100644 tests/test_native_suatom.py diff --git a/.github/workflows/source-regressions.yml b/.github/workflows/source-regressions.yml index 42bdf63f0..aeef7b07b 100644 --- a/.github/workflows/source-regressions.yml +++ b/.github/workflows/source-regressions.yml @@ -146,6 +146,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 13e80bcf4..c58137e98 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; @@ -1438,6 +1440,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 b61f843bd..d60718a5a 100644 --- a/tests/run_all.py +++ b/tests/run_all.py @@ -35,6 +35,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()