Базовая концепция CUDA

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

В настоящее время программирование на графических процессорах стало довольно распространенным явлением в мире, поскольку потребность в интенсивных вычислениях растет с каждым годом. Такие задачи, как машинное обучение, криптография и, очевидно, добыча монет, больше не могут эффективно выполняться одними только ЦП из-за большого количества данных, задействованных в обработке.

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

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

Почему CUDA работает только с графической картой Nvidia?

Что будет делать пользователь Radeon? Я пользователь Mac, а в Macbook только AMD или Intel в качестве графического процессора.

Ответ на первый вопрос очевиден после некоторого поиска в Google, поскольку CUDA создается Nvidia, они не будут инвестировать, чтобы сделать своего конкурента сильнее. Что касается второго вопроса, это такие инструменты, как OpenACC и OpenCL, которые позволяют выполнять параллельные вычисления на графической карте Nvidia или AMD. AMD также недавно выпустила еще один инструмент под названием HIP, который преобразует код CUDA, работающий только на графическом процессоре Nvidia, в переносимый код, который можно запускать на графических картах AMD.

На первый взгляд кажется, что у CUDA много ограничений. Почему бы вместо этого не использовать OpenACC или OpenCL для обучения?

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

OpenACC - компания PGI, которая создает более простой в использовании бесплатный компилятор для OpenACC, была куплена Nvidia согласно здесь, и теперь они отказались от поддержки графической карты AMD. Более или менее это также стало похоже на изучение CUDA в любом случае, потому что под капотом OpenACC для Nvidia использует CUDA. Это потребовало от нас понимания CUDA.

Установить CUDA в окно

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

Сначала вам необходимо установить Microsoft Visual Studio на свой компьютер, а затем загрузить CUDA Toolkit. Есть и другой способ заставить CUDA работать с окном, за исключением использования Visual Studio, но это самый простой способ заставить CUDA работать с окном.



Visual Studio IDE | Microsoft Visual Studio
Полнофункциональная IDE для Android, iOS, Windows, Интернета и облака. Узнайте больше об интегрированной разработке Visual Studio… www.visualstudio.com





Загрузить CUDA Toolkit 9.1
Перед установкой CUDA Toolkit в Linux убедитесь, что у вас установлена ​​последняя версия драйвера NVIDIA R390. … developer.nvidia.com



На момент написания этой статьи версия этих инструментов

  • Microsoft Visual Studio 2017 - версия 15.5.6
  • CUDA - версия 9.1.85

Однако эти две версии несовместимы друг с другом и кажутся многим людям такого рода проблемами здесь, здесь и здесь. Начиная с Visual Studio 15.5, он использовал более новую версию набора инструментов MSVC, чем 15.4, которую CUDA еще не поддерживает.

Чтобы заставить его работать, проще отсюда. Необходимо установить предыдущую версию набора инструментов.

И в каждом проекте CUDA установите набор инструментов платформы для использования этого набора инструментов.

Примечание. Visual Studio и дополнительный набор инструментов платформы должны быть установлены перед CUDA Toolkit, в противном случае необходимо переустановить CUDA Toolkit, чтобы CUDA вступил в силу на дополнительном наборе инструментов платформы.

Подтвердите установку

Чтобы проверить правильность работы CUDA, сначала создайте новый проект CUDA.

Этот шаблон создаст и настроит проект для использования компилятора CUDA для запуска приложения CUDA. После завершения создания в проекте будет файл с именем kernel.cu, который является файлом по умолчанию, созданным из шаблона CUDA.

Нажмите Ctrl+F5, чтобы запустить программу, результат в командной строке должен быть

{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}
Press any key to continue . . .

Примечание. Если не удается скомпилировать и отобразить ошибку типа «неподдерживаемая версия Microsoft Visual Studio!», это означает, что набор инструментов платформы все еще является последней версией.

Модель программирования CUDA

