CUDA计算直方图，采用线程一对一映射实现，怎么定义线程块和线程

[dreamcraneleaf]

__global__ void myhistKernel(const unsigned int const * d_hist_data,unsigned int * const d_bin_data)
{
__shared__ unsigned int d_bin_data_shared[256];
const unsigned int idx=(blockIdx.x * (blockDim.x*N)) + threadIdx.x;
const unsigned int idy=(blockIdx.y * blockDim.y) + threadIdx.y;
const unsigned int tid=idx+idy * (blockDim.x*N) * (gridDim.x);
d_bin_data_shared[threadIdx.x]=0;
__syncthreads();
const unsigned int value_u32 = d_hist_data[tid];

atomicAdd(&(d_bin_data_shared[ ((value_u32 & 0x000000FF)) ]),1);
atomicAdd(&(d_bin_data_shared[ ((value_u32 & 0x0000FF00) >>8 ) ]),1);
atomicAdd(&(d_bin_data_shared[ ((value_u32 & 0x00FF0000) >>16 ) ]),1);
atomicAdd(&(d_bin_data_shared[ ((value_u32 & 0xFF000000) >>24 ) ]),1);

__syncthreads();
atomicAdd(&(d_bin_data[threadIdx.x]),d_bin_data_shared[threadIdx.x]);
}