Динамическое матричное умножение с CUDA - PullRequest
2 голосов
/ 12 февраля 2012

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

Я хочу взять ввод x на x, в настоящее время я не собираюсь умножать два разных размера в данный момент.

Как бы вы, ребята, предложили мне сделать это?

Извините, мой вопрос не был достаточно ясным, я хочу изменить это ядро, чтобы оно могло обрабатывать матрицу любого размера (где x и y являются эквивалентами, чтобы упростить его). Вместо кратных 16.

Я не уверен, что вам понадобится мой текущий код, но вот код ядра:

// CUDA Kernel
__global__ void matrixMul( float* C, float* A, float* B, int wA, int wB,size_t block_size)
{
    int bx = blockIdx.x;
    int by = blockIdx.y;
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    int aBegin = wA * block_size * by;
    int aEnd   = aBegin + wA - 1;
    int aStep  = block_size;

    int bBegin = block_size * bx;

    int bStep  = block_size * wB;
    float Csub=0;

    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) 
    {
        extern __shared__ float As[];
        extern __shared__ float Bs[];
        extern __shared__ float smem[];

        smem[ty*block_size+tx] = A[a + wA * ty + tx];

        smem[block_size*block_size+ty*block_size+tx]  = B[b + wB * ty + tx];

        __syncthreads();

        for (int k = 0; k < block_size; ++k)
            Csub += smem[ty*block_size+k] * smem[block_size*block_size+k*block_size+tx] ;

        __syncthreads();
    }

    int c = wB * block_size * by + block_size * bx;
    C[c + wB * ty + tx] = Csub;


}

Обновление: я решил пойти с дополнением нулями. Однако я получаю неправильные ответы. Возьмите матрицу A 2x2, дополненную до 16x16:

5.000 0.000 9.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000

Матрица B, 2x2, дополненная до 16x16:

7.000 4.000 8.000 7.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000

Таким образом, результат для C, который я получаю, правильный:

35.000 20.000 40.000 35.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000

Однако если убрать нули, матрица должна быть: A:

5.000 0.000
9.000 0.000

B:

7.000 4.000
8.000 7.000

C Должно быть:

35.000 20.000
63.000 36.000

Однако две матрицы C не совпадают.

1 Ответ

7 голосов
/ 13 февраля 2012

Это не очень понятный вопрос, поэтому этот ответ является неким предположением, основанным на том, что вы ранее задавали в нескольких довольно похожих вопросах ранее.

Хорошая отправная точка для понимания того, как выполнять операции такого рода, - вернуться к началу и подумать о проблеме умножения матрицы на матрицу из первых принципов. Вас интересует код для вычисления точечного произведения двух матриц, C = AB . У вас есть ограничение: ядро, которое вы используете, может вычислять только произведения матриц, которые кратны некоторому внутреннему размеру блока. Так что вы можете сделать?

Один из способов взглянуть на проблему - представить, что матрицы A и B были блочными матрицами . Матрицу умножения можно записать так:

enter image description here

и полученная матрица C могут быть затем образованы комбинациями произведений восьми подматриц в A и B :

enter image description here

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

  1. У вас есть оптимальное ядро ​​умножения матриц, которое использует внутренний размер блока 32 и является правильным, только если матрицы представляют собой круглые кратные этому размеру блока.
  2. У вас есть пара матриц 1000x1000 для умножения.

Из этих первых фактов следует, что ваше ядро ​​может корректно решать только продукт 1024x1024 или продукт 992x992, но не необходимую операцию 1000x1000.

Если вы решите использовать продукт 1024x1024, вы можете использовать идею декомпозиции блока, чтобы сформулировать проблему следующим образом:

enter image description here

, где O nn обозначает матрицу подходящих размеров нулей. Теперь у вас есть пара 1024x1024 матриц, и их произведение приведет к

enter image description here

т. верхний блок слева представляет собой матрицу 1000x1000, содержащую AB . Это фактически заполнение нулями для достижения правильного результата. В этом примере это означает, что выполняется примерно на 7% больше вычислений, чем требуется. Чем это важно или нет, вероятно, зависит от приложения.

Второй подход состоит в том, чтобы использовать базовое ядро ​​для вычисления продукта 992x992, а затем выработать стратегию для работы с другими семью продуктами в разложенной на блоки версии расчета, что-то вроде этого:

enter image description here

с A 11 и B 11 , являющимися матрицами 992x992, и O nn - нулевые матрицы, как и раньше. При первом осмотре это не выглядит очень полезным, но стоит помнить, что все вычисления, чтобы сделать правую матрицу, содержат только около 1,2% всех вычислений, необходимых для вычисления продукта матрицы. Их можно легко сделать на центральном процессоре, пока графический процессор выполняет основные вычисления, а затем добавить к результату графического процессора для формирования окончательной матрицы. Поскольку CUDA API является асинхронным, большая часть вычислений хоста может быть полностью скрыта и фактически бесплатна.

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

Добро пожаловать на сайт PullRequest, где вы можете задавать вопросы и получать ответы от других членов сообщества.
...