CUDA сам по себе является платформой для параллельных вычислений и моделью программирования, которая обеспечивает расширение программирования на C / C ++. Концепция предоставления расширения вместо всего нового языка программирования понятна, потому что с этим небольшим набором расширений разработчики могут разрабатывать параллельный алгоритм на CUDA так же просто, как программировать на C / C ++, хотя я не думаю, что C / C ++ - это легко. Существующие программы на C / C ++ также можно легко преобразовать для использования параллельных вычислений на графическом процессоре для некоторой части алгоритмов вместо перезаписи всей программы. Прежде чем приступить к непосредственному выполнению, необходимо знать базовое понимание того, как работает CUDA, с целью максимального ускорения программы.

Модель программирования CUDA, конечно, имеет некоторое сходство с обычной моделью программирования процессора, но некоторые части отличаются, например:

  • Возможность организовать потоки или доступ к памяти на GPU в иерархической структуре.
  • Возможность напрямую управлять кешем на GPU, что не разрешено в CPU.
  • И т.п.

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

Структура программы CUDA

CUDA требует так называемых «гетерогенных вычислений» для выполнения программы CUDA. Итак, программы CUDA всегда состоят из 2 частей:

  1. Код, выполняемый на ЦП (код хоста)
  2. Код, выполняемый на GPU (код устройства)

Гетерогенные вычисления - это способ использования архитектуры набора процессоров для выполнения программы, а не только архитектуры одного процессора. (В этом случае CUDA использует как CPU, так и GPU)

Решение, какая часть кода будет выполняться на хосте или устройстве, является ответственностью разработчиков.

Из приведенных выше двух частей основной строительный блок программы CUDA обычно будет иметь:

  1. Получение входных данных, размещение этих данных в памяти устройства и запуск кода устройства (код хоста ЦП)
  2. Код устройства выполняет код, который обычно требует больших объемов вычислений. (Код устройства GPU)
  3. Копирование выходных данных из памяти устройства в память хоста для дальнейшей обработки (код хоста CPU).

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

На первый взгляд, этот рабочий процесс связан с большими накладными расходами, связанными с копированием данных из / в CPU и GPU. Если алгоритм небольшой и не обрабатывает много данных, то выполнение алгоритма только на ЦП может быть быстрее.

Следуя этому потоку, центральная часть всей программы CUDA находится на шаге 2, который можно назвать ядром, то есть блоком кода, выполняемым на GPU. Хотя ядро ​​ здесь может показаться пугающим и трудным для понимания, как какой-то внутренний модуль в операционной системе, это просто обычная функция C со специальным ключевым словом __global__ впереди. Это ключевое слово указывает компилятору CUDA, что эта функция будет выполняться на графическом процессоре. Чтобы запустить ядро ​​, необходимо использовать специальный синтаксис. Этот синтаксис представляет собой комбинацию обычного метода вызова функции C с <<<gridDim, blockDim>>>, который настраивает GPU, сколько потоков будет создано для выполнения функции.

Из этой запускающей сигнатуры количество gridDim и blockDim - это то, что будет определять, как иерархия потоков будет организована в CUDA. gridDim и blockDim также могут быть вектором до 3-х измерений, чтобы иерархия потоков могла быть напрямую сопоставлена ​​измерению данных (массив, матрица или объем). Это всего лишь представление измерения потока CUDA, внутренне оно все еще отображается в одномерный массив потоков.

Иерархия потоков

Иерархия потоков - это способ, которым CUDA позволяет вам управлять тем, как поток порожден графическим процессором. Он состоит из 3-х уровней от самого маленького до самого большого:

  1. Нить
  2. Блок - Группа потоков
  3. Сетка - Группа блоков

Обычно при запуске ядра все потоки создаются внутри одной сетки. Переменные gridDim и blockDim, указанные при запуске ядра, представляют собой соответственно количество блоков внутри сетки и потоков внутри блока. В качестве примера из ядра на изображении выше, gridDim является двухмерным (3, 2), а blockDim также двухмерным (5, 3).

