Row: `GFX1100-TG200`
T24 adds an LDS-buffered quant epilogue to the cooperative RmsNormRowKernel. When the norm-quant fused path is active, the cooperative kernel writes Q8_K blocks to LDS before the global store, reducing register pressure and improving throughput.
Also restores the `__half` Load helper and `hip_fp16.h` include that were lost when the file was restructured.
Depends on #2782 (F1), #2790 (T4a), #2792 (lever-C), and #2796 (T27).
Row: `GFX1100-TG200`
T24 adds an LDS-buffered quant epilogue to the cooperative RmsNormRowKernel. When the norm-quant fused path is active, the cooperative kernel writes Q8_K blocks to LDS before the global store, reducing register pressure and improving throughput.
Also restores the `__half` Load helper and `hip_fp16.h` include that were lost when the file was restructured.
Depends on #2782 (F1), #2790 (T4a), #2792 (lever-C), and #2796 (T27).