Kopieren eine struct mit Zeigern auf CUDA-Gerät

Arbeite ich an einem Projekt, wo ich mein CUDA-Gerät, um Berechnungen auf eine struct mit Zeigern.

typedef struct StructA {
    int* arr;
} StructA;

Wenn ich Speicher für das struct und dann kopieren Sie Sie auf das Gerät, wird es kopiert nur die Struktur und nicht der Inhalt der pointer. Jetzt arbeite ich, um dieses durch die Zuweisung der Zeiger zuerst, dann legen Sie die host-Struktur zu verwenden, dass die neuen Zeiger (die sich auf die GPU). Das folgende Codebeispiel beschreibt diesen Ansatz mit der struct von oben:

#define N 10

int main() {

    int h_arr[N] = {1,2,3,4,5,6,7,8,9,10};
    StructA *h_a = (StructA*)malloc(sizeof(StructA));
    StructA *d_a;
    int *d_arr;

    //1. Allocate device struct.
    cudaMalloc((void**) &d_a, sizeof(StructA));

    //2. Allocate device pointer.
    cudaMalloc((void**) &(d_arr), sizeof(int)*N);

    //3. Copy pointer content from host to device.
    cudaMemcpy(d_arr, h_arr, sizeof(int)*N, cudaMemcpyHostToDevice);

    //4. Point to device pointer in host struct.
    h_a->arr = d_arr;

    //5. Copy struct from host to device.
    cudaMemcpy(d_a, h_a, sizeof(StructA), cudaMemcpyHostToDevice);

    //6. Call kernel.
    kernel<<<N,1>>>(d_a);

    //7. Copy struct from device to host.
    cudaMemcpy(h_a, d_a, sizeof(StructA), cudaMemcpyDeviceToHost);

    //8. Copy pointer from device to host.
    cudaMemcpy(h_arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost);

    //9. Point to host pointer in host struct.
    h_a->arr = h_arr;
}

Meine Frage ist: Ist dies der Weg, es zu tun?

Es scheint wie eine Menge Arbeit, und ich erinnere Sie daran, dass dies eine sehr einfache Struktur. Wenn mein struct enthalten eine Menge von Zeigern oder Strukturen mit dem Pointer selbst, der code für die Zuweisung und Kopie ist Recht umfangreich und verwirrend.

Die Schritte 7 und 9 sind überflüssig, aber sonst ist ziemlich viel, wie es ist. Als Antwort unten sagt, Sie sind am besten gedient, durch die Vermeidung von komplexen, Zeiger-basierte Datenstrukturen auf der GPU. Die Leistung auf der GPU ist schlimmer, und die APIs sind wirklich nicht dafür geeignet.
Ich kann sehen, dass der Schritt 7 ist überflüssig, aber warum Schritt 9?
gut h_a ist (oder sollte sein) ein "Bild" der device-Struktur gehalten, die in dem host-Speicher. Die Zuordnung zu halten, einen Zeiger in dem host-Speicher ist wahrscheinlich eine Kombination von schlechte Praxis/falsch/device memory-leak-je nachdem, was Ihre wahren Absichten sind. Nachdem Sie kopiert den Inhalt d_a zurück h_a Sie haben "der Kreis" und sind wieder dort wo Sie begann.
Aber um die Kopie der struct korrekt an das Gerät muss ich den Zeiger von h_a zu d_arr (Schritt 4). Also, wenn ich kopieren Sie die Daten zurück, auch ich habe die Zeiger in h_a dem array habe ich nur kopiert. Ich bin damit einverstanden, dass der Schritt 7 ist überflüssig in meinem Beispiel oben, weil es keine anderen Daten in der Struktur, aber wenn es war, dass Schritt würde nicht überflüssig sein.. Oder bin ich komplett Irre?
Danke, tahatmat, für die Bereitstellung von uns mit diesem Muster kopieren von Strukturen über host-und device-Erinnerungen hin und her. Aber ich glaube, es ist nur erwähnenswert, einen zweiten Weg, der scheint etwas konsequenter und hilft zu vermeiden, die Umsetzung Schritt 9. Die spezifische Funktion cudaMemcpy() erlaubt tatsächlich die Dereferenzierung Gerät Zeiger in einen host-Codepunkt in einer solchen Art und Weise: Sie Schritt 4 überspringen und nach dem kopieren h_a zu d_a auf Schritt 5 kopieren Sie manuell jedes Gerät pointer-Adresse in d_a, wie diese: cudaMemcpy(&(d_a->arr) &(d_arr), sizeof(int*), cudaMemcpyHostToDevice). Wieder, "d_a->arr" legit ist

InformationsquelleAutor Thorkil Holm-Jacobsen | 2012-02-16

Edit: CUDA 6 führt Unified-Memory, das macht dieses "Tiefe Kopie" - problem viel einfacher. Sehen dieser Beitrag für mehr details.

Vergessen Sie nicht, dass Sie kann pass Strukturen von Wert-Kernel. Dieser code funktioniert:

//pass struct by value (may not be efficient for complex structures)
__global__ void kernel2(StructA in)
{
    in.arr[threadIdx.x] *= 2;
}