С моей точки зрения, причина, по которой CUDA разделяет поток на этот тип сложной иерархии. Один из них связан с количеством данных, которое обычно превышает количество потоков на GPU. Это делает наиболее оптимальное решение (и самый быстрый способ выполнения данных), которое отображает все индивидуальные данные в отдельный поток и делает невозможным их параллельное вычисление. Для этой цели количества потоков будет недостаточно, следовательно, некоторые потоки необходимо повторно использовать для вычисления большего, чем просто отдельных данных. Эта пара с архитектурой аппаратного обеспечения графического процессора Nvidia, называемой потоковыми мультипроцессорами (SM), которые могут одновременно выполнять несколько потоков с помощью одной инструкции, затрудняет управление потоком с плоской иерархией. По этой причине CUDA разделяет поток на группу независимых друг от друга блоков. Каждый блок будет выполняться по расписанию на доступных SM с помощью графического процессора.

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

Чтобы различать каждый поток внутри функции ядра, CUDA предоставляет встроенные переменные с именами blockIdx и threadIdx. Эти две переменные также имеют 3 измерения, которые могут использоваться как blockIdx.x, blockIdx.y и blockIdx.z. Как показано в примере ниже

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

Размер сетки - рассчитывается на основе размера данных, разделенного на размер блока.

Размер блока - это значение необходимо определять на основе характеристик ядра приложения и ограничений аппаратного обеспечения графического процессора (неправильное решение о размере блока может привести к снижению производительности).

Чтобы определить идеальный размер блока, необходимо знать некоторую часть архитектуры графического процессора Nvidia.

Обзор архитектуры GPU

Архитектура графического процессора Nvidia состоит из массива потоковых мультипроцессоров (SM). Каждый SM независим друг от друга, и с точки зрения разработчика, этот SM выполняет ту же роль, что и ядро ​​ЦП, но вместо этого использует графический процессор. После запуска ядра блоки потоков этого ядра будут распределены между доступными SM. После того, как потоки в блоке потока назначены SM, он останется на этом SM до конца выполнения этого блока потока. SM может содержать несколько блоков потоков в зависимости от имеющихся у него ресурсов (обычно доступной памяти). Затем потоки в каждом блоке делятся на группу Warp, которая всегда состоит из 32 потоков, которые будут выполняться вместе. Это 32 - особое число в CUDA, потому что оно происходит от количества ядер в SM, которые могут выполняться параллельно. Внутри SM есть 32 ядра на архитектуре Fermi, как показано на рисунке ниже. В более новой версии, такой как архитектура Pascal, количество ядер увеличилось до 64. Тем не менее, он по-прежнему группирует потоки в основу из 32 потоков. Это означает, что количество деформаций, которые могут выполняться за цикл, увеличивается с 1 до 2 от архитектуры Ферми к архитектуре Паскаля.

Деформация

Деформация - это сердце параллельного выполнения в архитектуре CUDA. Здесь, когда выдается одна инструкция, одновременно обрабатываются данные или, другими словами, SIMD (одна инструкция, несколько данных). Если посмотреть на код ядра, это означает, что когда выполняется одна строка кода, она будет выполняться вместе 32 потоками. Единственное, что отличает каждый поток, - это blockIdx и threadIdx, которые позволяют нам получать доступ к разным данным. В приведенном ниже коде проверяется, что в CUDA есть только деформация, которая может одновременно запускать не более 32 потоков с выдачей одной инструкции.

Приведенный выше фрагмент кода создает ядро ​​с блоком и 33 потоками на блок, что превышает количество потоков, доступных на деформацию. CUDA обрабатывает этот случай, создавая 2 деформации для этого блока: один имеет 32 потока, другой - поток. Порядок, в котором варп будет выполнен первым, неизвестен (согласно здесь). В приведенном выше случае, несмотря на то, что выходные данные показывают, что сначала выполняется вторая деформация, но нет способа гарантировать это при следующем выполнении.

