cuda GPUアクセラレーションコードに一貫性のない結果

Question

私はGPU上の画像に対してローカルバイナリパターンを計算しようとしていました。しかし、CPUとGPU上で同様のアルゴリズムを実行することによって得られる結果は、異なる結果を生み出しています。あなたは問題を理解するのを助けることができますか？

以下

は、私が実行しようとしていたコードの抜粋です：

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

は今、あなたはあなたが投稿元のカーネルコード内の紛れも明白な間違いを修正したことから、このコードを防止している2つの問題があります正しく動作します。

最初の、そして最も深刻なのは、カーネル内のメモリ競合です。このヒストグラムビンの更新：

histVec[mask]+=1 

は安全ではありません。複数のブロック内の複数のスレッドは、グローバルメモリ内の同じビンカウンタを同時に読み書きしようとします。このような状況では、CUDAは正確性または再現性を保証していません。

これまでの最も簡単な方法（ハードウェアによっては、パフォーマンスが最も高いとは限りません）は、アトミックメモリトランザクションを使用することです。これらはインクリメント操作がシリアル化されることを保証しますが、シリアル化はパフォーマンス上のペナルティを意味します。 CUDAは唯一あなたがhistVecの種類を確認する必要がありますので、32ビットおよび64ビットアトミックメモリトランザクションをサポートする互換性のある32ビットまたは64ビットであることを

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

注：あなたのようなものに更新コードを変更することでこれを行うことができます整数型。

これは、ビンカウンタベクトルをnumpy.uint8に定義したという2番目の問題につながります。つまり、メモリ競争がなかったとしても、カウントを格納するビット数は8ビットに過ぎず、意味のあるサイズのイメージではすぐにオーバーフローします。したがって、アトミックメモリトランザクションとの互換性とカウンタのロールオーバーを防止するためには、カウンタのタイプを変更する必要があります。

あなたが投稿したコードでこれらの変更を加えたところで、私はランダムな8ビット入力配列のGPUとホストコード計算ヒストグラムの間で正確な一致を得ることができました。

CUDAの基本的なヒストグラムの問題は非常によく説明されています。たとえば、hereのように、パフォーマンスを心配するときに勉強できるサンプルとコードベースがたくさんあります。