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