Skip to content

Conversation

@arai713
Copy link
Contributor

@arai713 arai713 commented Oct 17, 2025

Proposed changes

This PR is expanding the supported datatypes for Stream-K Gemm by adding examples for fp8 and bf8. Previously examples for fp16 and bf 16 were added through this PR. Currently these examples only support atomic reduction. Unit tests for fp8 and bf8 have been added through a test suite.

Checklist

Please put an x into the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask.

  • I have added tests relevant to the introduced functionality, and the unit tests are passing locally
  • I have added the test to REGRESSION_TESTS list defined at the top of CMakeLists.txt in tests/CMakeLists.txt, IF the test takes more than 30 seconds to run.
  • I have added inline documentation which enables the maintainers with understanding the motivation
  • I have removed the stale documentation which is no longer relevant after this pull request
  • (If this change is user-facing) I have added release notes which provide the end users with a brief summary of the improvement from this pull request
  • I have run clang-format on all changed files
  • Any dependent changes have been merged

Copy link
Collaborator

@cgmillette cgmillette left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Need to consolidate new fp8 tests into smoke tests with existing test harness.

@arai713 arai713 force-pushed the arai/ck_tile/sk_ex_fp8_bf8 branch from 3114a67 to 92f2fcc Compare October 24, 2025 21:46
@arai713 arai713 changed the title Stream-K Gemm Example for fp8 and bf8 [CK_TILE] Stream-K Gemm Example for fp8 and bf8 Oct 24, 2025
@arai713 arai713 force-pushed the arai/ck_tile/sk_ex_fp8_bf8 branch 2 times, most recently from b55587b to ffe459a Compare October 24, 2025 22:31
@arai713 arai713 requested review from a team and ddembeckAMD as code owners October 24, 2025 22:31
ecamartins
ecamartins previously approved these changes Oct 24, 2025
Copy link
Collaborator

@ecamartins ecamartins left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM! Great job on refactoring the example and tests :)

cgmillette
cgmillette previously approved these changes Oct 24, 2025
Copy link
Collaborator

@cgmillette cgmillette left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This looks great! Approved pending CI!

@cgmillette cgmillette requested a review from Copilot October 24, 2025 23:36
Copy link
Contributor

Copilot AI left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull Request Overview

This PR extends Stream-K GEMM support to fp8 and bf8 data types, building on the existing fp16 and bf16 implementation. The changes introduce type configuration headers, test suite infrastructure, and example code updates to handle the new 8-bit floating-point formats with atomic reduction support.

  • Adds fp8 and bf8 type definitions and layout configurations
  • Implements smoke tests covering multiple layout combinations (RRR, RCR, CRR, CCR) for both data types
  • Updates the basic example to support fp8 and bf8 with appropriate command-line interface changes

Reviewed Changes

Copilot reviewed 17 out of 17 changed files in this pull request and generated 4 comments.

Show a summary per file
File Description
test/ck_tile/gemm_streamk/test_gemm_streamk_types_fp8.hpp Defines F8Layouts struct and macros for fp8 test configurations
test/ck_tile/gemm_streamk/test_gemm_streamk_types_bf8.hpp Defines BF8Layouts struct and macros for bf8 test configurations
test/ck_tile/gemm_streamk/test_gemm_streamk_types.hpp Adds F8 and BF8 type aliases and includes new type headers
test/ck_tile/gemm_streamk/smoke_tests/*.cpp Eight new smoke test files covering fp8/bf8 with different layouts
test/ck_tile/gemm_streamk/CMakeLists.txt Adds compile options for OCP FP8 support and registers new test files
example/ck_tile/40_streamk_gemm/streamk_gemm_basic.cpp Adds fp8 and bf8 case handling in the example runner
example/ck_tile/40_streamk_gemm/gemm_utils.hpp Adds DataTypeTraits specializations and updates documentation
example/ck_tile/40_streamk_gemm/README.md Updates documentation to reflect fp8/bf8 support
example/ck_tile/40_streamk_gemm/CMakeLists.txt Adds OCP FP8 compile options for the example
example/ck_tile/CMakeLists.txt Adds trailing newline

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

@arai713 arai713 dismissed stale reviews from cgmillette and ecamartins via cd8d967 October 27, 2025 04:49
@arai713 arai713 force-pushed the arai/ck_tile/sk_ex_fp8_bf8 branch from ffe459a to cd8d967 Compare October 27, 2025 04:49
Refactored the unit tests for fp8/bf8 to utilize the test harness.
Implemented smoke tests with layouts: CCR, CRR, RCR, RRR for fp8/bf8.
The tests are using 128x128x32 for the tile configuration, as other
configurations revealed implementation gaps that are currently being
documented.
@arai713 arai713 force-pushed the arai/ck_tile/sk_ex_fp8_bf8 branch from cd8d967 to b0f39cb Compare October 27, 2025 16:26
Copy link
Collaborator

@cgmillette cgmillette left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

lgtm!

@arai713 arai713 merged commit 715395b into develop Oct 28, 2025
48 checks passed
@arai713 arai713 deleted the arai/ck_tile/sk_ex_fp8_bf8 branch October 28, 2025 02:29
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants