* \author Andrey Alekseenko <al42and@gmail.com>
*/
-/*! \brief Access mode to use for atomic accessors.
- *
- * Intel DPCPP compiler has \c sycl::atomic_ref, but has no \c sycl::atomic_fetch_add for floats.
- * However, \c atomic_ref can not be constructed from \c sycl::atomic, so we can not use
- * atomic accessors. Thus, we use \c mode::read_write accessors and \c atomic_ref.
- *
- * hipSYCL does not have \c sycl::atomic_ref, but has \c sycl::atomic_fetch_add for floats, which
- * requires using atomic accessors. Thus, we use \c mode::atomic accessors.
- *
- * The \ref atomicFetchAdd function could be used for doing operations on such accessors.
- */
-static constexpr auto mode_atomic = GMX_SYCL_DPCPP ? cl::sycl::access::mode::read_write :
- /* GMX_SYCL_HIPSYCL */ cl::sycl::access::mode::atomic;
+//! \brief Full warp active thread mask used in CUDA warp-level primitives.
+static constexpr unsigned int c_cudaFullWarpMask = 0xffffffff;
/*! \brief Convenience wrapper to do atomic addition to a global buffer.
+ */
+template<typename T, sycl_2020::memory_scope MemoryScope = sycl_2020::memory_scope::device>
+static inline void atomicFetchAdd(T& val, const T delta)
+{
+ sycl_2020::atomic_ref<T, sycl_2020::memory_order::relaxed, MemoryScope, cl::sycl::access::address_space::global_space> ref(
+ val);
+ ref.fetch_add(delta);
+}
+
+/*! \brief Convenience wrapper to do atomic loads from a global buffer.
+ */
+template<typename T, sycl_2020::memory_scope MemoryScope = sycl_2020::memory_scope::device>
+static inline T atomicLoad(T& val)
+{
+ sycl_2020::atomic_ref<T, sycl_2020::memory_order::relaxed, MemoryScope, cl::sycl::access::address_space::global_space> ref(
+ val);
+ return ref.load();
+}
+
+
+/*! \brief Issue an intra sub-group barrier.
*
- * Skips sub-normal values.
+ * Equivalent with CUDA's \c syncwarp(c_cudaFullWarpMask).
*
- * The implementation differences between DPCPP and hipSYCL are explained in \ref mode_atomic.
*/
-template<class IndexType>
-static inline void atomicFetchAdd(DeviceAccessor<float, mode_atomic> acc, const IndexType idx, const float val)
+template<int Dim>
+static inline void subGroupBarrier(const cl::sycl::nd_item<Dim> itemIdx)
{
-#if GMX_SYCL_DPCPP
- if (cl::sycl::isnormal(val))
- {
- sycl_2020::atomic_ref<float, sycl_2020::memory_order::relaxed, sycl_2020::memory_scope::device, cl::sycl::access::address_space::global_space>
- fout_atomic(acc[idx]);
- fout_atomic.fetch_add(val);
- }
-#elif GMX_SYCL_HIPSYCL
- if (std::isnormal(val)) // No sycl::isnormal in hipSYCL
- {
-# ifdef SYCL_DEVICE_ONLY
- /* While there is support for float atomics on device, the host implementation uses
- * Clang's __atomic_fetch_add intrinsic, that, at least in Clang 11, does not support
- * floats. Luckily, we don't want to run on host. */
- acc[idx].fetch_add(val);
-# else
- GMX_UNUSED_VALUE(acc);
- GMX_UNUSED_VALUE(idx);
-# endif
- }
+#if GMX_SYCL_HIPSYCL
+ cl::sycl::group_barrier(itemIdx.get_sub_group(), cl::sycl::memory_scope::sub_group);
+#else
+ itemIdx.get_sub_group().barrier();
#endif
}
namespace sycl_2020
{
#if GMX_SYCL_HIPSYCL
-__device__ static inline float shift_left(sycl_2020::sub_group, float var, sycl_2020::sub_group::linear_id_type delta)
+__device__ __host__ static inline float shift_left(sycl_2020::sub_group,
+ float var,
+ sycl_2020::sub_group::linear_id_type delta)
{
// No sycl::sub_group::shift_left / shuffle_down in hipSYCL yet
# ifdef SYCL_DEVICE_ONLY
# if defined(HIPSYCL_PLATFORM_CUDA) && defined(__HIPSYCL_ENABLE_CUDA_TARGET__)
- static const unsigned int sc_cudaFullWarpMask = 0xffffffff;
- return __shfl_down_sync(sc_cudaFullWarpMask, var, delta);
+ return __shfl_down_sync(c_cudaFullWarpMask, var, delta);
# elif defined(HIPSYCL_PLATFORM_ROCM) && defined(__HIPSYCL_ENABLE_HIP_TARGET__)
// Do we need more ifdefs? https://github.com/ROCm-Developer-Tools/HIP/issues/1491
return __shfl_down(var, delta);
return NAN;
# endif
}
-__host__ static inline float shift_left(sycl_2020::sub_group, float, sycl_2020::sub_group::linear_id_type)
-{
- // Should never be called
- assert(false);
- return NAN;
-}
#elif GMX_SYCL_DPCPP
static inline float shift_left(sycl_2020::sub_group sg, float var, sycl_2020::sub_group::linear_id_type delta)
{
#endif
#if GMX_SYCL_HIPSYCL
-__device__ static inline float shift_right(sycl_2020::sub_group,
- float var,
- sycl_2020::sub_group::linear_id_type delta)
+__device__ __host__ static inline float shift_right(sycl_2020::sub_group,
+ float var,
+ sycl_2020::sub_group::linear_id_type delta)
{
// No sycl::sub_group::shift_right / shuffle_up in hipSYCL yet
# ifdef SYCL_DEVICE_ONLY
# if defined(HIPSYCL_PLATFORM_CUDA) && defined(__HIPSYCL_ENABLE_CUDA_TARGET__)
- static const unsigned int sc_cudaFullWarpMask = 0xffffffff;
- return __shfl_up_sync(sc_cudaFullWarpMask, var, delta);
+ return __shfl_up_sync(c_cudaFullWarpMask, var, delta);
# elif defined(HIPSYCL_PLATFORM_ROCM) && defined(__HIPSYCL_ENABLE_HIP_TARGET__)
// Do we need more ifdefs? https://github.com/ROCm-Developer-Tools/HIP/issues/1491
return __shfl_up(var, delta);
return NAN;
# endif
}
-__host__ static inline float shift_right(sycl_2020::sub_group, float, sycl_2020::sub_group::linear_id_type)
+#elif GMX_SYCL_DPCPP
+static inline float shift_right(sycl_2020::sub_group sg, float var, sycl_2020::sub_group::linear_id_type delta)
{
+ return sg.shuffle_up(var, delta);
+}
+#endif
+
+#if GMX_SYCL_HIPSYCL
+/*! \brief Polyfill for sycl::isfinite missing from hipSYCL
+ *
+ * Does not follow GROMACS style because it should follow the name for
+ * which it is a polyfill. */
+template<typename Real>
+__device__ __host__ static inline bool isfinite(Real value)
+{
+ // This is not yet implemented in hipSYCL pending
+ // https://github.com/illuhad/hipSYCL/issues/636
+# ifdef SYCL_DEVICE_ONLY
+# if defined(HIPSYCL_PLATFORM_CUDA) && defined(__HIPSYCL_ENABLE_CUDA_TARGET__)
+ return isfinite(value);
+# elif defined(HIPSYCL_PLATFORM_ROCM) && defined(__HIPSYCL_ENABLE_HIP_TARGET__)
+ return isfinite(value);
+# else
+# error "Unsupported hipSYCL target"
+# endif
+# else
// Should never be called
assert(false);
- return NAN;
+ GMX_UNUSED_VALUE(value);
+ return false;
+# endif
}
#elif GMX_SYCL_DPCPP
-static inline float shift_right(sycl_2020::sub_group sg, float var, sycl_2020::sub_group::linear_id_type delta)
+template<typename Real>
+static inline bool isfinite(Real value)
{
- return sg.shuffle_up(var, delta);
+ return cl::sycl::isfinite(value);
}
+
#endif
+
+#if GMX_SYCL_HIPSYCL
+
+/*! \brief Polyfill for sycl::vec::load buggy in hipSYCL
+ *
+ * Loads from the address \c ptr offset in elements of type T by
+ * NumElements * offset, into the components of \c v.
+ *
+ * Can probably be removed when
+ * https://github.com/illuhad/hipSYCL/issues/647 is resolved. */
+template<cl::sycl::access::address_space AddressSpace, typename T, int NumElements>
+static inline void loadToVec(size_t offset,
+ cl::sycl::multi_ptr<const T, AddressSpace> ptr,
+ cl::sycl::vec<T, NumElements>* v)
+{
+ for (int i = 0; i < NumElements; ++i)
+ {
+ (*v)[i] = ptr.get()[offset * NumElements + i];
+ }
+}
+
+/*! \brief Polyfill for sycl::vec::store buggy in hipSYCL
+ *
+ * Loads from the address \c ptr offset in elements of type T by
+ * NumElements * offset, into the components of \c v.
+ *
+ * Can probably be removed when
+ * https://github.com/illuhad/hipSYCL/issues/647 is resolved. */
+template<cl::sycl::access::address_space AddressSpace, typename T, int NumElements>
+static inline void storeFromVec(const cl::sycl::vec<T, NumElements>& v,
+ size_t offset,
+ cl::sycl::multi_ptr<T, AddressSpace> ptr)
+{
+ for (int i = 0; i < NumElements; ++i)
+ {
+ ptr.get()[offset * NumElements + i] = v[i];
+ }
+}
+
+#elif GMX_SYCL_DPCPP
+
+/*! \brief Polyfill for sycl::vec::load buggy in hipSYCL
+ *
+ * Loads from the address \c ptr offset in elements of type T by
+ * NumElements * offset, into the components of \c v.
+ *
+ * Can probably be removed when
+ * https://github.com/illuhad/hipSYCL/issues/647 is resolved. */
+template<cl::sycl::access::address_space AddressSpace, typename T, int NumElements>
+static inline void loadToVec(size_t offset,
+ cl::sycl::multi_ptr<const T, AddressSpace> ptr,
+ cl::sycl::vec<T, NumElements>* v)
+{
+ v->load(offset, ptr);
+}
+
+/*! \brief Polyfill for sycl::vec::store buggy in hipSYCL
+ *
+ * Loads from the address \c ptr offset in elements of type T by
+ * NumElements * offset, into the components of \c v.
+ *
+ * Can probably be removed when
+ * https://github.com/illuhad/hipSYCL/issues/647 is resolved. */
+template<cl::sycl::access::address_space AddressSpace, typename T, int NumElements>
+static inline void storeFromVec(const cl::sycl::vec<T, NumElements>& v,
+ size_t offset,
+ cl::sycl::multi_ptr<T, AddressSpace> ptr)
+{
+ v.store(offset, ptr);
+}
+
+#endif
+
} // namespace sycl_2020
#endif /* GMX_GPU_UTILS_SYCL_KERNEL_UTILS_H */
biod.pnpi.spb.ru / alexxy / gromacs.git / blobdiff