Count asynchronous global copies as unsupported instead of dropping them

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(

src/ptxpass_hbf/transform.cpp

@@ -51,7 +51,24 @@ 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.
+        //
+        // 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;
@@ -60,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)
 {
@@ -137,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_$.]+))");
 
@@ -224,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';

tests/cpu/ptx_async_copy_coverage_test.cpp

@@ -0,0 +1,200 @@
+// `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 <algorithm>
+#include <cstdio>
+#include <string>
+
+#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;"));
+
+    // 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 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 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(
+        "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;"));
+    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;"));
+
+    // 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.
+    {
+        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;
+}