derek.holloway
7502bf64a4
CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
CK Tile Distribution Encoding Tutorial
Overview
Every load_tile and store_tile in CK needs to know which thread reads which data element.
This mapping is defined by a tile_distribution_encoding — a compile-time struct with 6 template
parameters:
tile_distribution_encoding<Rs, Hs, Ps_major, Ps_minor, Ys_major, Ys_minor>
Every level of Hs (hierarchical dimensions) is assigned to exactly one role:
| Role | Meaning |
|---|---|
| P (parallel) | Thread ID selects which slice — different threads get different data |
| Y (yield) | Each thread owns the entire range in its buffer |
| R (replicate) | Identical data broadcast to multiple thread groups |
Tutorials
These tutorials use the exact tile sizes from the naive GEMM tutorial
(01_naive_gemm/): MPerBlock=256, NPerBlock=128, KPerBlock=32, BlockSize=256, fp16.
| # | File | Matrix | Tile | Key Concept |
|---|---|---|---|---|
| 1 | tile_distribution_1.cpp |
A (DRAM load) | 256×32 | NDimP=2, warp_id→M1, lane_id→M2×K0 (coalesced) |
| 2 | tile_distribution_2.cpp |
B (DRAM load) | 128×32 | Same pattern as A, but N0=2 iterations (vs A's M0=4) due to smaller N |
| 3 | tile_distribution_3.cpp |
C (registers) | 256×128 | Warp-level MFMA output + block-level composition, standard vs transposed |
Tutorial 3 responds to CK_TILE_ENABLE_TRANSPOSED_C_DISTRIBUTION — rebuild with =0 or =1
to see both C register layouts.
Architecture note: All comments and concrete values assume CDNA (warp_size=64). On RDNA (warp_size=32), the thread-to-data mapping will differ.
Building
cd <repo-root>/projects/composablekernel/build
# Build all tutorials:
make tutorials -j
# or: ninja tutorials
# Or build individually:
make tile_tutorial_tile_distribution_1 -j
make tile_tutorial_tile_distribution_2 -j
make tile_tutorial_tile_distribution_3 -j
# Tutorial 3 with standard (non-transposed) C:
cmake -DCMAKE_CXX_FLAGS="-DCK_TILE_ENABLE_TRANSPOSED_C_DISTRIBUTION=0" ..
make tile_tutorial_tile_distribution_3 -j
Reference
- Encoding definition:
include/ck_tile/core/tensor/tile_distribution_encoding.hpp - Thread identity (NDimP):
include/ck_tile/core/tensor/tile_distribution.hpp - MFMA warp output layout:
include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp - Production A/B distributions:
include/ck_tile/ops/gemm/pipeline/gemm_pipeline_agmem_bgmem_creg_v1_default_policy.hpp - Naive GEMM tutorial:
tutorial/ck_tile/gemm/01_naive_gemm/