Студопедия — Процедура перемножения матриц на GPU SM4
Студопедия Главная Случайная страница Обратная связь

Разделы: Автомобили Астрономия Биология География Дом и сад Другие языки Другое Информатика История Культура Литература Логика Математика Медицина Металлургия Механика Образование Охрана труда Педагогика Политика Право Психология Религия Риторика Социология Спорт Строительство Технология Туризм Физика Философия Финансы Химия Черчение Экология Экономика Электроника

Процедура перемножения матриц на GPU SM4






Рассмотрим конкретную программу, реализующую перемножение матриц на GPU SM4. Эта программа может быть откомпилирована любыми средствами, поддерживающими библиотеки CUDA. Детали реализации алгоритма обсуждаются в комментариях.

 

// Файл matrixMul_kernel.cu

 

#ifndef _MATRIXMUL_KERNEL_H_

#define _MATRIXMUL_KERNEL_H_

 

#include < stdio.h>

#include " matrixMul.h"

 

/* Следующая процедура имеет модификатор __global__, т.е. является вычислительным ядром. Это ядро исполняет­ся одновременно на каждом мультипроцессоре и каждым потоком. Но всё-таки конкретные вычислительные процес­сы, исполняемые в различных «связках» и потоках, разли­чаются значениями таких системных переменных, как blockIdx и threadIdx. Это позволяет нам «привязать» конкретные ячейки матриц к конкретным потокам.

Параметрами процедуры являются указатели (float*) на области оперативной памяти компьютера, в которых распо­ложены матрицы A, B и C, а также целочисленные (int) значения wA и wB, задающие ширину входных матриц в единицах размера блока Block_Size */

 

__global__ void

matrixMul (float* C, float* A, float* B, int wA, int wB)

