From f30c38becccefbc8ec2dbdcbb5baf5f3c192bc2a Mon Sep 17 00:00:00 2001 From: Philip Maybank Date: Fri, 14 Nov 2025 10:44:25 +0000 Subject: [PATCH] fix file paths and add an index page with pseudo-code for each pipeline --- docs/reference/pipelines/ck_tile-v1.rst | 4 ++- docs/reference/pipelines/index.rst | 44 +++++++++++++++++++++++++ docs/reference/pipelines/xdl-v1.rst | 3 +- 3 files changed, 49 insertions(+), 2 deletions(-) create mode 100644 docs/reference/pipelines/index.rst diff --git a/docs/reference/pipelines/ck_tile-v1.rst b/docs/reference/pipelines/ck_tile-v1.rst index 223b7b75a1..076cbd7f5b 100644 --- a/docs/reference/pipelines/ck_tile-v1.rst +++ b/docs/reference/pipelines/ck_tile-v1.rst @@ -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 \ No newline at end of file +.. literalinclude:: ../../../include/ck_tile/ops/gemm/pipeline/gemm_pipeline_agmem_bgmem_creg_v1.hpp \ No newline at end of file diff --git a/docs/reference/pipelines/index.rst b/docs/reference/pipelines/index.rst new file mode 100644 index 0000000000..956b509397 --- /dev/null +++ b/docs/reference/pipelines/index.rst @@ -0,0 +1,44 @@ +Pipeline policies +==================== + +Xdl pipelines +-------------- + +GEMM Pipelines defined in ``include/ck/tensor_operation/gpu/block`` directory. + +:ref:`BlockwiseGemmXdlops_pipeline_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 ` + +.. code-block:: + + buffer_load(0) + lds_write(0) + loop: + buffer_load(i+1) + gemm(i) + lds_write(i+1) + +**TODO** + +- GemmPipelineAgBgCrMem +- GemmPipelineAgBgCrCompV3 \ No newline at end of file diff --git a/docs/reference/pipelines/xdl-v1.rst b/docs/reference/pipelines/xdl-v1.rst index fab2c3c5e1..b3bdfd96d2 100644 --- a/docs/reference/pipelines/xdl-v1.rst +++ b/docs/reference/pipelines/xdl-v1.rst @@ -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 \ No newline at end of file +.. literalinclude:: ../../../include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v1.hpp \ No newline at end of file