Реализация критического раздела в CUDA - PullRequest
13 голосов
/ 07 января 2010

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

#include <cuda_runtime.h>
#include <cutil_inline.h>
#include <stdio.h>

__global__ void k_testLocking(unsigned int* locks, int n) {
    int id = threadIdx.x % n;
    while (atomicExch(&(locks[id]), 1u) != 0u) {} //lock
    //critical section would go here
    atomicExch(&(locks[id]),0u); //unlock
}

int main(int argc, char** argv) {
    //initialize the locks array on the GPU to (0...0)
    unsigned int* locks;
    unsigned int zeros[10]; for (int i = 0; i < 10; i++) {zeros[i] = 0u;}
    cutilSafeCall(cudaMalloc((void**)&locks, sizeof(unsigned int)*10));
    cutilSafeCall(cudaMemcpy(locks, zeros, sizeof(unsigned int)*10, cudaMemcpyHostToDevice));

    //Run the kernel:
    k_testLocking<<<dim3(1), dim3(256)>>>(locks, 10);

    //Check the error messages:
    cudaError_t error = cudaGetLastError();
    cutilSafeCall(cudaFree(locks));
    if (cudaSuccess != error) {
        printf("error 1: CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
        exit(-1);
    }
    return 0;
}

Этот код, к сожалению, сильно замораживает мою машину на несколько секунд и, наконец, выходит, выводя сообщение:

fcudaSafeCall() Runtime API error in file <XXX.cu>, line XXX : the launch timed out and was terminated.

, что означает, что один из циклов while не возвращается, но, похоже, это должно сработать.

В качестве напоминания atomicExch(unsigned int* address, unsigned int val) атомарно устанавливает значение ячейки памяти, сохраненной в адресе, равным val и возвращает значение old. Итак, идея моего механизма блокировки заключается в том, что он изначально 0u, поэтому один поток должен пройти цикл while, а все остальные потоки должны ждать в цикле while, так как они будут читать locks[id] как 1u , Затем, когда поток завершается с критической секцией 1020 *, он сбрасывает блокировку на 0u, чтобы другой поток мог войти.

Чего мне не хватает?

Кстати, я компилирую с:

nvcc -arch sm_11 -Ipath/to/cuda/C/common/inc XXX.cu

Ответы [ 3 ]

20 голосов
/ 07 января 2010

Хорошо, я понял это, и это еще одна боль в парадигме.

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

Вот ядро, которое выполнит трюк без ошибок:

__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);
        }
    } 
}
5 голосов
/ 13 апреля 2017

Автор уже нашел ответ на свой вопрос. Тем не менее, в приведенном ниже коде я предоставляю общую структуру для реализации критической секции в 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. Соответственно, все потоки будут ждать, пока кто-нибудь разблокирует состояние, но ни один другой поток не сможет это сделать, и код застрянет в тупике.

3 голосов
/ 20 января 2010

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

...