3 #if __OPENCL_VERSION__ < 110
4 #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
7 #ifndef cl_khr_byte_addressable_store
8 #error This kernel requires byte-addressable store
14 __global uchar* restrict text, /* output text */
15 __global const packel* restrict enc, /* encoded input */
16 uint textsize, uint numpacks)
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);
23 /* number of decoded symbols */
24 size_t numsyms = SYM_PER_PACKEL*get_local_size(0);
26 /* offset into the global memory array, i.e. the number
27 of symbols decoded by groups with id lower than ours
29 size_t goff = numsyms*get_group_id(0);
31 uchar b[SYM_PER_PACKEL];
32 unpack_els(enc[gix % numpacks], b);
34 for (size_t i = 0; i < SYM_PER_PACKEL; ++i) {
35 /* gix here is the write global index */
38 text[goff+lix+i] = alphabet[ b[i] ];