Внедрение критического раздела в CUDA

Ответ 1

Хорошо, я понял это, и это еще одна-одна-из-куда-парадигма-боли.

Как знает любой хороший программист cuda (обратите внимание, что я не помню этого, что делает меня плохим программистом cuda, я думаю) все потоки в warp должны выполнять тот же код. Код, который я написал, будет работать отлично, если бы не этот факт. Однако, как бы то ни было, есть, вероятно, два потока в одном и том же варпе, доступ к одному и тому же замку. Если один из них получает блокировку, он просто забывает о выполнении цикла, но он не может продолжить цикл, пока все остальные потоки в его warp не завершили цикл. К сожалению, другой поток никогда не завершится, потому что он ждет, когда первый разблокируется.

Вот ядро, которое сделает трюк без ошибок:

__global__ void k_testLocking(unsigned int* locks, int n) {
    int id = threadIdx.x % n;
    bool leaveLoop = false;
    while (!leaveLoop) {
        if (atomicExch(&(locks[id]), 1u) == 0u) {
            //critical section
            leaveLoop = true;
            atomicExch(&(locks[id]),0u);
        }
    } 
}

Ответ 2

Кстати, нужно помнить, что глобальная память пишет и! чтения не завершены, где и пишут их в коде... так что для этого вам нужно добавить глобальное memfence, то есть __threadfence()

Ответ 3

Плакат уже нашел ответ на свой вопрос. Тем не менее, в приведенном ниже коде, я предоставляю общую структуру для реализации критического раздела в 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, существует вероятность того, что один из потоков в warp, например поток с индексом локального потока t≠0, заблокирует состояние. В этом случае другие потоки в одном и том же warp t, в том числе с threadIdx.x == 0, будут выполнять то же самое, что и оператор цикла, в виде потока t, являющийся выполнением потоков в том же самом ване, который выполняется в lockstep. Соответственно, все потоки будут ждать, пока кто-то разблокирует состояние, но ни один другой поток не сможет это сделать, и код застрянет в тупике.