WIP: perf(rocjitsu): streamline bulk vector memory transfers

[newling]

perf(rocjitsu): streamline bulk vector memory transfers

Overview

rocJITsu's functional global-memory path performs substantial host-side work
for wide vector loads and stores. This is especially visible while loading a
large code object: the ROCclr copy kernel repeatedly moves wide lane payloads,
and complete destination cache lines are read before being entirely
overwritten.

This change removes three sources of redundant work:

• generated flat stores acquire one observed VGPR region and marshal its data
  into lane-major storage, instead of repeating ownership lookup and plugin
  dispatch for every register/lane pair;
• ordinary dword load completion obtains one raw destination pointer per VGPR,
  after validating ownership for the complete range, instead of making one
  virtual write call per register/lane pair; and
• aligned complete-line L1/L2 store misses allocate and overwrite the new line
  without first fetching bytes that the store will replace.

These are generic functional-emulation paths. The implementation does not
recognize specific kernels, bypass instruction execution, or suppress plugin
observation. The complete-line path retains dirty-eviction publication and the
existing write-through behavior.

The generated flat-store path uses the wave-owned register-access boundary
introduced by #10030. The VGPR-region snapshot follows the same general pattern
as the scalar MFMA operand snapshot in #10706, while the cache work complements
the page-local MTYPE batching in #10583. PR #10898 optimizes the separate f16
and bf16 MFMA path and was present in both arms of the performance comparison
below.

Performance

The primary workload was one f16 1024x1024x1024 GEMM in functional emulation
with one simulator thread:

hipblaslt-bench --precision f16_r \
  --sizem 1024 --sizen 1024 --sizek 1024 \
  --initialization zero --iters 1 --cold_iters 0

Four balanced parent/candidate pairs measured the complete process without an
execution plugin:

Metric	Parent median	Candidate median	Speedup
Complete process	10.685 s	9.360 s	1.14x
hipBLASLt-reported interval	9.219 s	8.094 s	1.14x

Two balanced pairs with the throughput plugin enabled showed where the saving
comes from:

Metric	Parent	Candidate	Speedup
Complete process	11.505 s	9.740 s	1.18x
221 MiB module-copy dispatch	3.456 s	2.408 s	1.44x
GEMM dispatch	5.286 s	4.739 s	1.12x

The module-copy dispatch executed exactly 2,287,712 wave instructions in both
revisions. Its global-instruction handler interval fell from approximately
0.736 seconds to 0.250 seconds, a 2.95x improvement. A separate 25 MiB
hipMemset probe improved from approximately 0.416 seconds to 0.331 seconds,
or 1.26x.

The comparison used exact revision af74731cc as its parent so that #10898's
MFMA optimization was identical in both arms. The implementation was then
ported to current develop at 2d311cd6f, including adaptation to #10030's
wave-owned register-access API.

Validation

• Release build of rocjitsu_tests, rocjitsu_shared, and the throughput
  plugin
• 3,566 C++ tests: 3,564 passed, two skipped, one disabled
• 518 generator tests passed, one skipped
• 72 focused register/L1/L2 tests passed
• 14 global-store and race-observation tests passed
• full ISA regeneration for all supported AMDGPU families
• changed-file pre-commit and git diff --check
• exact ROCclr copy-kernel execution copied 221.3 MiB correctly
• nonzero f16 256x256x256 verification completed with norm_error=0

Issue Tracking

Related: #9577

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.