-1

同じためにPythonでcudaモジュールを利用して、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)
4

1 に答える 1

0

投稿した元のカーネル コードの明らかな間違いを修正したので、このコードが正しく動作するのを妨げている 2 つの問題があります。

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

histVec[mask]+=1

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

これに対する最も簡単な (ただし、ハードウェアによっては必ずしも最もパフォーマンスが高いとは限りません) ソリューションは、アトミック メモリ トランザクションを使用することです。これらは、インクリメント操作がシリアル化されることを保証しますが、もちろん、シリアル化はパフォーマンスの低下を意味します。これを行うには、更新コードを次のように変更します。

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

histVecCUDA は 32 ビットと 64 ビットのアトミック メモリ トランザクションのみをサポートするため、 の型が互換性のある 32 ビットまたは 64 ビットの整数型であることを確認する必要があることに注意してください。

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

あなたが投稿したコードでこれらのことを変更したとき (そして初期の欠落コードの問題を修正したとき)、ランダムな 8 ビット入力配列に対して GPU とホスト コードで計算されたヒストグラムの間で正確な一致を得ることができました。

根底にある並列ヒストグラムの問題は CUDA について非常によく説明されており、パフォーマンスについて心配するようになったときに学習できる多くの例とコードベースがあります。たとえば、こちら.

于 2017-11-15T15:19:14.140 に答える