Skip to content

Commit

Permalink
clang-format
Browse files Browse the repository at this point in the history 
Signed-off-by: Matthew Michel <[email protected]>
  • Loading branch information
@mmichel11
mmichel11 committed Dec 19, 2024
1 parent 505bdf3 commit 114924d
Show file tree
Hide file tree
Showing 4 changed files with 43 additions and 36 deletions.
28 changes: 17 additions & 11 deletions include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_for.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -89,8 +89,7 @@ struct __parallel_for_large_submitter<__internal::__optional_kernel_name<_Name..
// device.
static inline std::tuple<std::size_t, std::size_t, bool>
__stride_recommender(const sycl::nd_item<1>& __item, std::size_t __count, std::size_t __iters_per_work_item,
std::size_t __adj_elements_per_work_item,
std::size_t __work_group_size)
std::size_t __adj_elements_per_work_item, std::size_t __work_group_size)
{
if constexpr (oneapi::dpl::__internal::__is_spirv_target_v)
{
Expand All @@ -101,20 +100,27 @@ struct __parallel_for_large_submitter<__internal::__optional_kernel_name<_Name..
const std::size_t __work_group_id = __item.get_group().get_group_linear_id();

const std::size_t __sub_group_start_idx =
__iters_per_work_item * __adj_elements_per_work_item * (__work_group_id * __work_group_size + __sub_group_size * __sub_group_id);
__iters_per_work_item * __adj_elements_per_work_item *
(__work_group_id * __work_group_size + __sub_group_size * __sub_group_id);
const bool __is_full_sub_group =
__sub_group_start_idx + __iters_per_work_item * __adj_elements_per_work_item * __sub_group_size <= __count;
const std::size_t __work_item_idx = __sub_group_start_idx + __adj_elements_per_work_item * __sub_group_local_id;
return std::make_tuple(__work_item_idx, __adj_elements_per_work_item * __sub_group_size, __is_full_sub_group);
__sub_group_start_idx + __iters_per_work_item * __adj_elements_per_work_item * __sub_group_size <=
__count;
const std::size_t __work_item_idx =
__sub_group_start_idx + __adj_elements_per_work_item * __sub_group_local_id;
return std::make_tuple(__work_item_idx, __adj_elements_per_work_item * __sub_group_size,
__is_full_sub_group);
}
else
{
const std::size_t __work_group_start_idx =
__item.get_group().get_group_linear_id() * __work_group_size * __iters_per_work_item * __adj_elements_per_work_item;
const std::size_t __work_item_idx = __work_group_start_idx + __item.get_local_linear_id() * __adj_elements_per_work_item;
const std::size_t __work_group_start_idx = __item.get_group().get_group_linear_id() * __work_group_size *
__iters_per_work_item * __adj_elements_per_work_item;
const std::size_t __work_item_idx =
__work_group_start_idx + __item.get_local_linear_id() * __adj_elements_per_work_item;
const bool __is_full_work_group =
__work_group_start_idx + __iters_per_work_item * __work_group_size * __adj_elements_per_work_item <= __count;
return std::make_tuple(__work_item_idx, __work_group_size * __adj_elements_per_work_item, __is_full_work_group);
__work_group_start_idx + __iters_per_work_item * __work_group_size * __adj_elements_per_work_item <=
__count;
return std::make_tuple(__work_item_idx, __work_group_size * __adj_elements_per_work_item,
__is_full_work_group);
}
}

Expand Down
44 changes: 22 additions & 22 deletions include/oneapi/dpl/pstl/hetero/dpcpp/unseq_backend_sycl.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -116,12 +116,12 @@ template <typename... _Ranges>
class walk_vector_or_scalar_base
{
using _ValueTypes = std::tuple<oneapi::dpl::__internal::__value_t<_Ranges>...>;
constexpr static std::uint8_t __min_type_size =
oneapi::dpl::__internal::__min_nested_type_size<_ValueTypes>::value;
constexpr static std::uint8_t __min_type_size = oneapi::dpl::__internal::__min_nested_type_size<_ValueTypes>::value;
// Empirically determined 'bytes-in-flight' to maximize bandwidth utilization
constexpr static std::uint8_t __bytes_per_item = 16;
// Maximum size supported by compilers to generate vector instructions
constexpr static std::uint8_t __max_vector_size = 4;

public:
constexpr static bool __can_vectorize =
(oneapi::dpl::__ranges::__is_vectorizable_range<_Ranges>::value && ...) &&
Expand All @@ -138,9 +138,9 @@ template <typename... _Ranges>
class walk_scalar_base
{
using _ValueTypes = std::tuple<oneapi::dpl::__internal::__value_t<_Ranges>...>;
constexpr static std::uint8_t __min_type_size =
oneapi::dpl::__internal::__min_nested_type_size<_ValueTypes>::value;
constexpr static std::uint8_t __min_type_size = oneapi::dpl::__internal::__min_nested_type_size<_ValueTypes>::value;
constexpr static std::uint8_t __bytes_per_item = 16;

public:
constexpr static bool __can_vectorize = false;
// With no vectorization, the vector size is 1
Expand All @@ -163,8 +163,8 @@ struct walk1_vector_or_scalar : public walk_vector_or_scalar_base<_Range>
{
// This is needed to enable vectorization
auto __raw_ptr = __rng.begin();
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(
__is_full, __idx, __f, __raw_ptr);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(__is_full, __idx, __f,
__raw_ptr);
}

// _IsFull is ignored here. We assume that boundary checking has been already performed for this index.
Expand Down Expand Up @@ -213,8 +213,8 @@ struct walk2_vectors_or_scalars : public walk_vector_or_scalar_base<_Range1, _Ra
__is_full, __idx, oneapi::dpl::__par_backend_hetero::__lazy_store_transform_op<_F>{__f}, __rng1_vector,
__raw_ptr2);
// 3. Explicitly call destructor of lazy union type
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(__is_full, 0,
oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng1_vector);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(
__is_full, 0, oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng1_vector);
}

