Объявление переменных в CUDA ядра

Question

Допустим, вы объявляете новую переменную в ядре CUDA, а затем использовать его в несколько потоков, как:

__global__ void kernel(float* delt, float* deltb) { 
int i = blockIdx.x * blockDim.x + threadIdx.x; 
float a; 
a = delt[i] + deltb[i]; 
a += 1; 
} 

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

int threads = 200; 
uint3 blocks = make_uint3(200,1,1); 
kernel<<<blocks,threads>>>(d_delt, d_deltb); 

1. является "" хранится в стеке?
2. Является ли новый «a» создан для каждого потока, когда они инициализируются?
3. Или каждый поток будет независимо получать доступ к «а» в неизвестное время, потенциально испортить алгоритм?

Answer 1

Ничего из перечисленного. Компилятор CUDA достаточно умен и достаточно агрессивен с оптимизацией, что он может обнаружить, что a не используется, и полный код можно оптимизировать. Вы можете это подтвердить, компилируя ядро ​​с -Xptxas=-v в качестве опции и посмотрите на количество ресурсов, которое должно в основном не регистрируются и не имеют локальной памяти или кучи.

В менее тривиальном примере a, вероятно, будет храниться в каждом потоковом регистре или в локальной потоковой памяти потока, которая является вне DRAM.

Answer 2

Любая переменная (скаляр или массив), объявленная внутри функции ядра, без спецификатора extern, является локальной для каждого потока, то есть каждый поток имеет свою собственную «копию» этой переменной, при этом не будет происходить расхождение данных среди потоков!

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

Более подробную информацию о том, какие переменные можно найти в локальной памяти, можно найти в руководстве пользователя NVIDIA Cuda, глава 5.3.2.2