Одна вещь, которую я замечаю во время тестирования кода, - это результат, выводимый функцией printf(), потому что с точки зрения разработчика приложения (меня самого) обновление пользовательского интерфейса из фонового потока недопустимо. При дальнейшем исследовании printf() функция ядра устройства не совпадает с printf() функцией нормальной программы на языке C, хотя у них почти такая же сигнатура, просто взглянув. Это один из встроенных методов, который CUDA предоставляет для помощи в отладке, хотя он имеет побочный эффект в виде сообщения ниже в их руководстве. Это может нарушить порядок вывода на печать.

Внутри printf() используется общая структура данных, поэтому возможно, что вызов printf() может изменить порядок выполнения потоков. (Ссылка здесь)

Из приведенного выше примера возникает один вопрос. У второй основы есть только один поток, тогда что делают другие потоки, когда этот один поток выполняется?

Ответ: они ничего не делают. Приведенный выше параметр конфигурации ядра - пустая трата ресурсов. Чтобы проверить это, CUDA предоставляет команду nvprof, чтобы помочь приложению CUDA профилировать разработчика. Для вышеуказанного приложения, поскольку имя проекта CudaTest, в результате мы можем найти файл .exe. Файл создается из Visual Studio, после чего мы можем запустить команду профиля, которая

nvprof --metrics warp_execution_efficiency CudaTest

Эффективность составляет всего около 50%, что понятно, поскольку из возможных 64 потоков выполняется только 33 потока. Если мы изменим размер блока на 32, эффективность будет близка к 100%. Это не 100%, потому что инструкция используется в функции ядра printf. На мой взгляд, функция printf выполняет дополнительную внутреннюю синхронизацию, что может привести к остановке некоторых потоков. Тем не менее, изменение строки printf на что-то вроде int i = 0; приведет к тому, что эффективность станет 100% для размера блока 32. Из этого наблюдения можно сделать вывод, что для того, чтобы не тратить ненужные ресурсы, размер блока должен быть умножен на 32. Для multi- размерного блока, умножение x, y и z также должно приводить к умножению на 32, потому что внутренне весь размер измерения отображается в одномерный массив и выполняется так же, как и размер одномерного блока. Хотя данные реального мира не всегда имеют идеальный размер - 32, 64, 96, 128 и т. Д.… Единственный способ решить эту проблему - проверить blockIdx и threadIdx и определить индекс данных. Если этот индекс превышает последний индекс, то выполняется прямой возврат из ядра для этого конкретного потока, как показано в приведенном ниже коде.

вывод для threadIdx 2 индекса блока 1 составляет всего 3 строки, потому что индекс других потоков превышает размер данных. Наконец, количество блоков в сетке рассчитывается из количества потоков в блоке, которое было решено равным 32. Это наиболее распространенный способ определения конфигурации ядра.

Варп Дивергенция

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

Из-за этой проблемы графический процессор не подходит для выполнения сложных приложений с большим количеством условий. С другой стороны, он подходит для чего-то вроде обработки данных дампа, которая практически не требует потока управления. Чтобы избежать влияния расхождения деформации, нужно избегать потока управления внутри ядра. Однако в некоторых случаях это кажется невозможным, поэтому другой способ - переупорядочить входные данные так, чтобы все данные, имеющие один и тот же путь, находились в одной и той же деформации при выполнении, как показано ниже.

Приведенный выше код имеет 64 вывода данных, состоящих из 32 данных с 1 и 2 значениями. Первый метод warpDivergence делает это обычным способом с использованием модуля. Результат будет иметь 1 и 2 данных чередования. С другой стороны, второй метод avoidWarpDivergence предполагает, что первая деформация имеет все 1 значение, а вторая деформация - все 2 значения. Это требует переупорядочения вывода, если требуется чередование, подобное первому методу. Проверка branch_efficiency дает следующий результат.

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

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

