Блог

для функции ядра CUDA получите ветвящееся расхождение, показанное ниже, как его оптимизировать?

 int gx = threadIdx.x   blockDim.x * blockIdx.x; val = g_data[gx];  if (gx % 4 == 0)  val = op1(val); else if (gx % 4 == 1)  val = op2(val); else if (gx % 4 == 2)  val = op3(val); else if (gx % 4 == 3)  val = op4(val);  g_data[gx] = val;  

Если бы я программировал на CUDA, я бы, конечно, ничего этого не делал. Однако, чтобы ответить на ваш вопрос:

как избежать расхождения потоков в этом ядре CUDA?

Вы могли бы сделать что-то вроде этого:

 int gx = threadIdx.x   blockDim.x * blockIdx.x; val = g_data[gx];  int gx_bit_0 = gx amp; 1; int gx_bit_1 = (gx amp; 2) gt;gt; 1; val = (1-gx_bit_1)*(1-gx_bit_0)*op1(val)   (1-gx_bit_1)*(gx_bit_0)*op2(val)   (gx_bit_1)*(1-gx_bit_0)*op3(val)   (gx_bit_1)*(gx_bit_0)*op4(val);  g_data[gx] = val;  

Вот полный тестовый пример:

 $ cat t1914.cu #include lt;iostreamgt;  __device__ float op1(float val) { return val   1.0f;} __device__ float op2(float val) { return val   2.0f;} __device__ float op3(float val) { return val   3.0f;} __device__ float op4(float val) { return val   4.0f;}  __global__ void k(float *g_data){   int gx = threadIdx.x   blockDim.x * blockIdx.x;  float val = g_data[gx];   int gx_bit_0 = gx amp; 1;  int gx_bit_1 = (gx amp; 2) gt;gt; 1;  val = (1-gx_bit_1)*(1-gx_bit_0)*op1(val)   (1-gx_bit_1)*(gx_bit_0)*op2(val)   (gx_bit_1)*(1-gx_bit_0)*op3(val)   (gx_bit_1)*(gx_bit_0)*op4(val);   g_data[gx] = val; }  const int N = 32; int main(){   float *data;  cudaMallocManaged(amp;data, N*sizeof(float));  for (int i = 0; i lt; N; i  ) data[i] = 1.0f;  klt;lt;lt;1,Ngt;gt;gt;(data);  cudaDeviceSynchronize();  for (int i = 0; i lt; N; i  ) std::cout lt;lt; data[i] lt;lt; std::endl; } $ nvcc -o t1914 t1914.cu $ compute-sanitizer ./t1914 ========= COMPUTE-SANITIZER 2 3 4 5 2 3 4 5 2 3 4 5 2 3 4 5 2 3 4 5 2 3 4 5 2 3 4 5 2 3 4 5 ========= ERROR SUMMARY: 0 errors $  

Решение путем изменения работы для каждого потока

Лучшее решение с существующей компоновкой данных — позволить каждому потоку вычислять 4 последовательных значения. Лучше иметь меньше потоков, которые могут работать должным образом, чем больше, которые не могут.

 float* g_data; int gx = threadIdx.x   blockDim.x * blockIdx.x; g_data[4 * gx] = op1(g_data[4 * gx]); g_data[4 * gx   1] = op2(g_data[4 * gx   1]); g_data[4 * gx   2] = op3(g_data[4 * gx   2]); g_data[4 * gx   3] = op4(g_data[4 * gx   3]);  

Если размер g_data не кратен 4, поместите if вокруг операций с индексом. Если оно всегда кратно 4 и правильно выровнено, загрузите и сохраните 4 значения как a float4 для повышения производительности.

Решение путем изменения порядка работы

Как, возможно, следовало из всего моего рассказа о float4, ваши входные данные представляют собой некоторую форму 2D-структуры, в которой каждый из четырех элементов выполняет аналогичную функцию. Возможно, это массив структур или массив векторов-другими словами, матрица.

Для того, чтобы объяснить, что я имею в виду, я рассматриваю это как матрицу Nx4. Если вы перенесете это в матрицу 4xN и примените к ней ядро, большинство ваших проблем исчезнет. Потому что тогда записи, для которых необходимо выполнить одну и ту же операцию, помещаются рядом друг с другом в памяти, и это облегчает написание эффективного ядра. Что-то вроде этого:

 float* g_data; int rows_in_g; int gx = threadIdx.x   blockDim.x * blockIdx.x; int gy = threadIdx.y; floatamp; own_g = g_data[gx   rows_in_g * gy]; switch(gy) { case 0: own_g = op1(own_g); break; case 1: own_g = op2(own_g); break; case 2: own_g = op3(own_g); break; case 3: own_g = op4(own_g); break; default: break; }  

Начните это как 2D-ядро с размером блоков x=32, y=4 и размером сетки x=N/32, y=1.

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