This is substantial, careful work and I went through the kernel line by line rather than trusting the comments. The correctness holds up. I'm holding approval on one thing only — the copyright header — plus a question about the tuning data. Everything else below is a suggestion.

What I verified

The transpose is correct. Composing the two phases: the store reads tile[lx * pitch + rl], which the load wrote as a[(tile_r + rl) * stridea + (tile_c + lx) * lda], and writes it to b[(tile_c + lx) * strideb + (tile_r + rl) * ldb]. With R = tile_r + rl and C = tile_c + lx that is exactly b[C * strideb + R * ldb] = alpha * a[R * stridea + C * lda], matching the contract in your comment.

The tile bounds are in range. Max local index is (cols - 1) * pitch + (rows - 1) = cols * rows + cols - 2, against an allocation of cols * pitch = cols * rows + cols. Fits, with the odd pitch doing its bank-conflict job.

No uninitialized local reads, which is where these kernels usually break. The load guards on load_r < logical_m with load_r = tile_r + lx, and the store guards on r = tile_r + rl < logical_m — and the entry it reads was written by the load iteration where lx_load = rl, so the two conditions are the same predicate. Same story for the column guards against logical_n. Partial tiles at both edges are handled correctly.

The barrier is outside the divergent load branch, so every item in the group reaches it. Easy thing to get wrong; you didn't.

Conjugation is applied only on the transposing path, which is right — oneMath's transpose has no plain-conjugate value, so launch_nontrans correctly has nothing to do.

Every tile-size claim in your comments is arithmetically correct. I checked all of them: gfx90a complex<double> strided is 32 × 65 × 16 = 33,280 B, i.e. the 32.5 KB you cite; nvidia strided 16-byte is 16 × 65 × 16 = 16,640 B, inside your 17 KB bound; and the work-group sizes match your prose too — gfx950 8-byte is 16 × 64 = 1024 items, nvidia strided is 4 × 64 = 256, gfx942 8-byte is a 32×8 tile at 256 items. The static_asserts on cols <= rows and rows % block == 0 && cols % block == 0 hold for every geometry in the file. That level of internal consistency is a good sign, and it made the review much faster.

Please fix: the copyright header

src/blas/backends/omatcopy2_kernels.hpp is a new file, and its header reads:

*  Copyright (C) Codeplay Software Limited

I assume that came along with the boilerplate from a neighbouring file. Since this is new code you wrote, the attribution should reflect that. This is the one thing I'd like corrected before merge — it's a two-word change, but it's a licensing/provenance detail in a Linux Foundation project rather than a style nit, so I'd rather not wave it through.

Question: can the tuning data be reproduced?

The geometry tables carry a lot of specific empirical claims — 8.5% on MI210 for 16-byte strided, 4% on double and 2% on complex<float> for the NVIDIA work-group choice, 2.4% on gfx942, 0.4% agreement between gfx950 and A100, 2–231% versus geam when pipelined. That's measurements across gfx950, two MI210s, gfx942 and an A100.

Two things follow from that:

1. No reviewer can check any of it, and there's no benchmark in the PR to re-run.
2. The tables can't be regenerated. When gfx951 or Blackwell shows up, whoever maintains this will have no way to re-derive the numbers, and given how thin maintainer capacity on this project is, that matters more than usual.

Could you include the benchmark you used — even as a rough standalone file under tests/ or referenced from the header comment? That turns the tables from constants-of-unknown-provenance into something maintainable.

And to ask plainly, since you've mentioned leaning on AI tooling: were these figures produced by running on those four devices? I'm not doubting the code — the arithmetic all checks out — but percentage claims of that precision are only worth having if they came off real hardware, and if some are estimates I'd rather the comments said so.

Suggestions

Guard against the hardware limits you're brushing up against. Two geometries sit close to the edge: the gfx90a complex<double> strided tile asks for 32.5 KB of local memory, and the gfx950 8-byte config requests a 1024-item work-group — the maximum on both vendors. Neither is checked at runtime. On a device reporting a smaller local_mem_size or max_work_group_size, this fails at launch with an opaque backend error. A query against device.get_info<sycl::info::device::local_mem_size>() and max_work_group_size, falling back to the conservative geometry, would turn that into a graceful degradation.

Say out loud that misclassification is harmless. The MI210/MI250/MI300/MI308/MI325/gfx90a/gfx942 substring matching will inevitably rot as devices are added and renamed. The saving grace is that guessing wrong only picks a different tile shape — it cannot produce a wrong result. That's worth one explicit sentence in the comment, because the next maintainer looking at a stale device list needs to know whether they're staring at a performance issue or a correctness landmine. Your explanation of why the fallback exists (HIP adapter reports unknown architecture, AdaptiveCpp doesn't expose it) is already good.

The store phase idles lanes. You note it, and when cols < rows it's substantial — the NVIDIA 16-byte strided geometry leaves three quarters of each group idle during the store. Presumably you measured that it still wins; did you try a separate index mapping for the store phase so all lanes participate, and it lost? Worth a sentence either way.

Tests: good news here — tests/unit_tests/blas/extensions/omatcopy2.cpp and omatcopy2_usm.cpp already exist, so this stops being a path that throws unimplemented and starts actually being exercised. Worth confirming those cases cover conjtrans, non-unit element strides, and shapes that produce partial tiles in both dimensions, since the edge guards are the part most likely to regress under future tuning changes.

Fix the header, answer on the benchmarks, and I'll approve.