Loads/Stores

The load and store functions enable developers to assert that all work-items in a sub-group read/write from/to contiguous locations in memory. Such operations can be mapped directly to SIMD operations.

Member function Description
template <typename T,
  access::address_space Space>
T load(const multi_ptr<T,Space> src) const

Loads contiguous data from src. Returns one element per work-item, corresponding to the memory location at src + get_local_id().

The value of src must be the same for all work-items in the sub-group.

template <int N, typename T,
  access::address_space Space>
vec<T,N> load(const multi_ptr<T,Space> src) const

Loads contiguous data from src. Returns N elements per work-item, corresponding to the N memory locations at src + i * get_max_local_range() + get_local_id() for i between 0 and N.

The value of src must be the same for all work-items in the sub-group.

template <typename T,
  access::address_space Space>
void store(multi_ptr<T,Space> dst, const T& x) const

Stores contiguous data to dst. The value of x from each work-item is written to the memory location at dst + get_local_id().

The value of dst must be the same for all work-items in the sub-group.

template <int N, typename T,
  access::address_space Space>
void store(
    multi_ptr<T,Space> dst,
    const vec<T,N>& x
) const

Stores contiguous data to dst. The N elements from each work-item are written to the memory locations at dst + i * get_max_local_range() + get_local_id() for i between 0 and N.

The value of dst must be the same for all work-items in the sub-group.