Почему мое ядро ​​CUDA (Numba) по-разному ведет себя при последовательных вызовах с одинаковыми входными данными? - PullRequest
Новая
       153

Почему мое ядро ​​CUDA (Numba) по-разному ведет себя при последовательных вызовах с одинаковыми входными данными?

0 голосов
/ 06 августа 2020

У меня ошибка noob в numba + cuda в python. Версия Numba - 0.51, версия CUDA - 10.2. Приведенный ниже код дает очень разные результаты при многократном вызове с одними и теми же входами. 

import numpy as np
from numba import cuda, jit

@cuda.jit()
def writeToArray(vec, array_in, array_out):
    ''' vec is a 3x1 vector, array_in is a 3D array, array_out is a 3D array of the shape of array in'''
    i,j,k = cuda.grid(3)
    value = array_in[i,j,k] * vec[0] + array_in[i,j,k] * vec[0] + array_in[i,j,k] * vec[0]
    cuda.atomic.max(array_out,(i,j,k), value)
    # cuda.synchronize()

def test():
    
    threadsperblock = (8,8,8)
    blockspergrid_x = ( 17 + threadsperblock[0]) // threadsperblock[0]
    blockspergrid_y = ( 21 + threadsperblock[1]) // threadsperblock[1]
    blockspergrid_z = ( 5 + threadsperblock[2]) // threadsperblock[2]
    blockspergrid = (blockspergrid_x, blockspergrid_y, blockspergrid_z)
    array_in = np.random.rand(17,21,5).astype(np.float_)
    vec = np.array([1.0, -1.0, 1.0]).astype(np.float_)
    d_array_in = cuda.to_device(array_in)
    d_vec = cuda.to_device(vec)
    while True:
        array_out_1 = -999.999*np.ones_like(array_in)
        array_out_2 = -999.999*np.ones_like(array_in)
        d_array_out_1= cuda.to_device(array_out_2)
        d_array_out_2 = cuda.to_device(array_out_2)
        writeToArray[blockspergrid, threadsperblock](d_vec, d_array_in, d_array_out_1)
        writeToArray[blockspergrid, threadsperblock](d_vec, d_array_in, d_array_out_2)
        array_out_1_host = d_array_out_1.copy_to_host()
        array_out_2_host = d_array_out_2.copy_to_host()
        assert(np.allclose(array_out_1_host, array_out_2_host))

if __name__ == "__main__":
    test()

Это не должно нарушаться, но в конечном итоге, после примерно 10 итераций while l oop утверждение не выполняется. Что я делаю не так?

1 Ответ

1 голос
/ 06 августа 2020

Код вашего ядра делает незаконный доступ за пределами допустимого диапазона. Когда вы изменяете размер сетки таким образом: 

blockspergrid_x = ( 17 + threadsperblock[0]) // threadsperblock[0]
blockspergrid_y = ( 21 + threadsperblock[1]) // threadsperblock[1]
blockspergrid_z = ( 5 + threadsperblock[2]) // threadsperblock[2]

, вы запускаете возможность создания «дополнительных» потоков. Это потоки, индексы i, j, k которых находятся за пределами «формы» вашего входного массива. Вы не хотите, чтобы эти потоки действовали. Обычный подход - включить «проверку потока» в код ядра: 

$ cat t31.py
import numpy as np
from numba import cuda, jit

@cuda.jit()
def writeToArray(vec, array_in, array_out):
    ''' vec is a 3x1 vector, array_in is a 3D array, array_out is a 3D array of the shape of array in'''
    i,j,k = cuda.grid(3)
    if i < array_in.shape[0] and j < array_in.shape[1] and k < array_in.shape[2]:
        value = array_in[i,j,k] * vec[0] + array_in[i,j,k] * vec[0] + array_in[i,j,k] * vec[0]
        cuda.atomic.max(array_out,(i,j,k), value)
    # cuda.synchronize()

def test():

    threadsperblock = (8,8,8)
    blockspergrid_x = ( 17 + threadsperblock[0] -1) // threadsperblock[0]
    blockspergrid_y = ( 21 + threadsperblock[1] -1) // threadsperblock[1]
    blockspergrid_z = (  5 + threadsperblock[2] -1) // threadsperblock[2]
    blockspergrid = (blockspergrid_x, blockspergrid_y, blockspergrid_z)
    array_in = np.random.rand(17,21,5).astype(np.float_)
    vec = np.array([1.0, -1.0, 1.0]).astype(np.float_)
    d_array_in = cuda.to_device(array_in)
    d_vec = cuda.to_device(vec)
    i=0
    while i<20:
        array_out_1 = -999.999*np.ones_like(array_in)
        array_out_2 = -999.999*np.ones_like(array_in)
        d_array_out_1= cuda.to_device(array_out_2)
        d_array_out_2 = cuda.to_device(array_out_2)
        writeToArray[blockspergrid, threadsperblock](d_vec, d_array_in, d_array_out_1)
        writeToArray[blockspergrid, threadsperblock](d_vec, d_array_in, d_array_out_2)
        array_out_1_host = d_array_out_1.copy_to_host()
        array_out_2_host = d_array_out_2.copy_to_host()
        assert(np.allclose(array_out_1_host, array_out_2_host))
        i+=1
        print(i)

if __name__ == "__main__":
    test()
$ cuda-memcheck python t31.py
========= CUDA-MEMCHECK
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
========= ERROR SUMMARY: 0 errors
$

Обратите внимание на использование cuda-memcheck выше. Вы можете использовать его аналогичным образом с вашим исходным кодом, чтобы увидеть сообщения об ошибках, выходящих за границы. Такой доступ вне пределов приведет к сбою контекста CUDA, что означает, что каждый запуск ядра (и все другие действия CUDA) после первого сбоя вообще ничего не делает. 1008 * 1009 Точечное равенство будет работать во всех возможных сценариях ios, но, похоже, оно работает для этого конкретного случая. Вы можете wi sh читать this .

...