@@ -2024,16 +2024,17 @@ namespace main_krn {
2024
2024
template <class KernelName > struct NDRangeAtomic64 ;
2025
2025
} // namespace main_krn
2026
2026
} // namespace reduction
2027
+
2027
2028
// Specialization for devices with the atomic64 aspect, which guarantees 64 (and
2028
2029
// temporarily 32) bit floating point support for atomic add.
2029
2030
// TODO 32 bit floating point atomics are eventually expected to be supported by
2030
2031
// the has_fast_atomics specialization. Corresponding changes to
2031
2032
// IsReduOptForAtomic64Add, as prescribed in its documentation, should then also
2032
2033
// be made.
2033
2034
template <typename KernelName, typename KernelType, int Dims, class Reduction >
2034
- void reduCGFuncImplAtomic64 (handler &CGH, KernelType KernelFunc,
2035
- const nd_range<Dims> &Range, Reduction &,
2036
- typename Reduction::rw_accessor_type Out) {
2035
+ void reduCGFuncAtomic64 (handler &CGH, KernelType KernelFunc,
2036
+ const nd_range<Dims> &Range, Reduction &Redu) {
2037
+ auto Out = Redu. getReadWriteAccessorToInitializedMem (CGH);
2037
2038
static_assert (Reduction::has_atomic_add_float64,
2038
2039
" Only suitable for reductions that have FP64 atomic add." );
2039
2040
constexpr size_t NElements = Reduction::num_elements;
@@ -2058,20 +2059,6 @@ void reduCGFuncImplAtomic64(handler &CGH, KernelType KernelFunc,
2058
2059
});
2059
2060
}
2060
2061
2061
- // Specialization for devices with the atomic64 aspect, which guarantees 64 (and
2062
- // temporarily 32) bit floating point support for atomic add.
2063
- // TODO 32 bit floating point atomics are eventually expected to be supported by
2064
- // the has_fast_atomics specialization. Corresponding changes to
2065
- // IsReduOptForAtomic64Add, as prescribed in its documentation, should then also
2066
- // be made.
2067
- template <typename KernelName, typename KernelType, int Dims, class Reduction >
2068
- void reduCGFuncAtomic64 (handler &CGH, KernelType KernelFunc,
2069
- const nd_range<Dims> &Range, Reduction &Redu) {
2070
- auto Out = Redu.getReadWriteAccessorToInitializedMem (CGH);
2071
- reduCGFuncImplAtomic64<KernelName, KernelType, Dims, Reduction>(
2072
- CGH, KernelFunc, Range, Redu, Out);
2073
- }
2074
-
2075
2062
template <typename ... Reductions, size_t ... Is>
2076
2063
void associateReduAccsWithHandler (handler &CGH,
2077
2064
std::tuple<Reductions...> &ReduTuple,
0 commit comments