{

/* Присвоение локальным переменным значений координат текущей «связки» потоков */

int block_x = blockIdx.x;

int block_y = blockIdx.y;

 

/* Определение индексов (i, k), соответствующих конкретному потоку в составе «связки» */

int i = threadIdx.y;

int k = threadIdx.x;

 

/* Для понимания дальнейшего необходимо учесть, что в глобальной памяти графического процессора (Global memory на рис. 9.2) элементы исходных матриц A и B хранятся последовательно, в форме линейного массива, как это показано следующими формулами:

.

Как можно увидеть из рис. 9.4, элементы матрицы A, обрабатываемые одной «связкой», расположены в линейном массиве последовательно. Если нумеровать все элементы линейного массива подряд с нуля, то номер первого элемента последовательности, с которой работает заданная «связка» -

.

Вводим целочисленную константу, равную этому начальному номеру: */

 

int a_First = wA * Block_Size * block_y;

 

/* Как, опять же, видно из рис. 9.4, последователь­ность элементов матрицы A, обрабатываемая выделенной «связкой», проходит через несколько блоков размера Block_Size, которые загружаются в разделяемую память поочерёдно. Для организации цикла, внутри которого будут загружаться в память и обрабатываться эти блоки, нам потребуется номер 1-ого элемента в последнем из блоков, который вычисляется по формуле

.

Определяем соответствующую целочисленную константу: */

 

int a_Last = a_First + (wA – 1) * BLOCK_SIZE;

 

/* Для организации этого же цикла потребуется шаг, задающий изменение 1-ого номера (т.е. «левого верхнего» элемента в блоке) при переходе от текущего блока к следующему */

int a_Step = Block_Size;

 

/* Индекс первого блока матрицы B, обрабатываемого текущей связкой

int b_First = Block_Size * block_x;

 

/* Шаг итераций «сверху вниз» (см. рис. 9.3) по блокам матрицы B */

int b_Step = Block_Size * wB;

 

/* Величины шагов при переходах от одних блоков к другим в матрицах A и B проиллюстрированы на рис. 9.5.

Рис. 9.5. Переходы к следующим блокам при умножении матриц */

// Вводим переменную, в которую будет записываться
//значение искомого элемента матрицы Csub (рис. 9.4-5)

float Csub = 0;

 

// Внешний цикл по всем блокам матриц A и B

for (int a = a_First, b = b_First;

a < = a_Last;

a += a_Step, b += b_Step) {

 

/* Создаём в разделяемой памяти массив, куда из гло­бальной видеопамяти будет скопирован очередной блок матрицы A. Модификатор __shared__ как раз означает, что массив As будет создан в «быстрой» разделяемой памяти */

__shared__ float As[Block_Size][Block_Size];

 

/* Такой же массив создается для хранения блока матрицы B */

__shared__ float Bs[Block_Size][Block_Size];

 

// Копируем блоки матриц в разделяемую память:

/* каждый поток загружает 1 элемент матрицы A и 1 элемент матрицы B */

As[i][k] = A[a + wA * i + k];

BS[i][k] = B[b + wB * i + k];

/* Принцип определения номера загружаемого элемента в блоке показан на рис. 9.6 */

Рис. 9.6. Принцип расчёта линейного номера элемента с индексами (i, j) в текущем блоке: a – номер 1-го элемента в данном блоке, (i, j) – индексы вычислитель­ного потока, копирующего данный элемент; эти индексы отсчитываются не от начала матрицы, а от начала блока.

/* Синхронизируем потоки, чтобы гарантировать то, что блоки матриц загружены полностью */

 

__syncthreads();

 

/* На этом этапе все элементы обрабатываемых блоков загружены в разделяемую память. Создавать цикл для их загрузки не требуется, так как этот код будет параллельно исполняться стольким числом потоков, равным количеству элементов в блоке матрицы Block_Size× Block_Size. */

 

/* Перемножение блоков матриц A и B, загружен­ных в «быструю» память GPU. Каждый поток вычисляет 1 элемент результирующей матрицы Csub

.

Использование оператора += обеспечивает добавление рассчитываемого слагаемого ко всей сумме

. */

 

for (int j = 0; j < Block_Size; ++j)

Csub += As[i][j] * Bs[j][k];

 

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

 

/* Синхронизируем потоки. Это гарантирует, что все вычисления будут закончены до того, как начнётся загрузка других блоков на следующей итерации внешнего цикла */

__syncthreads();

} // Возвращение к началу внешнего цикла

 

/* Обратное копирование вычисленного блока результирующей матрицы C в глобальную память GPU, где данные будут доступны центральному процессору. Каждый поток сохраняет 1 элемент */

int c = wB * Block_Size * block_y + Block_Size * block_x;

C[c + wB * i + k] = Csub;

}

 

#endif // #ifndef _MATRIXMUL_KERNEL_H_

 







Дата добавления: 2014-12-06; просмотров: 613. Нарушение авторских прав; Мы поможем в написании вашей работы!



Шрифт зодчего Шрифт зодчего состоит из прописных (заглавных), строчных букв и цифр...

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

Практические расчеты на срез и смятие При изучении темы обратите внимание на основные расчетные предпосылки и условности расчета...

Функция спроса населения на данный товар Функция спроса населения на данный товар: Qd=7-Р. Функция предложения: Qs= -5+2Р,где...

Понятие о синдроме нарушения бронхиальной проходимости и его клинические проявления Синдром нарушения бронхиальной проходимости (бронхообструктивный синдром) – это патологическое состояние...

Опухоли яичников в детском и подростковом возрасте Опухоли яичников занимают первое место в структуре опухолей половой системы у девочек и встречаются в возрасте 10 – 16 лет и в период полового созревания...

Способы тактических действий при проведении специальных операций Специальные операции проводятся с применением следующих основных тактических способов действий: охрана...

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

ЛЕЧЕБНО-ПРОФИЛАКТИЧЕСКОЙ ПОМОЩИ НАСЕЛЕНИЮ В УСЛОВИЯХ ОМС 001. Основными путями развития поликлинической помощи взрослому населению в новых экономических условиях являются все...

МЕТОДИКА ИЗУЧЕНИЯ МОРФЕМНОГО СОСТАВА СЛОВА В НАЧАЛЬНЫХ КЛАССАХ В практике речевого общения широко известен следующий факт: как взрослые...

Studopedia.info - Студопедия - 2014-2024 год . (0.009 сек.) русская версия | украинская версия