#include "gmxsycl.h"
SYCL kernel helper functions.
- Author
- Andrey Alekseenko al42a.nosp@m.nd@g.nosp@m.mail..nosp@m.com
|
template<typename T , sycl_2020::memory_scope MemoryScope = sycl_2020::memory_scope::device> |
static void | atomicFetchAdd (T &val, const T delta) |
| Convenience wrapper to do atomic addition to a global buffer.
|
|
template<typename T , sycl_2020::memory_scope MemoryScope = sycl_2020::memory_scope::device> |
static T | atomicLoad (T &val) |
| Convenience wrapper to do atomic loads from a global buffer.
|
|
template<int Dim> |
static void | subGroupBarrier (const cl::sycl::nd_item< Dim > itemIdx) |
| Issue an intra sub-group barrier. More...
|
|
|
static constexpr unsigned int | c_cudaFullWarpMask = 0xffffffff |
| Full warp active thread mask used in CUDA warp-level primitives.
|
|
template<int Dim>
static void subGroupBarrier |
( |
const cl::sycl::nd_item< Dim > |
itemIdx | ) |
|
|
inlinestatic |
Issue an intra sub-group barrier.
Equivalent with CUDA's syncwarp(c_cudaFullWarpMask)
.