Tun, so dass bedeutet, dass Sie nur kopieren müssen Sie das array an das Gerät, nicht auf die Struktur:

int h_arr[N] = {1,2,3,4,5,6,7,8,9,10};
StructA h_a;
int *d_arr;

//1. Allocate device array.
cudaMalloc((void**) &(d_arr), sizeof(int)*N);

//2. Copy array contents from host to device.
cudaMemcpy(d_arr, h_arr, sizeof(int)*N, cudaMemcpyHostToDevice);

//3. Point to device pointer in host struct.
h_a.arr = d_arr;

//4. Call kernel with host struct as argument
kernel2<<<N,1>>>(h_a);

//5. Copy pointer from device to host.
cudaMemcpy(h_arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost);

//6. Point to host pointer in host struct 
//   (or do something else with it if this is not needed)
h_a.arr = h_arr;

InformationsquelleAutor harrism

Wie bereits von Mark Harris, Strukturen übergeben werden können, indem Sie Werte, um CUDA-Kernel. Jedoch, einige Sorgfalt gewidmet werden sollte, um ein geeignetes Destruktor, da der Destruktor aufgerufen wird bei der Ausfahrt aus dem kernel.

Betrachten Sie das folgende Beispiel

#include <stdio.h>

#include "Utilities.cuh"

#define NUMBLOCKS  512
#define NUMTHREADS 512 * 2

/***************/
/* TEST 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 (wrong version)
    //~Lock(void) { 
    // printf("Calling destructor\n");
    // gpuErrchk(cudaFree(d_state)); 
    //}

    //--- Destructor (correct version)
// __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 WITH LOCK */
/**********************************/
__global__ void blockCounterLocked(Lock lock, int *nblocks) {

    if (threadIdx.x == 0) {
        lock.lock();
        *nblocks = *nblocks + 1;
        lock.unlock();
    }
}

/********/
/* MAIN */
/********/
int main(){

    int h_counting, *d_counting;
    Lock lock;

    gpuErrchk(cudaMalloc(&d_counting, sizeof(int)));

    //--- Locked case
    h_counting = 0;
    gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice));

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

mit der Destruktor auskommentiert (nicht zahlen zu viel Aufmerksamkeit auf das, was der code tatsächlich tut). Wenn Sie ausführen, dass code, erhalten Sie die folgende Ausgabe

Calling destructor
Counting in the locked case: 512
Calling destructor
GPUassert: invalid device pointer D:/Project/passStructToKernel/passClassToKernel/Utilities.cu 37

Dann gibt es zwei Aufrufe des destruktors, sobald auf den kernel verlassen, und sobald am Hauptausgang. Die Fehlermeldung bezieht sich auf die Tatsache, dass, wenn die memory locations, die Spitzen, von d_state befreit sind, in den kernel verlassen, Sie können nicht freigegeben werden, nicht mehr an den Hauptausgang. Entsprechend der Destruktor müssen verschiedene host-und Geräte-Ausführungen. Dies wird erreicht, indem die kommentierte Destruktor in den obigen code.

InformationsquelleAutor JackOLantern

-3

struct-arrays ist ein Albtraum, der in cuda. Sie müssen auf jeder Kopie der Zeiger auf eine neue Struktur, die das Gerät verwenden können. Vielleicht haben Sie stattdessen verwenden könnte, die ein array von structs? Wenn nicht die einzige Möglichkeit, die ich gefunden habe, ist, um es anzugreifen, wie Sie es tun, das ist in keiner Weise ziemlich.

BEARBEITEN:
da kann ich nicht geben, Kommentare auf den oberen post: Schritt 9 ist redundant, da Sie sich ändern können Schritt 8 und 9 in
```
//8. Copy pointer from device to host.
cudaMemcpy(h->arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost);
```
Erste, diese Antwort ist gefährlich, weil es geht gegen die standard-Weisheit über AOS/SOA, parallel computing. Structure of Arrays (SOA) ist vorzuziehen Array of Structures (AOS) in allen parallel-computing, einschließlich multicore-CPUs mit SSE/AVX-Befehlssätzen. Der Grund dafür ist, dass SOA unterhält die Lokalität der Verweis auf mehrere threads (z.B. benachbarte Elemente von d_a.arr zugegriffen werden, indem benachbarte threads, die gleichzeitig ausgeführt werden). Eine Struktur mit einem Zeiger im es ist NICHT die gleiche wie die Struktur des Arrays. Zweitens vereinfachen Sie diesen code, indem die Struktur von Wert.
Warum ist ein Array von Structs nicht vorzuziehen cuda? Ich verstehe nicht, kannst du mir ein Beispiel oder einen link? Dank
href="http://stackoverflow.com/questions/18136785/kernel-using-aos-is-faster-than-using-soa/18137311#18137311">here eine Frage/Antwort, die beschreibt SOA vs. AOS mit Beispielen.
Danke für die Antwort Robert, aber ich habe bereits eine Frage über die es Here und die Antworten waren ziemlich eindeutig. 🙂

InformationsquelleAutor martiert

Schreibe einen Kommentar

Du musst angemeldet sein, um einen Kommentar abzugeben.