From 6236d3cb23ab3127cb0909e28b4828e941e3e3af Mon Sep 17 00:00:00 2001 From: Yanpeng hu Date: Sat, 29 Aug 2026 11:41:29 -0700 Subject: [PATCH 1/4] Count asynchronous global copies as unsupported instead of dropping them `cp.async` and the bulk tensor copy instructions read global memory without going through a register. The rewrite pattern in src/ptxpass_hbf/ptx_memory_op.cpp only matches `ld`/`st` `.global`, and the unsupported pattern in src/ptxpass_hbf/transform.cpp listed `atom`, `red`, non-global `ld`/`st`, `tex`, `suld`, `sust` and inline `asm`. An asynchronous copy matched neither, so an HBF address reached by one produced no entry of any kind: no modeled delay, and no unsupported-list entry either. The coverage record showed no trace the instruction had been there. Line 27 of docs/superpowers/specs/2026-08-09-hbfsim-hybrid-design.md sets the goal this violates: "Fail closed whenever an HBF address could reach an uninstrumented or unsupported memory operation." The two new alternatives require `.global` in the opcode on purpose. The same families carry pure synchronisation forms -- cp.async.commit_group, cp.async.wait_group, cp.async.bulk.wait_group, cp.async.mbarrier.arrive -- which touch no memory, and a shared-to-shared bulk copy names no global space either. None of those may be reported as unsupported memory operations. WHAT THIS CHANGES DOWNSTREAM. The unsupported list is filtered by hbf_relevant_unsupported_opcode in src/ptxpass_hbf/plugin.cpp, written into the module manifest, and consumed at src/cuda_runtime/coverage.cpp:379-396. There are two outcomes, and the second needs a decision: - On a timing-backed range under a non-strict policy, the launch is still allowed and is recorded as `opaque_unmodeled_timing` with `modeled: false`. This is a pure visibility gain: the access was already uncharged, and now it is uncharged and recorded. - Under a strict policy, and on a capacity-backed range, the launch is refused with reason `unsupported_operation`. A kernel that uses these instructions against a registered range and used to run will now be refused. That is what failing closed means here, but whether you want it turned on now is yours to decide, since Triton emits `cp.async` freely on recent architectures. tests/cpu/ptx_async_copy_coverage_test.cpp fails on the previous pattern at the first asynchronous copy and passes on this one. Each case is measured as an increment over the same kernel without the instruction, because the kernel needs a `ld.param` to load its pointer and `ld.param` is itself on the unsupported list; counting absolute totals would have reported every case as unsupported and hidden the defect. No case in the file rewrites an instruction, so the file does not need the embedded device helper and runs in a CPU-only build; the rewriting cases stay in ptx_transform_test. Test suite before and after: the same four tests fail (context_lifecycle, vmem_tuning, run_with_bpftime, mqsim_benchmark), all pre-existing in a CPU-only build. Co-Authored-By: Claude Opus 5 (1M context) --- CMakeLists.txt | 14 ++ src/ptxpass_hbf/transform.cpp | 7 +- tests/cpu/ptx_async_copy_coverage_test.cpp | 150 +++++++++++++++++++++ 3 files changed, 170 insertions(+), 1 deletion(-) create mode 100644 tests/cpu/ptx_async_copy_coverage_test.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 5fa1a53..0793209 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -251,6 +251,20 @@ set_tests_properties( ptx_transform PROPERTIES WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}" ) +add_executable( + ptx_async_copy_coverage_test tests/cpu/ptx_async_copy_coverage_test.cpp +) +target_link_libraries(ptx_async_copy_coverage_test PRIVATE hbfsim_core) +target_include_directories( + ptx_async_copy_coverage_test + PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/src/ptxpass_hbf" +) +add_test(NAME ptx_async_copy_coverage COMMAND ptx_async_copy_coverage_test) +set_tests_properties( + ptx_async_copy_coverage + PROPERTIES WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}" +) + add_executable(ptxpass_hbf src/ptxpass_hbf/main.cpp) target_link_libraries(ptxpass_hbf PRIVATE hbfsim_core) target_include_directories( diff --git a/src/ptxpass_hbf/transform.cpp b/src/ptxpass_hbf/transform.cpp index 81c6e06..afd636e 100644 --- a/src/ptxpass_hbf/transform.cpp +++ b/src/ptxpass_hbf/transform.cpp @@ -51,7 +51,12 @@ bool unsupported_memory_instruction(const std::string& line, std::string& opcode) { static const std::regex expression( - R"(^\s*(?:@!?%[A-Za-z0-9_$]+\s+)?((?:atom|red)\.global\S*|ld\.(?!global)\S*|st\.(?!global)\S*|tex\S*|suld\S*|sust\S*|asm\s*\().*;\s*(?://.*)?$)"); + // The `cp.async` and `cp.reduce.async` alternatives require `.global` + // in the opcode on purpose. The same families carry pure + // synchronisation forms -- cp.async.commit_group, cp.async.wait_group, + // cp.async.bulk.wait_group -- which touch no memory and must not be + // reported as unsupported memory operations. + R"(^\s*(?:@!?%[A-Za-z0-9_$]+\s+)?((?:atom|red)\.global\S*|ld\.(?!global)\S*|st\.(?!global)\S*|cp\.async\S*\.global\S*|cp\.reduce\.async\S*\.global\S*|tex\S*|suld\S*|sust\S*|asm\s*\().*;\s*(?://.*)?$)"); std::smatch match; if (!std::regex_match(line, match, expression)) { return false; diff --git a/tests/cpu/ptx_async_copy_coverage_test.cpp b/tests/cpu/ptx_async_copy_coverage_test.cpp new file mode 100644 index 0000000..62e5713 --- /dev/null +++ b/tests/cpu/ptx_async_copy_coverage_test.cpp @@ -0,0 +1,150 @@ +// `cp.async` and the bulk tensor copy instructions read global memory without +// going through a register. Before this test, neither the rewrite pattern in +// src/ptxpass_hbf/ptx_memory_op.cpp nor the unsupported pattern in +// src/ptxpass_hbf/transform.cpp matched them, so an HBF address reached by one +// of them produced no entry of any kind: no modeled delay, and no +// unsupported-list entry either. The design goal on line 27 of +// docs/superpowers/specs/2026-08-09-hbfsim-hybrid-design.md is to fail closed +// whenever an HBF address could reach an uninstrumented or unsupported memory +// operation, which needs the instruction to be visible first. +// +// The point this test pins down is narrow: an asynchronous copy that names +// .global has to be counted, and the synchronisation instructions of the same +// family, which touch no memory, must not be. + +#include "ptx_memory_op.hpp" +#include "transform.hpp" + +#include +#include +#include + +#define CHECK(condition) \ + do { \ + if (!(condition)) { \ + std::printf("failed at line %d: %s\n", __LINE__, #condition); \ + return __LINE__; \ + } \ + } while (false) + +namespace { + +// Wraps one instruction in the smallest kernel the pass will walk. +std::string kernel_with(const std::string& instruction) +{ + return ".version 8.0\n" + ".target sm_120\n" + ".address_size 64\n" + ".visible .entry probe(.param .u64 probe_param_0)\n" + "{\n" + " .reg .b32 %r<8>;\n" + " .reg .b64 %rd<8>;\n" + " .reg .f32 %f<8>;\n" + " ld.param.u64 %rd1, [probe_param_0];\n" + + std::string{" "} + instruction + "\n" + + " ret;\n" + "}\n"; +} + +hbfsim::ptx::TransformResult run(const std::string& instruction) +{ + hbfsim::ptx::TransformRequest request{}; + request.full_ptx = kernel_with(instruction); + return hbfsim::ptx::transform_ptx(request); +} + +// The kernel body needs a `ld.param` to load the pointer, and `ld.param` +// itself matches the unsupported pattern. Counting the absolute total would +// therefore report every instruction as unsupported, so each case is measured +// as the increment over the same kernel without the instruction under test. +std::uint64_t baseline_unsupported() +{ + static const auto value = + run("ret;").coverage.unsupported_instructions; + return value; +} + +std::uint64_t baseline_rewritten() +{ + static const auto value = run("ret;").coverage.rewritten_instructions; + return value; +} + +bool counted_unsupported(const std::string& instruction) +{ + return run(instruction).coverage.unsupported_instructions > + baseline_unsupported(); +} + +} // namespace + +int main() +{ + // Reads global memory into shared memory. Must be visible. + CHECK(counted_unsupported( + "cp.async.ca.shared.global [%r1], [%rd1], 4;")); + CHECK(counted_unsupported( + "cp.async.cg.shared.global [%r1], [%rd1], 16;")); + + // Bulk tensor copy from global. Must be visible. + CHECK(counted_unsupported( + "cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::" + "complete_tx::bytes [%r1], [tmap, {%r2,%r3}], [%r4];")); + CHECK(counted_unsupported( + "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes " + "[%r1], [%rd1], %r2, [%r3];")); + + // Reduction form that reads global. Must be visible. + CHECK(counted_unsupported( + "cp.reduce.async.bulk.global.shared::cta.bulk_group.add.u32 " + "[%rd1], [%r1], %r2;")); + + // Same family, but pure synchronisation: these touch no memory and must + // not be reported as unsupported memory operations. + CHECK(!counted_unsupported("cp.async.commit_group;")); + CHECK(!counted_unsupported("cp.async.wait_group 0;")); + CHECK(!counted_unsupported("cp.async.wait_all;")); + CHECK(!counted_unsupported("cp.async.bulk.commit_group;")); + CHECK(!counted_unsupported("cp.async.bulk.wait_group.read 0;")); + + // The instructions the pass already handled must keep their old + // classification: an ordinary global load is still rewritten and is not on + // the unsupported list, and a shared load is still unsupported. + // An ordinary global load must still be rewritten rather than counted + // here. That case is not exercised in this file: rewriting makes + // transform_ptx append the embedded device helper, which only exists in a + // CUDA build, so the assertion lives in ptx_transform_test instead. Every + // case in this file is chosen so that nothing is rewritten. + { + const auto result = run("ld.shared.u32 %r1, [%r2];"); + CHECK(result.coverage.rewritten_instructions == baseline_rewritten()); + CHECK(result.coverage.unsupported_instructions == + baseline_unsupported() + 1); + } + { + const auto result = run("atom.global.add.u32 %r1, [%rd1], 1;"); + CHECK(result.coverage.rewritten_instructions == baseline_rewritten()); + CHECK(result.coverage.unsupported_instructions == + baseline_unsupported() + 1); + } + + // An asynchronous copy that never names global memory is a shared-to-shared + // move and is not an HBF access. + CHECK(!counted_unsupported( + "cp.async.bulk.shared::cluster.shared::cta [%r1], [%r2], %r3;")); + + // The recorded opcode has to name the instruction, so the coverage record + // says which operation was refused. + { + const auto result = run("cp.async.ca.shared.global [%r1], [%rd1], 4;"); + const auto named = std::any_of( + result.coverage.unsupported_opcodes.begin(), + result.coverage.unsupported_opcodes.end(), + [](const std::string& opcode) { + return opcode.rfind("cp.async", 0) == 0; + }); + CHECK(named); + } + + return 0; +} From 9414191e6c8486de0a2d84e8e77e1346c0c0e647 Mon Sep 17 00:00:00 2001 From: Yanpeng hu Date: Sat, 29 Aug 2026 12:23:22 -0700 Subject: [PATCH 2/4] Leave the bulk tensor prefixes to the branch that models them The first version of this pattern matched every `cp.async` naming `.global`, including the bulk tensor forms. Branch feature/sm120-exact-stage1 models those forms rather than refusing them: parse_tma in src/ptxpass_hbf/ptx_async_op.cpp accepts exactly three prefixes -- cp.async.bulk.tensor., cp.reduce.async.bulk.tensor. and cp.async.bulk.prefetch.tensor. -- and gives the third of them its own direction, TmaDirection::Prefetch, with a different completion rule. Marking those unsupported here would refuse, once that branch merges, exactly the launches it can model. Two negative lookaheads now exclude the three prefixes. What remains is the gap that branch leaves open: its unsupported pattern is character-for-character the one on hybrid and carries no `cp.` alternative either, so the plain cp.async.ca/cg.shared.global form -- the one Triton emits freely on earlier architectures -- is matched by neither of its patterns. The two lines together now cover the family without overlapping. Test suite: the same four tests fail before and after (context_lifecycle, vmem_tuning, run_with_bpftime, mqsim_benchmark), all pre-existing in a CPU-only build. Co-Authored-By: Claude Opus 5 (1M context) --- src/ptxpass_hbf/transform.cpp | 12 ++++++++++- tests/cpu/ptx_async_copy_coverage_test.cpp | 25 +++++++++++++++++----- 2 files changed, 31 insertions(+), 6 deletions(-) diff --git a/src/ptxpass_hbf/transform.cpp b/src/ptxpass_hbf/transform.cpp index afd636e..2f6cbc9 100644 --- a/src/ptxpass_hbf/transform.cpp +++ b/src/ptxpass_hbf/transform.cpp @@ -56,7 +56,17 @@ bool unsupported_memory_instruction(const std::string& line, // synchronisation forms -- cp.async.commit_group, cp.async.wait_group, // cp.async.bulk.wait_group -- which touch no memory and must not be // reported as unsupported memory operations. - R"(^\s*(?:@!?%[A-Za-z0-9_$]+\s+)?((?:atom|red)\.global\S*|ld\.(?!global)\S*|st\.(?!global)\S*|cp\.async\S*\.global\S*|cp\.reduce\.async\S*\.global\S*|tex\S*|suld\S*|sust\S*|asm\s*\().*;\s*(?://.*)?$)"); + // + // The three bulk TENSOR prefixes are excluded by the lookaheads: + // cp.async.bulk.tensor., cp.reduce.async.bulk.tensor. and + // cp.async.bulk.prefetch.tensor. Those are the prefixes parse_tma + // accepts on branch feature/sm120-exact-stage1, where they are modeled + // rather than refused. Matching them here would refuse, once that + // branch merges, exactly the launches it can model. This pattern + // closes the gap that branch leaves open -- the plain + // cp.async.ca/cg.shared.global form, which its unsupported pattern + // does not match either -- and stays out of what it handles. + R"(^\s*(?:@!?%[A-Za-z0-9_$]+\s+)?((?:atom|red)\.global\S*|ld\.(?!global)\S*|st\.(?!global)\S*|cp\.async(?!\.bulk\.(?:tensor|prefetch\.tensor)\.)\S*\.global\S*|cp\.reduce\.async(?!\.bulk\.tensor\.)\S*\.global\S*|tex\S*|suld\S*|sust\S*|asm\s*\().*;\s*(?://.*)?$)"); std::smatch match; if (!std::regex_match(line, match, expression)) { return false; diff --git a/tests/cpu/ptx_async_copy_coverage_test.cpp b/tests/cpu/ptx_async_copy_coverage_test.cpp index 62e5713..c45d286 100644 --- a/tests/cpu/ptx_async_copy_coverage_test.cpp +++ b/tests/cpu/ptx_async_copy_coverage_test.cpp @@ -86,19 +86,34 @@ int main() CHECK(counted_unsupported( "cp.async.cg.shared.global [%r1], [%rd1], 16;")); - // Bulk tensor copy from global. Must be visible. - CHECK(counted_unsupported( - "cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::" - "complete_tx::bytes [%r1], [tmap, {%r2,%r3}], [%r4];")); + // Non-tensor bulk copy from global. Must be visible. CHECK(counted_unsupported( "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes " "[%r1], [%rd1], %r2, [%r3];")); - // Reduction form that reads global. Must be visible. + // Reduction form that reads global and is not a tensor copy. Must be + // visible. CHECK(counted_unsupported( "cp.reduce.async.bulk.global.shared::cta.bulk_group.add.u32 " "[%rd1], [%r1], %r2;")); + // The three bulk TENSOR prefixes are deliberately left alone, because + // parse_tma on branch feature/sm120-exact-stage1 models them properly + // rather than refusing them: cp.async.bulk.tensor., + // cp.reduce.async.bulk.tensor. and cp.async.bulk.prefetch.tensor. Marking + // them unsupported here would refuse, once that branch merges, exactly the + // launches it can model. This pass closes the gap that branch leaves open, + // and stays out of what it handles. + CHECK(!counted_unsupported( + "cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::" + "complete_tx::bytes [%r1], [tmap, {%r2,%r3}], [%r4];")); + CHECK(!counted_unsupported( + "cp.reduce.async.bulk.tensor.2d.global.shared::cta.add.tile." + "bulk_group [tmap, {%r2,%r3}], [%r1];")); + CHECK(!counted_unsupported( + "cp.async.bulk.prefetch.tensor.2d.L2.global.tile " + "[tmap, {%r2,%r3}];")); + // Same family, but pure synchronisation: these touch no memory and must // not be reported as unsupported memory operations. CHECK(!counted_unsupported("cp.async.commit_group;")); From 66f5fb76b6f3c0a9c7e7ab9a91e73e26d4447bc5 Mon Sep 17 00:00:00 2001 From: Yanpeng hu Date: Sun, 30 Aug 2026 03:23:57 -0700 Subject: [PATCH 3/4] Keep the bulk tensor families unsupported on hybrid, per review The previous commit excluded cp.async.bulk.tensor., cp.reduce.async.bulk.tensor. and cp.async.bulk.prefetch.tensor. so this pattern would not collide with parse_tma on feature/sm120-exact-stage1. The review is right that this was the wrong trade. hybrid has no parse_tma, no ptx_async_op and no TmaDirection, so on the only branch that exists today the exclusion left those instructions neither modeled nor reported as unsupported -- preserving the exact fail-open coverage hole this pattern exists to close, and encoding that behaviour in the test. The lookaheads are removed and the three families are counted again. The test asserts the positive case for all three. Both this file's test and ptx_transform pass. The merge collision is real but belongs at merge time: when feature/sm120-exact-stage1 lands, the same commit removes the three prefixes it models from this pattern, so the hole is never open in between. That instruction is now in the comment beside the pattern rather than in a branch that has not merged. Test suite: the same four tests fail before and after (context_lifecycle, vmem_tuning, run_with_bpftime, mqsim_benchmark), all pre-existing in a CPU-only build. Co-Authored-By: Claude Opus 5 (1M context) --- src/ptxpass_hbf/transform.cpp | 22 ++++++++++++---------- tests/cpu/ptx_async_copy_coverage_test.cpp | 20 ++++++++++---------- 2 files changed, 22 insertions(+), 20 deletions(-) diff --git a/src/ptxpass_hbf/transform.cpp b/src/ptxpass_hbf/transform.cpp index 2f6cbc9..1c5dd7a 100644 --- a/src/ptxpass_hbf/transform.cpp +++ b/src/ptxpass_hbf/transform.cpp @@ -57,16 +57,18 @@ bool unsupported_memory_instruction(const std::string& line, // cp.async.bulk.wait_group -- which touch no memory and must not be // reported as unsupported memory operations. // - // The three bulk TENSOR prefixes are excluded by the lookaheads: - // cp.async.bulk.tensor., cp.reduce.async.bulk.tensor. and - // cp.async.bulk.prefetch.tensor. Those are the prefixes parse_tma - // accepts on branch feature/sm120-exact-stage1, where they are modeled - // rather than refused. Matching them here would refuse, once that - // branch merges, exactly the launches it can model. This pattern - // closes the gap that branch leaves open -- the plain - // cp.async.ca/cg.shared.global form, which its unsupported pattern - // does not match either -- and stays out of what it handles. - R"(^\s*(?:@!?%[A-Za-z0-9_$]+\s+)?((?:atom|red)\.global\S*|ld\.(?!global)\S*|st\.(?!global)\S*|cp\.async(?!\.bulk\.(?:tensor|prefetch\.tensor)\.)\S*\.global\S*|cp\.reduce\.async(?!\.bulk\.tensor\.)\S*\.global\S*|tex\S*|suld\S*|sust\S*|asm\s*\().*;\s*(?://.*)?$)"); + // The bulk TENSOR families are matched here too, deliberately. Branch + // feature/sm120-exact-stage1 models them in parse_tma, and an earlier + // version of this pattern excluded them so the two would not collide + // on merge. That was the wrong trade: hybrid has no parse_tma, so + // excluding them left those instructions neither modeled nor reported + // -- the exact fail-open hole this pattern exists to close, preserved + // on the only branch that exists today. When + // feature/sm120-exact-stage1 merges, remove the three prefixes it + // handles (cp.async.bulk.tensor., cp.reduce.async.bulk.tensor. and + // cp.async.bulk.prefetch.tensor.) from this pattern in the same + // commit, so the coverage hole is never open in between. + R"(^\s*(?:@!?%[A-Za-z0-9_$]+\s+)?((?:atom|red)\.global\S*|ld\.(?!global)\S*|st\.(?!global)\S*|cp\.async\S*\.global\S*|cp\.reduce\.async\S*\.global\S*|tex\S*|suld\S*|sust\S*|asm\s*\().*;\s*(?://.*)?$)"); std::smatch match; if (!std::regex_match(line, match, expression)) { return false; diff --git a/tests/cpu/ptx_async_copy_coverage_test.cpp b/tests/cpu/ptx_async_copy_coverage_test.cpp index c45d286..e6863eb 100644 --- a/tests/cpu/ptx_async_copy_coverage_test.cpp +++ b/tests/cpu/ptx_async_copy_coverage_test.cpp @@ -97,20 +97,20 @@ int main() "cp.reduce.async.bulk.global.shared::cta.bulk_group.add.u32 " "[%rd1], [%r1], %r2;")); - // The three bulk TENSOR prefixes are deliberately left alone, because - // parse_tma on branch feature/sm120-exact-stage1 models them properly - // rather than refusing them: cp.async.bulk.tensor., - // cp.reduce.async.bulk.tensor. and cp.async.bulk.prefetch.tensor. Marking - // them unsupported here would refuse, once that branch merges, exactly the - // launches it can model. This pass closes the gap that branch leaves open, - // and stays out of what it handles. - CHECK(!counted_unsupported( + // The bulk TENSOR families are counted here too. Branch + // feature/sm120-exact-stage1 models them in parse_tma, and an earlier + // version of this test required them NOT to be counted so the two would + // not collide on merge. That was wrong: hybrid has no parse_tma, so the + // exclusion left them neither modeled nor reported, preserving the exact + // hole this file exists to close. They stay counted until that branch + // merges, and the merge commit removes them from both sides at once. + CHECK(counted_unsupported( "cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::" "complete_tx::bytes [%r1], [tmap, {%r2,%r3}], [%r4];")); - CHECK(!counted_unsupported( + CHECK(counted_unsupported( "cp.reduce.async.bulk.tensor.2d.global.shared::cta.add.tile." "bulk_group [tmap, {%r2,%r3}], [%r1];")); - CHECK(!counted_unsupported( + CHECK(counted_unsupported( "cp.async.bulk.prefetch.tensor.2d.L2.global.tile " "[tmap, {%r2,%r3}];")); From 936b8e548c0ad487975fea0de73938a1949bc156 Mon Sep 17 00:00:00 2001 From: Yanpeng hu Date: Mon, 31 Aug 2026 15:24:00 -0700 Subject: [PATCH 4/4] Scan whole PTX statements, not physical lines The review is right. A PTX statement may be written across several physical lines, and the unsupported scan matched one line at a time, so an asynchronous copy split across two lines matched neither half and was never reported. If the same kernel also held an ordinary global load, the module was still marked instrumented, and the launch proceeded with an access nothing had recorded -- the fail-open case this scan exists to prevent. Physical lines are now accumulated into one logical statement before the scan. A statement is treated as open until its text, with any trailing comment removed, ends in one of `;`, `{`, `}` or `:`. The rewrite path is untouched and still matches per line; only the scan sees joined statements, so no rewritten output changes. The new case in ptx_async_copy_coverage_test writes the same bulk tensor copy once across two lines and once on one line, and requires the same count from both. It fails on the previous pass at the equality assertion and passes on this one. Test suite: 88% passed, the same four failing as before (context_lifecycle, vmem_tuning, run_with_bpftime, mqsim_benchmark), all pre-existing in a CPU-only build. Co-Authored-By: Claude Opus 5 (1M context) --- src/ptxpass_hbf/transform.cpp | 49 +++++++++++++++++++++- tests/cpu/ptx_async_copy_coverage_test.cpp | 35 ++++++++++++++++ 2 files changed, 83 insertions(+), 1 deletion(-) diff --git a/src/ptxpass_hbf/transform.cpp b/src/ptxpass_hbf/transform.cpp index 1c5dd7a..3d14a08 100644 --- a/src/ptxpass_hbf/transform.cpp +++ b/src/ptxpass_hbf/transform.cpp @@ -77,6 +77,44 @@ bool unsupported_memory_instruction(const std::string& line, return true; } +std::string joined_statement(const std::string& pending, + const std::string& line) +{ + if (pending.empty()) { + return line; + } + auto trimmed = line; + const auto first = trimmed.find_first_not_of(" \t"); + if (first != std::string::npos) { + trimmed.erase(0, first); + } + return pending + " " + trimmed; +} + +// A PTX statement may be written across several physical lines. The rewrite +// path matches per line, which is enough for the forms it rewrites, but the +// unsupported scan must see whole statements: an asynchronous copy split +// across two lines matches neither line on its own, so scanning per line lets +// it through unreported, and if the same kernel also holds an ordinary +// ld.global the module is still marked instrumented. The launch then proceeds +// with an unreported access, which is the fail-open case the scan exists to +// prevent. +bool statement_is_open(const std::string& text) +{ + auto without_comment = text; + if (const auto comment = without_comment.find("//"); + comment != std::string::npos) { + without_comment.erase(comment); + } + const auto last = without_comment.find_last_not_of(" \t\r"); + if (last == std::string::npos) { + return false; + } + const auto character = without_comment[last]; + return character != ';' && character != '{' && character != '}' && + character != ':'; +} + std::string replace_address(const PtxMemoryOp& op, const std::string& scratch) { @@ -154,6 +192,7 @@ TransformResult transform_ptx(const TransformRequest& request) bool selected = false; int brace_depth = 0; std::uint64_t scratch_id = 0; + std::string pending_statement; static const std::regex function_expression( R"(\.(?:visible\s+)?(?:entry|func)\s+([A-Za-z0-9_$.]+))"); @@ -241,11 +280,19 @@ TransformResult transform_ptx(const TransformRequest& request) result.modified = true; continue; } + // Accumulate physical lines into one logical statement before + // scanning, so a statement split across lines is seen whole. + pending_statement = joined_statement(pending_statement, line); + if (statement_is_open(pending_statement)) { + output << line << '\n'; + continue; + } std::string opcode; - if (unsupported_memory_instruction(line, opcode)) { + if (unsupported_memory_instruction(pending_statement, opcode)) { ++result.coverage.unsupported_instructions; result.coverage.unsupported_opcodes.push_back(opcode); } + pending_statement.clear(); } output << line << '\n'; diff --git a/tests/cpu/ptx_async_copy_coverage_test.cpp b/tests/cpu/ptx_async_copy_coverage_test.cpp index e6863eb..87287b7 100644 --- a/tests/cpu/ptx_async_copy_coverage_test.cpp +++ b/tests/cpu/ptx_async_copy_coverage_test.cpp @@ -148,6 +148,41 @@ int main() CHECK(!counted_unsupported( "cp.async.bulk.shared::cluster.shared::cta [%r1], [%r2], %r3;")); + // A statement split across physical lines has to be seen whole. Neither + // half matches the pattern on its own, so scanning per line lets the + // access through unreported; and if the same kernel also holds an ordinary + // global load, the module is still marked instrumented and the launch + // proceeds with an access nothing recorded. + { + hbfsim::ptx::TransformRequest request{}; + request.full_ptx = + ".version 8.0\n.target sm_120\n.address_size 64\n" + ".visible .entry probe(.param .u64 probe_param_0)\n" + "{\n" + " .reg .b32 %r<8>;\n" + " .reg .b64 %rd<8>;\n" + " ld.param.u64 %rd1, [probe_param_0];\n" + " cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::" + "complete_tx::bytes\n" + " [%r1], [tmap, {%r2,%r3}], [%r4];\n" + " ret;\n" + "}\n"; + const auto split = hbfsim::ptx::transform_ptx(request); + + hbfsim::ptx::TransformRequest one_line{}; + one_line.full_ptx = request.full_ptx; + const auto position = one_line.full_ptx.find("bytes\n"); + one_line.full_ptx.replace(position + 5, 10, " "); + const auto joined = hbfsim::ptx::transform_ptx(one_line); + + // Written on one line or on two, the same statement must be counted + // the same number of times. + CHECK(split.coverage.unsupported_instructions == + joined.coverage.unsupported_instructions); + CHECK(split.coverage.unsupported_instructions > + baseline_unsupported()); + } + // The recorded opcode has to name the instruction, so the coverage record // says which operation was refused. {