Doc reorg2 (#1189)

* doc_reorg2 updated TOC

* doc_reorg2 updates

* fix conflicts, add grid

docs/reference/API_Reference_Guide.rst

@@ -0,0 +1,54 @@
+.. meta::
+  :description: Composable Kernel documentation and API reference library
+  :keywords: composable kernel, CK, ROCm, API, documentation
+
+.. _api-reference:
+
+********************************************************************
+API reference guide
+********************************************************************
+
+
+This document contains details of the APIs for the Composable Kernel (CK) library and introduces
+some of the key design principles that are used to write new classes that extend CK functionality.
+
+=================
+Using CK API
+=================
+
+This section describes how to use the CK library API.
+
+=================
+CK Datatypes
+=================
+
+-----------------
+DeviceMem
+-----------------
+
+.. doxygenstruct:: DeviceMem
+
+---------------------------
+Kernels For Flashattention
+---------------------------
+
+The Flashattention algorithm is defined in :cite:t:`dao2022flashattention`. This section lists
+the classes that are used in the CK GPU implementation of Flashattention.
+
+**Gridwise classes**
+
+.. doxygenstruct:: ck::GridwiseBatchedGemmSoftmaxGemm_Xdl_CShuffle
+
+**Blockwise classes**
+
+.. doxygenstruct:: ck::ThreadGroupTensorSliceTransfer_v4r1
+
+.. doxygenstruct:: ck::BlockwiseGemmXdlops_v2
+
+.. doxygenstruct:: ck::BlockwiseSoftmax
+
+**Threadwise classes**
+
+.. doxygenstruct:: ck::ThreadwiseTensorSliceTransfer_StaticToStatic
+
+.. bibliography::

docs/reference/Supported_Primitives_Guide.rst

@@ -0,0 +1,80 @@
+.. meta::
+  :description: Composable Kernel documentation and API reference library
+  :keywords: composable kernel, CK, ROCm, API, documentation
+
+.. _supported-primitives:
+
+********************************************************************
+Supported Primitives Guide
+********************************************************************
+
+This document contains details of supported primitives in Composable Kernel (CK). In contrast to the API Reference Guide, the Supported Primitives Guide is an introduction to the math which underpins the algorithms implemented in CK.
+
+------------
+Softmax
+------------
+
+For vectors :math:`x^{(1)}, x^{(2)}, \ldots, x^{(T)}` of size :math:`B` you can decompose the
+softmax of concatenated :math:`x = [ x^{(1)}\ | \ \ldots \ | \ x^{(T)} ]` as,
+
+.. math::
+   :nowrap:
+
+   \begin{align}
+      m(x) & = m( [ x^{(1)}\ | \ \ldots \ | \ x^{(T)} ] ) = \max( m(x^{(1)}),\ldots, m(x^{(T)}) )  \\
+      f(x) & = [\exp( m(x^{(1)}) - m(x) ) f( x^{(1)} )\ | \ \ldots \ | \ \exp( m(x^{(T)}) - m(x) ) f( x^{(T)} )] \\
+      z(x) & = \exp( m(x^{(1)}) - m(x) )\ z(x^{(1)}) + \ldots + \exp( m(x^{(T)}) - m(x) )\ z(x^{(1)}) \\
+      \operatorname{softmax}(x) &= f(x)\ / \ z(x)
+   \end{align}
+
+where :math:`f(x^{(j)}) = \exp( x^{(j)} - m(x^{(j)}) )` is of size :math:`B` and
+:math:`z(x^{(j)}) = f(x_1^{(j)})+ \ldots+ f(x_B^{(j)})` is a scalar.
+
+For a matrix :math:`X` composed of :math:`T_r \times T_c` tiles, :math:`X_{ij}`, of size
+:math:`B_r \times B_c` you can compute the row-wise softmax as follows.
+
+For :math:`j` from :math:`1` to :math:`T_c`, and :math:`i` from :math:`1` to :math:`T_r` calculate,
+
+.. math::
+   :nowrap:
+
+   \begin{align}
+      \tilde{m}_{ij}   &= \operatorname{rowmax}( X_{ij} ) \\
+      \tilde{P}_{ij}   &= \exp(X_{ij} - \tilde{m}_{ij} ) \\
+      \tilde{z}_{ij}   &= \operatorname{rowsum}( P_{ij} ) \\
+   \end{align}
+
+If :math:`j=1`, initialize running max, running sum, and the first column block of the output,
+
+.. math::
+   :nowrap:
+
+   \begin{align}
+      m_i            &= \tilde{m}_{i1} \\
+      z_i            &= \tilde{z}_{i1} \\
+      \tilde{Y}_{i1} &= \diag(\tilde{z}_{ij})^{-1} \tilde{P}_{i1}
+   \end{align}
+
+Else if :math:`j>1`,
+
+1. Update running max, running sum and column blocks :math:`k=1` to :math:`k=j-1`
+
+.. math::
+   :nowrap:
+
+   \begin{align}
+      m^{new}_i &= \max(m_i, \tilde{m}_{ij} ) \\
+      z^{new}_i &= \exp(m_i - m^{new}_i)\ z_i + \exp( \tilde{m}_{ij} - m^{new}_i )\ \tilde{z}_{ij}  \\
+      Y_{ik}    &= \diag(z^{new}_{i})^{-1} \diag(z_{i}) \exp(m_i - m^{new}_i)\ Y_{ik}
+   \end{align}
+
+2. Initialize column block :math:`j` of output and reset running max and running sum variables:
+
+.. math::
+   :nowrap:
+
+   \begin{align}
+      \tilde{Y}_{ij} &= \diag(z^{new}_{i})^{-1} \exp(\tilde{m}_{ij} - m^{new}_i ) \tilde{P}_{ij} \\
+      z_i            &= z^{new}_i \\
+      m_i            &= m^{new}_i \\
+   \end{align}

docs/reference/wrapper.rst

@@ -0,0 +1,93 @@
+.. meta::
+  :description: Composable Kernel documentation and API reference library
+  :keywords: composable kernel, CK, ROCm, API, documentation
+
+.. _wrapper:
+
+********************************************************************
+Wrapper
+********************************************************************
+
+-------------------------------------
+Description
+-------------------------------------
+
+.. note::
+
+    The wrapper is under development and its functionality is limited.
+
+
+The CK library provides a lightweight wrapper for more complex operations implemented in 
+the library. It allows indexing of nested layouts using a simple interface 
+(avoiding complex descriptor transformations) and memory access (using Tensor).
+
+Example:
+
+.. code-block:: c
+
+    const auto shape_4x2x4         = ck::make_tuple(4, ck::make_tuple(2, 4));
+    const auto strides_s2x1x8      = ck::make_tuple(2, ck::make_tuple(1, 8));
+    const auto layout = ck::wrapper::make_layout(shape_4x2x4, strides_s2x1x8);
+    
+    std::array<ck::index_t, 32> data;
+    auto tensor = ck::wrapper::make_tensor<ck::wrapper::MemoryTypeEnum::Generic>(&data[0], layout);
+
+    for(ck::index_t w = 0; w < size(tensor); w++) {
+        tensor(w) = w;
+    }
+
+    // slice() == slice(0, -1) (whole dimension)
+    auto tensor_slice = tensor(ck::wrapper::slice(1, 3), ck::make_tuple(ck::wrapper::slice(), ck::wrapper::slice()));
+    std::cout << "dims:2,(2,4) strides:2,(1,8)" << std::endl;
+    for(ck::index_t h = 0; h < ck::wrapper::size<0>(tensor_slice); h++)
+    {
+        for(ck::index_t w = 0; w < ck::wrapper::size<1>(tensor_slice); w++)
+        {
+            std::cout << tensor_slice(h, w) << " ";
+        }
+        std::cout << std::endl;
+    }
+
+Output::
+
+    dims:2,(2,4) strides:2,(1,8)
+    1 5 9 13 17 21 25 29 
+    2 6 10 14 18 22 26 30 
+
+
+Advanced examples:
+
+* `Image to column <https://github.com/ROCm/composable_kernel/blob/develop/client_example/25_wrapper/wrapper_img2col.cpp>`_
+
+-------------------------------------
+Layout
+-------------------------------------
+
+.. doxygenstruct:: Layout
+
+-------------------------------------
+Layout helpers
+-------------------------------------
+
+.. doxygenfile:: include/ck/wrapper/utils/layout_utils.hpp
+
+-------------------------------------
+Tensor
+-------------------------------------
+
+.. doxygenstruct:: Tensor
+
+-------------------------------------
+Tensor helpers
+-------------------------------------
+
+.. doxygenfile:: include/ck/wrapper/utils/tensor_utils.hpp
+
+.. doxygenfile:: include/ck/wrapper/utils/tensor_partition.hpp
+
+-------------------------------------
+Operations
+-------------------------------------
+
+.. doxygenfile:: include/ck/wrapper/operations/copy.hpp
+.. doxygenfile:: include/ck/wrapper/operations/gemm.hpp