git.oblomov.eu Git - babel/blob - babel.cl

   1 #include "babel_cl.h"
   2
   3 #if __OPENCL_VERSION__ < 110
   4 #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
   5 #endif
   6
   7 #ifndef cl_khr_byte_addressable_store
   8 #error This kernel requires byte-addressable store
   9 #endif
  10
  11 __kernel
  12 void
  13 decode_string(
  14         __global uchar* restrict text, /* output text */
  15         __global const packel* restrict enc, /* encoded input */
  16         uint textsize, uint numpacks)
  17 {
  18         /* each work-item decodes a packel */
  19         size_t gix = get_global_id(0);
  20         /* into three uchars */
  21         size_t lix = SYM_PER_PACKEL*get_local_id(0);
  22
  23         /* number of decoded symbols */
  24         size_t numsyms = SYM_PER_PACKEL*get_local_size(0);
  25
  26         /* offset into the global memory array, i.e. the number
  27            of symbols decoded by groups with id lower than ours
  28          */
  29         size_t goff = numsyms*get_group_id(0);
  30
  31         uchar b[SYM_PER_PACKEL];
  32         unpack_els(enc[gix % numpacks], b);
  33
  34         for (size_t i = 0; i < SYM_PER_PACKEL; ++i) {
  35                 /* gix here is the write global index */
  36                 gix = goff + lix + i;
  37                 if (gix < textsize)
  38                         text[goff+lix+i] = alphabet[ b[i] ];
  39         }
  40 }