c++ - 使用cuda创建共现矩阵

Question

//This is my kernel function

__global__ void createSCM(Pixel*pixelMat, //image
                          int imgRows, //image dimensions
                          int imgCols,
                          int*matrizSCM, //Coocurrence matrix
                          int numNiveles, //coocurrence matrix levels = 256
                          int delta_R, //value = {-1,0 or 1}
                          int delta_C) //value = {-1,0 or 1}
{
    int i = blockIdx.y*blockDim.y+threadIdx.y;
    int j = blockIdx.x*blockDim.x+threadIdx.x;

    int cols = numNiveles;

    int posx,posy;

    if ( (j + delta_C) < imgCols && (i + delta_R) < imgRows && 
       ((j + delta_C) >= 0) && ((i + delta_R) >= 0) )
    {
       posx = pixelMat[i*imgCols+j].channel_0;
       posy = pixelMat[(i + delta_R)*imgCols+(j + delta_C)].channel_0;

       matrizSCM[posx*cols+posy]++;
       matrizSCM[posy*cols+posx]++;
    }

}

struct Pixel {

    int channel_0;
};

我在共生矩阵中有计数错误，因为

pixelMat[i*imgCols+j] 和 pixelMat[(i + delta_R)*imgCols+(j + delta_C)]

正在访问同一个线程的不同位置。

这是我的内核调用

int Grid_Dim_x=imagenTest.rows, Grid_Dim_y=imagenTest.cols;
int Block_Dim_x=1, Block_Dim_y=1;   

dim3 Grid(Grid_Dim_x, Grid_Dim_y);  
dim3 Block(Block_Dim_x,Block_Dim_x);

createSCM<<<Grid,Block>>>(...)

每个 block 上只有一个线程，每个 block 代表一个像素

这个问题有好的解决方案吗？

谢谢:)

Answer 1

从不可变输入的不同存储单元读取不会导致您必须处理的并行危险。问题出在 matrizSCM 中，其中同一个内存单元可以由多个线程同时递增。

atomicAdd(addr,1) 是一个快速修复 --- 它应该使算法正确，但它可能相当慢。把它改正应该是第一步；然后你可以在网上查看直方图计算和并行缩减算法的可用示例，并检查它是否可以应用于你的问题。

最后，正如 Robert 在评论中指出的那样，在一个 block 中只启动一个线程是非常低效的。您需要 32 的倍数才能利用硬件 SIMD 单元，并且通常需要大约 256 个线程来隐藏各种内存延迟。

此外，如果您的图像很大并且您仍然需要数千个 256 线程 block ，您可以考虑启动更少的 block (大约 60-120 个)但让每个 block 按顺序处理多个像素。如果这样做，您可能能够将 matrixSCM 的拷贝放入共享内存中。这将为每个 block 创建一个单独的 matrixSCM 拷贝，从而减少 block 之间的原子冲突。显然，在内核的末尾，您的 block 仍然需要将部分结果“提交”到全局结果中，但那将是单步操作。