Далее, кто хочет уже опробывать CUDA, выкладываю ссылки на скачивание первой версии, 2й бета и версии для маков и линукс.

Качаем:

CUDA 1.0

Существенная информация из анонса делится на две кучки, "сейчас и потом". Сейчас:
Асинхронное выполнение одновременно с копированием порции данных доступно только на "архитектуре 1.1 (g84/g86/g92)".... получается, что на своей G80 (8800GTX) я этого счастья лишен и это аппаратное ограничение.
Пересылку данных в карту можно делать асинхронно.
Поддержка 64-битных Windows, чего я ждал 11 месяцев
Непонятная мне "Graphics interoperability with CUDA across Multiple GPUs is enabled", это имеется в виду, карты в SLI-режиме ?

Потом:
поддержка double precision hardware, которого формально еще нет и хочется надеяться, что это не будет только Tesla;
Vista (сроки не объявлены, по идее, все ограничивается драйвером);
Mac OS X, обещают бету к январю (а у меня в макбуке-про 8600M ;
развитые средства отладки и профайлинга;
3D-текстуры;
оптимизированная пересылка данных карта-карта (SLI-режим, мимо процессора?).

For Windows:
CUDA Toolkit version 1.0 for Windows XP - x86
CUDA SDK version 1.0 for Windows XP - x86
Windows Display Driver version 162.01 for CUDA Toolkit version 1.0 - Download
For Linux:
CUDA Toolkit version 1.0 for Redhat Enterprise Linux 3.x - x86 | x86-64
CUDA Toolkit version 1.0 for Redhat Enterprise Linux 4.x - x86 | x86-64
CUDA Toolkit version 1.0 for Redhat Enterprise Linux 5.x - x86 | x86-64
CUDA Toolkit version 1.0 for SUSE Linux Enterprise Desktop 10 - x86 | x86-64
CUDA Toolkit version 1.0 for OpenSUSE 10.1 - x86 | x86-64
CUDA Toolkit version 1.0 for OpenSUSE 10.2 - x86 | x86-64
CUDA SDK version 1.0 for Linux - Download
Linux Display Driver version 100.14 for CUDA Toolkit version 1.0 - x86 | x86-64
For MacOS:
CUDA for Rocks Cluster Management: Complete CUDA Rocks Roll with driver, toolkit, and SDK - Download


CUDA 1.1

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

Впрочем, общий драйвер, который для всех пользователей, но с поддержкой CUDA, все еще в бете, поэтому распространять собственный софт с CUDA все еще неудобно. Но ждать, по всей видимости, недолго.

Из интересных проектов с CUDA увидел два:
Texture Tools от Гугла, которые обещают офигенно быструю texture compression, аж в 12.5 раз быстрее. Гейм-девелоперам будет приятно.
PyStream - Python interface to CUDA. Жизнь там пока вялая, но идея интересная.

For Windows:
CUDA Toolkit version 1.1 for Windows XP - x86 | x86-64
CUDA SDK version 1.1 for Windows XP - x86 | x86-64
NVIDIA Driver for Microsoft Windows XP with CUDA Support (169.21) - x86 | x86-64
For Linux:
CUDA Toolkit version 1.1 for Fedora 7 - x86 | x86-64
CUDA Toolkit version 1.1 for Redhat Enterprise Linux 3.x - x86 | x86-64
CUDA Toolkit version 1.1 for Redhat Enterprise Linux 4.x - x86 | x86-64
CUDA Toolkit version 1.1 for Redhat Enterprise Linux 5.x - x86 | x86-64
CUDA Toolkit version 1.1 for SUSE Linux Enterprise Desktop 10-SP1 - x86 | x86-64
CUDA Toolkit version 1.1 for OpenSUSE 10.1 - x86 | x86-64
CUDA Toolkit version 1.1 for OpenSUSE 10.2 - x86 | x86-64
CUDA Toolkit version 1.1 for Ubuntu 7.04 - x86 | x86-64
CUDA SDK version 1.1 for Linux - Download
NVIDIA Driver for Linux with CUDA Support (169.09) - x86 | x86-64
NVIDIA Driver for Linux with CUDA (171.05) specifically for Tesla S870 1U System - x86 | x86-64
For MacOS:
Beta CUDA for Mac OS X (10.5.2) Toolkit - Download
Beta CUDA for Mac OS X (10.5.2) SDK - Download


CUDA 2.0 Beta

Из важного
Поддержка Vista (32 и 64 бита);
Нет поддержки GeForce 9800GTX (вышедшей на пару недель раньше этой беты), что довольно странно.
С двойной точностью какая-то непонятная совсем история:
В CuBLAS она заявлена в документации, символы в библиотеке имеются (собирать еще ничего не пробовал).
В документации (programming guide) слово double встречается 8 раз (на 99 страниц текста), что как-то безобразно мало.
Времена вычислений для double в соответствующей секции не описаны (но я подозреваю, что они другие, чем у float).
Таблица с описанием double-функций (на которую есть указание в тексте) - отсутствует.
Другими словами, работы ведутся и довольно скоро все может появиться.

For Windows:
NVIDIA Driver for Microsoft Windows XP with CUDA Support (174.55) - x86 | x86-64
NVIDIA Driver for Microsoft Windows Vista with CUDA Support (174.55) - x86 | x86-64
CUDA Toolkit version 2.0 for Windows XP - x86 | x86-64
CUDA Toolkit version 2.0 for Windows Vista - x86 | x86-64
CUDA SDK version 2.0 for Windows XP - x86 | x86-64
CUDA SDK version 2.0 for Windows Vista - x86 | x86-64
For Linux:
NVIDIA Driver for Redhat Enterprise Linux 4.x with CUDA Support (174.55) - x86 | x86064
NVIDIA Driver for Redhat Enterprise Linux 5.x with CUDA Support (174.55) - x86 | x86-64
CUDA Toolkit version 2.0 for Redhat Enterprise Linux 4.x - x86 | x86-64
CUDA Toolkit version 2.0 for Redhat Enterprise Linux 5.x - x86 | x86-64
CUDA SDK version 2.0 for Linux - Download

Один пример программы, которая использует технологию CUDA.

Video Memory stress Test версии 1.5

Вышла новая версия 1.5 сборка 112 теста видеопамяти Video Memory stress Test. В новой версии добавлен экспериментальный режим тестирования с большей нагрузкой на видеопамять посредством технологии универсальных вычислений NVIDIA CUDA, также приложение отныне перед началом теста временно деактивирует экранную заставку и энергосберегающие технологии воизбежание прерывания теста. Загрузить последнюю версию программы можно с нашего сайта (~ 0.6 Мб).


Да и ещё.
Программирующим на CUDA может быть интересно: NVidia начала раздавать исходники библиотек CUBLAS/CUFFT.

Я, правда, не очень понимаю статус этого дела:

• С одной стороны, все выложено на девелоперском сайте, куда нужна регистрация (и говорят, что стоит большая очередь желающих оной регистрации, хотя меня в прошлом году зарегистрировали за один день).
• С другой стороны, в девелоперской рассылке пришли ссылки на незапароленый сайт, бери кто хочет.
  Посему, ссылки не публикую, если кому-то нужно и нет терпения ждать (со временем все попадает в полностью открытый доступ, всегда так было) - пишите лично.

А вот что точно открыто всем желающим, так это визуальный профайлер (beta) для той же CUDA. Пока не смотрел, руки не дошли.

Приобщение к CUDA. Серия 1 - обзорная

Я давно обещал сделать сообщение с обзорным рассказом о технологиях чего-то вроде суперкомпьютинга на рабочем столе. Но так, что-то толком не вытанцовывается полноценный грамотный обзор с правильным сравнением, хорошим и подробным стилем изложения. Поэтому выложу то, что имеется и перейду к практическим вещам в виде CUDA от NVIDIA.


Вообще-то специальные устройства, позволяющие превратить свой настольный компьютьер во что-то куда более быстродействующее, чем он обычно является, известны довольно давно, ещё с середины 80-х годов. Такая платка от фирмы INMOS, называвшаяся транспьютером и вставленная в ISA-слот могла дать до 20-30 полноценных MIPS, что было довольно круто для того времени. Но в дальнейшем они стали отставать от роста быстродействия обычных процессоров, и в конце-концов INMOS свернула их производство, не сумев поддержать достойную производительность. Были и есть и другие устройства подобного рода, всякого рода специализированные ускорители, часто их сфера применения чрезвычайно узкая и ограничена единственной профессиональной программой. Мало кто их видел и мало где они продаются.

Но в последнее время наметилась совершенно явная тенденция к выпуску на массовый рынок более-менее универсальных устройств, в которых стоят многоядерные процессоры с параллельными векторными вычислениями, за счёт чего достигается реальная (а не маркетинговая) пиковая производительность около 200-300 GFlops. Внушительная величина, ещё всего 6 лет назад компьютеры с подобным быстродействием входили в Top-500 суперкомпьютеров. И хотя подобное быстродействие достигается не на всех задачах и имеется ряд других ограничений, тем не менее, сфера применения достаточно широкая и практически востребованная, чтобы впечатлиться. Особенно, если использовать не одно устройство, а сразу несколько.

Реально такими устройствами являются игровая приставка Sony Playstation-3, Xbox360 и более-менее совремённые видеокарточки от NVIDIA и от AMD (бывшие ATI). Нетрудно обратить внимание, что всё это в основном предназначено для игр и только для игр. На самом деле ничего удивительного, игры, пытающиеся моделировать окружающий мир, всегда были жадными до вычислительных возможностей. "Матрица" требует много вычислений ;-) Между прочим, в черновом варианте сценария к фильму, в Матрице люди использовались не как батарейки, а как элементы вычислительной системы. Жаль, что сценарий изменили, первоначальный вариант был, по-моему, логичнее и интереснее.

