Программирование с использованием технологий AMD Stream и Nvidia CUDA
Технологии вычислений с использованием видеокарты с каждым днем получают все большее распространение. Особенно часто их применение можно встретить в задачах обработки мультимедийных данных, таких как распознавание текста, обработка изображений и декодирование видеоданных. Поэтому в данной статье мы рассмотрим две наиболее распространенные на сегодняшний день технологии программирования для видеокарты - AMD Stream и Nvidia CUDA.
Несмотря на то, что эти технологии были разработаны в компаниях, которые тесно конкурируют между собой на рынке графических решений, они обе имеют общую концепцию разработки приложений. Более того, данная концепция используется также в недавно появившемся стандарте OpenCL 1.0. Однако перед тем как перейти к непосредственному рассмотрению программного кода, необходимо овладеть базовыми понятиями общей технологии вычислений на видеокартах GPGPU. Поэтому далее мы рассмотрим основные из них на примере технологии AMD Stream (при разработке и тестировании программного кода использовалась версия 1.2), а потом покажем реализацию аналогичного программного кода в Nvidia CUDA (версии 2.0). Потоковая модель программирования GPGPU подразумевает наличие потоков данных (streams), которые необходимо обрабатывать с помощью специальной функции. Такая функция в Brook+ называется ядром (kernel). Потоки данных являются ничем иным, как массивами из элементов одинакового типа данных. Brook+ поддерживаются следующие типы данных: int - 32-битное целое число; float - 32-битное вещественное; double - 64-битное вещественное. Кроме того, допускается использование структур данных. Для описания потоков используется комбинация символов < >. Например: int a<> - одномерный поток (вектор) произвольной длины из целых чисел; int a<100> - одномерный поток (вектор) из ста чисел; int a<100,100> - двумерный поток (матрица) целых чисел размером 100х100; int a<100,100,100> - трехмерный поток целых чисел размером 100х100x100. В AMD Stream версии 1.2 существуют ограничения на размер потоков. Официально поддерживаются двумерные потоки до 8192х8192 и одномерные до 223 элементов. На практике у меня возникла ошибка "Stack overflow" уже при использовании матрицы размером 300x300. Возможно, это из-за нехватки мощности видеокарты или из-за использования не самой последней версии AMD Stream и не очень свежих драйверов. Будем надеяться, что AMD в новой версии 2.0 программного обеспечения Stream уже исправила это (мне не удалось протестировать это из-за несовместимости моей видеокарты). Теперь разберемся с ядрами. Итак, ядро предназначено для проведения вычисления над потоками. Исходящие потоки маркируются словом out. Например, следующее ядро суммирует 2 целочисленных потока a и b и выводит результат в потоке z:
kernel void sum(int a<>, int b<>, out float z<>)
{ z = a + b; }
Версия Brook+, входящая в состав AMD Stream 1.2, поддерживает до 128 потоков на входе и до 8 на выходе. К элементам потоков внутри ядра можно также обращаться по индексу как в обычных массивах. Кроме того, существуют также ядра отсечения (Reduction kernels), в которых можно применять условные операторы. Подробнее о них и о многом другом написано в файле brookspec из прилагающейся к AMD Stream документации. Настоятельно рекомендую с ней ознакомиться, так как там содержатся полные описания языковых синтаксических ограничений для ядер, и не только. Ну а теперь перейдем к рассмотрению примера программы на Brook+, исходный код и запускной файл которого можно скачать здесь: сайт Нижеописанный пример суммирует два потока вещественных чисел a и b и выводит результат в потоке z.
Листинг примера 1 (Файл sum.br):
#include "stdio.h"
//Ядро
kernel void sum(float a<>, float b<>, out float z<>)//Описание входящих и исходящих потоков
{
z = a + b; //Суммирование каждого элемента потока a и b c записью результата в поток z
}
//Основная функция
int main(int argc, char** argv)
{
int i, j; // переменные-счетчик цикла, который заполняет потоки a и b значениями
float a<10, 10>; //поток a - матрица из вещественных чисел размером 10x10
float b<10, 10>; //поток b - матрица из вещественных чисел размером 10x10
float z<10, 10>; //поток z - матрица из вещественных чисел размером 10x10
float input_a[10][10]; //матрица 10x10 для хранения данных потока a в оперативной памяти
float input_b[10][10]; //матрица 10x10 для хранения данных потока b в оперативной памяти
float input_z[10][10]; //матрица 10x10 для хранения данных потока z в оперативной памяти
//заполняем потоки a и b значениями
for(i=0; i<10; i++)
{
for(j=0; j<10; j++)
{
input_a[i][j] = (float) i;
input_b[i][j] = (float) j;
}
}
//Копируем данные (значения элементов) потоков а и b из оперативной памяти в память видеокарты
streamRead(a, input_a);
streamRead(b, input_a);
//запускаем ядро (суммирование элементов a и b c записью результата в поток z)
sum(a, b, z);
//Извлекаем полученный в потоке z результат из памяти видеокарты обратно в оперативную память
streamWrite(z, input_z);
//Выводим полученный результат на экран
for(i=0; i<10; i++)
{
for(j=0; j<10; j++)
{
printf("%6.2f", input_z[i][j]);
}
printf("\n");
}
return 0;
}
Как видно из примера, потоки a, b, z являются матрицами из элементов типа float размером 10х10 элементов. Обратите внимание, что в функции main, в отличие от ядра, указывается размерность потоков. Кроме того, из-за особенностей работы видеокарты все значения элементов потоков перед началом суммирования в ядре должны быть сначала заданы в оперативной памяти компьютера, а затем, с помощью функций streamRead (для входящих потоков a и b) и StreamWrite (для исходящего потока z), помещены и во внутреннюю память видеокарты или считаны оттуда. В результате выполнения кода программы каждый элемент матрицы а будет просуммирован c соответствующим по индексу элементом матрицы b, а результат будет записан в соответствующую позицию матрицы z, после чего результат будет выведен на экран. Однако как же откомпилировать данную программу? Для этого можно использовать интегрированную среду разработки Microsoft Visual Studio 2005, и следующий абзац посвящен тому, как это удобнее всего сделать. Итак, перейдем к рассмотрению процесса компиляции Brook+ программ с помощью Visual Studio 2005. Надеюсь, что к этому моменту на вашем компьютере уже установлены AMD Stream SDK и сама Visual Studio 2005, и вы уже скачали исходный код примера sum.br. Для начала зайдем в папку, где находятся примеры программ Brook+, которая по умолчанию для AMD Stream версии 1.2 находится в C:\Program Files\AMD\AMD BROOK 1.2\samples. Эта папка имеет в своем составе папки apps и bin. Папка apps содержит исходные коды различных примеров, а в папке bin хранятся откомпилированные из этих исходных кодов программы. После каждой компиляции исполняемые файлы проекта записываются именно в папку bin, так как таковы настройки проектов примеров, входящих в состав AMD Stream SDK. Если вы не специалист по параметрам проекта, то рекомендую писать ваши приложения на базе готовых примеров, как это будет показано ниже.
Итак, приступим. Откроем первый пример binary_search из папки apps. Для открытия дважды щелкаем по файлу binary_serach.vcproj. Далее в Solution Explorer открываем файл binary_serch.br. Теперь открываем исходный код нашего примера в Блокноте (для тех, кто забыл, где он находится: Меню Пуск -> Все программы -> Стандартные) и с помощью горячих клавиш Ctrl+C и Ctrl+V заменяем им исходный код файла binary_serch.br. Затем запускаем с помощью кнопки "Run". Если после запуска на экране появилось (промелькнуло) такое же окно, как на рисунке, то наша программа Brook+ успешно откомпилировалась.
Исполняемый файл будет находиться в папке bin\xp_x86_32 (при условии использования системы Windows XP 32-bit) под именем binary_search_d.exe. Поэтому в случае необходимости (например, для просмотра результата) его можно запустить из командной строки прямо оттуда. Обратите внимание, что в результате выполнения сначала был динамически сгенерирован файл binary_search.cpp, и только после этого из этого сгенерированного файла был получен путем компиляции исполняемый файл binary_search_d.exe. Такова особенность работы Brook+. Однако если вам понадобится использовать код в каком-либо своем приложении, то для интеграции необходимо использовать код именно из cpp-файла. Кроме того, исполняемые файлы требуют для запуска наличие уже установленного AMD Stream SDK, так как в него входят необходимые для запуска библиотеки. Теперь, зная, как работает технология программирования AMD Stream, перейдем к рассмотрению аналогичной конкурирующей технологии Nvidia CUDA, которая, несмотря на общую концепцию программирования, имеет целый ряд отличий. Первым отличием от AMD Stream является то, что примеры программ и полезная документация содержатся в разных инсталляционных файлах - Nvidia CUDA SDK и CUDA Toolkit соответственно. При этом, для удобства, рекомендую произвести установку этого всего в одну и ту же папку. Примеры программного кода при этом будут содержаться в папке projects, а вся полезная документация — в папке doc. Открыв любой из проектов примеров программного кода, вы, скорее всего, заметите второе основное отличие от AMD Stream - описание ядра и основной код содержатся в двух разных файлах, которые имеют расширение '.cu'. Ну а самое главное отличие, конечно же, состоит в реализации программного кода. Поэтому перейдем к рассмотрению непосредственно кода, чтобы вы сами смогли оценить это. С целью экономии места приведу лишь самые важные строки кода, без полного листинга.
Как и в случае с AMD Stream, нижеописанный пример суммирует два потока вещественных чисел a и b и выводит результат в потоке z. Следующий код показывает описание ядра и процедуры его вызова из основной программы:
//Ядро (Обычно находится в отдельном файле)
__global__ void Sum(float A[10] [10], float B[10] [10], float Z[10] [10])
{
int i = threadldx.x;
int j = threadldx.y;
Z[i] [j] = A[i] [j] + B[i] [j]; //Суммирование каждого элемента потока a и b c записью результата в поток z
}
//Основная функция (Также обычно в отдельном файле)
int main()
{
dim3 dimBlock(10, 10);
//запускаем ядро (суммирование элементов a и b c записью результата в поток z)
Sum<<<l, dimBlock>>> (А, В, Z);
}
А теперь посмотрите, как в CUDA происходит работа с памятью:
// Выделение места в оперативной памяти для матриц a и b
unsigned int size_A = 10 * 10;
unsigned int mem_size_A = sizeof(float) * size_A;
float* h_A = (float*) malloc(mem_size_A);
unsigned int size_B = 10 * 10;
unsigned int mem_size_B = sizeof(float) * size_B;
float* h_B = (float*) malloc(mem_size_B);
// Инициализация матрицы a и b
randomInit(h_A, size_A);
randomInit(h_B, size_B);
// Выделение для матриц a и b места в памяти видеокарты
float* d_A;
CUDA_SAFE_CALL(cudaMalloc((void**) &d_A, mem_size_A));
float* d_B;
CUDA_SAFE_CALL(cudaMalloc((void**) &d_B, mem_size_B));
// Копируем данные (значения элементов) потоков а и b из оперативной памяти в память видеокарты
CUDA_SAFE_CALL(cudaMemcpy(d_A, h_A, mem_size_A,
cudaMemcpyHostToDevice) );
CUDA_SAFE_CALL(cudaMemcpy(d_B, h_B, mem_size_B,
cudaMemcpyHostToDevice) );
// Выделение памяти видеокарты для результирующей матрицы z
unsigned int size_Z = 10 * 10;
unsigned int mem_size_Z = sizeof(float) * size_Z;
float* d_Z;
CUDA_SAFE_CALL(cudaMalloc((void**) &d_Z, mem_size_Z));
// Выделение места в оперативной памяти для результирующей матрицы z
float* h_Z = (float*) malloc(mem_size_Z);
// Извлекаем полученный в потоке z результат из памяти видеокарты обратно в оперативную память
CUDA_SAFE_CALL(cudaMemcpy(h_C, d_C, mem_size_C,
cudaMemcpyDeviceToHost) );
// Высвобождение занятой памяти
free(h_A);
free(h_B);
free(h_C);
free(reference);
CUDA_SAFE_CALL(cudaFree(d_A));
CUDA_SAFE_CALL(cudaFree(d_B));
CUDA_SAFE_CALL(cudaFree(d_C));
Как видно из вышеописанного кода, концепция программирования в Nvidia CUDA полностью идентична таковой в AMD Stream. Однако в отличие от AMD Stream, в CUDA большую часть действий по управлению памятью программисту приходится брать на себя. В том числе и по ее очистке. И это, на мой взгляд, является серьезным недостатком, так как усложняет и замедляет процесс разработки приложений. В общем, выбор, как всегда, только за вами.
Кроме того, у этих двух технологий есть еще одна общая черта, которой является широкий набор математических функций, которые у Nvidia CUDA содержатся в библиотеках CUBLAS и CUFFT, а у AMD Stream в библиотеке ACML. Нельзя не отметить также, что разрекламированные технологии ускорения работы c HD-контетом Nvidia PureVideo и ATI Avivo HD, которые обеспечивают плавное и качественное воспроизведение видеоданных высокой четкости, тоже основаны на возможности использования видеокарты для вычислений.
И в заключение нельзя не сказать пару слов о том, как же быть владельцам встроенных видеокарт от Intel и VIA/S3, ведь ни одна из вышеописанных технологий не будет работать с этими видеокартами. В этом случае тем, у кого встроенная видеокарта аппаратно поддерживает DirectX 9, можно попробовать использовать возможности Brook, разработанного в Стэндфордском университете. А для оценки производительности таких видеокарт и выигрыша от ускорения вычислений по сравнению с центральным процессором, а также их совместимости, можно воспользоваться тестом, который доступен для скачивания в Интернете по адресу сайт Этот тест является немного модифицированной версией (был добавлен графический интерфейс и результаты нескольких видеокарт для сравнения) теста из Стэндфордского набора Brook и совместим практически со всеми видеокартами, аппаратно поддерживающими DirectX 9, в том числе с Nvidia GeForce и ATI Radeon.
Виталий Сороко
Несмотря на то, что эти технологии были разработаны в компаниях, которые тесно конкурируют между собой на рынке графических решений, они обе имеют общую концепцию разработки приложений. Более того, данная концепция используется также в недавно появившемся стандарте OpenCL 1.0. Однако перед тем как перейти к непосредственному рассмотрению программного кода, необходимо овладеть базовыми понятиями общей технологии вычислений на видеокартах GPGPU. Поэтому далее мы рассмотрим основные из них на примере технологии AMD Stream (при разработке и тестировании программного кода использовалась версия 1.2), а потом покажем реализацию аналогичного программного кода в Nvidia CUDA (версии 2.0). Потоковая модель программирования GPGPU подразумевает наличие потоков данных (streams), которые необходимо обрабатывать с помощью специальной функции. Такая функция в Brook+ называется ядром (kernel). Потоки данных являются ничем иным, как массивами из элементов одинакового типа данных. Brook+ поддерживаются следующие типы данных: int - 32-битное целое число; float - 32-битное вещественное; double - 64-битное вещественное. Кроме того, допускается использование структур данных. Для описания потоков используется комбинация символов < >. Например: int a<> - одномерный поток (вектор) произвольной длины из целых чисел; int a<100> - одномерный поток (вектор) из ста чисел; int a<100,100> - двумерный поток (матрица) целых чисел размером 100х100; int a<100,100,100> - трехмерный поток целых чисел размером 100х100x100. В AMD Stream версии 1.2 существуют ограничения на размер потоков. Официально поддерживаются двумерные потоки до 8192х8192 и одномерные до 223 элементов. На практике у меня возникла ошибка "Stack overflow" уже при использовании матрицы размером 300x300. Возможно, это из-за нехватки мощности видеокарты или из-за использования не самой последней версии AMD Stream и не очень свежих драйверов. Будем надеяться, что AMD в новой версии 2.0 программного обеспечения Stream уже исправила это (мне не удалось протестировать это из-за несовместимости моей видеокарты). Теперь разберемся с ядрами. Итак, ядро предназначено для проведения вычисления над потоками. Исходящие потоки маркируются словом out. Например, следующее ядро суммирует 2 целочисленных потока a и b и выводит результат в потоке z:
kernel void sum(int a<>, int b<>, out float z<>)
{ z = a + b; }
Версия Brook+, входящая в состав AMD Stream 1.2, поддерживает до 128 потоков на входе и до 8 на выходе. К элементам потоков внутри ядра можно также обращаться по индексу как в обычных массивах. Кроме того, существуют также ядра отсечения (Reduction kernels), в которых можно применять условные операторы. Подробнее о них и о многом другом написано в файле brookspec из прилагающейся к AMD Stream документации. Настоятельно рекомендую с ней ознакомиться, так как там содержатся полные описания языковых синтаксических ограничений для ядер, и не только. Ну а теперь перейдем к рассмотрению примера программы на Brook+, исходный код и запускной файл которого можно скачать здесь: сайт Нижеописанный пример суммирует два потока вещественных чисел a и b и выводит результат в потоке z.
Листинг примера 1 (Файл sum.br):
#include "stdio.h"
//Ядро
kernel void sum(float a<>, float b<>, out float z<>)//Описание входящих и исходящих потоков
{
z = a + b; //Суммирование каждого элемента потока a и b c записью результата в поток z
}
//Основная функция
int main(int argc, char** argv)
{
int i, j; // переменные-счетчик цикла, который заполняет потоки a и b значениями
float a<10, 10>; //поток a - матрица из вещественных чисел размером 10x10
float b<10, 10>; //поток b - матрица из вещественных чисел размером 10x10
float z<10, 10>; //поток z - матрица из вещественных чисел размером 10x10
float input_a[10][10]; //матрица 10x10 для хранения данных потока a в оперативной памяти
float input_b[10][10]; //матрица 10x10 для хранения данных потока b в оперативной памяти
float input_z[10][10]; //матрица 10x10 для хранения данных потока z в оперативной памяти
//заполняем потоки a и b значениями
for(i=0; i<10; i++)
{
for(j=0; j<10; j++)
{
input_a[i][j] = (float) i;
input_b[i][j] = (float) j;
}
}
//Копируем данные (значения элементов) потоков а и b из оперативной памяти в память видеокарты
streamRead(a, input_a);
streamRead(b, input_a);
//запускаем ядро (суммирование элементов a и b c записью результата в поток z)
sum(a, b, z);
//Извлекаем полученный в потоке z результат из памяти видеокарты обратно в оперативную память
streamWrite(z, input_z);
//Выводим полученный результат на экран
for(i=0; i<10; i++)
{
for(j=0; j<10; j++)
{
printf("%6.2f", input_z[i][j]);
}
printf("\n");
}
return 0;
}
Как видно из примера, потоки a, b, z являются матрицами из элементов типа float размером 10х10 элементов. Обратите внимание, что в функции main, в отличие от ядра, указывается размерность потоков. Кроме того, из-за особенностей работы видеокарты все значения элементов потоков перед началом суммирования в ядре должны быть сначала заданы в оперативной памяти компьютера, а затем, с помощью функций streamRead (для входящих потоков a и b) и StreamWrite (для исходящего потока z), помещены и во внутреннюю память видеокарты или считаны оттуда. В результате выполнения кода программы каждый элемент матрицы а будет просуммирован c соответствующим по индексу элементом матрицы b, а результат будет записан в соответствующую позицию матрицы z, после чего результат будет выведен на экран. Однако как же откомпилировать данную программу? Для этого можно использовать интегрированную среду разработки Microsoft Visual Studio 2005, и следующий абзац посвящен тому, как это удобнее всего сделать. Итак, перейдем к рассмотрению процесса компиляции Brook+ программ с помощью Visual Studio 2005. Надеюсь, что к этому моменту на вашем компьютере уже установлены AMD Stream SDK и сама Visual Studio 2005, и вы уже скачали исходный код примера sum.br. Для начала зайдем в папку, где находятся примеры программ Brook+, которая по умолчанию для AMD Stream версии 1.2 находится в C:\Program Files\AMD\AMD BROOK 1.2\samples. Эта папка имеет в своем составе папки apps и bin. Папка apps содержит исходные коды различных примеров, а в папке bin хранятся откомпилированные из этих исходных кодов программы. После каждой компиляции исполняемые файлы проекта записываются именно в папку bin, так как таковы настройки проектов примеров, входящих в состав AMD Stream SDK. Если вы не специалист по параметрам проекта, то рекомендую писать ваши приложения на базе готовых примеров, как это будет показано ниже.
Исполняемый файл будет находиться в папке bin\xp_x86_32 (при условии использования системы Windows XP 32-bit) под именем binary_search_d.exe. Поэтому в случае необходимости (например, для просмотра результата) его можно запустить из командной строки прямо оттуда. Обратите внимание, что в результате выполнения сначала был динамически сгенерирован файл binary_search.cpp, и только после этого из этого сгенерированного файла был получен путем компиляции исполняемый файл binary_search_d.exe. Такова особенность работы Brook+. Однако если вам понадобится использовать код в каком-либо своем приложении, то для интеграции необходимо использовать код именно из cpp-файла. Кроме того, исполняемые файлы требуют для запуска наличие уже установленного AMD Stream SDK, так как в него входят необходимые для запуска библиотеки. Теперь, зная, как работает технология программирования AMD Stream, перейдем к рассмотрению аналогичной конкурирующей технологии Nvidia CUDA, которая, несмотря на общую концепцию программирования, имеет целый ряд отличий. Первым отличием от AMD Stream является то, что примеры программ и полезная документация содержатся в разных инсталляционных файлах - Nvidia CUDA SDK и CUDA Toolkit соответственно. При этом, для удобства, рекомендую произвести установку этого всего в одну и ту же папку. Примеры программного кода при этом будут содержаться в папке projects, а вся полезная документация — в папке doc. Открыв любой из проектов примеров программного кода, вы, скорее всего, заметите второе основное отличие от AMD Stream - описание ядра и основной код содержатся в двух разных файлах, которые имеют расширение '.cu'. Ну а самое главное отличие, конечно же, состоит в реализации программного кода. Поэтому перейдем к рассмотрению непосредственно кода, чтобы вы сами смогли оценить это. С целью экономии места приведу лишь самые важные строки кода, без полного листинга.
Как и в случае с AMD Stream, нижеописанный пример суммирует два потока вещественных чисел a и b и выводит результат в потоке z. Следующий код показывает описание ядра и процедуры его вызова из основной программы:
//Ядро (Обычно находится в отдельном файле)
__global__ void Sum(float A[10] [10], float B[10] [10], float Z[10] [10])
{
int i = threadldx.x;
int j = threadldx.y;
Z[i] [j] = A[i] [j] + B[i] [j]; //Суммирование каждого элемента потока a и b c записью результата в поток z
}
//Основная функция (Также обычно в отдельном файле)
int main()
{
dim3 dimBlock(10, 10);
//запускаем ядро (суммирование элементов a и b c записью результата в поток z)
Sum<<<l, dimBlock>>> (А, В, Z);
}
А теперь посмотрите, как в CUDA происходит работа с памятью:
// Выделение места в оперативной памяти для матриц a и b
unsigned int size_A = 10 * 10;
unsigned int mem_size_A = sizeof(float) * size_A;
float* h_A = (float*) malloc(mem_size_A);
unsigned int size_B = 10 * 10;
unsigned int mem_size_B = sizeof(float) * size_B;
float* h_B = (float*) malloc(mem_size_B);
// Инициализация матрицы a и b
randomInit(h_A, size_A);
randomInit(h_B, size_B);
// Выделение для матриц a и b места в памяти видеокарты
float* d_A;
CUDA_SAFE_CALL(cudaMalloc((void**) &d_A, mem_size_A));
float* d_B;
CUDA_SAFE_CALL(cudaMalloc((void**) &d_B, mem_size_B));
// Копируем данные (значения элементов) потоков а и b из оперативной памяти в память видеокарты
CUDA_SAFE_CALL(cudaMemcpy(d_A, h_A, mem_size_A,
cudaMemcpyHostToDevice) );
CUDA_SAFE_CALL(cudaMemcpy(d_B, h_B, mem_size_B,
cudaMemcpyHostToDevice) );
// Выделение памяти видеокарты для результирующей матрицы z
unsigned int size_Z = 10 * 10;
unsigned int mem_size_Z = sizeof(float) * size_Z;
float* d_Z;
CUDA_SAFE_CALL(cudaMalloc((void**) &d_Z, mem_size_Z));
// Выделение места в оперативной памяти для результирующей матрицы z
float* h_Z = (float*) malloc(mem_size_Z);
// Извлекаем полученный в потоке z результат из памяти видеокарты обратно в оперативную память
CUDA_SAFE_CALL(cudaMemcpy(h_C, d_C, mem_size_C,
cudaMemcpyDeviceToHost) );
// Высвобождение занятой памяти
free(h_A);
free(h_B);
free(h_C);
free(reference);
CUDA_SAFE_CALL(cudaFree(d_A));
CUDA_SAFE_CALL(cudaFree(d_B));
CUDA_SAFE_CALL(cudaFree(d_C));
Как видно из вышеописанного кода, концепция программирования в Nvidia CUDA полностью идентична таковой в AMD Stream. Однако в отличие от AMD Stream, в CUDA большую часть действий по управлению памятью программисту приходится брать на себя. В том числе и по ее очистке. И это, на мой взгляд, является серьезным недостатком, так как усложняет и замедляет процесс разработки приложений. В общем, выбор, как всегда, только за вами.
Кроме того, у этих двух технологий есть еще одна общая черта, которой является широкий набор математических функций, которые у Nvidia CUDA содержатся в библиотеках CUBLAS и CUFFT, а у AMD Stream в библиотеке ACML. Нельзя не отметить также, что разрекламированные технологии ускорения работы c HD-контетом Nvidia PureVideo и ATI Avivo HD, которые обеспечивают плавное и качественное воспроизведение видеоданных высокой четкости, тоже основаны на возможности использования видеокарты для вычислений.
И в заключение нельзя не сказать пару слов о том, как же быть владельцам встроенных видеокарт от Intel и VIA/S3, ведь ни одна из вышеописанных технологий не будет работать с этими видеокартами. В этом случае тем, у кого встроенная видеокарта аппаратно поддерживает DirectX 9, можно попробовать использовать возможности Brook, разработанного в Стэндфордском университете. А для оценки производительности таких видеокарт и выигрыша от ускорения вычислений по сравнению с центральным процессором, а также их совместимости, можно воспользоваться тестом, который доступен для скачивания в Интернете по адресу сайт Этот тест является немного модифицированной версией (был добавлен графический интерфейс и результаты нескольких видеокарт для сравнения) теста из Стэндфордского набора Brook и совместим практически со всеми видеокартами, аппаратно поддерживающими DirectX 9, в том числе с Nvidia GeForce и ATI Radeon.
Виталий Сороко
Компьютерная газета. Статья была опубликована в номере 07 за 2010 год в рубрике программирование
