Суммирование по одному измерению трехмерного массива с использованием общей памяти - PullRequest
Новая
       70

Суммирование по одному измерению трехмерного массива с использованием общей памяти

3 голосов
/ 01 апреля 2012

Мне нужно сделать расчет как: 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++;
        }

Я бы хотел, чтобы каждый блок вычислял одну сумму, поскольку каждый блок имеет собственную разделяемую память.Чтобы избежать скачек данных, я использую atomicAdd, что-то вроде этого:

Часть кода в глобальной памяти: 

dim3 block (n, 1, 1);
dim grid (height, width, 1);

Ядро: 

atomicAdd( &(A[blockIdx.x + blockIdx.y*gridDim.y]), 
           B[blockIdx.x + blockIdx.y*gridDim.y+threadIdx.x*blockDim.x*blockDim.y] 
           + C[blockIdx.x + blockIdx.y*gridDim.y+threadIdx.x*blockDim.x*blockDim.y] );

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

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

1 Ответ

3 голосов
/ 01 апреля 2012

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

__global__ void kernel(float *A, const float *B, const float *C, 
        const int width, const int height, const int n)
{
    int tidx = threadIdx.x + blockDim.x * blockIdx.x;
    int tidy = threadIdx.y + blockDim.y * blockIdx.y;

    if ( (tidx < height) && (tidy < width) ) {
        int stride = width * height;
        int ipos = tidx + tidy * height;

        float * oval = A + ipos;
        float sum = 0.f;
        for(int z=0; z<n; z++, ipos+=stride) {
            sum += B[ipos] + C[ipos];
        }
        *oval = sum;
    }
}

Этот подход должен быть оптимальным для основных данных столбца с width * height >= n.Для этого нет никаких преимуществ с точки зрения производительности при использовании разделяемой памяти, и нет необходимости использовать атомарные операции с памятью.Если у вас возникла проблема, когда width * height << n, возможно, имеет смысл попробовать блочное параллельное уменьшение на сумму.Но вы не указали, каковы типичные размеры проблемы.Оставьте комментарий, если ваша проблема больше похожа на последнюю, и я могу добавить в ответ образец ядра на основе сокращения.

...