c – 用cuda创建共生矩阵

2019年8月4日 121次阅读 
//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 [(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>>>(...)

每个块上只有一个线程,每个块代表一个像素

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

谢谢 ：)

最佳答案 从不同输入的不同存储单元中读取不会产生您必须处理的并行危险.问题在于matrizSCM,其中同一个存储器单元可以一次增加多个线程.

atomicAdd(addr,1)是一个快速修复—它应该使算法正确,但它可能相当慢.做到正确应该是第一步;然后,您可以查看直方图计算和并行缩减算法网络上的可用示例,并检查它是否可以应用于您的问题.

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

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

点赞