Когда я должен использовать встроенный в CUDA warpSize, в отличие от моей собственной постоянной?

Question

Когда я должен использовать встроенный в CUDA warpSize, в отличие от моей собственной постоянной?

Код устройства nvcc имеет доступ к встроенному значению, warpSize, который установлен на размер основы устройства, выполняющего ядро (то есть 32 в обозримом будущем). Обычно вы не можете отличить это от константы — но если вы попытаетесь объявить массив длины warpSize, вы получите жалобу на его неконстантность … (с CUDA 7.5)

Таким образом, по крайней мере, для этого вы мотивированы, чтобы иметь что-то вроде (редактировать):

enum : unsigned int { warp_size  = 32 };

где-то в ваших заголовках. Но сейчас — что мне лучше и когда? : warpSize, или же warp_size?

Редактировать: warpSize по-видимому, постоянная времени компиляции в PTX. Тем не менее вопрос стоит.

-1

c++constants cuda gpu-warp

Решение

Другие решения

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

Значение доступно:

В CUDA C через warpSizeгде находится не постоянная времени компиляции (PTX WARP_SZ переменная испускается компилятором в таких случаях).
В PTX ассемблер через WARP_SZгде это непосредственная постоянная времени выполнения
Из API времени выполнения как устройства имущество

Не объявляйте свою собственную константу для размера основы, это просто напрашивается на неприятности. Обычный сценарий использования массива в ядре, размер которого кратен размеру деформации, — использование динамически распределенной разделяемой памяти. Вы можете прочитать размер деформации из API хоста во время выполнения, чтобы получить его. Если у вас есть статически объявленный в ядре, вам нужно измерить размер деформации, использовать шаблоны и выбрать правильный экземпляр во время выполнения. Последний может показаться ненужным театром, но это правильно для случая использования, который практически никогда не возникает на практике. Выбор за вами.

10

Источник

Accepted Answer

Вопреки ответу talonmies я нахожу warp_size постоянная вполне приемлемая. Единственная причина использовать warpSize состоит в том, чтобы сделать код напрямую совместимым с возможным будущим оборудованием, которое может иметь перекосы различного размера. Однако когда приходит такое оборудование, код ядра, скорее всего, потребует и других изменений, чтобы оставаться эффективным. CUDA не является аппаратно-независимым языком — напротив, это все еще довольно низкоуровневый язык программирования. В производственном коде используются различные встроенные функции, которые приходят и уходят со временем (например, __umul24).

В тот день, когда мы получим другой размер основы (например, 64), многое изменится:

warpSize придется корректировать, очевидно,
Многие встроенные функции уровня варпа нуждаются в корректировке своей подписи или создании новой версии, например, int __ballotи пока int не должен быть 32-битным, это является чаще всего так!
Итеративным операциям, таким как сокращение уровня деформации, потребуется скорректировать их количество итераций. Я никогда не видел, чтобы кто-нибудь писал:
```
for (int i = 0; i < log2(warpSize); ++i) ...
```
это было бы слишком сложно в чем-то, что обычно является критичным по времени фрагментом кода.
warpIdx а также laneIdx вычисление из threadIdx нужно будет отрегулировать. В настоящее время наиболее типичный код, который я вижу для этого:
```
warpIdx = threadIdx.x/32;
laneIdx = threadIdx.x%32;
```
что сводится к простым операциям сдвига вправо и маски. Однако, если вы замените 32 с warpSize это внезапно становится довольно дорогой операцией!

В то же время, используя warpSize в коде предотвращается оптимизация, поскольку формально это не известная константа времени компиляции.
Кроме того, если объем общей памяти зависит от warpSize это заставляет вас использовать динамически выделенный шмем (согласно ответу талонмий). Однако использовать синтаксис для этого неудобно, особенно если у вас есть несколько массивов — это заставляет вас самостоятельно выполнять арифметику указателей и вручную вычислять сумму всего использования памяти.

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

deviceFunction<warp_size>(params)

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

Мое предложение будет иметь один заголовок, который контролирует все константы, характерные для модели, например,

#if __CUDA_ARCH__ <= 600
//all devices of compute capability <= 6.0
static const int warp_size = 32;
#endif

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

2