Предисловие: Когда вы новичок в программировании на CUDA, многие люди задаются вопросом, как следует устанавливать параметры в трех угловых скобках при запуске ядра? Какие факторы ограничивают эти параметры? И как они влияют на производительность работающего ядра? Эта статья ссылается на официальную документацию CUDA, чтобы проанализировать, как должны быть установлены эти параметры.
Обычно мы видим в коде запуск ядра CUDA, используя:
cuda_kernel<<<grid_size, block_size, 0, stream>>>(...)
cuda_kernel — это идентификатор глобальной функции, а (...) — это параметр, соответствующий вызову cuda_kernel. Синтаксис обоих такой же, как у C++, и >> является CUDA для C++. Расширение, называемое «Конфигурация выполнения», см. во введении в Руководстве по программированию CUDA C++ (далее — «Руководство»)
The execution configuration is specified by inserting an expression of the form <<< Dg, Db, Ns, S >>> between the function name and the parenthesized argument list, where:
Dg is of type dim3 (see dim3) and specifies the dimension and size of the grid, such that Dg.x * Dg.y * Dg.z equals the number of blocks being launched;
Db is of type dim3 (see dim3) and specifies the dimension and size of each block, such that Db.x * Db.y * Db.z equals the number of threads per block;
Ns is of type size_t and specifies the number of bytes in shared memory that is dynamically allocated per block for this call in addition to the statically allocated memory; this dynamically allocated memory is used by any of the variables declared as an external array as mentioned in shared; Ns is an optional argument which defaults to 0;
S is of type cudaStream_t and specifies the associated stream; S is an optional argument which defaults to 0.
Dg представляет размер сетки, Db представляет размер блока, а тип dim3.Если это простая одномерная структура, то есть в дополнение к x, значения, соответствующие двум измерениям yz оба равны 1, а Dg и Db также могут использоваться напрямую.Число, соответствующее размеру x, заменяется, что является способом представления в начале статьи.Для более конкретного описания размерности сетки и размера блока см. пожалуйста, обратитесь к модели программирования, Далее давайте обсудим, какие значения эти два значения обычно должны принимать.
grid_size и block_size соответственно представляют количество блоков, соответствующих этому запуску ядра, и количество потоков в каждом блоке, поэтому очевидно, что оба значения должны быть больше 0.
K.1. Характеристики и технические характеристики в Руководстве указывает, что максимальное количество потоков на блок и максимальное x- или y-размерность блока равны 1024, поэтому максимальный размер блока может быть 1024.
В одном и том же блоке 32 последовательных потока образуют варп.Эти 32 потока каждый раз выполняют одну и ту же инструкцию, так называемую SIMT.Даже если количество эффективных потоков в последнем варпе меньше 32, одни и те же аппаратные ресурсы должен использоваться. , поэтому block_size предпочтительно должен быть целым числом, кратным 32.
блок иногда называют кооперативными массивами потоков, см.
The Parallel Thread Execution (PTX) programming model is explicitly parallel: a PTX program specifies the execution of a given thread of a parallel thread array. A cooperative thread array, or CTA, is an array of threads that execute a kernel concurrently or in parallel.
Threads within a CTA can communicate with each other. To coordinate the communication of the threads within the CTA, one can specify synchronization points where threads wait until all threads in the CTA have arrived.
Аппаратный уровень, соответствующий блоку, — SM. SM предоставляет аппаратные ресурсы, необходимые для связи и синхронизации для потоков в одном блоке. Соответствующая связь между SM не поддерживается, поэтому все потоки в блоке выполняются в одном и том же SM. SM, и поскольку потоки могут быть синхронизированы, как только блок начинает выполняться на SM, все потоки в блоке выполняются в одном и том же SM одновременно (параллельно, а не параллельно), что означает, что процесс планирования Блок СМ является атомарным. SM позволяет одновременно выполнять на нем более одного блока. Если незанятые ресурсы SM удовлетворяют выполнению блока, то блок может быть немедленно запланирован для выполнения на SM. Конкретные аппаратные ресурсы обычно включают регистры, разделяемую память , и Различные ресурсы, связанные с планированием, ресурсы, связанные с планированием, здесь обычно выражаются в виде двух конкретных ограничений: Максимальное количество резидентных блоков на SM и Максимальное количество резидентных потоков на SM, то есть максимальное количество одновременно выполняемых блоков на SM. и количество потоков. Поскольку GPU характеризуется высокой пропускной способностью и высокой задержкой, точно так же, как эскалатор может за одну минуту перевезти 60 человек на другой этаж, но один человек не может пройти по эскалатору на другой этаж за одну секунду, добраться до эскалатора можно транспортом. достаточно людей, чтобы убедиться, что на эскалаторе одновременно находится достаточно людей, что соответствует GPU, это попытаться обеспечить, чтобы на конвейере одновременно было достаточно инструкций.
Существует множество способов достижения этой цели. Один из самых простых способов — позволить как можно большему количеству потоков выполняться на SM одновременно Отношение количества одновременно выполняемых потоков на SM к максимальному количеству потоков Поддерживаемый на SM называется Occupancy.Более высокая занятость представляет собой потенциально более высокую производительность. Очевидно, что block_size ядра должен быть больше, чем отношение максимального количества потоков к максимальному количеству блоков на SM, иначе оно не сможет достичь 100% заполнения, что соответствует разным архитектурам, это соотношение не равно то же самое, для V100, A100, GTX 1080 Ti это 2048/32=64, для RTX 3090 это 1536/16=96, поэтому для адаптации к основной архитектуре, если статически установить block_size, он не должен быть меньше 96. Учитывая атомарность блочного планирования, размер_блока должен быть делителем максимального количества потоков SM, иначе он не сможет достичь 100% занятости.Соглашение о максимальном количестве потоков SM для графических процессоров с основной архитектурой: 512, а к делителям выше 96 относятся также 128 и 256, то есть пока необязательные значения block_size составляют всего 128/256/512 три значения.
Или, поскольку планирование блоков для SM является атомарным, SM должен соответствовать ресурсам, необходимым для запуска хотя бы одного блока. Ресурсы включают общую память и регистры. Общая память обычно явно контролируется разработчиком, и если количество потоков в блоке * Количество регистров, необходимых каждому потоку, превышает максимальное количество регистров на блок, поддерживаемое SM, и ядро не запустится. В текущей основной архитектуре максимальное количество регистров на блок, поддерживаемое SM, составляет 32 КБ или 64 КБ 32-битных регистров, каждый поток может использовать до 255 32-битных регистров, и компилятор не будет выделять больше регистров для потоков, поэтому с точки зрения регистров Например, каждый SM может поддерживать как минимум 128 или 256 потоков, а размер блока 128 может предотвратить сбои при запуске, вызванные количеством регистров, но лишь немногие ядра могут использовать такое количество регистров, и только 128 или 256 потоков могут выполняться одновременно на потоки SM., также могут быть потенциальные проблемы с производительностью. Но установка block_size на 128 не имеет никаких потерь по сравнению с 256 и 512. 128 очень подходит в качестве общего значения для block_size.
После определения block_size можно дополнительно определить grid_size, то есть определить общее количество потоков.Для общего поэлементного ядра общее количество потоков не должно быть больше, чем общее количество элементов, то есть thread обрабатывает по крайней мере один элемент, а grid_size также имеет Верхний предел — максимальное x-размерность сетки блоков потоков, которая в настоящее время составляет 2^31 - 1 в основных архитектурах, что является достаточно большим значением для многих ситуаций.
Иногда полезно создать поток для каждого элемента, потому что создание потока — это достаточно низкая операция на графическом процессоре, но если каждый поток содержит общую операцию, увеличение количества потоков также представляет эту часть. Стоимость становится больше, например
__global__
void kernel(const float* x, const float* v, float* y) {
const float sqrt_v = sqrt(*v);
const int idx = blockIdx* gridDim.x + threadIdx.x;
y[idx] = x[idx] * sqrt_v;
}
Обработка v в этом ядре является обычной, если мы уменьшим количество потоков и зациклимся на y и x, накладные расходы sqrt(*v) будут соответственно уменьшены, но значение grid_size не должно быть меньше числа СМ на GPU, иначе будет СМ в состоянии простоя.
Мы можем представить, что GPU может запланировать количество SM * максимальное количество блоков на SM за раз.Поскольку количество вычислений для каждого блока одинаково, все SM должны завершить вычисление этих блоков почти в одно и то же время, и затем обработайте следующую партию. Каждая партия называется волной. Представьте, если grid_size окажется на один блок больше, чем волна, потому что следующее ядро в потоке не начнет выполняться до тех пор, пока ядро не будет полностью выполнено, поэтому после завершения первой волны на GPU будет выполняться только один блок, и фактическое использование графического процессора будет очень низким.Эта ситуация называется хвостовым эффектом, и мы должны попытаться избежать этой ситуации. Установка grid_size ровно на одну волну может не избежать эффекта хвоста, потому что GPU может не быть эксклюзивным для текущего потока, и это обычное дело, например, выполнение NCCL, которое требует некоторого SM. Поэтому без особых обстоятельств можно установить grid_size в целочисленную волну с достаточным числом, и зачастую добиться лучшего результата.Если число достаточно велико, нецелочисленная волна часто малоэффективна.
Подводя итог, можно сказать, что в обычном поэлементном ядре или подобных случаях для block_size устанавливается значение 128, а для grid_size устанавливается достаточное количество волн для получения лучшего результата. Но более сложные ситуации требуют конкретного анализа конкретных проблем.Например, если SM может выполнять только несколько блоков одновременно из-за ограничения shared_memory, то увеличение block_size имеет возможность улучшить производительность.Если есть межпотоковое синхронизации в ядре, то слишком большой размер блока приводит к уменьшению фактического использования SM, что у нас есть возможность обсудить отдельно.
Добро пожаловать, чтобы загрузить и испытать новое поколение среды глубокого обучения OneFlow с открытым исходным кодом:GitHub.com/oneflow-Inc…