initial work on producing a reference that explains each pipeline policy class

This commit is contained in:
Philip Maybank
2025-11-13 17:16:14 +00:00
parent 8ac9550ac8
commit ea1ea41889
2 changed files with 54 additions and 0 deletions

View File

@@ -0,0 +1,19 @@
ck_tile - AGmemBGmemCReg - V1 - [SCHEDULER]
--------------------------------------------
**loop: buffer_load(i+1)**
.. code-block::
elementwise_As_res = load_tile_with_elementwise(as_copy_dram_window, a_element_func);
elementwise_Bs_res = load_tile_with_elementwise(bs_copy_dram_window, b_element_func);
The ``load_tile_with_elementwise`` function calls the ``load`` method on ``tile_window_with_static_distribution`` instances.
- The ``tile_window_with_static_distribution::load`` method calls ``make_static_distributed_tensor`` to create distributed tensors for storing loaded data.
- The ``static_distributed_tensor`` class uses ``make_naive_tensor_descriptor_packed`` from ``ck_tile/core/tensor/tensor_descriptor.hpp`` in its tensor slicing operations (via ``set_y_sliced_thread_data``).
**full class definition**
.. literalinclude:: ./gemm_pipeline_agmem_bgmem_creg_v1.hpp

View File

@@ -0,0 +1,35 @@
Xdl - v1 - Intrawave
----------------------
**loop: buffer_load(i+1)**
.. code-block::
static_for<0, KRepeat, 1>{}([&](auto k) {
static_for<0, MRepeat, 1>{}([&](auto m0) {
a_thread_copy_.Run(a_block_desc_m0_m1_m2_k,
make_tuple(m0, I0, I0, Number<k * AMmaKStride>{}),
a_block_buf,
a_thread_desc_,
make_tuple(m0, I0, k, I0),
a_thread_buf);
static_for<0, NRepeat, 1>{}([&](auto n0) {
b_thread_copy_.Run(b_block_desc_n0_n1_n2_k,
make_tuple(n0, I0, I0, Number<k * BMmaKStride>{}),
b_block_buf,
b_thread_desc_,
make_tuple(n0, I0, k, I0),
b_thread_buf);
});
});
});
The ``a_thread_copy_.Run`` function is a member of the ``ThreadwiseTensorSliceTransfer_v4`` class. The ``ThreadwiseTensorSliceTransfer_v4`` class uses,
- ``make_tensor_coordinate`` from ``tensor_descriptor.hpp``
- ``make_naive_tensor_descriptor_packed`` from ``tensor_descriptor_helper.hpp``
**full class definition**
.. literalinclude:: ./blockwise_gemm_pipeline_xdlops_v1.hpp