diff --git a/docs/conf.py b/docs/conf.py index e441ff1ced..e8617a09ef 100644 --- a/docs/conf.py +++ b/docs/conf.py @@ -45,3 +45,5 @@ for sphinx_var in ROCmDocs.SPHINX_VARS: extensions += ['sphinxcontrib.bibtex'] bibtex_bibfiles = ['refs.bib'] + +cpp_id_attributes = ["__global__", "__device__", "__host__"] diff --git a/docs/wrapper.rst b/docs/wrapper.rst index 39e2fd0bbd..190fbcd445 100644 --- a/docs/wrapper.rst +++ b/docs/wrapper.rst @@ -64,31 +64,31 @@ Advanced examples: Layout ------------------------------------- -.. doxygenstruct:: ck::wrapper::Layout +.. doxygenstruct:: Layout ------------------------------------- Layout helpers ------------------------------------- -.. doxygenfile:: layout_utils.hpp +.. doxygenfile:: include/ck/wrapper/utils/layout_utils.hpp ------------------------------------- Tensor ------------------------------------- -.. doxygenstruct:: ck::wrapper::Tensor +.. doxygenstruct:: Tensor ------------------------------------- Tensor helpers ------------------------------------- -.. doxygenfile:: tensor_utils.hpp +.. doxygenfile:: include/ck/wrapper/utils/tensor_utils.hpp -.. doxygenfile:: tensor_partition.hpp +.. doxygenfile:: include/ck/wrapper/utils/tensor_partition.hpp ------------------------------------- Operations ------------------------------------- -.. doxygenfile:: copy.hpp -.. doxygenfile:: gemm.hpp +.. doxygenfile:: include/ck/wrapper/operations/copy.hpp +.. doxygenfile:: include/ck/wrapper/operations/gemm.hpp diff --git a/include/ck/wrapper/layout.hpp b/include/ck/wrapper/layout.hpp index 71c512e136..5cd1f614e6 100644 --- a/include/ck/wrapper/layout.hpp +++ b/include/ck/wrapper/layout.hpp @@ -5,8 +5,11 @@ #include "ck/wrapper/utils/layout_utils.hpp" +// Disable from doxygen docs generation +/// @cond INTERNAL namespace ck { namespace wrapper { +/// @endcond /** * \brief Layout wrapper that performs the tensor descriptor logic. @@ -19,6 +22,8 @@ namespace wrapper { template struct Layout { + // Disable from doxygen docs generation + /// @cond INTERNAL private: static constexpr auto I0 = Number<0>{}; static constexpr auto I1 = Number<1>{}; @@ -246,6 +251,7 @@ struct Layout using Descriptor1dType = remove_cvref_t; using DefaultIdxsTupleType = remove_cvref_t; + /// @endcond public: using LayoutShape = Shape; @@ -457,6 +463,8 @@ struct Layout return unrolled_descriptor_; } + // Disable from doxygen docs generation + /// @cond INTERNAL private: // All dimensions are unrolled UnrolledDescriptorType unrolled_descriptor_; @@ -469,6 +477,7 @@ struct Layout // Descriptor1dType lengths: (8) // MergedNestsDescriptorType lengths: (4, 2) const Shape shape_; + /// @endcond }; } // namespace wrapper diff --git a/include/ck/wrapper/operations/copy.hpp b/include/ck/wrapper/operations/copy.hpp index 5f64031ebe..e8a919fdda 100644 --- a/include/ck/wrapper/operations/copy.hpp +++ b/include/ck/wrapper/operations/copy.hpp @@ -12,8 +12,11 @@ #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/tensor_description/tensor_space_filling_curve.hpp" +// Disable from doxygen docs generation +/// @cond INTERNAL namespace ck { namespace wrapper { +/// @endcond /** * \brief Perform optimized copy between two tensors partitions (threadwise copy). diff --git a/include/ck/wrapper/operations/gemm.hpp b/include/ck/wrapper/operations/gemm.hpp index e41cd5bd8a..42a70239ad 100644 --- a/include/ck/wrapper/operations/gemm.hpp +++ b/include/ck/wrapper/operations/gemm.hpp @@ -9,9 +9,14 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/tensor_operation/gpu/block/blockwise_gemm_xdlops.hpp" +// Disable from doxygen docs generation +/// @cond INTERNAL namespace ck { namespace wrapper { +/// @endcond +// Disable from doxygen docs generation +/// @cond INTERNAL namespace { namespace detail { /** @@ -45,6 +50,7 @@ __device__ constexpr auto GetBlockDescriptor() } // namespace detail } // namespace +/// @endcond /** * \brief Perform blockwise gemm xdl on tensors stored in lds. Result will be diff --git a/include/ck/wrapper/tensor.hpp b/include/ck/wrapper/tensor.hpp index 6946e79ea4..8dabb58451 100644 --- a/include/ck/wrapper/tensor.hpp +++ b/include/ck/wrapper/tensor.hpp @@ -7,9 +7,14 @@ #include "utils/tensor_partition.hpp" #include "utils/layout_utils.hpp" +// Disable from doxygen docs generation +/// @cond INTERNAL namespace ck { namespace wrapper { +/// @endcond +// Disable from doxygen docs generation +/// @cond INTERNAL namespace { namespace detail { /** @@ -189,6 +194,7 @@ __host__ __device__ constexpr auto GenerateSlicedDescriptor(const Tuple& } } // namespace detail } // namespace +/// @endcond /** * \brief Tensor wrapper that performs static and dynamic buffer logic. @@ -394,6 +400,8 @@ struct Tensor } private: + // Disable from doxygen docs generation + /// @cond INTERNAL using DynamicBufferType = DynamicBuffer struct Layout; diff --git a/include/ck/wrapper/utils/tensor_partition.hpp b/include/ck/wrapper/utils/tensor_partition.hpp index 141e0a58e5..69fd502d63 100644 --- a/include/ck/wrapper/utils/tensor_partition.hpp +++ b/include/ck/wrapper/utils/tensor_partition.hpp @@ -9,9 +9,14 @@ #include "ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp" #include "ck/tensor_description/cluster_descriptor.hpp" +// Disable from doxygen docs generation +/// @cond INTERNAL namespace ck { namespace wrapper { +/// @endcond +// Disable from doxygen docs generation +/// @cond INTERNAL namespace { namespace detail { @@ -236,6 +241,7 @@ __host__ __device__ constexpr auto CalculateThreadMultiIdx( } } // namespace detail } // namespace +/// @endcond /** * \brief Create local partition for thread (At now only packed partition diff --git a/include/ck/wrapper/utils/tensor_utils.hpp b/include/ck/wrapper/utils/tensor_utils.hpp index ee9e438a40..ccab99fac3 100644 --- a/include/ck/wrapper/utils/tensor_utils.hpp +++ b/include/ck/wrapper/utils/tensor_utils.hpp @@ -13,8 +13,11 @@ #include "ck/utility/amd_address_space.hpp" #include "ck/utility/multi_index.hpp" +// Disable from doxygen docs generation +/// @cond INTERNAL namespace ck { namespace wrapper { +/// @endcond /** * \brief Memory type, allowed members: @@ -27,7 +30,7 @@ namespace wrapper { using MemoryTypeEnum = AddressSpaceEnum; // Disable from doxygen docs generation -/// @cond +/// @cond INTERNAL // forward declarations template struct Layout;