Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
91 commits
Select commit Hold shift + click to select a range
ad748da
GEMM reference HIP implementation
matthiasdiener Dec 9, 2025
11e090b
blockwise amax
matthiasdiener Dec 11, 2025
9006224
Merge branch 'dev' into compute-ref-offload
matthiasdiener Dec 18, 2025
3ecea7f
Change to use Tensor arguments, combine mxfp8/non-mxfp8 paths
matthiasdiener Jan 13, 2026
cafee59
Merge remote-tracking branch 'origin/dev' into compute-ref-offload
matthiasdiener Jan 14, 2026
86fbbac
skip on SwizzleScale limitation on gfx950
matthiasdiener Jan 14, 2026
54de3db
Revert "skip on SwizzleScale limitation on gfx950"
matthiasdiener Jan 14, 2026
311ddfe
MXFP8 fix
matthiasdiener Jan 14, 2026
306e432
Merge remote-tracking branch 'origin/dev' into compute-ref-offload
matthiasdiener Jan 15, 2026
445e64f
correct scale_inv packing and exp2(biased−127) conversion
matthiasdiener Jan 15, 2026
462945f
cleanups
matthiasdiener Jan 15, 2026
e32fb3d
Merge branch 'dev' into compute-ref-offload
matthiasdiener Jan 19, 2026
7bf8adb
Merge remote-tracking branch 'origin/dev' into compute-ref-offload
matthiasdiener Jan 22, 2026
e11e400
use Tensor class for more device objects
matthiasdiener Jan 22, 2026
325ece6
Pass D Tensor into run_reference and move RefD allocation into Perfor…
matthiasdiener Jan 23, 2026
fc64b8c
[WIP] proof-of-concept: grouped GEMM with ck_tile
matthiasdiener Jan 26, 2026
134b350
Merge branch 'dev' into ck-grouped-gemm
matthiasdiener Jan 28, 2026
9091e6c
restructure and enable tests
matthiasdiener Jan 29, 2026
7435062
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Jan 29, 2026
a00a1c8
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Jan 30, 2026
4e9ead9
grid improvements
matthiasdiener Jan 30, 2026
259645c
restructure
matthiasdiener Feb 3, 2026
9986bd4
reduce code duplication & simplify
matthiasdiener Feb 4, 2026
355ec2f
make the code more similar to nv, check emopty gelu/bias
matthiasdiener Feb 4, 2026
df5e3ea
Merge branch 'dev' into ck-grouped-gemm
matthiasdiener Feb 4, 2026
a42f7ca
further simplify & make closer to nv
matthiasdiener Feb 4, 2026
fac7c11
add ck_tile reference
matthiasdiener Feb 4, 2026
71b97e0
rename in error messages
matthiasdiener Feb 4, 2026
dd3ed2f
allow flattened higher-D tensors
matthiasdiener Feb 4, 2026
7b0413e
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Feb 5, 2026
ebc005f
relax tolerance on gfx942
matthiasdiener Feb 5, 2026
c0bf502
enable more tests
matthiasdiener Feb 5, 2026
0b16287
return early when num_gemms<=0
matthiasdiener Feb 5, 2026
58b34e7
simplify normalization
matthiasdiener Feb 5, 2026
74f229a
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Feb 10, 2026
e28c801
run hipblaslt for num_gemms==1
matthiasdiener Feb 11, 2026
6151b96
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Feb 12, 2026
5c57d47
disable ck_tile when accumulate=true
matthiasdiener Feb 17, 2026
29d6ab7
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Feb 17, 2026
6e9aae4
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Feb 17, 2026
2e844d9
remove test file
matthiasdiener Feb 17, 2026
4aa8229
Merge branch 'dev' into ck-grouped-gemm
matthiasdiener Feb 23, 2026
f680d6a
fix copyright header
matthiasdiener Feb 23, 2026
6d85088
simplify calls in dispatch_grouped
matthiasdiener Feb 23, 2026
7910038
remove is_mi3*0_class
matthiasdiener Feb 23, 2026
e8ebb0e
disable unused constants
matthiasdiener Feb 23, 2026
deb7474
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Feb 24, 2026
e866bc6
add another fallback
matthiasdiener Feb 24, 2026
ee438fb
implement Primus-Turbo selection logic, persistent descs
matthiasdiener Feb 25, 2026
b65dbfa
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Feb 25, 2026
0cbf1cd
tighten tolerances
matthiasdiener Feb 25, 2026
98e0c66
use namespace, various cleanups
matthiasdiener Feb 25, 2026
36bd68e
avoid creating vector with Tensors
matthiasdiener Feb 26, 2026
070c58d
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Feb 26, 2026
c5d83a4
merge dispatch_grouped into ck_tile_grouped_gemm
matthiasdiener Feb 26, 2026
56afb04
Merge remote-tracking branch 'origin/dev' into ck-grouped-gemm
matthiasdiener Feb 27, 2026
26dfbb6
same tolerances for gfx950
matthiasdiener Feb 27, 2026
5a7eb69
feat(gemm): enable TensorQuant pipeline for FP8 on GFX942
matthiasdiener Feb 27, 2026
54da682
Include Float8 E4M3/E5M2 in is_supported_dtype and remove float8 from…
aris134 Mar 4, 2026
78a702f
forward pass ck_tile with matching FP8 data type inputs passing accur…
aris134 Mar 4, 2026
f198341
Support mixed FP8/BF8 grouped GEMM in CK backward path
aris134 Mar 6, 2026
6b24be2
include more descriptive comment regarding tensor normalization in ck…
aris134 Mar 6, 2026
a161c20
Refactor CK grouped GEMM: split FP8/FP16 implementations and introduc…
aris134 Mar 9, 2026
af95382
Add explicit template instantiations for CK grouped GEMM runners
aris134 Mar 9, 2026
9990db3
Split CK grouped GEMM implementation to reduce compile-time coupling
aris134 Mar 10, 2026
bff80fe
Split CK grouped GEMM explicit instantiations by operand dtype
aris134 Mar 10, 2026
e9cd6b8
Add runtime architecture dispatch for CK FP8 grouped GEMM (gfx942/gfx…
aris134 Mar 11, 2026
32f2ac3
Merge dev and fix FP8 logic/CMake
aris134 Mar 11, 2026
5275ac6
Fix dev merge conflicts in CMakeLists.txt
aris134 Mar 11, 2026
3940748
add copyright headers and add blank line at bottom of each new file. …
aris134 Mar 12, 2026
76f207a
Merge branch 'dev' into amartin/ck-grouped-gemm-fp8
aris134 Mar 19, 2026
670d2c4
Fix cudnn-frontend submodule pointer after dev merge
aris134 Mar 19, 2026
a1f4e37
Improve TE FP8 CK backward dX path by reformulating NN as NT on repac…
aris134 Mar 21, 2026
5b67d3d
refactor fp8 gfx942 tile configs
aris134 Mar 23, 2026
3194b17
refactor gfx950 tile configs
aris134 Mar 23, 2026
565ee11
reduce number of kernel instantiation files
aris134 Mar 23, 2026
b1721d6
Merge branch 'dev' into amartin/ck-grouped-gemm-fp8
aris134 Mar 24, 2026
619a741
remove explicit memory-operation template argument from ck epilogue f…
aris134 Mar 24, 2026
d70252c
update TileGemmQuantTraits signature to match change in updated CK
aris134 Mar 25, 2026
e130a8c
Refactor code to reduce file count.
aris134 Mar 26, 2026
a2b971b
remove no longer needed files
aris134 Mar 26, 2026
2e89fa4
re-enable accumulate support for FP16
aris134 Mar 26, 2026
899b102
update test_numerics
aris134 Mar 26, 2026
fe17d45
turn hard-fails into warnings to support fallbacks like is done in co…
aris134 Mar 27, 2026
0fb97f8
Merge branch 'dev' into amartin/ck-grouped-gemm-fp8
aris134 Mar 27, 2026
c8d0e02
Remove unnecessary template instantiations
aris134 Mar 27, 2026
d1caa82
address reviewer PR comment in detect_gpu_arch
aris134 Mar 27, 2026
25b6f2f
address additional reviewer comments
aris134 Mar 27, 2026
8195d63
add EOL to test_numerics.py
aris134 Mar 27, 2026
7739254
remove extaneous build_log.txt
aris134 Mar 27, 2026
1591653
remap gfx942-compatible tile config to gfx950-friendly config to avoi…
aris134 Mar 27, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 6 additions & 2 deletions tests/pytorch/test_numerics.py
Original file line number Diff line number Diff line change
Expand Up @@ -2171,13 +2171,17 @@ def test_grouped_linear_accuracy(
@pytest.mark.parametrize("num_gemms", [3, 6])
@pytest.mark.parametrize("bs", batch_sizes)
@pytest.mark.parametrize("model", ["126m"])
@pytest.mark.parametrize("fp8_model_params", all_boolean)
@pytest.mark.parametrize("recipe", fp8_recipes + [None])
@pytest.mark.parametrize("fuse_wgrad_accumulation", all_boolean)
@pytest.mark.parametrize("delay_wgrad_compute", all_boolean)
def test_grouped_linear_accuracy_cutlass(
dtype,
num_gemms,
bs,
model,
recipe,
fp8_model_params,
fuse_wgrad_accumulation,
delay_wgrad_compute,
):
Expand All @@ -2187,8 +2191,8 @@ def test_grouped_linear_accuracy_cutlass(
num_gemms,
bs,
model,
None,
False,
recipe,
fp8_model_params,
fuse_wgrad_accumulation,
False,
delay_wgrad_compute,
Expand Down
4 changes: 3 additions & 1 deletion transformer_engine/common/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -254,7 +254,9 @@ else()
list(APPEND transformer_engine_cpp_sources
fused_attn_rocm/fused_attn.cpp
gemm/rocm_gemm.cu
gemm/ck_grouped_gemm.cpp
gemm/ck_grouped_gemm/ck_grouped_gemm.cpp
gemm/ck_grouped_gemm/ck_grouped_gemm_fp8.cpp
gemm/ck_grouped_gemm/ck_grouped_gemm_fp16.cpp
amd_detail/system.cpp)
list(APPEND transformer_engine_cuda_sources
fused_attn_rocm/fused_attn_aotriton.cpp
Expand Down
338 changes: 0 additions & 338 deletions transformer_engine/common/gemm/ck_grouped_gemm.cpp

This file was deleted.

Loading
Loading