Код общей памяти с Cuda Python - PullRequest
Новая
       40

Код общей памяти с Cuda Python

0 голосов
/ 30 марта 2020

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

import numpy as np
from numba import cuda, types
m = 128
n = 32
a = np.arange(m*n).reshape(m,n).astype(np.int32)
b = np.arange(m*n).reshape(n,m).astype(np.int32)
c = np.zeros((m, n)).astype(np.int32)

d_a = cuda.to_device(a)
d_b = cuda.to_device(b)
d_c = cuda.to_device(c)

block_size = (m,n)
grid_size = (int(m/n),int(m/n))


@cuda.jit
def mm(a, b, c):
    column, row = cuda.grid(2)
    sum = 0

    # `a_cache` and `b_cache` are already correctly defined
    a_cache = cuda.shared.array(block_size, types.int32)
    b_cache = cuda.shared.array(block_size, types.int32)


    a_cache[cuda.threadIdx.y, cuda.threadIdx.x] = a[row, column]
    b_cache[cuda.threadIdx.x, cuda.threadIdx.y] = b[column, row]
    cuda.syncthreads()
    for i in range(a.shape[1]):
        sum += a_cache[row][i] * b_cache[i][column]
    c[row][column] = sum

, а тестирование 

mm[grid_size, block_size](d_a, d_b, d_c)
solution = a@b
output = d_c.copy_to_host()

продолжает приводить к следующей ошибке: 

CudaAPIError: [700] Call to cuMemcpyDtoH results in UNKNOWN_CUDA_ERROR

После общения с поставщиком одного ответа я обновил функцию. Но все еще не может сделать эту работу. Таким образом, для вычисления суммы для каждого элемента в выводе c нам нужно l oop по столбцам A и строкам B, используя i в качестве индекса. Поэтому у нас есть n * n продуктов. Я думаю, что мы правильно в сумме, но я не могу получить правильный индекс для строки и столбца a и b в выражении для суммы. 

import numpy as np
from numba import cuda, types
@cuda.jit
def mm_shared(a, b, c):
    column, row = cuda.grid(2)
    sum = 0

    # `a_cache` and `b_cache` are already correctly defined
    a_cache = cuda.shared.array(block_size, types.int32)
    b_cache = cuda.shared.array(block_size, types.int32)


    a_cache[cuda.threadIdx.x, cuda.threadIdx.y] = a[row, column]
    b_cache[cuda.threadIdx.x, cuda.threadIdx.y] = b[row, column]

    cuda.syncthreads()


    for i in range(a.shape[1]):

        sum += a_cache[cuda.threadIdx.x, i] * b_cache[i, cuda.threadIdx.y]

    c[row][column] = sum

1 Ответ

1 голос
/ 31 марта 2020

Ваш размер блока неверен. Устройства CUDA имеют ограничение из 1024 потоков на блок. Когда я запускаю ваш код, я вижу это: 

/opt/miniconda3/lib/python3.7/site-packages/numba/cuda/cudadrv/driver.py in _check_error(self, fname, retcode)
    327                     _logger.critical(msg, _getpid(), self.pid)
    328                     raise CudaDriverError("CUDA initialized before forking")
--> 329             raise CudaAPIError(retcode, msg)
    330 
    331     def get_device(self, devnum=0):

CudaAPIError: [1] Call to cuLaunchKernel results in CUDA_ERROR_INVALID_VALUE

Когда я исправляю это, я вижу это: 

$ cuda-memcheck python somethingsometing.py

========= CUDA-MEMCHECK
========= Invalid __shared__ read of size 4
=========     at 0x000008b0 in cudapy::__main__::mm$241(Array<int, int=2, A, mutable, aligned>, Array<int, int=2, A, mutable, aligned>, Array<int, int=2, A, mutable, aligned>)
=========     by thread (15,11,0) in block (3,2,0)
=========     Address 0x00000ec0 is out of bounds

Почему это довольно очевидно: 

for i in range(a.shape[1]):
    sum += a_cache[row][i] * b_cache[i][column]

row и column являются измерениями в сетке выполнения, а не плиткой локальной общей памяти, и аналогично i ограничен формой a, а не формой a_cache (обратите внимание, что вы, похоже, пропустите синтаксис индексации двумерного массива в стиле C примерно на полпути кода, что является потенциальной ошибкой, если вы не понимаете разницу между ними в Python).

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

Также обратите внимание, что

  • Размеры, выбранные вами для c, неверны (должно быть mxm)
  • Размер сетки, на которой вы запускаете ядро, равен также неправильно, потому что размеры C неверны, и поэтому ваш код никогда не сможет рассчитать всю матрицу
  • Даже после исправления всего этого, вполне вероятно, что результаты умножения будут неверными при любых других значениях, кроме тривиальных размеров, потому что целочисленного переполнения.
...