[ESIMD] Doxygen update part I - fix structure, filter non-public APIs. (#5383)

Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>
Co-authored-by: Alexander Batashev <alexbatashev@outlook.com>

Author: kbobrovs

sycl/doc/doxygen.cfg.in

@@ -1118,7 +1118,7 @@ COLS_IN_ALPHA_INDEX    = 4
 # while generating the index headers.
 # This tag requires that the tag ALPHABETICAL_INDEX is set to YES.
 
-IGNORE_PREFIX          = cl::sycl::
+IGNORE_PREFIX          = cl::sycl:: cl::sycl::ext::intel::experimental::esimd
 
 #---------------------------------------------------------------------------
 # Configuration options related to the HTML output
@@ -2115,7 +2115,10 @@ PREDEFINED             = "__SYCL_INLINE_NAMESPACE(X)=namespace X" \
 # definition found in the source code.
 # This tag requires that the tag ENABLE_PREPROCESSING is set to YES.
 
-EXPAND_AS_DEFINED      =
+EXPAND_AS_DEFINED      = \
+  __ESIMD_DEF_SIMD_OBJ_IMPL_BIN_OP __ESIMD_BITWISE_OP_FILTER \
+  __ESIMD_SHIFT_OP_FILTER __ESIMD_UNARY_INTRINSIC_DEF __ESIMD_EMATH_COND \
+  __ESIMD_BINARY_INTRINSIC_DEF __ESIMD_INTRINSIC_DEF
 
 # If the SKIP_FUNCTION_MACROS tag is set to YES then doxygen's preprocessor will
 # remove all references to function-like macros that are alone on a line, have

sycl/include/sycl/ext/intel/experimental/esimd.hpp

@@ -10,7 +10,43 @@
 
 #pragma once
 
-/// \defgroup sycl_esimd DPC++ Explicit SIMD API
+/// @defgroup sycl_esimd DPC++ Explicit SIMD API
+/// This is a low-level API providing direct access to Intel GPU hardware
+/// features. ESIMD overview can be found
+/// [here](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md).
+
+///@{
+/// @ingroup sycl_esimd
+
+/// @defgroup sycl_esimd_core ESIMD core.
+/// Core APIs defining main vector data types and their interfaces.
+
+/// @defgroup sycl_esimd_memory Memory access API.
+/// ESIMD APIs to access memory via accessors, USM pointers, perform per-element
+/// atomic operations.
+
+/// @defgroup sycl_esimd_math ESIMD math operations.
+/// Defines math operations on ESIMD vector data types.
+
+/// @defgroup sycl_esimd_bitmanip Bit and mask manipulation APIs.
+
+/// @defgroup sycl_esimd_conv Explicit conversions.
+/// @ingroup sycl_esimd
+/// Defines explicit conversions (with and without saturation), truncation etc.
+/// between ESIMD vector types.
+
+/// @defgroup sycl_esimd_misc Miscellaneous ESIMD convenience functions.
+
+/// The main components of the API are:
+///   - @ref sycl_esimd_core - core API defining main vector data types and
+///   their
+///     interfaces.
+///   - @ref sycl_esimd_memory
+///   - @ref sycl_esimd_math
+///   - @ref sycl_esimd_bitmanip
+///   - @ref sycl_esimd_conv
+///   - @ref sycl_esimd_misc
+///@}
 
 #include <sycl/ext/intel/experimental/esimd/alt_ui.hpp>
 #include <sycl/ext/intel/experimental/esimd/common.hpp>

sycl/include/sycl/ext/intel/experimental/esimd/alt_ui.hpp

@@ -18,13 +18,15 @@ namespace intel {
 namespace experimental {
 namespace esimd {
 
+/// @{
+/// @ingroup sycl_esimd_misc
+
 /// "Merges" elements of the input vectors according to the merge mask.
 /// @param a the first vector
 /// @param b the second vector
 /// @param m the merge mask
 /// @return a vector, where each element equals to corresponding element from
 ///    \c a (if corresponding mask element is zero) or \c b (otherwise)
-/// \ingroup sycl_esimd
 template <class T, int N>
 __ESIMD_API simd<T, N> merge(simd<T, N> a, simd<T, N> b, simd_mask<N> m) {
   a.merge(b, m);
@@ -53,6 +55,8 @@ __ESIMD_API auto merge(simd_view<BaseT1, RegionT1> v1,
   return merge(v1.read(), v2.read(), m);
 }
 
+/// @} sycl_esimd_misc
+
 } // namespace esimd
 } // namespace experimental
 } // namespace intel

sycl/include/sycl/ext/intel/experimental/esimd/common.hpp

@@ -74,6 +74,9 @@ namespace intel {
 namespace experimental {
 namespace esimd {
 
+/// @{
+/// @ingroup sycl_esimd_core
+
 using uchar = unsigned char;
 using ushort = unsigned short;
 using uint = unsigned int;
@@ -255,6 +258,8 @@ using EsimdSbarrierType = split_barrier_action;
 /// identified by its "binding table index" - surface index.
 using SurfaceIndex = unsigned int;
 
+/// @} sycl_esimd_core
+
 } // namespace esimd
 } // namespace experimental
 } // namespace intel

sycl/include/sycl/ext/intel/experimental/esimd/detail/atomic_intrin.hpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 #pragma once
 
+/// @cond ESIMD_DETAIL
+
 #include <CL/sycl/exception.hpp>
 
 // This function implements atomic update of pre-existing variable in the
@@ -19,3 +21,5 @@ template <typename Ty> Ty atomic_add_fetch(Ty *ptr, Ty val) {
   return __atomic_add_fetch(ptr, val, __ATOMIC_RELAXED);
 #endif
 }
+
+/// @endcond ESIMD_DETAIL

sycl/include/sycl/ext/intel/experimental/esimd/detail/elem_type_traits.hpp

@@ -85,6 +85,8 @@
 
 #include <CL/sycl/half_type.hpp>
 
+/// @cond ESIMD_DETAIL
+
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace __SEIEED {
 
@@ -580,7 +582,6 @@ template <typename T>
 static inline constexpr bool is_generic_floating_point_v =
     element_type_traits<T>::is_floating_point;
 
-// @{
 // Get computation type of a binary operator given its operand types:
 // - if both types are arithmetic - return CPP's "common real type" of the
 //   computation (matches C++)
@@ -653,8 +654,6 @@ template <class T1, class T2 = T1>
 using computation_type_t =
     typename computation_type<remove_cvref_t<T1>, remove_cvref_t<T2>>::type;
 
-// @}
-
 ////////////////////////////////////////////////////////////////////////////////
 // sycl::half traits
 ////////////////////////////////////////////////////////////////////////////////
@@ -723,3 +722,5 @@ inline std::istream &operator>>(std::istream &I, sycl::half &rhs) {
 
 } // namespace __SEIEED
 } // __SYCL_INLINE_NAMESPACE(cl)
+
+/// @endcond ESIMD_DETAIL

sycl/include/sycl/ext/intel/experimental/esimd/detail/host_util.hpp

@@ -10,6 +10,8 @@
 
 #pragma once
 
+/// @cond ESIMD_DETAIL
+
 #ifndef __SYCL_DEVICE_ONLY__
 
 #include <assert.h>
@@ -458,3 +460,5 @@ template <> struct dwordtype<unsigned int> { static const bool value = true; };
 } // __SYCL_INLINE_NAMESPACE(cl)
 
 #endif // #ifndef __SYCL_DEVICE_ONLY__
+
+/// @endcond ESIMD_DETAIL

sycl/include/sycl/ext/intel/experimental/esimd/detail/intrin.hpp

@@ -11,6 +11,8 @@
 
 #pragma once
 
+/// @cond ESIMD_DETAIL
+
 #include <sycl/ext/intel/experimental/esimd/common.hpp>
 #include <sycl/ext/intel/experimental/esimd/detail/types.hpp>
 #include <sycl/ext/intel/experimental/esimd/detail/util.hpp>
@@ -343,3 +345,5 @@ __esimd_wrindirect(__SEIEED::vector_type_t<T, N> OldVal,
 }
 
 #endif // __SYCL_DEVICE_ONLY__
+
+/// @endcond ESIMD_DETAIL

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp

@@ -11,6 +11,8 @@
 
 #pragma once
 
+/// @cond ESIMD_DETAIL
+
 #include <CL/sycl/builtins.hpp>
 #include <sycl/ext/intel/experimental/esimd/common.hpp>
 #include <sycl/ext/intel/experimental/esimd/detail/elem_type_traits.hpp>
@@ -1199,3 +1201,5 @@ __ESIMD_INTRIN __ESIMD_raw_vec_t(T, N)
 #endif // #ifdef __SYCL_DEVICE_ONLY__
 
 #undef __ESIMD_raw_vec_t
+
+/// @endcond ESIMD_DETAIL

sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp

@@ -9,6 +9,8 @@
 // the SIMD classes objects.
 //===----------------------------------------------------------------------===//
 
+/// @cond ESIMD_DETAIL
+
 #pragma once
 
 #include <CL/sycl/detail/accessor_impl.hpp>
@@ -924,3 +926,5 @@ __ESIMD_INTRIN void __esimd_raw_send2_noresult(
   throw cl::sycl::feature_not_supported();
 }
 #endif // __SYCL_DEVICE_ONLY__
+
+/// @endcond ESIMD_DETAIL