Implementando una sección crítica en CUDA

synchronization locking

Question

Estoy tratando de implementar una sección crítica en CUDA usando instrucciones atómicas, pero me encontré con algunos problemas. Creé el programa de prueba para mostrar el problema:

#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;
}

Este código, desafortunadamente, congela mi máquina por varios segundos y finalmente sale, imprimiendo el mensaje:

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

lo que significa que uno de esos ciclos while no está regresando, pero parece que esto debería funcionar.

Como recordatorio, atomicExch(unsigned int* address, unsigned int val) establece atómicamente el valor de la ubicación de memoria almacenada en la dirección en val y devuelve el valor anterior. Entonces la idea detrás de mi mecanismo de bloqueo es que inicialmente es 0u , por lo que un hilo debería pasar el ciclo while y todos los otros hilos deberían esperar en el ciclo while ya que leerán locks[id] como 1u . Luego, cuando el hilo termina con la sección crítica , restablece el bloqueo a 0u para que otro hilo pueda ingresar.

¿Qué me estoy perdiendo?

Por cierto, estoy compilando con:

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

Answer 1

De acuerdo, lo descubrí, y este es todavía otro más de los cuda-paradigma-dolores.

Como cualquier buen programador de cuda sabe (nótese que no recuerdo esto, lo que me convierte en un mal programador de cuda, creo) todos los hilos en un warp deben ejecutar el mismo código. El código que escribí funcionaría perfectamente si no fuera por este hecho. Sin embargo, tal como es, es probable que haya dos subprocesos en la misma urd que accedan al mismo bloqueo. Si uno de ellos adquiere el bloqueo, simplemente se olvida de ejecutar el ciclo, pero no puede continuar más allá del ciclo hasta que todos los otros hilos en su warp hayan completado el ciclo. Lamentablemente, el otro subproceso nunca se completará porque está esperando que se desbloquee el primero.

Aquí hay un kernel que hará el truco sin error:

__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);
        }
    } 
}

Answer 2

por cierto tienes que recordar que la memoria global escribe y! las lecturas no se completan cuando las escribes en el código ... así que para que esto sea una práctica necesitas agregar una memfence global ie __threadfence ()

Answer 3

El póster ya ha encontrado una respuesta a su propio problema. Sin embargo, en el siguiente código, proporciono un marco general para implementar una sección crítica en CUDA. Más en detalle, el código realiza un conteo de bloques, pero es fácilmente modificable para alojar otras operaciones que se realizarán en una sección crítica . A continuación, también informo algunas explicaciones del código, con algunos errores "típicos" en la implementación de secciones críticas en CUDA.

EL CÓDIGO

#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));
}

CÓDIGO EXPLICACIÓN

Las secciones críticas son secuencias de operaciones que deben ser ejecutadas secuencialmente por los hilos CUDA.

Supongamos que se construye un núcleo que tiene la tarea de calcular el número de bloques de hilos de una cuadrícula de hilos. Una posible idea es permitir que cada hilo en cada bloque tenga threadIdx.x == 0 aumentar un contador global. Para evitar condiciones de carrera, todos los aumentos deben ocurrir de forma secuencial, por lo que deben incorporarse en una sección crítica.

El código anterior tiene dos funciones de kernel: blockCountingKernelNoLock y blockCountingKernelLock . El primero no usa una sección crítica para aumentar el contador y, como se puede ver, arroja resultados incorrectos. Este último encapsula el aumento del contador dentro de una sección crítica y, por lo tanto, produce resultados correctos. Pero, ¿cómo funciona la sección crítica?

La sección crítica se rige por un estado global d_state . Inicialmente, el estado es 0 . Además, dos métodos __device__ , lock y unlock , pueden cambiar este estado. Los métodos de lock y unlock solo pueden invocarse mediante un único hilo dentro de cada bloque y, en particular, mediante el hilo que tiene el índice de hilo local threadIdx.x == 0 .

Aleatoriamente durante la ejecución, uno de los subprocesos tiene un índice de subproceso local threadIdx.x == 0 e índice de subproceso global, por ejemplo, t será el primero en invocar el método de lock . En particular, lanzará atomicCAS(d_state, 0, 1) . Como inicialmente d_state == 0 , entonces d_state se actualizará a 1 , atomicCAS devolverá 0 y el hilo saldrá de la función de lock , pasando a la instrucción de actualización. Mientras tanto, tal subproceso realiza las operaciones mencionadas, todos los demás subprocesos de todos los demás bloques que tengan threadIdx.x == 0 ejecutarán el método de lock . Sin embargo, encontrarán un valor de d_state igual a 1 , de modo que atomicCAS(d_state, 0, 1) no realizará ninguna actualización y devolverá 1 , por lo que dejará que estos subprocesos ejecuten el ciclo while. Después de que thread t realiza la actualización, entonces ejecuta la función de unlock , es decir, atomicExch(d_state, 0) , restaurando así d_state a 0 . En este punto, al azar, otro de los hilos con threadIdx.x == 0 volverá a bloquear el estado.

El código anterior también contiene una tercera función del kernel, es decir, blockCountingKernelDeadlock . Sin embargo, esta es otra implementación incorrecta de la sección crítica, lo que lleva a bloqueos. De hecho, recordamos que las urdimbres operan al unísono y se sincronizan después de cada instrucción. Entonces, cuando ejecutamos blockCountingKernelDeadlock , existe la posibilidad de que uno de los hilos en un warp, digamos un hilo con índice de hilo local t≠0 , bloqueará el estado. Bajo esta circunstancia, los otros hilos en la misma urdimbre de t , incluyendo eso con threadIdx.x == 0 , ejecutarán la misma sentencia de ciclo while que thread t , siendo la ejecución de hilos en la misma urdimbre realizada en lockstep. En consecuencia, todos los hilos esperarán a que alguien desbloquee el estado, pero ningún otro hilo podrá hacerlo y el código quedará atrapado en un punto muerto.