В настоящее время и приставки PS3 и видеокарточки позволят проводить вычисления общего назначения, а не только играть в игры. Исключение - Xbox, Microsoft тщательно оберегает покупателя приставки от использования её не по игровому назначению. Остальные производители, кто в большей, кто в меньшей степени, не только не препятствуют проводить произвольные вычисления, но даже поощряют подобный интерес.
Хотя, всё-таки, не без ограничений. Sony не даёт напрямую воспользоваться всеми компонентами PS3, Nvidia не собирается раскрывать низкоуровневые детали взаимодействия с процессором видеоплаты, а AMD хотя выполняя обещания открыть техническую документацию по программированию своих чипов, сделала крупные шаги к раскрытию спецификаций, всё ещё довольно много ключевых документов сохраняет закрытыми.

Вычислительными возможностями видеокарточек и PS3 уже воспользовалось немало людей, вот несколько примеров:


Кластер на видеокартах для астрофизических расчётов
Кластер на основе PS3 и тоже для астрофизических нужд
Рендеринг сложной трёхмерной сцены в реальном времени с помощью трёх приставок PS3
Использование для томографии четырёх двухпроцессорных Geforce 9800GX2 вместо суперкомпьютера за $4.6 миллиона долларов"

Во всех этих случаях были достигнуты результаты, раннее требовавшие суперкомпьютерных ресурсов и стоимости.
Кроме того, для PS3 и AMD Radeon серии 2xxx и 3xxx существует клиент для распределённых вычислений Folding@Home
В скором времени, его обещают написать и для карточек от nvidia.

Каким же образом обладатель приставки или видеокарточек может их запрограммировать для своих нужд?

Наиболее продвинутым в этом плане выглядит Sony Playstation-3, работающая на чипе Cell, производимом консорциумом из Sony, Toshiba и IBM. На сайте DevelopersWorks IBM размещено немало материалов по этому процессору и подробно объясняющих как писать для него программы, в том числе и для PS3. Хотя основное предназначение Cell не PS3, а высокопроизводительные сервера.
Начать знакомиться можно, например, с серии статей Программирование высокопроизводительных приложений на процессоре Cell BE, Часть 1: Введение в Linux на PLAYSTATION 3

В тоже время, я встречал ряд статей в интернете, где рассказывалось о проблемах в PS3 с производительностью вычислений, что выглядит это не настолько радужно как кажется поначалу. К сожалению, ссылка теперь битая, но в статье рассказывалось, что из-за проблем с хэшем, латентностью и объёмом памяти (256 Мб) действительно быстрые результаты получаются на относительно не очень больших объёмах данных.


