__global__  void  histo_kernel(unsigned char* buffer, long size, unsigned int* histo){
    __shared__ unsigned int temp[256];
    tmp[threadIdx.x] = 0;
    __syncThreads();

    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int offset = blockDim.x * gridDim.x;
    while(i<size){
    
        atomicAdd( &temp[buffer[i]], 1);
        i += offset;
    }
    __syncthreads();
    atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x]);
}