先写个错误的version,折腾了我一天呀
__global__ void CalHistKernel(int*imgData,int*bins,int datasize)
{
__shared__ int _bins[3];
int tx=threadIdx.x;
int idx=blockIdx.x*blockDim.x+threadIdx.x;
if(tx<3)
{
_bins[tx]=0;
}
__syncthreads();
if(idx<datasize)
{
atomicAdd((int*)&_bins[imgData[idx]],1);
}
__syncthreads();
for(int i=0;i<3;i++)
atomicAdd((int*)&bins[i],_bins[i]);
}
void CalHist()
{
const int datasize=96;
int imgData[datasize];
for(int i=0;i<datasize;i++)
{
imgData[i]=rand()%3;
}
int bins[3]={0};
int*d_imgData;
int*d_bins;
cudaMalloc((void**)&d_imgData,sizeof(int)*datasize);
cudaMalloc((void**)&d_bins,sizeof(int)*3);
cudaMemcpy(d_imgData,imgData,sizeof(int)*datasize,cudaMemcpyHostToDevice);
cudaMemcpy(d_bins,bi