Для видеокарточек несколько лет назад зародилась целая технология, названная GPGPU General-Purpose computation on GPUs
Первоначально эксплуатировалась идея, что пиксельные шейдеры в видеокарточках стали очень умными и быстрыми и неплохо бы их как-то заставить рассчитывать не только изображение в играх (или CAD'ах), но и что-нибудь поинтереснее. Поначалу использовались разные извращения с OpenGL и DirectX, когда в видеокарточку под видом текстур передаются данные, а под видом шейдеров функции для их обработки. Недостатком оказывается повышенная сложность программирования и невысокая скорость обмена данными между видеокарточкой и основной системой в направлении от видеокарточки к системе. Ведь в играх, фактически нет нужды передавать данные обратно.

В настоящее время сами производители Nvidia и AMD предлагают поддержку вычислений на всех уровнях: аппаратном, сделав процессоры видеокарточек более универсальными и ускорив обмен данными, на уровне драйверов, предоставив универсальные механизмы не завязанные на OpenGL или DirectX и на пользовательском уровне, предоставив библиотеки, компиляторы и SDK с примерами программ и документацией.

У Nvidia это технология CUDA (Compute Unified Device Architecture) http://www.nvidia.com/object/cuda_home.html, доступная на чипах G80 и G92 в GeForce 8-ой серии и выше. В настоящее время, используется CUDA v1.1, но одновремённо с ожидаемым в июне этого года выпуском новых более производительных и с новыми возможностями видеокарт на чипе GT280 и GT260 (названия рабочие), выйдет и CUDA v2.0 с новыми возможностями. В частности, предполагается поддержка вещественного типа c двойной точностью (64 бит, double), пока что поддерживается только одинарная точность (32 бит, float). Версии CUDA существуют для Linux, Windows и MacOS X, как для 32-битных, так и для 64-битных систем.

Аналогичная технология и даже появившаяся немного раньше, чем у Nvidia, имеется и у AMD (ATI). Это CTM (загадочно названная Close to Metall) или AMD Stream Computing. Работает на видеокартах Radeon HD 2-ой, 3-й и выше серии. Готовые версии существуют только для 32-х и 64-х битных Windows, хотя недавно была выпущена бета-версия и для Linux. По сравнению с CUDA может быть быстрее в некоторых случаях и поддерживает 64-битный вещественный тип, хотя я встречал высказывания, что в целом программировать в рамках AMD Stream Computing немного сложнее, чем CUDA. Однако подробностей я не знаю и вообще, так вышло, в том числе и из-за плохой поддержки Linux, что моё внимание более привлекла CUDA.

Немного об аппаратном строении.

Cell в PS3. Состоит из управляющего PPE + 8 SPE.

PPE (Power Processor Element), RISC-процессор, фактически PowerPC, работающий на частоте 3.2 Ггц. Одно ядро, но с двумя конвейерами, представляющими аналог HyperThreading но, вроде бы более эффективный. Исполняет в среднем одну команду за такт и вследствие этого, общая производительность на фоне совремённых x86 процессоров скромная. Соответствует примерно 1 - 1.5 Ггц Pentium-III или Athlon. Написание программы с учётом особенностей HyperThreading может увеличить скорость работы где-то на 50%

SPE - так называемые, синергические обрабатывающие элементы. В PS3 один из восьми заблокирован с целью увеличения выхода годных изделий (Cell) и снижения себестоимости PS3. Другой - используется в гипервизоре, не пускающем обычного программиста куда с точки зрения Sony, ему не надо лезть. Например, в видеопамять (но не только). Таким образом, для обычного программиста доступен PPE и 6 SPE. SPU - представляет собой векторный процессор (SPU), объединённый с очень быстрой локальной памятью 256 Кб, у каждого SPU - своя.

G80 (в GeForce 8800)

Очень кратко, состоит из

1) 128-ми потоковых (thread) процессоров, также называющихся программируемыми пиксельными шейдерами.
Каждый такой процессор, содержит блок вычислений с плавающей запятой (FPU) и 1024 32-битных регистра.
FPU позволяет проводить вычисления только с одинарной точностью (32-битный вещественный тип float в Си)

8 потоковых процессоров объединены в кластер. Всего имеется 16 таких кластеров.
Каждый кластер имеет локальную общую память, объёмом 16 Кб, поддерживающую параллельный доступ к данным.

2) Устройства аппаратного распределения исполняемых потоков(thread), которое автоматически раскидывает
потоки по процессорам, избавляя от необходимости писать подробный управляющий код.
Максимальное количество, конкурирующих за процессоры, потоков 12288.

3) Общей памяти. 768 Мб у GeForce 8800 GTX и Ultra.

4) Текстурного кэша и так сказать обвязки,

В других картах, поддерживающих CUDA может быть другое количества процессоров, памяти и разрядность шины памяти.

(c)

Приобщение к CUDA. Серия 2 - выбор, установка видеокарты и результаты Specviewperf

Выбор

Итак мне захотелось поиграться с чем-то быстросчитающим. И впрямом смысле поиграться тоже

Встал вопрос выбора. Что взять: приставку PS3 с Cell, карточки Radeon для AMD Stream Processing или GeForce для CUDA? И какие именно? Каждый вариант имеет свои плюсы и минусы. Мысль решить всё самым интересным способом: взять все три вещи, и потом сравнивать в своё удовольствие, была отброшена по причине нежелания расставаться с потребной для этого суммой.



Плюсы Sony Playstation 3. Сама технология Cell, активно продвигаемая IBM, выглядит многообещающей. Именно на Cell (который на PS3, впрочем в усеченном виде) сейчас начинают строиться суперкомпьютеры, предполагающие преодоление порога в петафлопс (квадриллион операций в секунду с плавающей запятой) Кстати, сейчас в России проводится конкурс по программированию на Cell. Cell более универсальный процессор, чем в видеокарточках. PS3 - это редкая сейчас после перехода Mac'ов на платформу Intel возможность приобрести на бытовом рынке новый процессор PowerPC архитектуры. Напомню, что Cell - фактически сочетание PowerPC-совместимого (по крайней мере на уровне команд) процессора и синергических модулей. Кроме того, это ещё и дешёвый плеер Blue-Ray Disc. PS3 может использоваться и как отдельный универсальный компьютер и как, разумеется, игровая приставка - главная цель выпуска.

