Confusion about task vs NDRange implementaions

[Altera_Forum]

I'm attempting to implement a simple barrel shifter. I'm able to get it to work using a task based approach but am unable to do so using an NDRange based kernel. I've posted both kernels below. I'd appreciate some insight into how the implementation of the NDRange approach differs from the task. 

 __kernel  __attribute__((reqd_work_group_size(360,1,1))) __attribute__((max_work_group_size(360))) void barrelShifter_nd(  __global unsigned char * restrict dataIn,  const uint shift,  __global unsigned char * restrict dataOut ) {  char __attribute((register)) din[720];    uint gid = get_global_id(0);    din[gid] = dataIn[gid] & 0x1;  din[360+gid] = dataIn[gid] & 0x1;      dataOut[gid] = din[shift+gid];   }      __kernel  __attribute__((task))  void barrelShifter_task(  __global unsigned char * restrict dataIn,  const uint shift,  __global unsigned char * restrict dataOut ) {  char __attribute((register)) din[720];    for(uint i=0; i<360; i++)  {  din = datain & 0x1;  din[360+i] = dataIn & 0x1;  }    # pragma unroll  for(uint i=0; i<360; i++)  {  dataout = din[shift+i];  }   }

[Altera_Forum]

It is not possible to infer shift registers of any type in NDRange kernels. As per OpenCL's specification, there is no guarantee of thread ordering in this kernel type, and local memory consistency is only guaranteed at barriers. It is not possible to describe a shift register in NDRange kernels, while conforming to both of these conditions.