@@ -565,7 +565,7 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
565
565
using rw_accessor_type = accessor<T, accessor_dim, access::mode::read_write,
566
566
access::target::device, is_placeholder,
567
567
ext::oneapi::accessor_property_list<>>;
568
- static constexpr bool has_atomic_add_float64 =
568
+ static constexpr bool has_float64_atomics =
569
569
IsReduOptForAtomic64Op<T, BinaryOperation>::value;
570
570
static constexpr bool has_fast_atomics =
571
571
IsReduOptForFastAtomicFetch<T, BinaryOperation>::value;
@@ -645,7 +645,7 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
645
645
// / require initialization with identity value, then return user's read-write
646
646
// / accessor. Otherwise, create global buffer with 'num_elements' initialized
647
647
// / with identity value and return an accessor to that buffer.
648
- template <bool HasFastAtomics = (has_fast_atomics || has_atomic_add_float64 )>
648
+ template <bool HasFastAtomics = (has_fast_atomics || has_float64_atomics )>
649
649
std::enable_if_t <HasFastAtomics, rw_accessor_type>
650
650
getReadWriteAccessorToInitializedMem (handler &CGH) {
651
651
if constexpr (is_rw_acc) {
@@ -2046,8 +2046,9 @@ template <typename KernelName, typename KernelType, int Dims, class Reduction>
2046
2046
void reduCGFuncAtomic64 (handler &CGH, KernelType KernelFunc,
2047
2047
const nd_range<Dims> &Range, Reduction &Redu) {
2048
2048
auto Out = Redu.getReadWriteAccessorToInitializedMem (CGH);
2049
- static_assert (Reduction::has_atomic_add_float64,
2050
- " Only suitable for reductions that have FP64 atomic operations." );
2049
+ static_assert (
2050
+ Reduction::has_float64_atomics,
2051
+ " Only suitable for reductions that have FP64 atomic operations." );
2051
2052
constexpr size_t NElements = Reduction::num_elements;
2052
2053
using Name =
2053
2054
__sycl_reduction_kernel<reduction::main_krn::NDRangeAtomic64, KernelName>;
0 commit comments