Минусы Sony Playstation-3. Более высокая цена, по сравнению с видеокарточками. Базовой комплектации с 60Гб винчестером уже не продаётся, а в 40Гб кое-что выкинуто для удешевления, в том числе некоторые неигровые возможности. Винчестер, кстати, всё-равно для неигровых целей пришлось бы ставить более вместительный. К сожалению, в отличие от винчестера, нарастить оперативную память сверх лимита в 256Мб нереально, потому что это не обычная память DDR2 или DDR3, а Rambus XDR, более новая разработка, чем та что была в эпоху ранних Pentium-4 и на рынке просто так не продаётся. Игры для PS3 менее доступны и разнообразны, чем для простых компьютеров, даже если и там и там приобретать исключительно лицензионную продукцию. Программиста работающий гипервизор всё-таки стесняет, в том числе и в быстродействии. Как я уже говорил в предыдущей части, имеются некоторые проблемы с реальной производительностью на PS3. То есть, выжать под 200Gflops всё-таки можно, но с более неприятными ограничениями, чем казалось раньше.

Плюсы видеокарточек. Дешевле PS3, больше памяти, причём и скорость работы с памятью больше и латентность меньше, да и скорость обмена данными между основной памятью в компьютере и памятью в видеокарточке составляет от 1 до 5 Гб/сек, в зависимости от вариантов. Производительность видеокарточек совсем не хуже, чем у PS3 и оказывается, может быть даже больше. Тем более, готовящиеся к выпуску новые видеоплаты от Nvidia и AMD уже точно будут существенно быстрее устаревающего PS3. Правда не исключено, что Sony выпустит какую-нибудь новую более быструю модификацию, но подобный шаг маловероятен, поскольку вряд ли вписывается в рыночную модель использования приставок. Скорее уж следует ожидать PS4, но для нового поколения PS3 слишком мало времени продавалась.

Минусы видеокарточек. Много выделяют тепла В отличие от PS3 видеокарточки не самодостаточны, круг решаемых задач уже чем у Cell - на них нельзя сделать компьютер, нужен отдельный общий процессор и для работы необходимо дополнительное ПО, программирование более закрытое, чем у Cell. В отличие от PS3 в них не имеется ограничивающего гипервизора, но зато для работы требуются закрытые модули, транслирующие промежуточный код в коды непосредственно микропроцессора карточки и описание работы на низком уровне отсутствует у Nvidia и пока не раскрыто (если они вообще собираются раскрывать) у AMD. Правда, благодаря фирменным SDK и архитектуре, само воплощение на них параллельных алгоритмов смотрится проще.

Ещё о PS3. Стоит сама приставка у нас от 15-16 тыс. рублей (примерно $660), к ней ещё нужен или хороший большой телевизор или монитор с поддержкой HDCP или специальный переходник. Переходник, который назло копирастам разводит HDMI с HDCP на обычные аналоговые или DVI (без шифрования) видеокабели и позволяет подключить в том числе и обычный монитор, стоит существенно дешевле нового телевизора, но всё равно требует примерно 6 тыс. руб. А без переходника нельзя получить изображение с разрешением более чем 720x576 точек, с которым работать на мониторе не слишком комфортно. Итого с приставкой примерно 22 тысячи выходит. И при наличии более дешёвых альтернатив и других минусов приставки её приобретение показалось менее интересным среди других вариантов. Хотя может ещё и возьму

Поэтому от PS3 я отказался и практически выбирал между Radeon HD 3870, GeForce 8800 (GTX,GT,GTS g80 и GTS от G92), GeForce 9800 GTX и ещё не выпущенными, но анонсируемыми в июне, новыми поколениями видеокарточек как от Nvidia, так и от AMD(ATI).

Ждать выхода новых видеокарточек я не стал, несмотря на все обещаемые вкусности. Потому что в Москве они поначалу будут продаваться вряд ли раньше августа, а поработать с ними хочется уже сейчас, да ещё и как водится продаваться будут, скорее всего, по невменяемым ценам, которым ещё несколько месяцев предстоит опускаться до адекватного уровня.

Взял то, что сильно захотелось: GeForce 8800 GTX с заводским разгоном Должен отметить, что AMD/ATI-шные карточки слегка (и то вопрос) проигрывая в играх топам от nvidia, показывают на ряде тестов закономерно более высокую общую вычислительную производительность, например, у Radeon HD 3870 имеется 320 потоковых процессоров и 512-битная шина против 128 потоковых и 384 бит у GeForce 8800 GTX и Ultra, хотя потоковые процессоры у nvidia и ati и нельзя напрямую сравнивать из-за их разной производительности.

Правда уже после покупки карточки возникла мысль, что возможно более удачным приобретением была бы GeForce 8800 GTS 512 Мб или 8800 GT с теми же 512 Мб на новом чипе G92 (8800 GTX на G80). Первая на 5-10% быстрее, а вторая слегка отставая, существенно дешевле. И обе поддерживают некоторые дополнительные, хотя и не принципиально важные, функции (AtomicFunction). Имеют слегка меньшие размеры и тепловыделение. С другой стороны, у них меньше памяти (512Мб вместо 768Мб) и уже шина: 256 бит вместо 384. А GeForce 9800 GTX мне показалось, что брать особого смысла не имеет, при более высокой цене, чем 8800 GTX, он фактически не более быстр, чем новый 8800 GTS.

Установка

Насколько мне известно, различий между GeForce 8800 GTX от разных производителей почти нет. Все они сделаны из референсной платы от NVIDIA и отличаются, в основном, дополнительным содержимым и размером коробки в которой поставляются, а также раскраской самой платы. Ещё наличием или отсутствием заводского разгона и есть модели с заменённой системой охлаждения. Поэтому размер в 27 с лишним сантиметров и два занимаемых слота - это их родовой признак (говорят, что новые платы не будут короче). Влезут не во всякий корпус и требуют хорошего блока питания, как указывает nvidia, не менее 450 честных ватт и суммарных 30A на двух линиях 12В, которые вставляются в два шестиштырьковых разъёма на плате.

Вот как она выглядит, вставленной в корпус моего компьютера:



IDE-шнурок ни к чему не подключён, но я вынужден его так оставить из-за расположения разъёма на плате и его ориентации штырьками вдоль платы, а не вверх. Вследствие чего, если я его сейчас выну, то захотев подключить IDE-винчестер, обратно его уже не вставлю без разбора понатыканного.
Сзади во втором слоте PCI-E находится GeForce 7600GS от Gigabyte. Первоначально я планировал оставить её в первом слоте, а новую плату поставить во второй, тем более, что учитывая вырезы на боковой стенке корпуса, там она, по-моему и лучше бы охлаждалась.
Блока питания Chieftec 550Вт хватило для платы, но вот в корпус она едва влезла:



