//统计
   for (int i=0; i<SIZE; i++)
       histo[buffer[i]]++;

//声明变量
 	unsigned int *dev_histo;
    HANDLE_ERROR( cudaMalloc( (void**)&dev_histo,
                              256 * sizeof( int ) ) );
    //代码块内的变量一定要手动初始化
	HANDLE_ERROR( cudaMemset( dev_histo, 0,
                              256 * sizeof( int ) ) );
...............
__global__ void histo_kernel( unsigned char *buffer,
                              long size,
                              unsigned int *histo ) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    while (i < size) {
        atomicAdd( &histo[buffer[i]], 1 );
        i += stride;
    }
}
..............
    
    // kernel launch - 2x the number of mps gave best timing
    cudaDeviceProp  prop;
    HANDLE_ERROR( cudaGetDeviceProperties( &prop, 0 ) );
    int blocks = prop.multiProcessorCount;
    histo_kernel<<<blocks*2,256>>>( dev_buffer, 
                                    SIZE, 
                                   dev_histo );

#define SIZE    (100*1024*1024)


__global__ void histo_kernel( unsigned char *buffer,
                              long size,
                              unsigned int *histo ) {

    __shared__  unsigned int temp[256];//声明一个共享缓冲区
    temp[threadIdx.x] = 0;             //将清除内存，每个线程写一次，由于我们在核函数设置启动线程中
    								   //为每个block分配了256个线程，所以很容易清除累计缓冲区temp。
    __syncthreads();

    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;//因为线程数没有数据多所以要设定步长（步长为分配的线程数目）
    while (i < size) {
        atomicAdd( &temp[buffer[i]], 1 );
        i += stride;
    }

    __syncthreads();
    atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x] );
}

·············
    cudaDeviceProp  prop;
    HANDLE_ERROR( cudaGetDeviceProperties( &prop, 0 ) );
    int blocks = prop.multiProcessorCount;
    histo_kernel<<<blocks*2,256>>>( dev_buffer,
                                    SIZE, 
                                   dev_histo );