Динамическое умножение матриц с помощью CUDA

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

Я ищу входные данные 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

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

35.000 20.000
63.000 36.000

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


matrix cuda matrix-multiplication
person Dan    schedule 12.02.2012    source источник
comment
Какой у Вас вопрос? Вы спрашиваете, как получить информацию от пользователя?   -  person harrism    schedule 13.02.2012
comment
Если я правильно понял ваши предыдущие вопросы, ваш реальный вопрос заключается в том, как изменить этот код ядра (сам по себе очень слегка модифицированная версия примера умножения матриц CUDA SDK), чтобы его можно было использовать для умножения матриц произвольных size, в отличие от округления, кратного размеру блока ядра. Не могли бы вы отредактировать свой вопрос, чтобы отразить это? На данный момент очень неясно, что вы на самом деле спрашиваете.   -  person talonmies    schedule 13.02.2012
comment
@talonmies, ты прав. Это именно то, что я ищу   -  person Dan    schedule 13.02.2012
comment
Привет Дэн, ты нашел правильный код для своей реализации выше?   -  person    schedule 18.04.2012
comment
@JohnSmith Спасибо. Но я думаю, что это ограничено 16 блоками. Я просто просматриваю это, чтобы увидеть, могу ли я это изменить.   -  person Dan    schedule 19.04.2012
comment
Не могли бы вы обновить код выше? Мне любопытно узнать, что такое полный код, включая стратегию, предложенную talonmies. Спасибо   -  person    schedule 19.04.2012
comment
@JohnSmith Вы имеете в виду заполнение нулями для работы с блоком из 16?   -  person Dan    schedule 19.04.2012

Ответы (1)


arrow_upward
7
arrow_downward

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

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

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

введите здесь описание изображения

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

введите здесь описание изображения

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

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

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

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

введите здесь описание изображения

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

введите здесь описание изображения

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

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

введите здесь описание изображения

где A11 и B11 представляют собой матрицы 992 x 992, а O nn — нулевые матрицы, как и раньше. На первый взгляд это не кажется очень полезным, но стоит помнить, что все вычисления для построения правой части матрицы содержат только около 1,2% всех вычислений, необходимых для вычисления матричного произведения. Их можно легко выполнить на центральном процессоре, пока графический процессор выполняет основные вычисления, а затем добавить к результату графического процессора, чтобы сформировать окончательную матрицу. Поскольку API CUDA является асинхронным, большая часть вычислений хоста может быть полностью скрыта и практически бесплатна.

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

person talonmies    schedule 13.02.2012
comment
Спасибо за хорошо оформленное объяснение. Я иду с мягким решением ради времени. Мне любопытно, где вы придумали на 7% больше вычислений. - person Dan; 14.02.2012
comment
Я пробовал заполнять его нулями, но это дает неправильные результаты. Я потратил некоторое время, пытаясь понять это, но кажется, что во время заполнения он лежит в 0. Пожалуйста, смотрите мой отредактированный OP для обновления. - person Dan; 14.02.2012
comment
То, как вы сделали заполнение нулями в своем примере, неверно. А цифра 7% получается из разницы между 2*1024^3 и 2*1000^3, которые представляют собой количество операций скалярных произведений для двух размеров. - person talonmies; 14.02.2012
comment
Привет, talonmies, спасибо за отличный ответ. Вы упомянули, что для стратегии потребуется изменить одну строку кода. Так что речь идет не о создании матрицы внутри матрицы, а просто о наличии оператора if для проверки границ, верно? Учитывая приведенный выше код Дэна, как мне определить границы? Спасибо - person ; 17.04.2012