Пришлось ножницами по металлу откромсать выступающую часть над отсеком для винчестеров, я не совсем понял зачем она вообще была предусмотрена, но возможно для 2.5 дюймовых винчестеров. Этого оказалось недостаточно. Двухслотовая плата уперлась и вот что-то у задней стенки корпуса. Оказалось виновато оригинальное крепление: сверху над задними панелями для их фиксации предусмотрен ряд пластмассовых защёлок, которые поворачиваются, поднимаются вверх и после установки платы, снова опускаются и закрываются, фиксируя планку. Так вот между слотами сверху расположены пластмассовые перегородки. Пришлось и их выкусить и только тогда видеокарточка влезла в компьютер.

Охлаждение

Видеокарточка недаром называется ПЕЧ, если GTX набрать в кириллической раскладке Она довольно таки много потребляет энергии даже в простое и доходит до 150 Вт и даже более при полной загрузке. На ней стоит штатный кулер, выбрасывающий поток тёплового воздуха за пределы корпуса компьютера, через радиаторную решётку во втором слоте. Разумеется, часть тепла всё-равно остаётся внутри корпуса и пространство вокруг платы заметно нагревается. В общем, в плохо проветриваемом корпусе эту видеокарточку не стоит держать.

Добавляет проблем и поведение штатной системы охлаждения. По умолчанию, она настроена на автоматическое изменение скорости вращения встроенного вентилятора зависимости от её загруженности. Где-то я читал и подтвердил на собственном опыте, что эта скорость вращения установлена в 60% от максимальной, что соответствует примерно 1720 об/мин. и при разогреве почти не увеличивается и только при температуре за 86 градусов (Цельсия) обороты реально растут, что выглядит не допустимым. Положение осложняется наличием заводского разгона и как следствие, повышенным нагревом.

К счастью, скоростью вращения вентилятора платы можно управлять вручную. В Linux для этой цели можно использовать консольную программу nvclock. nvclock создавался для разгона или тонокой подстройки видеоплат от Nvidia и управления такими параметрами как частота работы памяти, GPU и др., умеет и устанавливать частоту вращения кулера.

Следует знать, что страница man к этой программе (по крайней мере в Debian) не полная и в ней не документирован параметр -F, задающий процент мощности работы кулера.

Как видно, в режиме простоя, от повышения мощности кулера c 60% до 82% температура снизилась примерно на 7 градусов. Хотя шум от кулера увеличился, но не слишком сильно.
Тем не менее, для интенсивных вычислений надо ставить мощность сразу на 100% и всё-равно у меня температура подбирается к 80 градусам.
Экспериментируя с кулером необходимо помнить, что при его остановке видеокарточка может очень быстро сильно нагреться, особенно в ходе работы.
Можно посоветовать записать вызов nvclock с нужными параметрами в стартовые скрипты, чтобы после включения компьютера вентилятор сразу вращался на повышенных оборотах.

Драйвера и производительность

Приобретение новой видеокарточки совпало с выпуском NVIDIA новых драйверов для Linux 173.14.05. Тем не менее, поначалу я использовал старые драйвера 169.12.
Приятной неожиданностью стало то, что установка новой видеокарты и переезд старой в другой слот, система Debian GNU/Linux Testing (Lenny) восприняла спокойно, не потребовав никаких изменений настроек и внешне вообще не заметно, если не считать увеличения fps в играх и тестах. Попугаи в glxgears увеличились с 6500 до 18000 с чем-то. Интересно как в Windows бы прошло? Драйвера понадобилось бы переустанавливать или нет? Хотя, вообще-то в Debian приходится переустанавливать драйвер от nvidia практически каждый раз как меняется ядро или версия x.org
Но раз новые драйвера вышли я их всё-таки поставил.
Результат оказался интересным. Попугаи в glxgears уменьшились сразу на 2000 до 16500 примерно.
Протестировал я старую и новую видеокарточки и тестом Specviewperf, ориентированным на оценку быстродействия в профессиональных CAD-системах. Благо, что недавно вышла бета-версия SpecviewPerf 10 alpha 12 для Linux.



Очень забавно, но на некоторых тестах Geforce 7600GS оказалась даже чуть-чуть быстрее, чем 8800 GTX!
Практически получается, что для многих CAD-систем переход на 8800 GTX не оправдан, потому что они или не получают заметного прироста
или несмотря на прирост fps, все-равно слишком медленно работают и для них необходимо приобретать профессиональную серию Quadro.

Правда я сильно подозреваю, что особых различий в железе между игровыми платами и дорогими профессиональными не существует, если не считать увеличенного объёма памяти. Уж основной чип точно тот же самый и скорее всего, можно было бы заставить работать 8800 GTX в CAD'ах на уровне Quadro, если подкорректировать настройки в BIOS и в драйверах. Эти настройки в игровых платах явно ориентированы на функции OpenGL, более применяемые в игровых приложениях, а в Quadro - на компьютерное моделирование.

В тоже время, похоже, что CUDA в этом смысле универсальна, потому что не имеет заточки под какие-то конкретные функции.

(c)

Приобщение к CUDA. Серия 3 - установка и настройка Nvidia Toolkit и SDK в Debian

Чтобы иметь возможность создавать новые программы, использующие технологию CUDA,
необходимо установить как минимум NVIDIA CUDA Toolkit, состоящий из набора
фирменных библиотек и фирменного компилятора nvcc. nvcc компилирует код,
непосредственно исполняющийся на видеокарточке.



Крайне желательно установить ещё и NVIDIA CUDA SDK, в котором содержится немало
готовых примеров с исходными текстами, имеющими ценность как для изучения
программирования для CUDA, так и представляющими самостоятельный интерес для
встраивания в свои программы. Надо только не забывать внимательно читать
условия лицензии, носящей схожий с BSD смысл, но имеющей свои особенности.

И Toolkit и SDK свободно доступны на сайте Nvidia в
разделе Get Cuda
Перед установкой, в целях уменьшить количество возможных вопросов, крайне
желательно скачать и внимательно прочитать файлы:
CUDA_Toolkit_Release_Notes.txt и CUDA_SDK_release_notes_linux.txt.

