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

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

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

@@ -18,50 +18,93 @@
 #include <sycl/ext/intel/experimental/esimd/simd.hpp>
 #include <sycl/ext/intel/experimental/esimd/simd_view.hpp>
 
-// Table of contents:
-//
-// simd_obj_impl/simd/simd_mask global operators
-//   bitwise logic and arithmetic operators
-//     simd_obj_impl BINOP simd_obj_impl
-//     simd_obj_impl BINOP SCALAR
-//     SCALAR BINOP simd_obj_impl
-//   comparison operators
-//     simd_obj_impl CMPOP simd_obj_impl
-//     simd_obj_impl CMPOP SCALAR
-//     SCALAR CMPOP simd_obj_impl
-// simd_view global operators
-//   bitwise logic and arithmetic operators
-//     simd_view BINOP simd_view
-//     simd* BINOP simd_view<simd*...>
-//     simd_view<simd*...> BINOP simd*
-//     SCALAR BINOP simd_view
-//     simd_view BINOP SCALAR
-//   comparison operators
-//     simd_view CMPOP simd_view
-//     simd_view CMPOP simd_obj_impl
-//     simd_obj_impl CMPOP simd_view
-//     simd_view CMPOP SCALAR
-//     SCALAR CMPOP simd_view
-//
-// Some operations are enabled only for particular element and simd object type
-// (simd or simd_mask):
-// - bitwise logic operations - for integral element types (both simd and
-//   simd_mask)
-// - bit shift operations and and '%' - for the simd type (not for simd_mask)
-//   with integral element types.
-// - arithmetic binary operations - for the simd type (not for simd_mask)
-// In all cases, when an operation has a simd_view and a simd_obj_impl's
-// subclass objects as operands, it is enabled only when:
-// - simd_view's base type matches the simd object operand. I.e. only
-//   { simd_view<simd, ...>, simd } and { simd_view<simd_mask,...>, simd_mask }
-//   pairs are enabled (with any order of operand types).
-// - simd_view's value length matches the length of the simd object operand
-
-// Put operators into the ESIMD namespace to make argument-dependent lookup find
-// these operators instead of those defined in e.g. sycl namespace (which would
-// stop further lookup, leaving just non-viable sycl::operator < etc. on the
-// table).
-namespace __SEIEED {
+// Put operators into the ESIMD detail namespace to make argument-dependent
+// lookup find these operators instead of those defined in e.g. sycl namespace
+// (which would stop further lookup, leaving just non-viable sycl::operator <
+// etc. on the table).
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace intel {
+namespace experimental {
+namespace esimd {
+namespace detail {
+// clang-format off
+/// @ingroup sycl_esimd_core
+/// @{
+/// @defgroup sycl_esimd_core_binops C++ binary operators overloads for ESIMD.
+///
+/// Standard C++ binary operators overloads applicable to \c simd_obj_impl
+/// derivatives - \c simd , \c simd_mask , \c simd_view and their combinations.
+/// The following overloads are defined:
+///
+/// - \c simd_obj_impl global operators:
+///     + bitwise logic and arithmetic operators
+///          * \c simd_obj_impl BINOP \c simd_obj_impl
+///          * \c simd_obj_impl BINOP SCALAR
+///          * SCALAR BINOP \c simd_obj_impl
+///     + comparison operators
+///          * \c simd_obj_impl CMPOP \c simd_obj_impl
+///          * \c simd_obj_impl CMPOP SCALAR
+///          * SCALAR CMPOP \c simd_obj_impl
+/// - \c simd_view global operators
+///     + bitwise logic and arithmetic operators
+///          * \c simd_view BINOP \c simd_view
+///          * \c simd* BINOP \c simd_view<simd*...>
+///          * \c simd_view<simd*...> BINOP \c simd*
+///          * SCALAR BINOP \c simd_view
+///          * \c simd_view BINOP SCALAR
+/// - comparison operators
+///     * \c simd_view CMPOP \c simd_view
+///     * \c simd_view CMPOP \c simd_obj_impl
+///     * \c simd_obj_impl CMPOP \c simd_view
+///     * \c simd_view CMPOP SCALAR
+///     * SCALAR CMPOP \c simd_view
+///
+/// Some operations are enabled only for particular element type and/or simd
+/// object type (simd or simd_mask):
+/// - bitwise logic operations - for integral element types (both simd and
+///   simd_mask)
+/// - bit shift operations and and '%' - for the simd type (not for simd_mask)
+///   with integral element types
+/// - arithmetic binary operations - for the simd type (not for simd_mask)
+/// In all cases, when an operation has a simd_view and a simd_obj_impl's
+/// subclass objects as operands, it is enabled only when:
+/// - simd_view's base type matches the simd object operand. I.e. only
+///   { simd_view<simd, ...>, simd } and { simd_view<simd_mask,...>, simd_mask }
+///   pairs are enabled (with any order of operand types).
+/// - simd_view's value length matches the length of the simd object operand
+///
+/// The tables below provides more details about supported overloads.
+///
+/// Binary operators:
+/// |              |simd/simd_view (integer)|simd/simd_view (floating point)|simd_mask|
+/// |--------------|:----------------------:|:-----------------------------:|:-------:|
+/// | <tt>+   </tt>|      +                 |          +                    |         |
+/// | <tt>-   </tt>|      +                 |          +                    |         |
+/// | <tt>*   </tt>|      +                 |          +                    |         |
+/// | <tt>/   </tt>|      +                 |          +                    |         |
+/// | <tt>%   </tt>|      +                 |                               |         |
+/// | <tt>\<\<</tt>|      +                 |                               |         |
+/// | <tt>\>\></tt>|      +                 |                               |         |
+/// | <tt>^   </tt>|      +                 |                               |    +    |
+/// | <tt>\|  </tt>|      +                 |                               |    +    |
+/// | <tt>\&  </tt>|      +                 |                               |    +    |
+/// | <tt>\|\|</tt>|                        |                               |    +    |
+/// | <tt>\&\&</tt>|                        |                               |    +    |
+///
+/// Comparison operators
+/// |              |simd/simd_view (integer)|simd/simd_view (floating point)|simd_mask|
+/// |--------------|:----------------------:|:-----------------------------:|:-------:|
+/// | <tt>== </tt> |      +                 |         +                     |    +    |
+/// | <tt>!= </tt> |      +                 |         +                     |    +    |
+/// | <tt>\< </tt> |      +                 |         +                     |         |
+/// | <tt>\> </tt> |      +                 |         +                     |         |
+/// | <tt>\<=</tt> |      +                 |         +                     |         |
+/// | <tt>\>=</tt> |      +                 |         +                     |         |
+/// @}
+// clang-format on
 
 ////////////////////////////////////////////////////////////////////////////////
 // simd_obj_impl global operators
@@ -124,6 +167,7 @@ namespace __SEIEED {
     }                                                                          \
   }
 
+// TODO add doxygen for individual overloads.
 #define __ESIMD_BITWISE_OP_FILTER                                              \
   std::is_integral_v<T1> &&std::is_integral_v<T2>
 __ESIMD_DEF_SIMD_OBJ_IMPL_BIN_OP(^, BinOp::bit_xor, __ESIMD_BITWISE_OP_FILTER)
@@ -237,9 +281,7 @@ __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(||, BinOp::log_or,
                                  __SEIEED::is_simd_mask_type_v<SimdTx>)
 
 #undef __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP
-} // namespace __SEIEED
 
-namespace __SEIEED {
 ////////////////////////////////////////////////////////////////////////////////
 // simd_view global operators
 ////////////////////////////////////////////////////////////////////////////////
@@ -434,4 +476,10 @@ __ESIMD_DEF_SIMD_VIEW_CMP_OP(>=, __SEIEED::is_simd_type_v<SimdT1>)
 
 #undef __ESIMD_DEF_SIMD_VIEW_CMP_OP
 
-} // namespace __SEIEED
+} // namespace detail
+} // namespace esimd
+} // namespace experimental
+} // namespace intel
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)