Numba cuda: почему сумма 1D-массива не верна? - PullRequest
Новая
       19

Numba cuda: почему сумма 1D-массива не верна?

0 голосов
/ 03 октября 2019

Я занимаюсь программированием нумба и куда. Я попытался суммировать с Cuda множество из них. Сумма не верна. Я думаю, что должно быть что-то с синхронизацией и сбором в конце данных правильно. 

 @cuda.jit
def my_kernel(const_array, res_array):

    sbuf = cuda.shared.array(512, float32)

    # Thread id in a 1D block
    tx = cuda.threadIdx.x
    # Block id in a 1D grid
    ty = cuda.blockIdx.x
    # Block width, i.e. number of threads per block
    bw = cuda.blockDim.x
    # Compute flattened index inside the array
    pos = tx + ty * bw

    sbuf[tx] = 0

    if pos < const_array.shape[0]:

        sbuf[tx] = const_array[pos] # do the computation

    cuda.syncthreads()
    if cuda.threadIdx.x == 0:
        for i in range(bw):
            res_array[0] += sbuf[i] 


    return


data_size = 10000000
res = numpy.zeros(1, dtype=numpy.float64)
const_array = numpy.ones(data_size, dtype=numpy.int8)

threadsperblock = 512
blockspergrid = math.ceil(data_size / threadsperblock)

my_kernel[blockspergrid, threadsperblock](const_array, res)

print(res)

Каждый раз, когда я запускаю этот код, он получает другое значение, например, 28160.0, но, конечно, это должно быть 10 м.

А намек?

Ответы [ 2 ]

1 голос
/ 04 октября 2019

проблема в том, что вы не суммируете весь набор блоков. У вас есть векторное измерение 10000000 и 512 потоков, что означает, что вам нужно сложить по всем блокам 19532 блоков. Это достигается в стандартном языке CUDA путем запуска нескольких ядер (в основном для старых устройств) или с помощью атомарных операций. В частности, ваша проблема в этой части вашего кода: 

if pos < const_array.shape[0]:
    sbuf[tx] = const_array[pos] # do the computation    cuda.syncthreads()
if cuda.threadIdx.x == 0:
    for i in range(bw):
        res_array[0] += sbuf[i]

В первых двух строках вы копируете данные из глобальной памяти в общую память массива sbuf. Но затем все потоки в разных блоках одновременно пытаются добавить свои локальные данные в глобальный адрес памяти по адресу res_array, это не является последовательным, и разные потоки могут просто считывать одни и те же данные дважды и давать неправильные результаты. Решение состоит в том, чтобы сначала выполнить частичную сумму в разделяемой памяти, а затем выполнить атомарную сумму, чтобы избежать асинхронных операций чтения-записи 

if cuda.threadIdx.x == 0:
    sum = 0
    for i in range(bw):
        sum  += sbuf[i] 
    cuda.atomic.add(res_array, 0, sum)

, которые должны решить вашу проблему.

С уважением.

0 голосов
/ 04 октября 2019

Во-первых, логика суммирования вообще не имеет смысла и очень неэффективна. Проблема заключается в том, что вы пытаетесь записать данные в одну ячейку памяти из разных потоков в разных блоках, и это приводит к состоянию гонки. Вы должны использовать cuda.atomic.add, чтобы избежать состояния гонки. Вы можете прочитать больше в руководстве по программированию CUDA: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomic-functions

...