Предварительная подготовка

Итак задача: на 64-битную систему Debian Lenny (Testing) по состоянию на начало
июня этого года необходимо поставить Nvidia Toolkit, SDK, откомпилировать и запустить
примеры из SDK.

На сайте Nvidia предлагается выбор различных пакетов для разных дистрибутивов.
Debian'а среди них нет, поэтому я выбрал пакет для Ubuntu 7.04 x86_64 как
наиболее близкий к Debian'у. Обобщая свой опыт и результат поиска информации
по разным форумам, могу сказать следующее.

Перед установкой надо убедиться, что в системе установлены:

• драйвера 169.09 или более старшие. Причём я находил намёки, хотя и не
  проверял, что установка вместо оригинальных пакетов с драйверами от Nvidia,
  устанавливаемых из репозиториев дистрибутивов, может привести к проблемам
• gcc-4.1. Для Nvidia Toolkit нужна именно версия 4.1, если стоит другая
  можно, не удаляя других версий, дополнительно поставить gcc-4.1. Остальные
  два пакета нужны для компиляции примеров в SDK
• пакет build-essential. Вообще-то он нужен для сборки deb-пакетов, но похоже
  содержит ещё что-то полезное для сборки и без него кажется компиляция правильно
  не работает компиляция SDK
• пакет libglut3-dev, по зависимостям вытягивает freeglut3-dev.

К сожалению, в компиляторе nvcc жёстко прописан вызов именно /usr/bin/gcc ,
поэтому правок Makefile и прочих конфигурационных файлов недостаточно для
правильной работы. Можно вручную с помощью ln -s создать нужные символические
ссылки /usr/bin/gcc (а также g++ и cpp) на файлы версии 4.1, но я использовал
для этой цели механизм update-alternatives, который позволяет централизованно
управлять символьными ссылками, в том числе и группами ссылок.

Теперь командой

можно выбрать нужную версию gcc

Есть 2 альтернатив, которые предоставляют `gcc'.

Выбор Альтернатива
-----------------------------------------------
1 /usr/bin/gcc-4.1
*+ 2 /usr/bin/gcc-4.2

Нажмите enter, чтобы сохранить значение по умолчанию[*], или введите выбранное число: 1
Используется `/usr/bin/gcc-4.1' для предоставления `gcc'.

Теперь gcc -v показывает, что используется версия 4.1

Маленькое пояснение: командами update-alternatives для группового
наименования gcc устанавливается ссылка /usr/bin/gcc, указывающая на одну из
альтернатив /usr/bin/gcc-4.1 имеющую приоритет 413 или /usr/bin/gcc-4.2 с
приоритетом 424, каждая из которых имеет подчинённые групповые наименования
g++ и cpp с соответствующими ссылками, изменяющимися одновремённо с изменением
основной gcc. То есть версии g++ и cpp меняются одновремённо с gcc.
Сразу автоматически устанавливается версия с большим приоритетом, в качестве
которого был выбран номер версии. В дальнейшем можно легко переключать версии.

Установка и настройка Nvidia Toolkit

После запуска на исполнение файла NVIDIA_CUDA_Toolkit_1.1_Ubuntu7_x86_64.run
установщик, являющийся фактически просто самораспаковывающимся архивом, спросит
каталог для распаковки содержимого. В указанном каталоге будет автоматически
создан подкаталог cuda со всем содержимым.
По умолчанию предлагается путь /usr/local, если его выбрать Toolkit будет
установлен в /usr/local/cuda. Я советую не менять путь по умолчанию, иначе в
дальнейшем частенько придётся править конфигурационные файлы в проектах в
которых этот путь жёстко задан.
Я вначале этого не понял и установил в итоге в /usr/share/cuda

Затем, как написано в Release notes, необходимо поместить в PATH каталог bin
из cuda, то есть, в моём случае, /usr/share/cuda/bin и добавить библиотеки
из /usr/share/cuda/lib в путь для поиска библиотек при компиляции программ.

В SDK Release Notes для этой цели советуют добавить в bash_profile строки

или для библиотек модифицировать /etc/ld.so.conf.d и выполнить ldconfig

Я сделал немножко по-другому. В .bashrc появилась строчка
export PATH="$PATH:/usr/share/cuda/bin"

