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,22 @@ 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 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;

tests/cpu/ptx_async_copy_coverage_test.cpp

@@ -0,0 +1,165 @@
+// `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 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;"));
+    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;
+}