HIP : add shared-memory tiled transpose fast path for F32 concat on RDNA3.5

[Casten-Wang]

Overview

On AMD RDNA3.5 (gfx1151, HIP backend), concat on dim 0 with a transposed F32
src1 view falls back to the generic non-contiguous, per-element copy path. The
concrete shape seen during transformer prefill input processing is:

src0 : [3,    8192]   (contiguous)
src1 : [2048, 8192]   (transposed F32 view)
dst  : [2051, 8192]   = concat(src0, src1) on dim 0

Because src1 is a transposed view, logically adjacent elements are not
physically adjacent, so global-memory access is poorly coalesced — the result is
correct, only slow.

This PR adds a dedicated kernel concat_f32_transpose for exactly this case:

• New kernel. 256 threads per block (32 × 8), one logical 32×32 tile per
  block. src1 is staged through __shared__ float tile[32][33], transposed in
  shared memory, then written out coalesced. The [32][33] padding (one extra
  column) shifts each row's start so transposed reads don't collide on the same
  shared-memory bank — the standard fix for transpose bank conflicts. src0 is
  contiguous and copied into the first ne00 elements of each channel only when
  blockIdx.y == 0, avoiding redundant copies across vertical tiles.
• Gate concat_try_f32_transpose. Triggers only when dim == 0;
  src0/src1/dst are all F32; the device is RDNA3.5; src1->ne[0] is
  512/1024/2048; src0/dst are contiguous; src1 has strict transpose
  strides; and dst->ne[2] * dst->ne[3] <= 65535. Otherwise it returns false and
  the original implementation runs unchanged — no existing path is deleted or
  rewritten.

Only two files change: ggml/src/ggml-cuda/concat.cu and
tests/test-backend-ops.cpp (+132 / -1). Related to #21284.

Additional information

Why only prefill. The matching shape (ne00=3, ne10=2048, ne1=8192, ne2=1, ne3=1) only appears in prefill input processing; decode never matches and
falls back:

workload	fast-path hits
Q4_K_M prefill	60
BF16 prefill	60
Q4_K_M decode	0
BF16 decode	0

Tests (test-backend-ops.cpp). New transposed b layout (1 << 4, i.e.
v & 16) built with a real ggml_transpose, matching the model's layout.
Transposed cases are checked with std::memcmp for bit-identical output (the
kernel only copies/reorders F32 — no arithmetic or conversion). Fast-path shapes:
widths 512/1024/2048 × channel counts 31/32/33/64 (tile boundaries) plus a
ne2 = 2 case. Fallback shapes: non-whitelisted widths 1/31/32/511/513/1537, a
contiguous F32 case, and a transposed F16 case. Result: concat 198/198 pass;
transposed F32 output bit-identical.

Performance (Windows 11, ROCm 7.1, Radeon 8060S / gfx1151). A = upstream
master, B = master + this change; A/B build artifacts differ only in
ggml-hip.dll. Full GPU offload, Flash Attention, HIP Graph, p2048, b4096,
t16. Order A/B/B/A then extended; 10 samples per version per scenario; medians:

model / stage	ubatch	A (tok/s)	B (tok/s)	Δ
Q4_K_M prefill	512	985.611	1106.050	+12.22%
Q4_K_M prefill	1024	1218.855	1287.850	+5.66%
Q4_K_M prefill	2048	1293.875	1426.610	+10.26%
BF16 prefill	512	485.047	526.406	+8.53%
BF16 prefill	1024	440.108	465.546	+5.78%
BF16 prefill	2048	581.229	604.940	+4.08%
Q4_K_M decode	—	56.672	56.649	-0.04%
BF16 decode	—	27.077	27.037	-0.14%

All prefill scenarios improve in the same direction (+4.08% .. +12.22%); decode
changes are < 1% (run-to-run noise, not a regression). No thermal throttling; A/B
GPU clock, memory clock, power and temperature ranges overlap. Fixed-input output
hashes match A/B (Q4_K_M 5A2BC481...A45867F846, BF16
E8524416...EBB75C683).

Known limitation. The gate restricts width but not channel count
(dst->ne[1]), so some small-channel shapes can still enter the fast path and
regress. At width 2048, ne1 = 33 regresses ~ -4.25% and it turns positive from
ne1 = 64 (the real model width ne1 = 8192 improves clearly), because for very
small workloads the kernel's launch/sync/tiling overhead can exceed the generic
kernel. This first PR deliberately keeps the narrow 512/1024/2048 whitelist —
those widths are directly validated by the real model with minimal code surface.
A workload/channel gate to widen the range can be added separately if maintainers
prefer.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES — this change was produced with substantial AI
  assistance (code analysis, implementation drafting, and test/benchmark
  automation). I have reviewed the diff and understand the transpose indexing,
  the [32][33] shared-memory tiling, the blockIdx.y == 0 single copy of
  src0, the tail-block bounds checks, the fallback gating, and the known
  small-channel regression; I take responsibility for maintaining this change.

[ggml-gh-bot]

Hi @Casten-Wang, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[pwilkin]

What is "model" in this case? :)

[Casten-Wang]

Qwen3.6-35B-A3B (MoE, 35B total / 3B activated)