Операции на основе срезов PyCUDA GPUArray

В документации по PyCUDA мало примеров для тех из нас, кто не является гуру, но меня интересуют операции, доступные для операций с массивами на gpuarrays, т.е. если бы я хотел использовать gpuarray в этом цикле;

m=np.random.random((K,N,N))
a=np.zeros_like(m)
b=np.random.random(N) #example
for k in range(K):
    for x in range(N):
        for y in range(N):
            a[k,x,y]=m[k,x,y]*b[y]

Обычное сокращение Python для первой остановки для этого будет выглядеть примерно так:

for k in range(K):
    for x in range(N):
        a[k,x,:]=m[k,x,:]*b

Но я не вижу никакого простого способа сделать это с помощью GPUArray, кроме написания собственного поэлементного ядра, и даже тогда с этой проблемой в ядре должны быть циклические конструкции, и в этой точке сложности я, вероятно, лучше я просто пишу свое собственное полноценное ядро ​​​​SourceModule.

Кто-нибудь может подсказать мне?


python gpu multidimensional-array cuda pycuda
person Bolster    schedule 18.04.2011    source источник

Ответы (2)


arrow_upward
2
arrow_downward

Это, вероятно, лучше всего сделать с вашим собственным ядром. Хотя класс PyCUDA gpuarray представляет собой действительно удобную абстракцию памяти графического процессора во что-то, что можно использовать взаимозаменяемо с массивами numpy, нет необходимости кодировать для графического процессора что-либо, кроме стандартной линейной алгебры и параллельных операций редукции.

Тем не менее, это довольно тривиальное маленькое ядро ​​для написания. Настолько тривиально, что это будет связано с пропускной способностью памяти - вы можете посмотреть, сможете ли вы «объединить» несколько подобных операций вместе, чтобы немного улучшить соотношение FLOPS к транзакциям памяти.

Если вам нужна помощь с ядром, оставьте комментарий, и я могу расширить ответ, включив в него грубый прототип.

person talonmies    schedule 18.04.2011
comment
Я всегда ценю руководство, поэтому мне было бы очень интересно посмотреть, как кто-то еще возьмется за это, но моя маленькая экспериментальная реализация немного ленива и разбивает k как blockIdx.x, а x и y как их соответствующие threadid (N is никогда не будет близко к 512, так что я согласен с этим), с соответствующими размерами блока и сетки. Если я на правильном пути, тогда игнорируйте это, но любое дополнительное понимание было бы здорово. - person Bolster; 19.04.2011
comment
В случае упорядоченных массивов строк размерность y должна быть той размерностью, в которой операции чтения должны быть объединены для достижения максимальной пропускной способности. Так что я бы сделал внутренний цикл внутри блока, а внешние измерения развернуты по сетке. - person talonmies; 19.04.2011
comment
Итак (просто чтобы быть уверенным), значения k и x - это идентификаторы потоков, а y - идентификаторы блоков? - person Bolster; 19.04.2011
comment
Нет, по-другому. Имея потоки в одном блоке (и деформации в одном блоке), считываемые по y (третьему) измерению, операции чтения памяти должны быть объединены для упорядоченного хранения строк. Этот способ работы полностью ограничен полосой пропускания, поэтому оптимизация доступа к памяти является ключом к получению приемлемой производительности. - person talonmies; 19.04.2011

arrow_upward
2
arrow_downward

Вы также можете использовать метод memcpy_dtod() и функцию нарезки gpuarrays. Странно, что обычное задание не работает. set() не работает, так как предполагает передачу от хоста к устройству (используя memcpy_htod()).

    for k in range(K):
        for x in range(N):
            pycuda.driver.memcpy_dtod(a[k,x,:].gpudata, (m[k,x,:]*b).gpudata, a[k,x,:].nbytes)
person yanlend    schedule 01.08.2013