Поведение ядра CUDA в зависимости от количества блоков и потоков, даже если они не используются
У меня есть следующее ядро:
__global__ void myKernel(int k, int inc, int width, int* d_Xco, int* d_Xnum, bool*
Xvalid, float* d_X)
{
int i, k1;
i = threadIdx.x + blockIdx.x * blockDim.x;
//k1 = threadIdx.y + blockIdx.y * blockDim.y;
if( (i < k) ){
for(k1 = 0; k1 < inc; k1++){
int mul = (d_X[i*inc + k1] >= 2e2);
d_X[i*inc + k1] *= (float)(!mul);
d_Xco[i*width + k1] = k*mul;
d_Xnum[i] += mul;
d_Xvalid[i*inc + k1] = (!mul) ;
}
}// of if
}
который называется так:
int bx = (int)(k/32)+1;
int by = (int)(inc/32)+1;
dim3 b(bDim, 1);
dim3 t(tDim, 1);
cmyKernel<< b, t >>>( k, inc, width, d_Xco, d_Xnum, d_Xvalid, d_X );
cudaThreadSynchronize();
k около 9000 и inc составляет около 5000, поэтому я уверен, что я не превышаю количество блоков. Если myKernel вызывается с 1thread / 1block в y Ядро, кажется, работает нормально, просто меняя количество потоков и блоков в y Например, размерность до 10 дает неправильный вывод, даже если в ядре я не использую потоки и блоки в y, В идеале хотелось бы избавиться от for() с помощью k = threadIdx.y + blockIdx.y * blockDim.y
Решение
Если вы запускаете ядро с размером y = 10, чем вы их используете. Только потому, что вы не используете идентификатор потока threadIdx.y а также blockIdx.y не означает, что потоки не запускаются. Когда вы запустите ядро с размером y = 10, у вас будет 10 потоков с i = 0, 10 потоков с i = 1 и т. Д.
Скажем, вы запускаете 2×2 темы для простоты. У вас будут темы (0,0) (0,1) (1,0) (1,1). В вашем коде переменная i равна 0 для двух потоков (0,0) и (0,1), но threadIdx.y это отличается. Это означает, что оба потока будут оценивать код для одной и той же переменной i и вызывать условия гонки.
Вам необходимо разрешить зависимость между итерациями (d_Xnum [i] + = mul). Один из способов сделать это — использовать atomicAdd(..), Раскомментируйте k1замените цикл на if(k1 < inc) и добавить atomicAdd, Это должно дать вам правильное поведение.
Другие решения
Как уже говорилось в комментариях, ваше текущее решение запускает несколько потоков, каждый из которых будет применять свою работу к одному и тому же пространству памяти. Это происходит из-за того, что порождается несколько потоков, которые все имеют threadIdx.x с одинаковым значением, а значение threadIdx.y отличается. Это означает, что у вас будет несколько потоков, читающих и пишущих в одно и то же пространство памяти одновременно, это имеет много потенциальных проблем, вот краткое описание.
Чтобы избежать этого, есть несколько шагов, которые вы можете предпринять. Например, вы можете использовать синхронизированный доступ к данным (что приведет к значительному замедлению, поскольку потоки ждут, пока другие завершат доступ к данным). Если вы хотите, чтобы каждый поток имел дело с одним элементом ячейки, вам нужно удалить цикл for и вместо этого использовать k1, как и раньше, однако вы должны внимательно рассмотреть чтение и запись в память, поскольку любая часть процесса может быть выполнена до или после любого другие в другой теме!
Основное понимание здесь заключается в том, что вы никогда не можете полагаться на последовательность операций между потоками!
Когда дело доходит до доступа к данным, это помогает думать о всех ваших структурах данных как о сетках, где каждый поток должен получать доступ и изменять данные только в своих собственных координатах, например (3,2) для потока с threadIdx.x == 3 и threadIdx.y == 2. Таким образом, вы можете легко визуализировать поведение потоков и возможные условия гонки. Самый простой способ использовать это — создать одну запись в сетке для каждого элемента ваших выходных данных, поэтому, если у вас есть матрица из 9000×5000 элементов, вы можете потенциально порождать такое количество потоков для начала и оптимизировать их оттуда. Это, конечно, приведет к тому, что графический процессор придется выполнять на всех своих устройствах несколько раз, но это хорошая отправная точка.
Среди прочего, в университете в Осло есть курс для выпускников по этой теме. Вы можете найти эти слайды очень важно для дальнейшего вашего понимания. В частности, смотрите раздел, касающийся дозирования резьбы, сеток и блоков.