Проблема с суммированием записей с использованием поэлементного ядра в Cupy - PullRequest
Новая
       43

Проблема с суммированием записей с использованием поэлементного ядра в Cupy

0 голосов
/ 27 сентября 2019

В первом примере кода (kernel_conv) я запрограммировал простую свертку, и она работала с ожидаемым результатом [1,1,2,1,1].

Затем я использовал поэлементное ядро ​​для суммирования всех записей вектора.Однако, если я запускаю второй пример (kernel_sum), я получаю результат [3,0,0], но ожидаю [6,0,0].

В чем разница между этими двумя примерами?Почему переменная y в первом примере обновляется, а во втором она кажется перезаписанной? 

import numpy as np 
import cupy as cp 

kernel_conv = cp.ElementwiseKernel(
    'raw float32 x', 'raw float32 y',
    ''' int idx = i*2 + 1;
        for(size_t j=0;j<3;j++){
          y[idx - 1 + j] += x[j];
        }
    ''', 'conv')

x = cp.asarray(np.array([1,1,1]),dtype=np.float32)
y = cp.zeros((5,),dtype=np.float32)
z = kernel_conv(x,y,size=2)
print(z)

kernel_sum = cp.ElementwiseKernel(
  'raw float32 x', 'raw float32 y',
  ''' 
      y[0] += x[i]
  ''', 'summe')

x = cp.asarray(np.array([1, 2, 3]), dtype=np.float32)
y = cp.zeros((3,),dtype=np.float32)
z = kernel_sum(x,y,size=3)
print(z)

1 Ответ

1 голос
/ 28 сентября 2019

Неверный результат kernel_sum связан с гонкой данных.В этом случае 3 потока пытаются выполнить запись в один и тот же адрес глобальной памяти (y[0]) одновременно.Чтобы избежать гонки данных, вы должны: 1) использовать atomicAdd или 2) использовать cupy.ReductionKernel для сокращения.

На самом деле, kernel_conv также имеет гонку данных.Первый поток, выполняющий y[2] += x[2], может конфликтовать со вторым потоком, выполняющим y[2] += x[0].Поскольку первый из них слегка отставал от фактического выполнения, на результат это не повлияло, но это проблема синхронизации и вообще не гарантируется *.Чтобы исправить это, вы можете снова использовать atomicAdd или изменить способ запуска вычислений несколькими потоками (например, запустить 5 потоков, каждый из которых вычисляет отдельный элемент y).

* Действительно, в приведенном выше примере kernel_conv я предполагаю, что корректность гарантируется, когда все потоки работают в одной и той же деформации, т. Е. Число потоков не превышает 32. Это потому, что все потоки вта же самая деформация работает синхронно до расхождения по потокам управления.Если для числа потоков установлено большее значение, на границе деформации может произойти гонка данных.

...