Skip to content

CK MXFP8 Group Gemm gfx1250 Enablement #613

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
aris134 wants to merge 50 commits into
base: dev
Choose a base branch
from
Open

CK MXFP8 Group Gemm gfx1250 Enablement #613

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
50 commits
Select commit Hold shift + click to select a range
1f707d7
initial commit for CK Tile MXFP8 integration for gfx1250
aris134 Apr 25, 2026
e102f00
ck mxfp8 gfx1250 integration builds successfully
aris134 Apr 25, 2026
52a2887
add entrypoint to ck mx group gemm in caller
aris134 Apr 25, 2026
8022777
temporary hacky change to test_numerics for bringup testing
aris134 Apr 25, 2026
bc6253d
add warning print to confirm we are in fallback
aris134 Apr 25, 2026
d26f52e
MXFP8 grouped fwd/bwd now reaches CK path and runs without fallback/c…
aris134 May 2, 2026
e295e74
add cpp test for ck tile group mxfp8 gemm forward
aris134 May 2, 2026
1784045
Fix MXFP8 grouped GEMM scale handling for NN/TN/NT
aris134 May 3, 2026
fe99bf3
update ck mxfp8 group gemm gtest to exercise mixed dtypes
aris134 May 3, 2026
e7159c4
include renamed test file
aris134 May 3, 2026
972cea3
clean up code
aris134 May 3, 2026
c0fabff
Update cublaslt_gemm.cu
aris134 May 6, 2026
3db2e5a
address pr comments
aris134 May 11, 2026
910d30f
fix ck group mxfp8 dispatch
aris134 May 17, 2026
1b66d29
update CMakeLists.txt
aris134 May 17, 2026
23b505f
Add direct ROCm libraries dependency for CK grouped GEMM
aris134 May 19, 2026
746afea
Remove redundant MXFP8 env override from grouped linear test
aris134 May 19, 2026
175855d
factor out common definitions from mxfp8 ck ggemm
aris134 May 20, 2026
f00fb7f
add pr comments
aris134 May 21, 2026
45343f1
add MXFP8 pre-swizzling for gfx1250 GEMM (#568)
matthiasdiener May 21, 2026
a67bbe9
CK Tile Group GEMM gfx1250 (#576)
aris134 May 21, 2026
74744db
Merge remote-tracking branch 'origin/gfx1250' into amartin/ck-mxfp8-g…
aris134 May 21, 2026
7c3f499
Add sparse rocm-libraries submodule for Composable Kernel
aris134 Jun 8, 2026
7b5ba68
update submodule name
aris134 Jun 8, 2026
926701a
Merge branch 'dev' into amartin/ck-mxfp8-group-gemm-gfx1250-clean
aris134 Jun 8, 2026
508613c
override CK_ROOT
aris134 Jun 8, 2026
12461ee
fix util
aris134 Jun 8, 2026
e656341
add runtime guard for arch
aris134 Jun 8, 2026
cb1614a
Restore unrelated files from dev
aris134 Jun 8, 2026
a8bb950
Restore FP8 grouped GEMM source from dev
aris134 Jun 8, 2026
baeba44
Restore unrelated CK grouped GEMM files from dev
aris134 Jun 8, 2026
b18099f
update dispatch
aris134 Jun 8, 2026
d69f40c
Remove rocm_libraries submodule
aris134 Jun 8, 2026
88bb3dd
Add standalone Composable Kernel submodule
aris134 Jun 8, 2026
e077670
update gitmodules
aris134 Jun 8, 2026
1f764d2
minor fixes
aris134 Jun 8, 2026
4a262d5
address PR comments
aris134 Jun 8, 2026
669c4cc
address PR comments
aris134 Jun 9, 2026
f3ecda3
address pr comments: fix gfx1250 arch name and convert if-else to swi…
aris134 Jun 9, 2026
94b0126
use workspace for ck group gemm mxfp8 scales
aris134 Jun 10, 2026
479c509
add comment to ck gfx1250 mxfp8 scale swizzle
aris134 Jun 10, 2026
5b4b7fe
change random generation in test_ck_grouped_mxfp8.cu to use pre-exist…
aris134 Jun 10, 2026
68ed32a
address nits in ck_grouped_gemm_common.h
aris134 Jun 10, 2026
2e74a63
stylistic changes
aris134 Jun 10, 2026
f4c97ca
address pr comment: explicitly mention purpose of ck gfx1250 swizzle …
aris134 Jun 10, 2026
bdc6b4e
address PR comments
aris134 Jun 10, 2026
457bbc1
inline mxfp8_gemm bool into if statement
aris134 Jun 10, 2026
19151f4
address nit
aris134 Jun 10, 2026
b4b36c7
Merge remote-tracking branch 'origin/dev' into amartin/ck-mxfp8-group…
aris134 Jun 10, 2026
231c916
add warn fallback for mxfp8 ck
aris134 Jun 10, 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
5 changes: 5 additions & 0 deletions .gitmodules
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,3 +26,8 @@
[submodule "3rdparty/QoLA"]
path = 3rdparty/QoLA
url = https://github.com/ROCm/QoLA.git
[submodule "3rdparty/composable_kernel"]
path = 3rdparty/composable_kernel
url = https://github.com/ROCm/composable_kernel.git
branch = develop

1 change: 1 addition & 0 deletions 3rdparty/composable_kernel
Open in desktop
Submodule composable_kernel added at 0d18f4
3 changes: 2 additions & 1 deletion tests/cpp/operator/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,7 +47,8 @@ if(USE_ROCM)
list(APPEND test_cuda_sources
test_dequantize_nvfp4.cu
test_cublaslt_gemm.cu
test_cast_mxfp4_transpose.cu)
test_cast_mxfp4_transpose.cu
test_ck_grouped_mxfp8.cu)
TE_GetHipifiedSources("${test_cuda_sources}" ${CMAKE_CURRENT_SOURCE_DIR} test_hip_sources)
TE_AddHipifyDeps("${test_cuda_sources}" ${CMAKE_CURRENT_SOURCE_DIR})
message("${message_line}")
Expand Down
Loading
Loading