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..3d14a08 100644 --- a/src/ptxpass_hbf/transform.cpp +++ b/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'; 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..87287b7 --- /dev/null +++ b/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 +#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;")); + + // 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; +}