Сумма 3D матрица куда

Question

Мне нужно сделать расчет как: A [x] [y] = сумма {от z = 0 до z = n} {B [x] [y] [z] + C [x] [y] [z]}, где матрица A имеет размеры [высота] [ширина] и матрица B, C имеет размеры [высота] [ширина] [n].

Значения отображаются в памяти с помощью чего-то вроде:

index = 0;
for (z = 0; z<n; ++z)
    for(y = 0; y<width; ++y)
        for(x = 0; x<height; ++x) {
            matrix[index] = value;
            index++;
        }

Q1: нормально ли это ядро ​​Cuda?

idx = blockIdx.x*blockDim.x + threadIdx.x;
idy = blockIdx.y*blockDim.y + threadIdx.y;

for(z=0; z<n; z++){
    A[idx*width+idy] += B[idx*width+idy+z*width*height] + C[idx*width+idy+z*width*height];
}

Q2: Это более быстрый способ вычисления?

idx = blockIdx.x*blockDim.x + threadIdx.x;
idy = blockIdx.y*blockDim.y + threadIdx.y;
idz = blockIdx.z*blockDim.z + threadIdx.z;

int  stride_x = blockDim.x * gridDim.x;
int  stride_y = blockDim.y * gridDim.y;
int  stride_z = blockDim.z * gridDim.z;

while ( idx < height && idy < width && idz < n ) {
    atomicAdd( &(A[idx*width+idy]), B[idx*width+idy+idz*width*height] + C[idx*width+idy+idz*width*height] );
    idx += stride_x;
    idy += stride_y;
    idz += stride_z;
} 

Answer 1

Первое ядро ​​в порядке. Но мы не объединили доступ к матрице B и C.

Что касается второй функции ядра. У вас есть гонки данных, потому что не только один поток может писать в A[idx*width+idy] адресах. Вам нужна дополнительная синхронизация типа AttomicAdd

Что касается общего вопроса: Я думаю, что эксперименты показывают, что это лучше. Это зависит от типичных размеров матрицы, которые у вас есть. Помните, что максимальный размер блока нитей на Fermi <1024, и если матрицы имеют большой размер, вы получаете много блоков нитей. Обычно это медленнее (иметь много потоковых блоков). </p>

Answer 2

Очень просто в ArrayFire :

array A = randu(nx,ny,nz);
array B = sum(A,2); // sum along 3rd dimension
print(B);

Answer 3

Q1: проверить это с помощью матриц, где вы знаете ответ

Примечание: у вас могут возникнуть проблемы при использовании очень больших матриц. Используйте цикл while с соответствующими приращениями. Cuda by Example - это, как обычно, справочник.

Пример реализации вложенного цикла можно найти здесь: Для вложенных циклов с CUDA . Там реализован цикл while.

marina.k прав насчет состояния гонки. Это будет способствовать подходу 1, так как атомарные операции имеют тенденцию замедлять код.