КОМПЬЮТЕРЫ

Что нужно, чтобы начать

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

Набор доступных вам средств программирования зависит от установленной видеокарты - от ее серии, а также производителя. В таблицах ниже приведены серии видеокарт, а также рекомендуемые для них средства программирования.

Заметим, что в таблице приведены не обязательные, а рекомендуемые инструменты. На ГПУ старших моделей любого из производителей, естественно, доступны графические технологии программирования, однако они не являются самыми удобными для программирования общих вычислений. Visual Studio тоже не является обязательной - можно использовать любую другую интегрированную среду разработки или даже просто компилятор C++.

Итак, средства программирования мы скачали и поставили. Дальше следует начать с запуска примеров (samples), которые есть в каждом из приведенных выше инструментов. Во-первых, их запуск позволит вам убедиться, что вы ничего не напутали и выбрали подходящее средство программирования для вашей видеокарты. Во-вторых, вы посмотрите, как это работает, как это быстро или красиво. А в-третьих, вы увидите, сколько кода нужно написать, чтобы это все работало?

Итак, примеры работают нормально? Можно идти дальше. Что делать - это уже каждый решает сам. Можно, например, погрузиться в чтение руководства (а к каждому средству программирования оно прилагается, и не одно). Можно играться с примерами. А можно попытаться что-то сделать самому. Для тех, кто после долгих раздумий все-таки выбрал последнее, предназначены следующие два раздела. Они рассказывают о том, как адаптировать под ГПУ два классических алгоритма: умножение матриц и фильтрацию изображений. В качестве технологий программирования используются CUDA для NVidia и CAL для AMD.

Классический алгоритм: умножение матриц

Умножение плотных матриц является одним из наиболее известных классических алгоритмов, используемых в высокопроизводительных вычислениях. Если посмотреть на требования к задачам, которые эффективно реализуются на ГПУ (а также на сам алгоритм), то можно понять, что он им всем удовлетворяет. И действительно, умножение матриц на современных ГПУ позволяет достигать производительности 40 - 60% пика.

Начнем с классической записи алгоритма на C. В приводимом коде матрицы хранятся по строкам, A имеет размер m x k, B - k x n, а матрица С, в которую будет записан результат, имеет размер m x n.

void cpp_matmult(float* a, float* b, float* c, int m, int n, int k) {
 for(int i = 0; i < m; i++)
   for(int j = 0; j < n; j++) {
      float r = 0;
      for(int p = 0; p < k; p++)
         r += a[m * i + p] * b[k * p + j];
      c[m * i + j] = r;
   }  // end of for()
}  // end of cpp_matmult()

Можно начинать адаптировать его для ГПУ. Прежде чем делать это, вспомним, что ГПУ вычисляет параллельно все элементы некоторого массива (одномерного или двумерного). Поэтому нужно решить, какие именно из этих циклов будут отображены на измерения этого массива. В любой задаче логичнее всего отобразить на них те циклы, которые соответствуют измерениям вычисляемого массива элементов. Для карт AMD это вызвано ограничениями на возможности записи, для карт NVidia - необходимостью породить большое количество потоков для покрытия накладных расходов на доступ к памяти. Итак, циклы по i и j отображаем на измерения массива, вычисляемого на ГПУ. Остается цикл по p. Его можно отобразить либо на множество проходов, в каждом из которых будет выполняться одна итерация цикла, либо на цикл внутри шейдера. Поскольку работа с памятью на ГПУ является дорогой операцией, а первый вариант предполагает намного меньше операций с памятью, выбираем именно его. Итак, циклы по i и j отображаются на измерения вычисляемого массива, а по p - на цикл внутри шейдера.

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

void cuda_matmul(float* a, float* b, float*c, int m, int n, int k) {
 int i = threadIdx.y, j = threadIdx.x;
  float r = 0;
   for(int p = 0; p < k; p++)
     r += a[m * i + p] * b[k * p + j];
   c[m * i + j] = r;
 }  // end of cuda_matmul()

Обратите внимание, что здесь приведен только код ядра. Внешние измерения цикла исчезли (поскольку они задаются запуском ядра на двумерной области), остались только внутренние. Для ясности сохранены все названия переменных. Для AMD Brook+ тот же самый код будет выглядеть следующим образом (опять же, приведено только ядро):

kernel void cal_matmult(float a[][], float b[][], out float c<>, int m, int n, int k) {
 float2 ji = indexof(c).xy;
 float i = ji.y, j = ji.x;
 float2 ia = float2(0.0, j), ib = float2(i, 0.0);
 float r = 0;
 for(int p = 0; p < N; p++) {
   r += a[ia] * b[ib];
   ia += float2(1, 0); ib += float2(0, 1);
 }
 c = r;
}  // end of cal_matmult()

Код не совсем привычный, но разобраться можно. Обратите внимание, что хотя на ГПУ AMD и поддерживаются целочисленные операции, для индексирования массивов в памяти они используют вещественные числа, что находит отражение в коде на Brook. Как только это становится понятным, в алгоритме можно узнать ту же структуру, что и в оригинальной программе на C++.

Безусловно, программы в таком виде на обеих архитектурах работать будут, однако их производительность крайне низкая. На CUDA это 1 - 2 ГФлоп/c, что меньше 1% от пика. На AMD ситуация несколько лучше: примерно 5-7% пиковой производительности (в зависимости от карты) т.е. при производительности карточки в 250 ГФлоп/c это составит около 12 ГФлоп/c, что также очень мало.

На этом простое и прямолинейное программирование заканчивается. Далее для каждой архитектуры начинается свой путь оптимизации. Для AMD это блочный алгоритм умножения матрицы, при этом размер блока зависит от типа карточки и составляет 16 или 32 элемента матрицы. За счет блочности задействуются операции с четверками вещественных чисел. Для NVidia потенциал оптимизации лежит опять же в блочном алгоритме, однако с обязательным использованием статической памяти. Программы получаются намного длиннее и по этой причине здесь не приводятся. Их можно посмотреть в CTM SDK и CUDA SDK, которые свободно скачиваются с сайтов AMD и NVidia. Отметим, что реально получаемая при этом производительность составляет уже порядка 100 ГФлоп/c на обоих типах архитектур, что составляет порядка 20% от их пика.

RapidMind: адаптация программы программы

Технология RapidMind предполагает трехэтапный процесс адаптации программ для ГПУ:

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

Рассмотрим следующий код на C++, умножающий двумерный массив на скаляр и прибавляющий его к другому массиву (двумерный массив развернут в одномерный).

void main () {
 float f;
 int w = 512, h = 512;
 float a[w*h*3], b[w*h*3];
 for(int y = 0 y < h; y++)
  for(int x = 0; x < w; x++)
   for(int e = 0; e < 3; e++)
    a[(y * w + x) * 3 + e] += b[(y * w + x) * 3 + e] * f;
}  // end of main()

При адаптации его для RapidMind фрагмент примет следующий вид (опять же, на C++):

using namespace rapidmind;
void main() {
 Value1f f;
 rapidmind::init();
 int w = 512, h = 512;
 Array<2, Value3f> a(w, h), b(w, h);
 Program mul_add = BEGIN {
  In i1, i2;
  Out o;
  o = i1 + i2 * f;
 } END;
 a = mul_add(a, b);
}

© Лаборатория Параллельных информационных технологий НИВЦ МГУ