Что хуже всего в состоянии дел с CUDA, так это то, что:
- ни в одном из "крутых" источников (что в мэнюэлах с NVIDIA, что в книжках названных и достаточно бездарных) нет толком и 2-х слов о том, а). как собрать самое простенькое приложение CUDA, опции команд, что на что влияет; б). перечисления частей, входящих в CUDA tools и что для чего нужно; в). т.е. технологического описания "из чего это складывается и как это всё крутится"...
- а форумы, обсуждения etc. CUDA (в большинстве названных выше) использующие Windows описывают это ("с истошным криком идиота"
, как это назвал один мэтр рок-н-рола
) так: "... после установки запускаю Visual C 2010, нажимаю кнопку ... и всё получилось!" ... или: "... и ничего не получилось!"
Olej писал(а):
Итак, собираю себе ссылки источников по CUDA: а). обсуждения программистов, которые делятся опытом + б). русскоязычные (пока) + в). досаточно свежих (не позже 2-й половины 2010).
2. (1-я была ссылка цитируемая выше
)
http://www.bog.pp.ru/work/cuda.html
Bog BOS: Использование модели массового параллелизма CUDA для разработки программ
Последнее изменение файла: 2010.09.01
Вот здесь хоть кой-что:
Основной программой среды разработки является компилятор nvcc, который управляет всей цепочкой преобразований. Исходная программа может быть на языке C (.c), C++ (.cc, .cxx, .cpp), расширении языка C (.cu; .cup для препроцессированных файлов). В начале, стандартный cpp осуществляет макроподстановки, затем утилита cudafe в 2 прохода разделяет общий код программы, написанный на расширении языка C (.cu; .cup для препроцессированных файлов), на части, предназначенные к исполнению на хосте (.c, компилируется обычным gcc в режиме C++ в .o), и части, предназначенные к исполнению на GPU (.gpu, C). Простенький развёрнутый crypt-des (25 итераций по 16 раундов по 8 S-блоков по 10 переменных) обрабатывался 15 минут. Последние преобразуются nvopencc в код виртуальной машины (.ptx). nvopencc имеет ограничение по умолчанию - не более 26000 регистров каждого типа (а их препроцессор генерирует бесжалостно). После снятия ограничения на оптимизацию потребил (в фазе be) 6GB памяти (реально используются!) и 90 минут CPU, после чего всё равно упал. Код виртуальной машины компилируется оптимизирующим ассемблером ptxas (распределение виртуальных регистров в реальные, имеет свои ограничения) в двоичный блоб (.cubin), который утилита fatbin в сочетании с кодовым именем устройства либо размещает во внешний репозитарий (.fatbin), либо добавляет к коду хостовой части (.cu.c), из которого стандартый компилятор gcc (версии от 3.4 до 4.2; кстати, в Fedora gcc 4.3!) компилирует и собирает программу.
К хостовой программе в дополнение к остальным библиотекам (.a, .so) присоединяется библиотека функций (высокого уровня cudart и/или низкого уровня cuda; cudart использует cuda; приложение может использовать только одну из них), обеспечивающих управление и доступ к GPU со стороны хоста. К программе, выполняемой на GPU, присоединяется библиотека функций, специфических для устройства. Имеется также общая часть - подмножество библиотеки C, реализованное как на хосте, так и на GPU. nvcc генерирует код, предназначенный для работы с cudart, при этом инициализация, управление контекстами и модулями скрыты от программиста.
С помощью макро __CUDACC__ можно определить, компилируется ли программа с помощью nvcc.
Несколько хостовых потоков могут запускать ядра на одном GPU (ресурсы одного потока недоступны другому). Один хостовый поток не может использовать несколько GPU.
Это всё скорее "заметки на полях" от человека более-менее (скорее менее
) повозившимся с CUDA, но это и есть самое вкусное
:
Хостовые подпрограммы компилируются в режиме C++, чтобы компилировать в режиме C требуется указать ключ "--host-compilation=c".
Расширения языка C:
квалификатор типа функции, указывающий где должна выполняться функция и откуда её можно вызывать
__device__ (функция выполняется на GPU, вызывается из GPU; рекурсия запрещена; статические переменные запрещены; фиксированное число параметров; указатель на функцию недопустим)
__global__ (это ядро и должно вызываться специальной директивой запуска ядра; вызов является асинхронным; параметры передаются через разделяемую память (до 256 байт); функция выполняется на GPU, вызывается из хоста; рекурсия запрещена; статические переменные запрещены; фиксированное число параметров; тип функции только void)
__host__ (функция выполняется на хосте, вызывается из хоста, по умолчанию)
__device__ __host__ (код функции компилируется и для хоста, и для GPU)
квалификатор типа переменной, указывающий в памяти какого типа её размещать (можно использовать только в функциях, исполняемых на GPU):
__device__ (переменная размещается в глобальной памяти GPU; существует пока существует приложение; доступна из любого потока на GPU и из хоста через библиотеку)
__constant__ (переменная размещается в памяти констант GPU; существует пока существует приложение; доступна из любого потока на GPU и из хоста через библиотеку)
__shared__ (переменная размещается в разделяемой памяти блока, существует пока существует блок; доступна из любого потока блока; запись становится гарантированно видима остальным потокам блока только после вызова __syncthreads() )
внутренние переменные исполняемой на GPU функции без квалификатора обычно хранятся в регистре, но если оптимизатор из ptxas посчитает, что их слишком много (более 124), то в локальной нерегистровой памяти MP (по времени доступа неотличима от глобальной); узнать, что он туда перенёс, невозможно (в .ptx всё в регистрах, а в блоб не заглянешь); массивы с достпом по неконстантному индексу тоже попадают в локальную память
директива запуска ядра: вызов из хоста ядра (функции с квалификатором __global__) оформляется специальным образом - между именем функции и списком параметров в тройных угловых скобках ( "<<<...>>>" ) записываются через запятую размерность и размер решётки (тип dim3), размерность и размер блока (тип dim3), количество дополнительной динамически выделяемой разделяемой памяти на блок (опционально), номер задачи (stream, опционально); конфигурация ядра вычисляется до вычисления параметров и передаётся через разделяемую память
встроенные переменные, доступные из функции, работающих на GPU (только чтение, указывать на встроенные переменные недопустимо):
gridDim (тип dim3) - размерность и размер решётки
blockIdx (тип uit3) - индекс блока в решётке
blockDim (тип dim3) - размерность и размер блока
threadIdx (тип uit3) - индекс потока в блоке
warpSize (тип int) - размер порции
__noinline__ (по умолчанию, все функции, работающие на устройстве, компилируются, как встроенные (inline); nvcc может проигнорировать это указание, если список параметров слишком велик или при передаче ссылки)
компилятор по умолчанию развёртывает маленькие циклы с известным числом шагов; непосредственно перед циклом можно явно указать, что делать с этим циклом директивой "#pragma unroll [число-шагов]"