Die Implementierung eines kritischen Abschnitt in CUDA

Question

Ich versuche ein kritische Abschnitt in CUDA mit atomaren Befehlen zu implementieren, aber ich lief in einige Schwierigkeiten. Ich habe das Testprogramm erstellt, das Problem zu zeigen:

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

Dieser Code leider hart friert meine Maschine für einige Sekunden und schließlich verlässt, um die Nachricht auszudrucken:

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

Das bedeutet, dass einer von dem, während Schleifen nicht zurückkehren, aber es scheint, wie dies funktionieren soll.

Zur Erinnerung atomicExch(unsigned int* address, unsigned int val) atomar den Wert der in der Adresse gespeichert Speicherplatz stellt den val Wert old und zurückgibt. So ist die Idee hinter meinem Verriegelungsmechanismus ist, dass es zunächst 0u ist, so ein Thread vorbei an der while Schleife erhalten soll, und alle anderen Threads auf der while Schleife warten sollte, da sie locks[id] als 1u lesen. Dann, wenn der Faden mit dem kritischen Abschnitt erfolgt , setzt er das Schloss 0u so kann ein anderen Thread ein.

Was bin ich fehlt?

By the way, ich bin Kompilieren mit:

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

Answer 1

Okay, ich es herausgefunden, und dies ist noch-another-one-of-the-CUDA-Paradigma-Schmerzen.

Wie jeder gute CUDA Programmierer weiß (man beachte, dass habe ich nicht daran erinnern, was mich ein schlechter CUDA-Programmierer macht, glaube ich) alle Fäden in einer Kette müssen den gleichen Code ausführen. Der Code, den ich geschrieben würde perfekt funktionieren, wenn nicht für diese Tatsache. Wie es ist, gibt es jedoch wahrscheinlich die gleiche Verriegelungs Zugriff zwei Threads in der gleichen Kette sein. Wenn einer von ihnen die Sperre erwirbt, es vergisst, nur um die Schleife ausgeführt wird, aber es kann nicht über die Schleife fortgesetzt, bis alle anderen Threads in seiner Kette die Schleife abgeschlossen haben. Leider füllen Sie bitte den anderen Thread nie, weil es für die erste wartet zu entsperren.

Hier ist ein Kernel, der den Trick ohne Fehler tun:

__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

Das Plakat hat bereits eine Antwort auf seine eigene Frage. Dennoch in der unten stehenden Code, ich bin die Bereitstellung eines allgemeinen Rahmens eine kritischen Abschnitt in CUDA zu implementieren. Genauer gesagt, führt der Code eine Blockzählung, aber es ist leicht modifyiable andere Operationen Gastgeber in einem kritischen Abschnitt durchgeführt werden. Unten, ich bin eine Erklärung des Codes auch berichtet, mit einigen, „typischen“ Fehlern bei der Umsetzung der kritischen Abschnitte in CUDA.

der Code

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

Code Erläuterung

Kritische Abschnitte sind Sequenzen von Operationen, die sequentiell durch den CUDA Threads ausgeführt werden müssen.

Nehmen wir einen Kernel zu konstruieren, die die Aufgabe der Berechnung der Anzahl der Gewindeblöcke eines Gewindes Gitter hat. Eine mögliche Idee ist jeder Thread in jedem Block mit threadIdx.x == 0 läßt einen globalen Zähler erhöhen. Um zu verhindern, Rennbedingungen, alle Erhöhungen müssen nacheinander auftreten, so müssen sie in einem kritischen Abschnitt eingebracht werden.

Der obige Code hat zwei Kernfunktionen: blockCountingKernelNoLock und blockCountingKernelLock. Der ehemalige keinen kritischen Abschnitt verwenden, um die Zähler und zu erhöhen, wie man sehen kann, gibt falsche Ergebnisse. Letztere kapselt die Zähler Erhöhung innerhalb eines kritischen Abschnitt und erzeugt so korrekte Ergebnisse. Aber wie funktioniert der kritische Abschnitt arbeiten?

Der kritische Abschnitt wird von einem globalen Zustand d_state geregelt. Anfangs ist der Zustand 0. Darüber hinaus können zwei __device__ Methoden, lock und unlock kann diesen Zustand ändern. Die lock unlock und Methoden können nur innerhalb eines jeden Blockes und insbesondere durch einen einzelnen Thread aufgerufen werden, durch den Faden lokalen Thread-Index threadIdx.x == 0 aufweist.

Randomly während der Ausführung einer der Fäden lokalen Thread-Index threadIdx.x == 0 und globalen Thread Index, sagen wir, t wird der erste Aufruf des lock Methode. Insbesondere wird es starten atomicCAS(d_state, 0, 1). Da anfänglich d_state == 0, dann wird zu d_state 1 aktualisiert wird, wird atomicCAS 0 zurückzukehren, und der Faden wird die lock Funktion verlassen, an die Anweisungsaktualisierungsgeben. In der Zwischenzeit ein solcher Faden führt die genannten Operationen alle anderen Threads aller anderen Blöcke threadIdx.x == 0 aufweist, wird die lock Verfahren auszuführen. Sie werden jedoch einen Wert von d_state gleich finden 1, so dass atomicCAS(d_state, 0, 1) kein Update durchführen und wird 1 zurückkehren, so diese Fäden verlassen die while-Schleife ausgeführt wird. Danach Gewinde t das Update durchführt, dann führt es die unlock Funktion, nämlich atomicExch(d_state, 0), wodurch die Wiederherstellung zu d_state 0. Zu diesem Zeitpunkt zufällig, ein anderer der Fäden mit threadIdx.x == 0 rastet wieder den Zustand.

Der obige Code enthält auch eine dritte Kernfunktion, nämlich blockCountingKernelDeadlock. Dies ist jedoch eine andere falsche Umsetzung des kritischen Abschnitt, was zu Deadlocks. Tatsächlich erinnern wir uns, dass Verwerfungen in Lockstep betrieben werden und sie synchronisieren nach jeder Anweisung. Wenn wir also blockCountingKernelDeadlock auszuführen, besteht die Möglichkeit, dass eines der Gewinde in einer Kett, sagen wir ein Gewinde mit lokalen Thread-Index t≠0, wird den Zustand verriegeln. Unter diesem Umstand sind die anderen Fäden in der gleichen Kette von t, einschließlich der mit threadIdx.x == 0, werden dieselben ausführen, während Schleifenanweisung als Gewinde t, die Ausführung von Threads in der gleichen Kette in Lockstep durchgeführt wird. Dementsprechend werden alle Fäden für jemanden warten, um den Zustand zu entriegeln, aber kein anderer Thread, dies zu tun in der Lage, und der Code wird in einem D steckeneadlock.

Answer 3

durch die Art und Weise u haben, sich daran zu erinnern, dass die globale Speicher schreibt und! nicht abgeschlossen liest, wo u sie in den Code schreiben ... so für das Sie eine globale memfence dh __threadfence () hinzufügen müssen sein üben