cuda GPU加速代码的结果不一致

Question

我试图在我的 GPU 上为图像计算本地二进制模式，同样使用 python 中的 cuda 模块。但是在 CPU 和 GPU 上执行类似算法所产生的结果却产生了不同的结果。你能帮我找出问题吗？

下面是我尝试执行的代码的 sn-p：

from __future__ import division
from skimage.io import imread, imshow
from numba import cuda
import time
import math
import numpy

# CUDA Kernel
@cuda.jit
def pointKernelLBP(imgGPU, histVec, pos) :
    ''' Computes Point Local Binary Pattern '''
    row, col = cuda.grid(2)
    if row+1 < imgGPU.shape[0] and col+1 < imgGPU.shape[1] and col-1>=0 and row-1>=0 :
        curPos = 0
        mask = 0
        for i in xrange(-1, 2) :
            for j in xrange(-1, 2) :
                if i==0 and j==0 :
                    continue
                if imgGPU[row+i][col+j] > imgGPU[row][col] :
                    mask |= (1<<curPos)     
                curPos+=1
        histVec[mask]+=1


#Host Code for computing LBP 
def pointLBP(x, y, img) :
    ''' Computes Local Binary Pattern around a point (x,y),
    considering 8 nearest neighbours '''
    pos = [0, 1, 2, 7, 3, 6, 5, 4]  
    curPos = 0
    mask = 0
    for i in xrange(-1, 2) :
        for j in xrange(-1, 2) :
            if i==0 and j==0 :
                continue
            if img[x+i][y+j] > img[x][y] :
                mask |= (1<<curPos)         
            curPos+=1
    return mask                 

def LBPHistogram(img, n, m) :
    ''' Computes LBP Histogram for given image '''
    HistVec = [0] * 256 
    for i in xrange(1, n-1) :
        for j in xrange(1, m-1) :
            HistVec[ pointLBP(i, j, img) ]+=1
    return HistVec

if __name__ == '__main__' :

    # Reading Image
    img = imread('cat.jpg', as_grey=True)
    n, m = img.shape

    start = time.time() 
    imgHist = LBPHistogram(img, n, m)
    print "Computation time incurred on CPU : %s seconds.\n" % (time.time() - start)    

    print "LBP Hisogram Vector Using CPU :\n"
    print imgHist

    print type(img)

    pos = numpy.ndarray( [0, 1, 2, 7, 3, 6, 5, 4] )

    img_global_mem = cuda.to_device(img)
    imgHist_global_mem = cuda.to_device(numpy.full(256, 0, numpy.uint8))
    pos_global_mem = cuda.to_device(pos)

    threadsperblock = (32, 32)
    blockspergrid_x = int(math.ceil(img.shape[0] / threadsperblock[0]))
    blockspergrid_y = int(math.ceil(img.shape[1] / threadsperblock[1]))
    blockspergrid = (blockspergrid_x, blockspergrid_y)

    start = time.time() 
    pointKernelLBP[blockspergrid, threadsperblock](img_global_mem, imgHist_global_mem, pos_global_mem)
    print "Computation time incurred on GPU : %s seconds.\n" % (time.time() - start)

    imgHist = imgHist_global_mem.copy_to_host()

    print "LBP Histogram as computed on GPU's : \n"
    print imgHist, len(imgHist)

Answer 1

现在您已经修复了您发布的原始内核代码中明显的错误，有两个问题阻止了此代码正常工作。

第一个也是最严重的，是内核中的内存竞争。本次直方图箱更新：

histVec[mask]+=1

不安全。多个块中的多个线程将尝试同时在全局内存中读取和写入相同的 bin 计数器。在这种情况下，CUDA 不保证正确性或可重复性。

对此的最简单（但不一定是最高性能，具体取决于您的硬件）解决方案是使用原子内存事务。这些确实保证了增量操作将被序列化，但当然，序列化意味着一些性能损失。您可以通过将更新代码更改为以下内容来做到这一点：

cuda.atomic.add(histVec,mask,1)

请注意，CUDA 仅支持 32 位和 64 位原子内存事务，因此您需要确保 histVec 的类型是兼容的 32 位或 64 位整数类型。

这导致了第二个问题，即您已将 bin 计数器向量定义为 numpy.uint8。这意味着即使您没有内存竞争，您也只有 8 位来存储计数，并且对于任何有意义大小的图像，它们会很快溢出。因此，为了与原子内存事务兼容并防止计数器翻转，您需要更改计数器的类型。

当我在您发布的代码中更改了这些内容（并修复了早期丢失代码的问题）时，我可以得到 GPU 和主机代码计算的随机 8 位输入数组的直方图之间的精确一致性。

对 CUDA 的底层并行直方图问题进行了很好的描述，当您担心性能时，可以研究大量示例和代码库，例如 here。