Replace use of anonymous namespace in headers

Using anonymous namespace in headers is against best C++ practices
due to internal linkage of entities in that namespace.

Author: oleksandr-pavlyk

dpctl/tensor/libtensor/include/kernels/accumulators.hpp

@@ -90,7 +90,7 @@ template <typename BinOpT, typename T> struct can_use_inclusive_scan_over_group
     static constexpr bool value = sycl::has_known_identity<BinOpT, T>::value;
 };
 
-namespace
+namespace detail
 {
 template <typename T> class stack_t
 {
@@ -141,7 +141,7 @@ template <typename T> class stack_strided_t
     std::size_t get_local_stride() const { return local_stride_; }
 };
 
-} // end of anonymous namespace
+} // end of namespace detail
 
 // Iterative cumulative summation
 
@@ -627,7 +627,7 @@ sycl::event inclusive_scan_iter_1d(sycl::queue &exec_q,
             throw std::bad_alloc();
         }
 
-        std::vector<stack_t<outputT>> stack{};
+        std::vector<detail::stack_t<outputT>> stack{};
 
         // inclusive scans over blocks
         n_groups_ = n_groups;
@@ -867,7 +867,7 @@ sycl::event inclusive_scan_iter(sycl::queue &exec_q,
             throw std::bad_alloc();
         }
 
-        std::vector<stack_strided_t<outputT>> stack{};
+        std::vector<detail::stack_strided_t<outputT>> stack{};
 
         // inclusive scans over blocks
         acc_groups_ = acc_groups;

dpctl/tensor/libtensor/include/kernels/boolean_advanced_indexing.hpp

@@ -212,7 +212,7 @@ struct MaskedPlaceStridedFunctor
 
 // ======= Masked extraction ================================
 
-namespace
+namespace detail
 {
 
 template <std::size_t I, std::size_t... IR>
@@ -234,7 +234,7 @@ std::size_t get_lws(std::size_t n)
     return _get_lws_impl<lws0, lws1, lws2>(n);
 }
 
-} // end of anonymous namespace
+} // end of namespace detail
 
 template <typename MaskedDstIndexerT, typename dataT, typename indT>
 class masked_extract_all_slices_contig_impl_krn;
@@ -278,7 +278,7 @@ sycl::event masked_extract_all_slices_contig_impl(
 
     const std::size_t masked_extent = iteration_size;
 
-    const std::size_t lws = get_lws(masked_extent);
+    const std::size_t lws = detail::get_lws(masked_extent);
 
     const std::size_t n_groups = (iteration_size + lws - 1) / lws;
 
@@ -357,7 +357,7 @@ sycl::event masked_extract_all_slices_strided_impl(
 
     const std::size_t masked_nelems = iteration_size;
 
-    const std::size_t lws = get_lws(masked_nelems);
+    const std::size_t lws = detail::get_lws(masked_nelems);
 
     const std::size_t n_groups = (masked_nelems + lws - 1) / lws;
 
@@ -452,7 +452,7 @@ sycl::event masked_extract_some_slices_strided_impl(
 
     const std::size_t masked_extent = masked_nelems;
 
-    const std::size_t lws = get_lws(masked_extent);
+    const std::size_t lws = detail::get_lws(masked_extent);
 
     const std::size_t n_groups = ((masked_extent + lws - 1) / lws);
     const std::size_t orthog_extent = static_cast<std::size_t>(orthog_nelems);

dpctl/tensor/libtensor/include/utils/sycl_alloc_utils.hpp

@@ -142,7 +142,7 @@ smart_malloc_host(std::size_t count,
     return smart_malloc<T>(count, q, sycl::usm::alloc::host, propList);
 }
 
-namespace
+namespace detail
 {
 template <typename T> struct valid_smart_ptr : public std::false_type
 {
@@ -172,7 +172,7 @@ struct all_valid_smart_ptrs<Arg, RestArgs...>
     static constexpr bool value = valid_smart_ptr<Arg>::value &&
                                   (all_valid_smart_ptrs<RestArgs...>::value);
 };
-} // namespace
+} // end of namespace detail
 
 template <typename... Args>
 sycl::event async_smart_free(sycl::queue &exec_q,
@@ -184,7 +184,7 @@ sycl::event async_smart_free(sycl::queue &exec_q,
         n > 0, "async_smart_free requires at least one smart pointer argument");
 
     static_assert(
-        all_valid_smart_ptrs<Args...>::value,
+        detail::all_valid_smart_ptrs<Args...>::value,
         "async_smart_free requires unique_ptr created with smart_malloc");
 
     std::vector<void *> ptrs;

dpctl/tensor/libtensor/include/utils/sycl_utils.hpp

@@ -132,7 +132,7 @@ std::size_t choose_workgroup_size(const std::size_t nelems,
     return wg;
 }
 
-namespace
+namespace detail
 {
 
 template <typename LocAccT, typename OpT>
@@ -158,7 +158,7 @@ void _fold(LocAccT &local_mem_acc,
     }
 }
 
-} // namespace
+} // end of namespace detail
 
 template <typename T, typename GroupT, typename LocAccT, typename OpT>
 T custom_reduce_over_group(const GroupT &wg,
@@ -184,7 +184,8 @@ T custom_reduce_over_group(const GroupT &wg,
         for (std::uint32_t sz = high_sz; sz >= low_sz; sz >>= 1) {
             if (n_witems >= sz) {
                 const std::uint32_t n_witems_ = (n_witems + 1) >> 1;
-                _fold(local_mem_acc, lid, n_witems - n_witems_, n_witems_, op);
+                detail::_fold(local_mem_acc, lid, n_witems - n_witems_,
+                              n_witems_, op);
                 sycl::group_barrier(wg, sycl::memory_scope::work_group);
                 n_witems = n_witems_;
             }
@@ -196,7 +197,7 @@ T custom_reduce_over_group(const GroupT &wg,
         for (std::uint32_t sz = high_sz; sz >= low_sz; sz >>= 1) {
             if (n_witems >= sz) {
                 n_witems >>= 1;
-                _fold(local_mem_acc, lid, n_witems, op);
+                detail::_fold(local_mem_acc, lid, n_witems, op);
                 sycl::group_barrier(wg, sycl::memory_scope::work_group);
             }
         }