libclc/generic/lib/async/async_work_group_strided_copy.inc

   1
   2 #define STRIDED_COPY(dst, src, num_gentypes, dst_stride, src_stride)       \
   3   size_t size = get_local_size(0) * get_local_size(1) * get_local_size(2); \
   4   size_t id = (get_local_size(1) * get_local_size(2) * get_local_id(0)) +  \
   5               (get_local_size(2) * get_local_id(1)) +                      \
   6               get_local_id(2);                                             \
   7   size_t i;                                                                \
   8                                                                            \
   9   for (i = id; i < num_gentypes; i += size) {                              \
  10     dst[i * dst_stride] = src[i * src_stride];                             \
  11   }
  12
  13
  14 _CLC_OVERLOAD _CLC_DEF event_t async_work_group_strided_copy(
  15     local __CLC_GENTYPE *dst,
  16     const global __CLC_GENTYPE *src,
  17     size_t num_gentypes,
  18     size_t src_stride,
  19     event_t event) {
  20
  21   STRIDED_COPY(dst, src, num_gentypes, 1, src_stride);
  22   return event;
  23 }
  24
  25 _CLC_OVERLOAD _CLC_DEF event_t async_work_group_strided_copy(
  26     global __CLC_GENTYPE *dst,
  27     const local __CLC_GENTYPE *src,
  28     size_t num_gentypes,
  29     size_t dst_stride,
  30     event_t event) {
  31
  32   STRIDED_COPY(dst, src, num_gentypes, dst_stride, 1);
  33   return event;
  34 }