количество конвейеров cuda что это
Нужны ли графические ядра Nvidia CUDA для игр?
Ядра CUDA являются эквивалентом процессорных ядер Nvidia. Они оптимизированы для одновременного выполнения большого количества вычислений, что очень важно для современной графики. Естественно, на графические настройки больше всего повлияло количество ядер CUDA в видеокарте, и они требуют больше всего от графического процессора, то есть теней и освещения, среди прочего.
CUDA долгое время была одной из самых выдающихся записей в спецификациях любой видеокарты GeForce. Однако не все до конца понимают, что такое ядра CUDA и что конкретно они означают для игр.
В этой статье дан краткий и простой ответ на этот вопрос. Кроме того, мы кратко рассмотрим некоторые другие связанные вопросы, которые могут возникнуть у некоторых пользователей.
Что такое ядра видеокарты CUDA?
CUDA является аббревиатурой от одной из запатентованных технологий Nvidia: Compute Unified Device Architecture. Его цель? Эффективные параллельные вычисления.
Одиночное ядро CUDA аналогично ядру ЦП, основное отличие в том, что оно менее изощренное, но реализовано в большем количестве. Обычный игровой процессор имеет от 2 до 16 ядер, но количество ядер CUDA исчисляется сотнями, даже в самых низких современных видеокартах Nvidia GeForce. Между тем, у высококлассных карт сейчас их тысячи.
Что делают ядра CUDA в играх?
Обработка графики требует одновременного выполнения множества сложных вычислений, поэтому такое огромное количество ядер CUDA реализовано в видеокартах. И учитывая, как графические процессоры разрабатываются и оптимизируются специально для этой цели, их ядра могут быть намного меньше, чем у гораздо более универсального CPU.
И как ядра CUDA влияют на производительность в игре?
По сути, любые графические настройки, которые требуют одновременного выполнения вычислений, значительно выиграют от большего количества ядер CUDA. Наиболее очевидными из них считается освещение и тени, но также включены физика, а также некоторые типы сглаживания и окклюзии окружающей среды.
Ядра CUDA или потоковые процессоры?
Там, где у Nvidia GeForce есть ядра CUDA, у их основного конкурента AMD Radeon есть потоковые процессоры.
Ядра CUDA лучше оптимизированы, поскольку аппаратное обеспечение Nvidia обычно сравнивают с AMD, но нет никаких явных различий в производительности или качестве графики, о которых вам следует беспокоиться, если вы разрываетесь между приобретением Nvidia или AMD GPU.
Сколько ядер CUDA вам нужно?
И вот сложный вопрос. Как часто бывает с бумажными спецификациями, они просто не являются хорошим индикатором того, какую производительность вы можете ожидать от аппаратного обеспечения.
Многие другие спецификации, такие как пропускная способность VRAM, более важны для рассмотрения, чем количество ядер CUDA, а также вопрос оптимизации программного обеспечения.
Для общего представления о том, насколько мощен графический процессор, мы рекомендуем проверить UserBenchmark. Однако, если вы хотите увидеть детальное и всестороннее тестирование, есть несколько надежных сайтов, таких как GamersNexus, TrustedReviews, Tom’s Hardware, AnandTech и ряд других.
Вывод
Надеемся, что это помогло пролить некоторый свет на то, чем на самом деле являются ядра CUDA, что они делают и насколько они важны. Прежде всего, мы надеемся, что помогли развеять любые ваши заблуждения по этому поводу.
CUDA: Как работает GPU
Внутренняя модель nVidia GPU – ключевой момент в понимании GPGPU с использованием CUDA. В этот раз я постараюсь наиболее детально рассказать о программном устройстве GPUs. Я расскажу о ключевых моментах компилятора CUDA, интерфейсе CUDA runtime API, ну, и в заключение, приведу пример использования CUDA для несложных математических вычислений.
Вычислительная модель GPU:
При использовании GPU вы можете задействовать грид необходимого размера и сконфигурировать блоки под нужды вашей задачи.
CUDA и язык C:
Дополнительные типы переменных и их спецификаторы будут рассмотрены непосредственно в примерах работы с памятью.
CUDA host API:
Перед тем, как приступить к непосредственному использованию CUDA для вычислений, необходимо ознакомиться с так называемым CUDA host API, который является связующим звеном между CPU и GPU. CUDA host API в свою очередь можно разделить на низкоуровневое API под названием CUDA driver API, который предоставляет доступ к драйверу пользовательского режима CUDA, и высокоуровневое API – CUDA runtime API. В своих примерах я буду использовать CUDA runtime API.
Понимаем работу GPU:
Как было сказано, нить – непосредственный исполнитель вычислений. Каким же тогда образом происходит распараллеливание вычислений между нитями? Рассмотрим работу отдельно взятого блока.
Задача. Требуется вычислить сумму двух векторов размерностью N элементов.
Нам известна максимальные размеры нашего блока: 512*512*64 нитей. Так как вектор у нас одномерный, то пока ограничимся использованием x-измерения нашего блока, то есть задействуем только одну полосу нитей из блока (рис. 3). 
Рис. 3. Наша полоса нитей из используемого блока.
Заметим, что x-размерность блока 512, то есть, мы можем сложить за один раз векторы, длина которых N // Функция сложения двух векторов
__global__ void addVector( float * left, float * right, float * result)
<
//Получаем id текущей нити.
int idx = threadIdx.x;
Таким образом, распараллеливание будет выполнено автоматически при запуске ядра. В этой функции так же используется встроенная переменная threadIdx и её поле x, которая позволяет задать соответствие между расчетом элемента вектора и нитью в блоке. Делаем расчет каждого элемента вектора в отдельной нити.
Пишем код, которые отвечает за 1 и 2 пункт в программе:
#define SIZE 512
__host__ int main()
<
//Выделяем память под вектора
float * vec1 = new float [SIZE];
float * vec2 = new float [SIZE];
float * vec3 = new float [SIZE];
//Инициализируем значения векторов
for ( int i = 0; i //Указатели на память видеокарте
float * devVec1;
float * devVec2;
float * devVec3;
…
dim3 gridSize = dim3(1, 1, 1); //Размер используемого грида
dim3 blockSize = dim3(SIZE, 1, 1); //Размер используемого блока
Теперь нам остаеться скопировать результат расчета из видеопамяти в память хоста. Но у функций ядра при этом есть особенность – асинхронное исполнение, то есть, если после вызова ядра начал работать следующий участок кода, то это ещё не значит, что GPU выполнил расчеты. Для завершения работы заданной функции ядра необходимо использовать средства синхронизации, например event’ы. Поэтому, перед копированием результатов на хост выполняем синхронизацию нитей GPU через event.
Код после вызова ядра:
//Выполняем вызов функции ядра
addVector >>(devVec1, devVec2, devVec3);
//Хендл event’а
cudaEvent_t syncEvent;
cudaEventCreate(&syncEvent); //Создаем event
cudaEventRecord(syncEvent, 0); //Записываем event
cudaEventSynchronize(syncEvent); //Синхронизируем event
Рассмотрим более подробно функции из Event Managment API.
Рис. 4. Синхронизация работы основоной и GPU прграмм.
На рисунке 4 блок «Ожидание прохождения Event’а» и есть вызов функции cudaEventSynchronize.
Ну и в заключении выводим результат на экран и чистим выделенные ресурсы.
cudaFree(devVec1);
cudaFree(devVec2);
cudaFree(devVec3);
Думаю, что описывать функции высвобождения ресурсов нет необходимости. Разве что, можно напомнить, что они так же возвращают значения cudaError_t, если есть необходимость проверки их работы.
Заключение
Надеюсь, что этот материал поможет вам понять, как функционирует GPU. Я описал самые главные моменты, которые необходимо знать для работы с CUDA. Попробуйте сами написать сложение двух матриц, но не забывайте об аппаратных ограничениях видеокарты.
FAQ по видеокартам GeForce: что следует знать о графических картах?
Страница 4: GPU
Что скрывается за потоковым процессором, блоком шейдеров или ядром CUDA?
Потоковый процессор обрабатывает непрерывный поток данных, которых насчитываются многие сотни, причем они выполняются параллельно на множестве потоковых процессоров. Современные GPU оснащаются несколькими тысячами потоковых процессоров, они отлично подходят для задач с высокой степенью параллельности. Это и рендеринг графики, и научные расчеты. Что, кстати, позволило GPU закрепиться в серверном сегменте в качестве вычислительных ускорителей.
Еще одним шагом дальше можно назвать интеграцию ядер Tensor в архитектуру NVIDIA Ampere, которые способны эффективно вычислять менее сложные числа INT8 и INT4, но об этом мы поговорим чуть позже.
В составе GPU GA102 имеются семь кластеров Graphics Processing Clusters (GPC) с 12 потоковыми мультипроцессорами Streaming Multiprocessors (SM) каждый. Но на видеокартах GeForce RTX 3090 и GeForce RTX 3080 активны не все SM. GA102 GPU теоретически содержит 10.752 блоков FP32 (7 GPC x 12 SM x 128 блоков FP32). Но у GeForce RTX 3090 два SM отключены, поэтому видеокарта предлагает «всего» 10.496 блоков FP32. Такой подход повышает выход годных чипов NVIDIA, поскольку наличие одного-двух дефектных SM не приводит к отбраковке кристалла.
В случае GeForce RTX 3080 один кластер GPC полностью отключен, поэтому на GA102 GPU остаются шесть GPC, но только четыре из них содержат полные 12 SM, два ограничены десятью SM. Что дает в сумме 8.704 блока FP32 в составе 68 SM.
NVIDIA масштабирует архитектуру Ampere с видеокарты GeForce RTX 3060 вплоть до GeForce RTX 3090. Ниже представлен обзор видеокарт GeForce RTX 30:
| GeForce RTX 3090 | GeForce RTX 3080 Ti | GeForce RTX 3080 | GeForce RTX 3070 Ti | |
| GPU | Ampere (GA102) | Ampere (GA102) | Ampere (GA102) | Ampere (GA104) |
| Число транзисторов | 28 млрд. | 28 млрд. | 28 млрд. | 17,4 млрд. |
| Техпроцесс | 8 нм | 8 нм | 8 нм | 8 нм |
| Площадь кристалла | 628,4 мм² | 628,4 мм² | 628,4 мм² | 392,5 мм² |
| Число FP32 ALU | 10.496 | 10.240 | 8.704 | 6.144 |
| Число INT32 ALU | 5.248 | 5.120 | 4.352 | 3.072 |
| Число SM | 82 | 80 | 68 | 48 |
| Ядра Tensor | 328 | 320 | 272 | 192 |
| Ядра RT | 82 | 80 | 68 | 48 |
| Базовая частота | 1.400 МГц | 1.365 МГц | 1.440 МГц | 1.580 МГц |
| Частота Boost | 1.700 МГц | 1.665 МГц | 1.710 МГц | 1.770 МГц |
| Емкость памяти | 24 GB | 12 GB | 10 GB | 8 GB |
| Тип памяти | GDDR6X | GDDR6X | GDDR6X | GDDR6X |
| Частота памяти | 1.219 МГц | 1.188 МГц | 1.188 МГц | 1.188 МГц |
| Ширина шины памяти | 384 бит | 384 бит | 320 бит | 256 бит |
| Пропускная способность памяти | 936 Гбайт/с | 912 Гбайт/с | 760 Гбайт/с | 608 Гбайт/с |
| TDP | 350 Вт | 350 Вт | 320 Вт | 290 Вт |
| GeForce RTX 3070 | GeForce RTX 3060 Ti | GeForce RTX 3060 | |
| GPU | Ampere (GA104) | Ampere (GA104) | Ampere (GA106) |
| Число транзисторов | 17,4 млрд. | 17,4 млрд. | 12 млрд. |
| Техпроцесс | 8 нм | 8 нм | 8 нм |
| Площадь кристалла | 392,5 мм² | 392,5 мм² | 276 мм² |
| Число FP32 ALU | 5.888 | 4.864 | 3.584 |
| Число INT32 ALU | 2.944 | 2.432 | 1.792 |
| Число SM | 46 | 38 | 28 |
| Ядра Tensor | 184 | 152 | 112 |
| Ядра RT | 46 | 38 | 28 |
| Базовая частота | 1.500 МГц | 1.410 МГц | 1.320 МГц |
| Частота Boost | 1.730 МГц | 1.665 МГц | 1.780 МГц |
| Емкость памяти | 8 GB | 8 GB | 12 GB |
| Тип памяти | GDDR6 | GDDR6 | GDDR6 |
| Частота памяти | 1.725 МГц | 1.750 МГц | 1.875 МГц |
| Ширина шины памяти | 256 бит | 256 бит | 192 бит |
| Пропускная способность памяти | 448 Гбайт/с | 448 Гбайт/с | 360 Гбайт/с |
| TDP | 220 Вт | 200 Вт | 170 Вт |
Одновременное выполнение операций с целыми числами и числами с плавающей запятой
Как мы уже упоминали, вычислительные блоки FP32 могут работать в режиме 2x FP16, то же самое касается INT16. Чтобы увеличить вычислительную производительность и сделать ее более гибкой, в архитектуре NVIDIA Turing появилась возможность одновременного расчета чисел с плавающей запятой и целых чисел. Конечно, подобная возможность сохранилась и в архитектуре Ampere. NVIDIA проанализировала данные вычисления в конвейере рендеринга в десятках игр, обнаружив, что на каждые 100 расчетов FP выполняется примерно треть вычислений INT. Впрочем, значение среднее, на практике оно меняется от 20% до 50%. Конечно, если вычисления FP и INT будут выполняться одновременно, то конвейеру придется иногда «подтормаживать» в случае взаимных связей.
Соотношение 1/3 INT32 и 2/3 FP32 отражено в структуре Ampere Streaming Multiprocessor (SM), составляющем элементе архитектуры Ampere. NVIDIA удвоила число вычислительных блоков FP32 на каждый SM. Вместо 64 блоков FP32 на SM, их теперь насчитывается 128. Плюс 64 блока INT32. Теперь на квадрант SM насчитывается два пути данных, некоторые могут работать параллельно. Один из путей данных содержит 16 блоков FP32, то есть может выполнять 16 вычислений FP32 за такт. Второй путь данных содержит по 16 блоков FP32 и INT32. Каждый из квадрантов SM может выполнять либо 32 операции FP32, либо по 16 операций FP32 и INT32 за такт. Если же брать SM целиком, то возможно выполнение 128 операций FP32 или по 64 операции FP32 и INT32 за такт.
Параллельное выполнение продолжается и на других блоках. Например, ядра RT и Tensor могут работать параллельно в конвейере рендеринга, что снижает время, требующееся на рендеринг кадра.
Под термином «потоковые процессоры» сегодня подразумевают количество вычислительных блоков GPU, хотя следует помнить, что сложность вычислений бывает разной. Поэтому термин используется гибко, но обычно все равно описывает вычислительные блоки.
Текстурные блоки
Действительно, для рендеринга объекта простых текстур уже недостаточно, использование нескольких слоев позволяет, например, получить 3D-эффект вместо плоской текстуры. Раньше объекты приходилось рассчитывать на конвейере несколько раз, и каждый проход текстурный блок накладывал текстуру, сегодня достаточно одного процесса рендеринга, текстурный блок может получать данные объекта для многократной обработки из буфера.
Контроллер памяти
Помимо изменений в SM, новая архитектура NVIDIA получила оптимизированную структуру конвейеров растровых операций (ROP), а также соединения ROP и контроллера памяти. До поколения Turing ROP всегда подключались к интерфейсу памяти. И на каждый 32-битный контроллер памяти приходилось восемь ROP. Если число контроллеров памяти и ширина шины менялись, то же самое касалось и ROP. В архитектуре Ampere ROP перенесены в GPC. Используются два раздела ROP на GPC, каждый раздел содержит восемь ROP.
Что дает иную формулу вычисления ROP на GeForce RTX 3080. Шесть GPC с 2x 8 ROP на каждом дают 96 ROP. У GeForce RTX 3090 работают семь GPC с 2x 8 ROP, что дает 112 ROP. NVIDIA намеренно интегрировала ROP глубже, чтобы задняя часть конвейера рендеринга меньше зависела от интерфейса памяти. Например, видеокарта GeForce RTX 3080 использует 320-битный интерфейс памяти, но содержит 96 ROP, а не 80 ROP.
Интерфейс памяти разделен на 32-битные блоки. В зависимости от желаемой ширины интерфейса памяти или емкости, их можно набирать в произвольном количестве.
Ядра Tensor и RT
Ядра Tensor третьего поколения
С архитектурой Turing NVIDIA представила два новых вычислительных блока, ранее на GPU не использовавшихся. Конечно, ядра Tensor знакомы нам по архитектуре Volta, но там они использовались для научных расчетов. В случае GPU Ampere ядра Tensor перешли уже на третье поколение.
Ядра Tensor ранее использовались только для вычислений INT16 и FP16, но в третьем поколении они могут работать с FP32 и FP64. Что особенно важно для сегмента HPC с высокой точностью. Для игровых GPU GeForce намного важнее меньшая точность.
Ядра Tensor архитектуры Turing могут выполнять 64 операции FP16 Fused Multiply-Add (FMA) каждое. В случае Ampere число операций увеличено до 128 у GA102 GPU и до 256 у GA100 GPU с плотными матрицами. Если же используются разреженные матрицы, число операций FMA FP16 увеличивается до 256 у GA102 GPU и до 512 у GA100 GPU. Ядра Tensor архитектуры Turing разреженные матрицы не поддерживают.
Ядра RT второго поколения
Все они опираются на тот принцип, что удаленные от луча примитивы не могут с ним пересекаться. Следовательно, и смысла их просчитывать нет. Число лучей на сценах растет экспоненциально, поэтому на каждый луч следует обрабатывать как можно меньшее число примитивов, чтобы не увеличивать вычислительную нагрузку.
Поскольку NVIDIA не изменила число ядер RT на SM в архитектуре Ampere, количество блоков SM на GPU по-прежнему определяет производительность RT. Но в ядрах RT есть другие оптимизации.
Одна из проблем с расчетом пересечений при трассировке лучей связана с движущимися объектами, особенно если используется эффект размытия движения (motion blur). Для ядер RT в архитектуре Turing такой сценарий является «узким местом». Но второе поколение ядер RT уже лучше справляется с интерполяцией эффекта размытия движения. Пересечения просчитываются с упреждением, в итоге трассировка лучей рассчитываются только для тех областей, где она необходима.
Кэши L1 и L2
Между функциональными блоками (потоковые процессоры, ядра RT и Tensor) и видеопамятью располагаются еще два уровня хранения данных, без которых GPU не смог бы выдавать высокий уровень производительности. Цель этих кэшей заключается в том, чтобы хранить информацию как можно ближе к функциональным блокам. Данные передаются из видеопамяти сначала в кэш L2, а затем и в кэш L1.
NVIDIA с архитектурой Ampere вновь увеличила кэш L1 с 96 до 128 кбайт. Скорость работы L1 была вновь удвоена. NVIDIA реализовала такую же меру ранее при переходе с Pascal на Turing. Число 32-битных регистров не изменилось и осталось на уровне 16.384. То же самое касается числа блоков чтения/записи.
CUDA: с места в карьер
NB: Статья — краткое введение, покрыть все ньюансы программирования под CUDA в одной статье вряд ли возможно 🙂
О железе
CUDA работает на видеокартых начиная с 8400GS и выше. Разные видеокарты имеют разые возможности. В целом, если вы видите что в видеокарте например 128 SP(Streaming Processor) — это значит что там 8 SIMD MP (multiprocessor), каждый из которых делает одновременно 16 операций. На один MP есть 16кб shared memory, 8192 штуки 4-хбайтных регистров (В картах серии GTX2xx значения больше). Также есть 64кб констант общие для всех MP, они кешируются, при непопадании в кеш — достаточно большая задержка (400-600 тактов). Есть глобальная память видеокарты, доступ туда не кешируется, и текстуры (кешируется, кеш оптимизирован для 2D выборок). Для использования нескольких видеокарт нужно во первый отключать SLI в дровах, а во вторых — на каждую видеокарту запускать по потоку, и вызывать cudaSetDevice().
С чего начать?
Его вы можете использовать во всех своих проектах, только вместо «../../common/inc » можно указать абсолютный путь (или переменную окружения).
nvcc — это и есть великий и ужасный компилатор CUDA. На выходе он генерирует объектный файл, в котором уже включена откомпилированная программа для видеокарты.
Обратите внимение на описание интерфейса в Mandelbrot_kernel.h — тут руками приходится описывать kernel-ы которые мы собираемся вызывать из основной С++ программы (впрочем их обычно не много, так что это не страшно).
После того как вам удалось запустить пример SDK, можно рассмотреть, чем же CUDA программа отличается от обычной.
Определение функций
Определение данных
__constant__ — задает переменную в константной памяти. Следует обратить внимание, что значения для констант нужно загружать функцией cudaMemcpyToSymbol. Константы доступны из всех тредов, скорость работы сравнима с регистрами(когда в кеш попадает).
__shared__ — задает переменную в общей памяти блока тредов (т.е. и значение будет общее на всех). Тут нужно подходить с осторожностью — компилятор агрессивно оптимизирует доступ сюда(можно придушить модификатором volatile), можно получать race condition, нужно использовать __syncthreads(); чтобы данные гарантированно записались. Shared memory разделена на банки, и когда 2 потока одновременно пытаются обратиться к одному банку, возникает bank conflict и падает скорость.
Все локальные переменные которые вы определеили в ядре (__device__) — в регистрах, самая высокая скорость доступа.
Как поток узнает над чем ему работать
Основая идея CUDA в том, что для решения вашей задачи вы запускаете тысячи и тысячи потоков, поэтому не стоит пугаться того что тут будет дальше написано 🙂
Допустим, надо сделать какую-то операцию над картинкой 200×200. Картинка разбивается на куски 10×10, и на каждый пиксел такого кусочка запускаем по потоку. Выглядить это будет так:
dim3 threads(10, 10);//размер квардатика, 10*10
dim3 grid(20, 20);//сколько квадратиков нужно чтобы покрыть все изображение
your_kernel >>(image, 200,200);//Эта строка запустит 40’000 потоков (не одновременно, одновременно работать будет 200-2000 потоков примерно).
В отличии от Brook+ от AMD, где мы сразу определяем какому потоку над какими данными работать, в CUDA все не так: передаваеиые kernel-у параметры одинаковые для всех потоков, и поток должен сам получить данные для себя, чтобы сделать это, потоку нужно вычислить, в каком месте изображения он находится. В этом помогают магические переменные blockDim, blockIdx.
const int ix = blockDim.x * blockIdx.x + threadIdx.x;
const int iy = blockDim.y * blockIdx.y + threadIdx.y;
В ix и iy — координаты, с помощью которых можно получить исходные данные из массива image, и записать результат работы.
Оптимизация
Не получается?
В первую очередь следует прочитать документацию вместе с SDK (NVIDIA_CUDA_Programming_Guide, CudaReferenceManual, ptx_isa), после этого можно спросить на официальном форуме — там даже девелоперы nVidia часто отписываются, да и вообще много умных людей. По русски можно спросить у меня на форуме например, где отвечу я 🙂 Также много людей обитает на gpgpu.ru.
Надеюсь это введение поможет людям, решившим попробовать программирование для видеокарт. Если есть проблемы/вопросы — буду рад помочь. Ну а в переди нас ждет введение в Brook+ и SIMD x86