а в каталоге /etc/ld.so.conf.d/ добавился файл cuda.conf
(в Debian'е в /etc/ld.so.conf уже прописан include /etc/ld.so.conf.d/*.conf)

и вызвал ldconfig для построения кэша и привязок для библиотек.

Установка и компиляция Nvidia CUDA SDK

Также как и Toolkit файл NVIDIA_CUDA_SDK_1.1_Linux.run является фактически
самораспаковывающимся архивом, чьё содержимое предлагается по умолчанию
установить внутри домашнего каталога, создав каталог NVIDIA_CUDA_SDK и
всё распаковав туда. Именно туда я всё и поставил.

Для успешной компиляции пришлось только подправить файл
~/NVIDIA_CUDA_SDK/common/common.mk заменив в строчке
CUDA_INSTALL_PATH ?= /usr/local/cuda слово local на share, из-за
установки Toolkit не в каталог по умолчанию.

Теперь можно перейти в каталог ~/NVIDIA_CUDA_SDK/ и откомпилировать
все примеры. Команда make emu=1 соберёт в каталоге
~/NVIDIA_CUDA_SDK/bin/linux/emurelease программы, работающие в режиме эмуляции
CUDA, т.е., без реального использования процессора видеокарточки, а make без
параметров соберёт в каталог NVIDIA_CUDA_SDK/bin/linux/release реально
использующие ресурсы видеокарточки.

Пара скриншотов



На данном скриншоте изображена визуализация расчёта программой nbody задачи
движения N (здесь N=16384) тел в гравитационном поле (и не только гравитационном).



Скриншот не видео и не передаёт движение, но здесь изображено рассчитываемое движение
жидкости из 512x512 элементов. Можно мышкой как бы дать толчок, в том числе
закрутить жидкость и наблюдать перемещающие вихри.

(c)

О пирамидальном сложении на параллельной архитектуре

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

Для второго шага reduce обычно используется пирамидальная схема: сначала в N/2 потоков сложим N результатов попарно, затем сложим N/2 в N/4 и так далее. Число итераций равно, очевидно, log2N.

Возникает вопрос, «сколько данных складывать на каждой итерации?» Ведь можно складывать в N/4-N/16-N/256 кучек, можно по 1/8-64-512 и так далее. Из общих соображений, складывать по несколько лучше чем по два. Конечно, потоков получается меньше, но меньше и оверхед на создание-завершение потока.

Для NVidia CUDA идея "делать не по 2", выбирая все динамически, оказалась очень плохой. Да, с одной стороны мы действительно имеем оптимум при сложении по 8 или по 16. С другой стороны, код для вычисления содержит больше условного исполнения, отчего все ухудшается:

• С одной стороны, "умный" код для сложения по 8 примерно втрое быстрее, чем при сложении им же "по 2".
• С другой стороны, код рассчитанный только на сложение по 2 - в полтора раза быстрее чем умный складывает по восемь

Приведу на всякий случай быстрый код, взятый из примера scalarprod (CUDA SDK)

Здесь вложенный for - это на самом деле такой IF, который для рассматриваемого случая, когда N == blockDim.x, удовлетворяется только для части потоков, а выполняется тело цикла для этих потоков только один раз. Свой "умный" код не привожу, слишком уж умный.

Впрочем, есть и специальные случаи. Если известно, что количество складываемых элементов является степенью четверки, то ручное написание кода (без лишних IF) дает еще примерно полуторакратный рост производительности. Вот код:

Да, дальнейший unroll дает падение производительности.

(c) Алексей Тутубалин

CUDA: Как работает GPU

Внутренняя модель nVidia GPU – ключевой момент в понимании GPGPU с использованием CUDA. В этот раз я постараюсь наиболее детально рассказать о программном устройстве GPUs. Я расскажу о ключевых моментах компилятора CUDA, интерфейсе CUDA runtime API, ну, и в заключение, приведу пример использования CUDA для несложных математических вычислений.

Приступим.

Вычислительная модель GPU:

Рассмотрим вычислительную модель GPU более подробно.

1. Верхний уровень ядра GPU состоит из блоков, которые группируются в сетку или грид (grid) размерностью N1 * N2 * N3. Это можно изобразить следующим образом:


Рис. 1. Вычислительное устройство GPU.

Размерность сетки блоков можно узнать с помощь функции cudaGetDeviceProperties, в полученной структуре за это отвечает поле maxGridSize. К примеру, на моей GeForce 9600M GS размерность сетки блоков: 65535*65535*1, то есть сетка блоков у меня двумерная (полученные данные удовлетворяют Compute Capability v.1.1).

2. Любой блок в свою очередь состоит из нитей (threads), которые являются непосредственными исполнителями вычислений. Нити в блоке сформированы в виде трехмерного массива (рис. 2), размерность которого так же можно узнать с помощью функции cudaGetDeviceProperties, за это отвечает поле maxThreadsDim.


Рис. 2. Устройство блока GPU.

При использовании GPU вы можете задействовать грид необходимого размера и сконфигурировать блоки под нужды вашей задачи.

CUDA и язык C:

Сама технология CUDA (компилятор nvcc.exe) вводит ряд дополнительных расширений для языка C, которые необходимы для написания кода для GPU:
1. Спецификаторы функций, которые показывают, как и откуда буду выполняться функции.
2. Спецификаторы переменных, которые служат для указания типа используемой памяти GPU.
3. Спецификаторы запуска ядра GPU.
4. Встроенные переменные для идентификации нитей, блоков и др. параметров при исполнении кода в ядре GPU .
5. Дополнительные типы переменных.

Как было сказано, спецификаторы функций определяют, как и откуда буду вызываться функции. Всего в CUDA 3 таких спецификатора:

1. __host__ — выполнятся на CPU, вызывается с CPU (в принципе его можно и не указывать).
2. __global__ — выполняется на GPU, вызывается с CPU.
3. __device__ — выполняется на GPU, вызывается с GPU.

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

где

1. gridSize – размерность сетки блоков (dim3), выделенную для расчетов,
2. blockSize – размер блока (dim3), выделенного для расчетов,
3. sharedMemSize – размер дополнительной памяти, выделяемой при запуске ядра,
4. cudaStream – переменная cudaStream_t, задающая поток, в котором будет произведен вызов.

Ну и конечно сама myKernelFunc – функция ядра (спецификатор __global__). Некоторые переменные при вызове ядра можно опускать, например sharedMemSize и cudaStream.

Так же стоит упомянуть о встроенных переменных:

1. gridDim – размерность грида, имеет тип dim3. Позволяет узнать размер гридa, выделенного при текущем вызове ядра.
2. blockDim – размерность блока, так же имеет тип dim3. Позволяет узнать размер блока, выделенного при текущем вызове ядра.
3. blockIdx – индекс текущего блока в вычислении на GPU, имеет тип uint3.
4. threadIdx – индекс текущей нити в вычислении на GPU, имеет тип uint3.
5. warpSize – размер warp’а, имеет тип int (сам еще не пробовал использовать).

Кстати, gridDim и blockDim и есть те самые переменные, которые мы передаем при запуске ядра GPU, правда, в ядре они могут быть read only.

Дополнительные типы переменных и их спецификаторы будут рассмотрены непосредственно в примерах работы с памятью.

CUDA host API:

Перед тем, как приступить к непосредственному использованию CUDA для вычислений, необходимо ознакомиться с так называемым CUDA host API, который является связующим звеном между CPU и GPU. CUDA host API в свою очередь можно разделить на низкоуровневое API под названием CUDA driver API, который предоставляет доступ к драйверу пользовательского режима CUDA, и высокоуровневое API – CUDA runtime API. В своих примерах я буду использовать CUDA runtime API.

В CUDA runtime API входят следующие группы функций:

• Device Management – включает функции для общего управления GPU (получение инфор-мации о возможностях GPU, переключение между GPU при работе SLI-режиме и т.д.).
• Thread Management – управление нитями.
• Stream Management – управление потоками.
• Event Management – функция создания и управления event’ами.
• Execution Control – функции запуска и исполнения ядра CUDA.
• Memory Management – функции управлению памятью GPU.
• Texture Reference Manager – работа с объектами текстур через CUDA.
• OpenGL Interoperability – функции по взаимодействию с OpenGL API.
• Direct3D 9 Interoperability – функции по взаимодействию с Direct3D 9 API.
• Direct3D 10 Interoperability – функции по взаимодействию с Direct3D 10 API.
• Error Handling – функции обработки ошибок.

Понимаем работу GPU:

Как было сказано, нить – непосредственный исполнитель вычислений. Каким же тогда образом происходит распараллеливание вычислений между нитями? Рассмотрим работу отдельно взятого блока.

Задача. Требуется вычислить сумму двух векторов размерностью N элементов.

Нам известна максимальные размеры нашего блока: 512*512*64 нитей. Так как вектор у нас одномерный, то пока ограничимся использованием x-измерения нашего блока, то есть задействуем только одну полосу нитей из блока (рис. 3).


Рис. 3. Наша полоса нитей из используемого блока.

Заметим, что x-размерность блока 512, то есть, мы можем сложить за один раз векторы, длина которых N <= 512 элементов. В прочем, при более массивных вычислениях, можно использовать большее число блоков и многомерные массивы. Так же я заметил одну интересную особенность, возможно, некоторые из вас подумали, что в одном блоке можно задействовать 512*512*64 = 16777216 нитей, естественно это не так, в целом, это произведение не может превышать 512 (по крайней мере, на моей видеокарте).

В самой программе необходимо выполнить следующие этапы:

1. Получить данные для расчетов.
2. Скопировать эти данные в GPU память.
3. Произвести вычисление в GPU через функцию ядра.
4. Скопировать вычисленные данные из GPU памяти в ОЗУ.
5. Посмотреть результаты.
6. Высвободить используемые ресурсы.

Переходим непосредственно к написанию кода:

Первым делом напишем функцию ядра, которая и будет осуществлять сложение векторов:

Таким образом, распараллеливание будет выполнено автоматически при запуске ядра. В этой функции так же используется встроенная переменная threadIdx и её поле x, которая позволяет задать соответствие между расчетом элемента вектора и нитью в блоке. Делаем расчет каждого элемента вектора в отдельной нити.

Пишем код, которые отвечает за 1 и 2 пункт в программе:

Для выделения памяти на видеокарте используется функция cudaMalloc, которая имеет следующий прототип:

где

1. devPtr – указатель, в который записывается адрес выделенной памяти,
2. count – размер выделяемой памяти в байтах.

Возвращает:

1. cudaSuccess – при удачном выделении памяти
2. cudaErrorMemoryAllocation – при ошибке выделения памяти

Для копирования данных в память видеокарты используется cudaMemcpy, которая имеет следующий прототип:
cudaError_t cudaMemcpy(void* dst, const void* src ,size_t count, enum cudaMemcpyKind kind), где

1. dst – указатель, содержащий адрес места-назначения копирования,
2. src – указатель, содержащий адрес источника копирования,
3. count – размер копируемого ресурса в байтах,
4. cudaMemcpyKind – перечисление, указывающее направление копирования (может быть cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyHostToHost, cudaMemcpyDeviceToDevice).

Возвращает:

1. cudaSuccess – при удачном копировании
2. cudaErrorInvalidValue – неверные параметры аргумента (например, размер копирования отрицателен)
3. cudaErrorInvalidDevicePointer – неверный указатель памяти в видеокарте
4. cudaErrorInvalidMemcpyDirection – неверное направление (например, перепутан источник и место-назначение копирования)

Теперь переходим к непосредственному вызову ядра для вычисления на GPU.

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

Теперь нам остаеться скопировать результат расчета из видеопамяти в память хоста. Но у функций ядра при этом есть особенность – асинхронное исполнение, то есть, если после вызова ядра начал работать следующий участок кода, то это ещё не значит, что GPU выполнил расчеты. Для завершения работы заданной функции ядра необходимо использовать средства синхронизации, например event’ы. Поэтому, перед копированием результатов на хост выполняем синхронизацию нитей GPU через event.

Код после вызова ядра:

Рассмотрим более подробно функции из Event Managment API.

Event создается с помощью функции cudaEventCreate, прототип которой имеет вид:

где

1. *event – указатель для записи хендла event’а.

Возвращает:

1. cudaSuccess – в случае успеха
2. cudaErrorInitializationError – ошибка инициализации
3. cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
4. cudaErrorInvalidValue – неверное значение
5. cudaErrorMemoryAllocation – ошибка выделения памяти

Запись event’а осуществляется с помощью функции cudaEventRecord, прототип которой имеет вид:

где

1. event – хендл хаписываемого event’а,
2. stream – номер потока, в котором записываем (в нашем случае это основной нулевой по-ток).

Возвращает:

1. cudaSuccess – в случае успеха
2. cudaErrorInvalidValue – неверное значение
3. cudaErrorInitializationError – ошибка инициализации
4. cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
5. cudaErrorInvalidResourceHandle – неверный хендл event’а

Синхронизация event’а выполняется функцией cudaEventSynchronize. Данная функция ожидает окончание работы всех нитей GPU и прохождение заданного event’а и только потом отдает управление вызывающей программе. Прототип функции имеет вид:

где

1. event – хендл event’а, прохождение которого ожидается.

Возвращает:

1. cudaSuccess – в случае успеха
2. cudaErrorInitializationError – ошибка инициализации
3. cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
4. cudaErrorInvalidValue – неверное значение
5. cudaErrorInvalidResourceHandle – неверный хендл event’а

Понять, как работает cudaEventSynchronize, можно из следующей схемы:


Рис. 4. Синхронизация работы основоной и GPU прграмм.

На рисунке 4 блок «Ожидание прохождения Event’а» и есть вызов функции cudaEventSynchronize.

Ну и в заключении выводим результат на экран и чистим выделенные ресурсы.

Думаю, что описывать функции высвобождения ресурсов нет необходимости. Разве что, можно напомнить, что они так же возвращают значения cudaError_t, если есть необходимость проверки их работы.

Заключение

Надеюсь, что этот материал поможет вам понять, как функционирует GPU. Я описал самые главные моменты, которые необходимо знать для работы с CUDA. Попробуйте сами написать сложение двух матриц, но не забывайте об аппаратных ограничениях видеокарты.

Ув Змей!
... изменилось ли что-нибуть в мире Cuda за эти полгода ???

отличная статья..... о CUDA наслышан норм нвидиа сделали ток херова что на слабых не пойдет