4 
//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>>>(...)

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

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

谢谢 ：）

4

1 回答 1

2

从不可变输入的不同存储单元中读取不会产生您必须处理的并行风险。问题在于matrizSCM同一内存单元可以同时被多个线程递增。

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

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

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

于 2012-11-27T09:16:45.890 回答