Skip to content

Conversation

@AviralGoelAMD
Copy link
Collaborator

Proposed changes

In example/ck_tile/38_block_scale_gemm for quant_mode=aquant the current memory pipeline only supports one scheduling mechanism (intrawave) where each wave works independently of other waves to compute resultant wave tile for matrix C.

In this PR, we add interwave scheduling mechanism that orchestrates all the waves to work in lockstep and over finer grained chunks to compute resultant wave tile.

INTRAWAVE SCHEDULING (Independent Waves)
=========================================
Each wave loads full K dimension independently and computes its warp tile

Wave 0: [════════ Load K=0→32 ════════] → [══════ Compute ══════]
Wave 1: [════════ Load K=0→32 ════════] → [══════ Compute ══════]
Wave 2: [════════ Load K=0→32 ════════] → [══════ Compute ══════]
Wave 3: [════════ Load K=0→32 ════════] → [══════ Compute ══════]
         
        Time →→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→


INTERWAVE SCHEDULING (Lockstep Waves)
======================================
All waves synchronized, process K in chunks (MAC clusters)

         ┌─── MAC Cluster 0 ───┐  ┌─── MAC Cluster 1 ───┐
         │                      │  │                      │
Wave 0: [Load K=0→8] [Compute] ╫ [Load K=8→16] [Compute] ╫ ...
Wave 1: [Load K=0→8] [Compute] ╫ [Load K=8→16] [Compute] ╫ ...
Wave 2: [Load K=0→8] [Compute] ╫ [Load K=8→16] [Compute] ╫ ...
Wave 3: [Load K=0→8] [Compute] ╫ [Load K=8→16] [Compute] ╫ ...
                               ↑                          ↑
                          __barrier__               __barrier__
                          
        Time →→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→→

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

Discussion

If this is a relatively large or complex change, feel free to start a discussion by explaining why you chose the solution you did and what alternatives you considered

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.

2 participants