diff --git a/docs/reference/pipelines/ck_tile-v1.rst b/docs/reference/pipelines/ck_tile-v1.rst new file mode 100644 index 0000000000..223b7b75a1 --- /dev/null +++ b/docs/reference/pipelines/ck_tile-v1.rst @@ -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 \ No newline at end of file diff --git a/docs/reference/pipelines/xdl-v1.rst b/docs/reference/pipelines/xdl-v1.rst new file mode 100644 index 0000000000..fab2c3c5e1 --- /dev/null +++ b/docs/reference/pipelines/xdl-v1.rst @@ -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{}), + 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{}), + 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 \ No newline at end of file