// _IsFull is ignored here. We assume that boundary checking has been already performed for this index.
Expand Down Expand Up @@ -269,10 +269,10 @@ struct walk3_vectors_or_scalars : public walk_vector_or_scalar_base<_Range1, _Ra
__is_full, __idx, oneapi::dpl::__par_backend_hetero::__lazy_store_transform_op<_F>{__f}, __rng1_vector,
__rng2_vector, __raw_ptr3);
// 3. Explicitly call destructors of lazy union type
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(__is_full, 0,
oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng1_vector);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(__is_full, 0,
oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng2_vector);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(
__is_full, 0, oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng1_vector);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(
__is_full, 0, oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng2_vector);
}

// _IsFull is ignored here. We assume that boundary checking has been already performed for this index.
Expand Down Expand Up @@ -358,8 +358,8 @@ struct walk_adjacent_difference : public walk_vector_or_scalar_base<_Range1, _Ra
if (__idx == 0)
__rng2[0] = __rng1_vector[0].__v;
// 3. Delete temporary storage
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(__is_full, 0,
oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng1_vector);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(
__is_full, 0, oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng1_vector);
}
template <typename _IsFull, typename _ItemId>
void
Expand Down Expand Up @@ -1195,10 +1195,10 @@ struct __reverse_functor : public walk_vector_or_scalar_base<_Range>
oneapi::dpl::__par_backend_hetero::__lazy_store_transform_op<oneapi::dpl::__internal::__pstl_assign>{},
__rng_right_vector, __rng_pointer);
// 4. Call destructors of temporary storage
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(__is_full, 0,
oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng_left_vector);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(__is_full, 0,
oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng_right_vector);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(
__is_full, 0, oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng_left_vector);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(
__is_full, 0, oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng_right_vector);
}
template <typename _IsFull, typename _Idx>
void
Expand Down Expand Up @@ -1268,8 +1268,8 @@ struct __reverse_copy : public walk_vector_or_scalar_base<_Range1, _Range2>
__rng2_pointer[__output_start + __i] = __rng1_vector[__i].__v;
}
// 3. Cleanup
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(__is_full, 0,
oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng1_vector);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(
__is_full, 0, oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng1_vector);
}
template <typename _IsFull, typename _Idx>
void
Expand Down Expand Up @@ -1323,8 +1323,8 @@ struct __rotate_copy : public walk_vector_or_scalar_base<_Range1, _Range2>
oneapi::dpl::__par_backend_hetero::__lazy_store_transform_op<oneapi::dpl::__internal::__pstl_assign>{},
__rng1_vector, __rng2_pointer);
// 3. Delete temporary storage
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(__is_full, 0,
oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng1_vector);
oneapi::dpl::__par_backend_hetero::__vector_walk<__base_t::__preferred_vector_size>{__n}(
__is_full, 0, oneapi::dpl::__internal::__lazy_ctor_storage_deleter{}, __rng1_vector);
}
template <typename _IsFull, typename _Idx>
void
Expand Down
2 changes: 1 addition & 1 deletion include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@
#include <iterator>
#include <type_traits>
#if _ONEDPL_CPP20_RANGES_PRESENT && _ONEDPL_CPP20_CONCEPTS_PRESENT
#include <ranges> // std::ranges::contiguous_range
# include <ranges> // std::ranges::contiguous_range
#endif

#include "../../utils_ranges.h"
Expand Down
5 changes: 3 additions & 2 deletions include/oneapi/dpl/pstl/utils.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -785,11 +785,12 @@ union __lazy_ctor_storage
}
};

// Utility to explicitly call the destructor of __lazy_ctor_storage as a callback functor
// Utility to explicitly call the destructor of __lazy_ctor_storage as a callback functor
struct __lazy_ctor_storage_deleter
{
template <typename _Tp>
void operator()(__lazy_ctor_storage<_Tp> __storage) const
void
operator()(__lazy_ctor_storage<_Tp> __storage) const
{
__storage.__destroy();
}
Expand Down

0 comments on commit 114924d

Please sign in to comment.