ViperEkura
7dd184a4e5
- split gemm.cuh into gemm/{policy,load,scheduler,mainloop,epilogue}.cuh (humming/CUTLASS-style layering, files 28-336 lines); the umbrella keeps the kernel orchestrator, host planning and the gemm<> entry so ops.cu and the C tests build unchanged
- move the measured design essays (swizzle derivation, ring-depth barrier invariant, launch crossovers, NN swap) into an FP8 design-notes section in docs/developer/cuda_kernels.md, leaving one-line constraints at each symbol
- refresh the doc's FP8 file table and layout tree (fix stale mm.cu / fp8_mma_test.cu names)
- structure-only change: extension rebuilds identical, C tests all pass, tests/extension 65 passed, quantize layouts byte-exact, NT routing torch.equal, e2e M=8192 530.6ms / 1.26x unchanged
29 lines
1.1 KiB
Plaintext
29 lines
1.1 KiB
Plaintext
#pragma once
|
|
// Tile scheduler: the linear CTA id maps to (block_m, block_n) in grouped
|
|
// (L2-friendly) raster — consecutive CTAs share one B column stripe — or
|
|
// plain N-fastest raster (kRasterGroup=0, the measured best for dX's
|
|
// crosswise-B layouts where grouping was neutral).
|
|
|
|
namespace astrai {
|
|
namespace fp8 {
|
|
|
|
template <int kRasterGroup>
|
|
struct Fp8GemmTileScheduler {
|
|
static __device__ int2 tile(const uint3& block, const dim3& blocks) {
|
|
if constexpr (kRasterGroup > 0) {
|
|
constexpr int kGroupM = kRasterGroup;
|
|
const int bid = int(block.y) * int(blocks.x) + int(block.x);
|
|
const int group_first_m = (bid / (kGroupM * int(blocks.x))) * kGroupM;
|
|
const int group_rows =
|
|
min(int(blocks.y) - group_first_m, kGroupM); // M-tail group is short
|
|
return int2{group_first_m + bid % group_rows,
|
|
(bid % (kGroupM * int(blocks.x))) / group_rows};
|
|
} else {
|
|
return int2{int(block.y), int(block.x)};
|
|
}
|
|
}
|
|
};
|
|
|
|
} // namespace fp8
|
|
} // namespace astrai
|