Skip to content

Commit d079d1c

Browse files
authored
[SYCLomatic] Fix the bug in the group utils help function
Signed-off-by: intwanghao <hao3.wang@intel.com>
1 parent ff19999 commit d079d1c

File tree

1 file changed

+11
-10
lines changed

1 file changed

+11
-10
lines changed

clang/runtime/dpct-rt/include/dpct/group_utils.hpp

Lines changed: 11 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -800,18 +800,18 @@ load_direct_blocked(const ItemT &item, InputIteratorT input_iter,
800800
/// work-group, guarded by range.
801801
///
802802
/// \tparam T The data type to load.
803+
/// \tparam DefaultT The type of default value to assign out-of-bound items.
803804
/// \tparam ElementsPerWorkItem The number of consecutive elements partitioned
804805
/// onto each work-item.
805806
/// \tparam InputIteratorT The random-access iterator type for input \iterator.
806-
/// \tparam DefaultT The type of default value to assign out-of-bound items.
807807
/// \tparam ItemT The sycl::nd_item index space class.
808808
/// \param item The calling work-item.
809809
/// \param input_iter The work-group's base input iterator for loading from.
810810
/// \param data Data to load.
811811
/// \param valid_items Number of valid items to load
812812
/// \param default_value Default value to assign out-of-bound items.
813-
template <typename T, int ElementsPerWorkItem, typename InputIteratorT,
814-
typename DefaultT, typename ItemT>
813+
template <typename T, typename DefaultT, int ElementsPerWorkItem,
814+
typename InputIteratorT, typename ItemT>
815815
__dpct_inline__ void
816816
load_direct_striped(const ItemT &item, InputIteratorT input_iter,
817817
T (&data)[ElementsPerWorkItem], int valid_items,
@@ -1241,23 +1241,24 @@ class group_load {
12411241
T (&data)[ElementsPerWorkItem], int valid_items,
12421242
DefaultT default_value) {
12431243
if constexpr (LoadAlgorithm == group_load_algorithm::blocked) {
1244-
load_direct_blocked<T, ElementsPerWorkItem, InputIteratorT, DefaultT,
1244+
load_direct_blocked<T, DefaultT, ElementsPerWorkItem, InputIteratorT,
12451245
ItemT>(item, input_iter, data, valid_items,
12461246
default_value);
12471247
} else if constexpr (LoadAlgorithm == group_load_algorithm::striped) {
1248-
load_direct_striped<T, ElementsPerWorkItem, InputIteratorT, DefaultT,
1248+
load_direct_striped<T, DefaultT, ElementsPerWorkItem, InputIteratorT,
12491249
ItemT>(item, input_iter, data, valid_items,
12501250
default_value);
12511251
} else if constexpr (LoadAlgorithm == group_load_algorithm::transpose) {
1252-
load_direct_striped<T, ElementsPerWorkItem, InputIteratorT, ItemT>(
1253-
item, input_iter, data, valid_items, default_value);
1252+
load_direct_striped<T, DefaultT, ElementsPerWorkItem, InputIteratorT,
1253+
ItemT>(item, input_iter, data, valid_items,
1254+
default_value);
12541255
dpct::group::exchange<T, ElementsPerWorkItem>(_local_memory)
12551256
.striped_to_blocked(item, data, data);
12561257
} else if constexpr (LoadAlgorithm ==
12571258
group_load_algorithm::sub_group_transpose) {
1258-
load_direct_sub_group_striped<T, ElementsPerWorkItem, InputIteratorT,
1259-
ItemT>(item, input_iter, data, valid_items,
1260-
default_value);
1259+
load_direct_sub_group_striped<T, DefaultT, ElementsPerWorkItem,
1260+
InputIteratorT, ItemT>(
1261+
item, input_iter, data, valid_items, default_value);
12611262
dpct::group::exchange<T, ElementsPerWorkItem>(_local_memory)
12621263
.sub_group_striped_to_blocked(item, data, data);
12631264
}

0 commit comments

Comments
 (0)