
#define __1D_WORK_GROUP
#define __1D_GRID
#include <opencl.h>

//__axiom(get_global_size(0) <= 1024);

__axiom(get_local_size(0) == 64);
__axiom(get_num_groups(0) == 16);



#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable

#if mykernel == 1
__kernel void read_kernel ( volatile __global uint4 *in,
                            volatile __global uint  *out,
                                     uint   ni,
                                     uint   val,
                                     uint   nk )
{
   if( nk == 0 ) return;
   
           uint pcount = 0;
   __local uint lcount;
           uint i, idx;

   if( get_local_id(0) == 0)
       lcount=0;

   barrier( CLK_LOCAL_MEM_FENCE );

   for(int n=0; n<nk; n++)
   {
      for( i=0, idx=get_global_id(0); i<ni; i++, idx+=get_global_size(0) )
      {
         if(in[idx].x == val) pcount++;
         if(in[idx].y == val) pcount++;
         if(in[idx].z == val) pcount++;
         if(in[idx].w == val) pcount++;
      }
	}
      
     (void) atomic_add( &lcount, pcount );

     barrier( CLK_LOCAL_MEM_FENCE );

     if( get_local_id(0) == 0 )
        out[get_group_id(0)] = lcount/nk;
}

#endif
#if mykernel == 2

__kernel void write_kernel ( volatile __global uint  *in,
                             volatile __global uint4 *out,
                                      uint  ni,
                                      uint  val,
                                      uint  nk )
{
   if( nk == 0 ) return;
   
   uint i, idx;
   uint4 pval = (uint4) (val, val, val, val);

   for(int n=0; n<nk; n++)
   {
      for( i=0, idx=get_global_id(0); i<ni; i++, idx+=get_global_size(0) )
	  for( i=0; i<ni; i++ )
      {
         out[idx] = pval;
      }
	}
}
#endif
