cuda直方图计算——利用shared memory统计直方图

最新推荐文章于 2025-09-14 13:23:20 发布

原创

最新推荐文章于 2025-09-14 13:23:20 发布 · 2.5k 阅读

2 ·

CC 4.0 BY-SA版权

本文探讨了在CUDA中计算直方图时遇到的问题，初始版本由于每个线程独立加法导致结果是CPU的32倍。解决方法是通过利用共享内存进行优化。同时指出，对于计算能力1.0的设备，不支持atomic操作。要设置计算能力，需在项目属性中将CUDA C/C++ Device的Code Generation设为compute_20,sm_20。" 114966540,8764207,SQL学习：SELECT INTO创建新表及实例解析,"['SQL', '数据库', '数据操作']

先写个错误的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,bins,sizeof(int)*3,cudaMemcpyHostToDevice);
    int blocksize