#if __OPENCL_VERSION__ <= CL_VERSION_1_0
    #define atomic_add(p, v) atom_add(p, v)
    #define atomic_inc(p) atom_inc(p)
#endif

kernel void double(constant int input[], global int output[])
{
    const int i = get_global_id(0);
    output[i] = input[i] * 2;
}

#define CAPACITY(store) store[0]
#define NTH_LIST(store, index) (store + 1 + index * (CAPACITY(store) + 1))
#define LIST_SIZE(list) list[0]
#define LIST_GET(list, index) list[1 + index]
#define LIST_SET(list, index, val) list[1 + index] = val
#define LIST_APPEND(list, val) LIST_SET(list, atomic_inc(list), val)

kernel void fill(global int * store, int size)
{
    global int * list = NTH_LIST(store, 1);
    LIST_APPEND(list, get_global_id(0));
}