cuda GPU 가속 코드의 결과가 일치하지 않습니다.

Question

저는 GPU에서 이미지 용 로컬 바이너리 패턴을 계산하려고했는데, 파이썬에서 동일한 방식으로 cuda 모듈을 사용했습니다. 그러나 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

을 이제 당신은 당신이 게시 원래 커널 코드에서 눈부시게 명백한 실수를 해결 한 것을에서이 코드를 방해하는 두 가지 문제가있다 올바르게 작동합니다.

가장 중요한 것은 커널에서 메모리 경쟁입니다. 이 히스토그램 빈의 업데이트 :

histVec[mask]+=1 

은 안전하지 않습니다. 여러 블록의 다중 스레드는 전역 메모리에서 동일한 저장소 카운터를 동시에 읽고 쓰려고 시도합니다. CUDA는 이러한 상황에서 정확성 또는 반복성을 보장하지 않습니다.

가장 간단한 방법 (하드웨어에 따라 가장 중요한 것은 아니지만)은 원자 메모리 트랜잭션을 사용하는 것입니다. 이것들은 인크 리먼 트 조작이 직렬화되는 것을 보증 합니다만, 직렬화는 약간의 퍼포먼스 패널티를 의미합니다. CUDA는 당신이 histVec의 종류를 확인해야합니다, 그래서 32 비트 및 64 비트 원자 메모리 트랜잭션을 지원하는 호환 32 또는 64 비트입니다

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

참고 : 당신은 같은에 업데이트 코드를 변경하여이 작업을 수행 할 수 있습니다 정수 유형.

두 번째 문제는 빈 카운터 벡터를 numpy.uint8으로 정의한 것입니다. 즉, 메모리 경쟁이 없다고해도 카운트를 저장하는 데 8 비트 밖에 없으며 의미있는 크기의 이미지로 인해 오버플로가 발생합니다. 따라서 원자 메모리 트랜잭션과의 호환성과 카운터 롤오버 방지를 위해 카운터 유형을 변경해야합니다.

내가 게시 한 코드에서 이러한 사항을 변경하고 초기 누락 된 코드 문제를 해결했을 때 임의의 8 비트 입력 배열에 대해 GPU와 호스트 코드 계산 막대 그래프 사이에 정확히 일치 할 수있었습니다.

기본 병렬 히스토그램 문제는 CUDA에 대해 매우 잘 설명되어 있습니다. 예를 들어 here과 같이 성능에 대해 걱정할 필요가있을 때 많은 예제와 코드베이스를 공부할 수 있습니다.