gfx942-only build bugfix and kittens refactor

[alextmagro]

Fixes a bug where gfx942 only build runs into compilation failure due to kittens symbol issues with grouped gemm

Refactors cmake file to modularize architecture macros passed to TE
Removes arch checks within TE, deferring to within the kittens library
Modularizes mxfp8 to match class structure introduced for blockwise fp8 PR -- 1 folder per arch with pure virtual class

[github-actions]

Moving these TUs from anonymous namespaces into a shared named namespace gives their contents external linkage and puts mxfp8_gemm.cpp and cdna4/blockwise_fp8_gemm.cpp into the same scope. They now both declare namespace-scope names that mean different things:

• NUM_WARPS, WARPS_ROW, WARPS_COL, HALF_ROW, HALF_COL, REG_M, REG_N, NUM_THREADS, G
• GemmEpilogue — this file defines te_kittens::cdna4::GemmEpilogue {DEFAULT, BIAS, GELU_AUX, GELU_AUX_BIAS} (line 56), while blockwise_fp8_gemm_helper.cuh pulls te_kittens::blockwise::GemmEpilogue {DEFAULT, BIAS, GELU_AUX, BETA, BIAS_BETA, GELU_AUX_BETA} into te_kittens::cdna4 via a using namespace directive.

This links today only by luck: the constants are constexpr (internal linkage), the aliases have no linkage, and none of the __global__/__device__ inline names happen to collide. But __global__ templates and inline/__device__ functions emit weak COMDAT symbols — the day two same-named-but-different definitions appear here, the linker silently keeps one and the miscompile is invisible. The anonymous namespaces made that impossible. Note this file already #includes ../kittens_kernel_common.cuh, so te_kittens::blockwise::GemmEpilogue is already declared in this TU; only the absence of a using namespace directive keeps the two apart.

Cheap fix that keeps the refactor's benefit: wrap each file's internals in a nested anonymous namespace inside te_kittens::cdna4, or give each file its own sub-namespace (te_kittens::cdna4::mxfp8 / ::blockwise), leaving only the *Cdna4 backend class visible to the factory at the bottom.

[github-actions]

This is the one thing the new kittens_add_arch abstraction doesn't cover, and it sits outside the arch guard: -ffast-math gets applied unconditionally, including on gfx942-only builds where cdna4/blockwise_fp8_gemm.cpp is never compiled. Harmless today (a source property on an unused file is a no-op), but it means the source list now lives in two places and has to be kept in sync by hand — exactly the coupling the refactor is trying to remove.

Verified the flags themselves are equivalent to before: -gline-tables-only moved from a per-source COMPILE_FLAGS to FLAGS (target-wide) and now covers both cdna4 sources, which matches the old behavior since both had it. Only -ffast-math is genuinely per-source.

Consider a SOURCE_FLAGS pair argument on kittens_add_arch so the per-source case is handled inside the guarded region too, or at minimum move this call inside the if(gfx950 IN_LIST ...) scope.

[github-actions]

Claude review

Reviewed the full diff (11 files, ~700 lines) against dev — first review on this PR, no prior Claude or human review comments to dedupe against.

Verdict: approach is sound, no blocking correctness issues found. Two non-blocking findings posted inline.

The build fix is correct and minimal. try_kittens_grouped_mxfp8_gemm/_wgrad are called from cublaslt_gemm.cu:1220 under only USE_HIPKITTENS_GEMM, but previously depended on declarations that lived behind #ifdef KITTENS_HAVE_CDNA4 — hence the gfx942-only compile failure. Routing them through the unconditional inline wrappers in kittens_common.h, with MXFP8GemmBackend::get() returning nullptr when KITTENS_HAVE_CDNA4 is undefined, fixes it without leaving an undefined get_cdna4() reference at link time.

Things I checked that are fine, in case they look alarming in the diff:

• Dropping the host-side m%256==0 && n%256==0 && k%128==0 && k>=256 guard from cublas_gemm is not a behavior change — mxfp8_gemm_impl applies exactly the same predicate via check_tn_constraints (mxfp8_gemm.cpp:1252) and returns false, and use_hipkittens is reassigned from that return value, so unsupported shapes still fall through to hipBLASLt.
• CanonicalizeGemmInput now runs before the shape check, but it is pure pointer/dtype selection with no kernel launches or workspace writes, and the fallback (hipblaslt_gemm) canonicalizes identically. try_mxfp8_non_tn_transpose_to_tn is gated on sm_arch()==125, disjoint from the gfx950 kittens path.
• kittens_align_up(x, 256) is bit-identical to the three removed align_up_pow2ws/hk_align_up/hk_align256 helpers.
• CMake list accumulation via set(... ${...} ... PARENT_SCOPE) inside kittens_add_arch accumulates correctly across both calls, and KITTENS_HAVE_DEFS cannot be empty when USE_HIPKITTENS_GEMM is ON.

Upstream compatibility: clean. Every change is in ROCm-only files (rocm_gemm.cu, gemm/kittens/**) or inside the else() # USE_ROCM branch of common/CMakeLists.txt. No CUDA-reachable code path is touched.

Test coverage: none added, which is reasonable for a build-configuration fix — but the failure mode this PR fixes is only caught by actually configuring a single-arch (CMAKE_HIP_ARCHITECTURES=gfx942) build. Worth confirming CI exercises that, otherwise the same class of regression can reappear silently.

Copyright headers: OK — all 10 in-scope files carry AMD lines ending in 2026, the new kittens_kernel_common.cuh is correctly AMD-only, and no NVIDIA year range was altered.