Непоследовательные результаты в ускоренном коде cuda GPU

Question

Я пытался вычислить локальные двоичные шаблоны для изображения на моем GPU, используя для этого модуль cuda в python. Но результаты, полученные при выполнении аналогичного алгоритма на 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.atomic.add(histVec,mask,1)

Обратите внимание, что CUDA поддерживает только 32- и 64-битные атомарные транзакции памяти, поэтому вам нужно будет указать тип histVec совместимый 32- или 64-битный целочисленный тип.

Это приводит ко второй проблеме, которая заключается в том, что вы определили вектор счетчика бина как numpy.uint8, Это означает, что даже если бы у вас не было гонки памяти, у вас было бы только 8 бит для хранения счетчиков, и они быстро переполнились бы для изображений любого значимого размера. Таким образом, для совместимости с атомарными транзакциями памяти и для предотвращения опрокидывания счетчика вам необходимо изменить тип ваших счетчиков.

Когда я изменил эти вещи в коде, который вы опубликовали (и исправил проблему с отсутствующим ранее кодом), я смог получить точное согласие между GPU и вычисленными гистограммами кода хоста для случайного 8-битного входного массива.

Основная проблема параллельной гистограммы очень хорошо описана для CUDA, и есть много примеров и кодовых баз, которые вы можете изучить, когда приступите к проблеме производительности, например, здесь.