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