#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));
  }