подматрица

Question

проблема:

У меня есть 4 матрицы (64x64) чисел одинарной точности.нужно сделать расчет как:

R = A * sin(B) + C * cos(D)

идея:

для ускорения расчетов используйте разделяемую память.поскольку каждый блок потоков имеет (в случае моего графического процессора) 16 КБ разделяемой памяти и размер с плавающей запятой равен 4, можно хранить 4000 чисел с плавающей запятой в общей памяти.поэтому для каждой матрицы используйте 1000 элементов, что составляет 31 элемент на измерение.

поэтому каждая матрица должна быть разделена на 16 подматриц (16x16).

dim3 dimBlock(16, 16, 1)
dim3 dimGrid(4, 4, 1)

kernel:

int Tx = threadIdx.x;
int Ty = threadIdx.y;

int Bx = blockIdx.x;
int By = blockIdx.y;

int idx = Bx * blockDim.x + Tx;
int idy = By * blockDim.y + Ty;

__shared__ float s_A[16*16];
__shared__ float s_B[16*16];
__shared__ float s_C[16*16];
__shared__ float s_D[16*16];

// I am not sure how to write this part

s_A[(Tx * blockDim.x + Ty + By) + Bx] = A[idx * 64 + idy];
s_B[(Tx * blockDim.x + Ty + By) + Bx] = B[idx * 64 + idy];
s_C[(Tx * blockDim.x + Ty + By) + Bx] = C[idx * 64 + idy];
s_D[(Tx * blockDim.x + Ty + By) + Bx] = D[idx * 64 + idy];

R[idx * 64 + idy] = s_A[(Tx * blockDim.x + Ty + By) + Bx] * sin(s_B[(Tx * blockDim.x + Ty + By) + Bx]) + s_C[(Tx * blockDim.x + Ty + By) + Bx] * cos(s_D[(Tx * blockDim.x + Ty + By) + Bx]);

Как разделить исходную матрицу на подматрицы, чтобы каждый блок имелвладеть 4 подматрицами и рассчитывать по ним.

Answer 1

Если я неправильно истолковал ваш вопрос, вам не нужно и не следует использовать общую память для этой операции.Общая память полезна для совместного использования и возобновления данных между потоками в одном блоке, а также для облегчения доступа к объединенной памяти.Похоже, что ваша операция не требует ни одной из этих вещей для правильной работы.Использование общей памяти так, как вы предлагаете, вероятно, будет медленнее , чем простое чтение из глобальной памяти напрямую.Кроме того, поскольку вас беспокоят только поэлементные операции, схема индексации вашего ядра может быть значительно упрощена - тот факт, что A, B, C и D являются "матрицами", не имеет отношения квычисления, как я понимаю ваш вопрос.

В результате, почти оптимальная версия вашего ядра может быть записана так же просто, как

__global__ void kernel(const float *A, const float *B, const float *C, 
                        const float *D, const int n, float *R)
{
    int tidx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    while(tidx < n) {
        R[tidx] = A[idx] * sinf(B[idx]) + C[idx]*cosf(D[idx]);
        tidx += stride
    }
}

В этом коде вы запустите столько блоков, сколькодостигнет пиковой пропускной способности вашего графического процессора, и каждый поток обработает более одного значения ввода / вывода, если размер массива превысит размер оптимальной 1D сетки, которую вы запустили.Конечно, это довольно академично, если вы обрабатываете всего 4096 элементов - это, вероятно, на 2 порядка меньше, чтобы получить какую-либо выгоду от использования графического процессора.

Answer 2

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

Решение с общей памятью обычно лучше всего, когда некоторые данные вызываются повторно из глобальной памяти. Вместо многократной загрузки этих данных из глобальной памяти с низкой пропускной способностью и высокой задержкой, вы загружаете их один раз оттуда и выполняете последующие загрузки из общей памяти с более высокой пропускной способностью и меньшей задержкой. Обратите внимание, что выше и ниже , а не максимум и минимум . По-прежнему снижается производительность при использовании общей памяти.

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

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

Вы можете даже получить приличные результаты при многопоточности на процессоре.