Автор уже нашел ответ на свой вопрос. Тем не менее, в приведенном ниже коде я предоставляю общую структуру для реализации критической секции в CUDA. Более подробно, код выполняет подсчет блоков, но его легко модифицировать для размещения других операций, выполняемых в критической секции . Ниже я также сообщаю о некотором объяснении кода с некоторыми «типичными» ошибками в реализации критических разделов в CUDA.
КОД
#include <stdio.h>
#include "Utilities.cuh"
#define NUMBLOCKS 512
#define NUMTHREADS 512 * 2
/***************/
/* LOCK STRUCT */
/***************/
struct Lock {
int *d_state;
// --- Constructor
Lock(void) {
int h_state = 0; // --- Host side lock state initializer
gpuErrchk(cudaMalloc((void **)&d_state, sizeof(int))); // --- Allocate device side lock state
gpuErrchk(cudaMemcpy(d_state, &h_state, sizeof(int), cudaMemcpyHostToDevice)); // --- Initialize device side lock state
}
// --- Destructor
__host__ __device__ ~Lock(void) {
#if !defined(__CUDACC__)
gpuErrchk(cudaFree(d_state));
#else
#endif
}
// --- Lock function
__device__ void lock(void) { while (atomicCAS(d_state, 0, 1) != 0); }
// --- Unlock function
__device__ void unlock(void) { atomicExch(d_state, 0); }
};
/*************************************/
/* BLOCK COUNTER KERNEL WITHOUT LOCK */
/*************************************/
__global__ void blockCountingKernelNoLock(int *numBlocks) {
if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; }
}
/**********************************/
/* BLOCK COUNTER KERNEL WITH LOCK */
/**********************************/
__global__ void blockCountingKernelLock(Lock lock, int *numBlocks) {
if (threadIdx.x == 0) {
lock.lock();
numBlocks[0] = numBlocks[0] + 1;
lock.unlock();
}
}
/****************************************/
/* BLOCK COUNTER KERNEL WITH WRONG LOCK */
/****************************************/
__global__ void blockCountingKernelDeadlock(Lock lock, int *numBlocks) {
lock.lock();
if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; }
lock.unlock();
}
/********/
/* MAIN */
/********/
int main(){
int h_counting, *d_counting;
Lock lock;
gpuErrchk(cudaMalloc(&d_counting, sizeof(int)));
// --- Unlocked case
h_counting = 0;
gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice));
blockCountingKernelNoLock << <NUMBLOCKS, NUMTHREADS >> >(d_counting);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost));
printf("Counting in the unlocked case: %i\n", h_counting);
// --- Locked case
h_counting = 0;
gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice));
blockCountingKernelLock << <NUMBLOCKS, NUMTHREADS >> >(lock, d_counting);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost));
printf("Counting in the locked case: %i\n", h_counting);
gpuErrchk(cudaFree(d_counting));
}
КОД ОБЪЯСНЕНИЕ
Критические секции - это последовательности операций, которые должны выполняться последовательно потоками CUDA.
Предположим, для создания ядра, задача которого состоит в вычислении количества блоков потоков в сетке потоков. Одна из возможных идей - позволить каждому потоку в каждом блоке, имеющем threadIdx.x == 0, увеличить глобальный счетчик. Чтобы предотвратить гонки, все увеличения должны происходить последовательно, поэтому они должны быть включены в критическую секцию.
Приведенный выше код имеет две функции ядра: blockCountingKernelNoLock и blockCountingKernelLock. Первый не использует критическую секцию для увеличения счетчика и, как видно, возвращает неверные результаты. Последний инкапсулирует увеличение счетчика в критической секции и, таким образом, дает правильные результаты. Но как работает критический раздел?
Критическая секция управляется глобальным состоянием d_state. Первоначально состояние 0. Кроме того, два __device__ метода, lock и unlock, могут изменить это состояние. Методы lock и unlock могут вызываться только одним потоком в каждом блоке и, в частности, потоком с индексом локального потока threadIdx.x == 0.
Случайно во время выполнения, один из потоков, имеющий локальный индекс потока threadIdx.x == 0 и глобальный индекс потока, скажем, t, будет первым, вызывающим метод lock. В частности, запустится atomicCAS(d_state, 0, 1). Поскольку изначально d_state == 0, то d_state будет обновлено до 1, atomicCAS вернет 0 и поток выйдет из функции lock, перейдя к инструкции обновления. Тем временем такой поток выполняет упомянутые операции, все другие потоки всех других блоков, имеющих threadIdx.x == 0, будут выполнять метод lock. Однако они найдут значение d_state, равное 1, так что atomicCAS(d_state, 0, 1) не будет выполнять обновление и вернет 1, поэтому эти потоки будут работать в цикле while. После того, как этот поток t завершит обновление, он выполнит функцию unlock, а именно atomicExch(d_state, 0), восстановив, таким образом, d_state до 0. В этот момент случайным образом другой поток с threadIdx.x == 0 снова заблокирует состояние.
Приведенный выше код также содержит третью функцию ядра, а именно blockCountingKernelDeadlock. Однако это еще одна неправильная реализация критического раздела, приводящая к тупикам. В самом деле, мы помним, что деформации работают в режиме блокировки и синхронизируются после каждой инструкции. Таким образом, когда мы выполняем blockCountingKernelDeadlock, существует вероятность, что один из потоков в деформации, скажем, поток с индексом локального потока t≠0, заблокирует состояние. При этом обстоятельстве другие потоки в той же деформации t, в том числе с threadIdx.x == 0, будут выполнять ту же инструкцию цикла while, что и нить t, являясь выполнением потоков в той же деформации, выполненной в lockstep. Соответственно, все потоки будут ждать, пока кто-нибудь разблокирует состояние, но ни один другой поток не сможет это сделать, и код застрянет в тупике.