fix file paths and add an index page with pseudo-code for each pipeline

This commit is contained in:
Philip Maybank
2025-11-14 10:44:25 +00:00
parent ea1ea41889
commit f30c38becc
3 changed files with 49 additions and 2 deletions

View File

@@ -1,3 +1,5 @@
.. _ck_tile-v1:
ck_tile - AGmemBGmemCReg - V1 - [SCHEDULER]
--------------------------------------------
@@ -16,4 +18,4 @@ The ``load_tile_with_elementwise`` function calls the ``load`` method on ``tile_
**full class definition**
.. literalinclude:: ./gemm_pipeline_agmem_bgmem_creg_v1.hpp
.. literalinclude:: ../../../include/ck_tile/ops/gemm/pipeline/gemm_pipeline_agmem_bgmem_creg_v1.hpp

View File

@@ -0,0 +1,44 @@
Pipeline policies
====================
Xdl pipelines
--------------
GEMM Pipelines defined in ``include/ck/tensor_operation/gpu/block`` directory.
:ref:`BlockwiseGemmXdlops_pipeline_v1 <xdl-v1>`
.. code-block::
buffer_load(0)
lds_write(0)
loop:
buffer_load(i+1)
gemm(i)
lds_write(i+1)
**TODO**
- BlockwiseGemmXdlops_pipeline_v2
- BlockwiseGemmXdlops_pipeline_v3
CK_TILE GEMM pipelines
-------------------------
Pipelines defined in ``include/ck_tile/ops/gemm/pipeline`` directory.
:ref:`GemmPipelineAGmemBGmemCRegV1 <ck_tile-v1>`
.. code-block::
buffer_load(0)
lds_write(0)
loop:
buffer_load(i+1)
gemm(i)
lds_write(i+1)
**TODO**
- GemmPipelineAgBgCrMem
- GemmPipelineAgBgCrCompV3

View File

@@ -1,3 +1,4 @@
.. _xdl-v1:
Xdl - v1 - Intrawave
----------------------
@@ -32,4 +33,4 @@ The ``a_thread_copy_.Run`` function is a member of the ``ThreadwiseTensorSliceTr
**full class definition**
.. literalinclude:: ./blockwise_gemm_pipeline_xdlops_v1.hpp
.. literalinclude:: ../../../include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v1.hpp