Синхронизация

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

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

Другой метод синхронизации - это синхронизация на уровне блоков с использованием функции __syncthreads(). Этот метод гарантирует, что все потоки внутри одного блока завершат выполнение всех инструкций до __syncthreads(), пока какой-либо поток не сможет продолжить выполнение инструкций после __syncthreads(). В приведенном ниже коде показано, как важно использовать __syncthreads() для связи между потоками внутри одного блока.

Код использует общую память, которая будет доступна для всех потоков внутри одного блока (как указано с помощью __shared__). Он печатает valAfter для текущего потока как значение следующего потока. Приведенный выше экземпляр кода правильно печатает результат из-за __syncthreads(). Если эта функция будет удалена, некоторые потоки будут иметь valAfter как 0, потому что поток, ответственный за этот результат, прибыл в строку оператора 13, в то время как следующий поток еще не завершил строку оператора 9. Число потоков на блок здесь составляет 1024 максимально возможное на Архитектура Pascal выбрана потому, что для имитации эффекта __syncthreads() она требует большого количества перекосов в блоке. Небольшое количество перекосов, например 2–3 перекоса на блок, не вызовет проблемы синхронизации. Потому что в списке планировщика SM не так много деформаций, ожидающих выполнения. Таким образом, SM может выбирать деформацию по порядку. И наоборот, при большом количестве деформаций в списке планировщика SM будет много деформаций, ожидающих выполнения, тогда проблема возникает здесь, если SM выберет неправильный порядок, что, вероятно, произойдет, потому что нет ограничений на порядок или информации о том, как SM выберет следующую готовую основу.

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

Управление памятью

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

Копирование данных с хоста на устройство также разделено на 2 части. Первая часть выделяет место в памяти на устройстве. Другой - скопировать данные в память, выделенную первой частью. С другой стороны, после завершения выполнения ядра данные результатов также необходимо скопировать обратно в код хоста. Те же шаги применимы и здесь, но в обратном порядке. Сначала выделите память в коде хоста для хранения данных результатов из кода устройства. Затем скопируйте данные результата по коду устройства в область памяти, выделенную на первом шаге.

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

  • malloc - ›cudaMalloc — выделять память на GPU при возврате указателя туда, где он находится.
  • memcpy - ›cudaMemcpy — копировать данные из одного указателя в другой указатель, можно указать вариант перехода (cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost и cudaMemcpyDeviceToDevice)
  • free - ›cudaFree — освободить память на GPU

В приведенном ниже коде показан пример завершенной базовой структуры добавления массива.

Код довольно прост, исходя из базовой структуры программы. Данные результата можно использовать непосредственно в printData без cudaDeviceSynchronize(), потому что cudaMemcpy - это вызываемая синхронизация, которая блокирует поток ЦП до тех пор, пока не завершится копирование данных памяти. Внутренне я думаю, что метод cudaMemcpy просто вызывает cudaDeviceSynchronize() перед копированием данных. Поскольку копирование памяти запускается только после завершения работы ядра, гарантируется, что после вызова cudaMemcpy данные будут доступны для кода ЦП. В приведенном выше примере память выделяется в глобальной памяти, которая является наиболее распространенной и будет доступна всем SM на устройстве.

Типы памяти в CUDA

В CUDA есть много типов памяти, которые позволяют разработчикам использовать.

  • Регистры - самые быстрые и самые маленькие. Любые обычные переменные, объявленные внутри ядра без дополнительных квалификаторов, считаются хранимыми и доступными в регистре, значения, хранящиеся здесь, также являются частными для каждого потока. Регистры также используются всеми активными варпами на SM как состояние здесь. Однако регистры - это такой ограниченный ресурс, поэтому, если регистров недостаточно для хранения всех этих данных, они будут перетекать в глобальную память (разлив регистров), что приведет к более медленному времени доступа.
  • Совместно используемая память - программируемый кэш (быстрее, чем глобальная память, медленнее, чем регистры). Это можно рассматривать как программируемый кеш, который мы можем контролировать, как он хранит или выпускает, в отличие от ЦП, который будет выполнять весь кеш L1 / L2 для разработчиков, но который также имеет причину меньшей оптимизации. CUDA вместо автоматического управления кешем позволяет разработчику управлять ими вручную. Чтобы объявить общую память в ядре, требуется квалификатор __share__ перед переменной, чтобы указать, что она будет выделена в разделяемой памяти. Совместно используемая память также разделена между всеми блоками потоков на одном SM, поэтому она имеет то же ограничение, что и регистр. Несмотря на разлив подобных регистров, вместо этого будет уменьшено количество активных деформаций.
  • Глобальная память - самая большая и самая медленная. Это единственный тип памяти, который можно выделить в коде хоста или ЦП.

Хотя управление памятью - это гораздо больше, чем передача памяти между хостом и устройством или решение, использовать ли глобальную или общую память, это своего рода продвинутая тема и требует глубоких знаний о том, как работает шаблон доступа к памяти. Это выходит за рамки базовых знаний и только тогда, когда требуется серьезная оптимизация. Одним из примеров этого случая является неправильный доступ к памяти. Данные в глобальной памяти всегда делятся на блок из 32 элементов, это сделано намеренно и соответствует количеству потоков в деформации. Когда активная деформация запрашивает память, если память выровнена, требуется только одна транзакция памяти. Если память запросов не выровнена, пропускная способность памяти будет потрачена впустую из-за выполнения нескольких транзакций с памятью, как показано ниже, и некоторые данные в этой транзакции никогда не будут использоваться.

Как видно из рисунка, во втором случае требуется 2 транзакции с памятью, поскольку требуемые данные находятся между первым и вторым блоками. Первая половина первого блока и вторая половина второго блока данных будут потрачены впустую, потому что они не используются потоками. Чтобы показать влияние этой проблемы, измените строку 3 в arrayAddition вышеприведенного примера на int i = threadIdx.x + 0;. Изменяет число 0 для имитации доступа к невыровненной памяти как 0, 1, 31, 32 и 33 и измеряет время выполнения, используя nvprof каждой конфигурации. Для использования доступа к памяти используйте nvprof --metrics gld_efficiency -metrics gst_efficiency <execution file name> для проверки операций загрузки и сохранения.

+--------+----------+----------------+----------------+
| offset | time     | gld_efficiency | gst_efficiency |
+--------+----------+----------------+----------------+
| 0      | 2.9760us | 100.00%        | 100.00%        |
| 1      | 3.2640us | 87.50%         | 87.50%         |
| 31     | 3.2960us | 68.75%         | 68.75%         |
| 32     | 2.9120us | 100.00%        | 100.00%        |
| 33     | 3.2320us | 96.88%         | 96.88%         |
+--------+----------+----------------+----------------+

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

Заключение

CUDA - это расширение для C, которое обеспечивает возможность легко улучшить время выполнения алгоритма в 10, 100 или более раз с небольшими усилиями по сравнению с другими технологиями, возможно, кроме OpenACC. Что касается аппаратного обеспечения, из-за того, что базовый модуль выполняет деформацию, программирование на графическом процессоре Nvidia CUDA также имеет иной образ мышления, чем программирование на обычной программе CPU C (выполнение одного действия против выполнения нескольких действий). Эта концепция деформации позволяет проводить много оптимизаций, особенно для вычисления данных. CUDA также предоставляет разработчику способ взаимодействия с оборудованием и организации потоков и памяти вручную. Несмотря на то, что это более подвержено ошибкам и труднее исправить это, но это дает больше возможностей выжать из карты все до последней капли энергии.

Ссылка

[1] https://www.amazon.com/Professional-CUDA-Programming-John-Cheng-ebook/dp/B00NGK1LE0

[2] http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

[3] http://15418.courses.cs.cmu.edu/spring2013/article/11