Kopiowanie struktury zawierającej wskaźniki do urządzenia CUDA

Question

Kopiowanie struktury zawierającej wskaźniki do urządzenia CUDA

Pracuję nad projektem, w którym potrzebuję mojego urządzenia CUDA do wykonywania obliczeń na strukturze zawierającej wskaźniki.

typedef struct StructA {
    int* arr;
} StructA;

Kiedy przydzielę pamięć dla struktury, a następnie skopiuję ją do urządzenia, skopiuje tylko strukturę, a nie zawartość wskaźnika. W tej chwili pracuję nad tym, przydzielając wskaźnik najpierw, a następnie ustawiając strukturę hosta, aby używała tego nowego wskaźnika (który znajduje się na GPU). Poniższy przykład kodu opisuje to podejście przy użyciu struktury z powyżej:

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

Moje pytanie brzmi: czy tak to się robi?

Wygląda to na strasznie dużo pracy i przypominam, że jest to bardzo prosta struktura. Jeśli moja struktura zawiera wiele wskaźników lub struktur z samymi wskaźnikami, kod alokacji i kopiowania będzie dość obszerny i mylący.

25

pointers struct device cuda host

Author: Guru Swaroop, 2012-02-16

Source

3 answers

Jak zauważył Mark Harris, struktury mogą być przekazywane przez wartości do jąder CUDA. Należy jednak poświęcić trochę uwagi ustawieniu WŁAŚCIWEGO destruktora, ponieważ Destruktor jest wywoływany przy wyjściu z jądra.

Rozważ następujący przykład

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

Z nieautoryzowanym destruktorem(nie zwracaj zbyt dużej uwagi na to, co faktycznie robi kod). Jeśli uruchomisz ten kod, otrzymasz następujące wyjście

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

Są wtedy dwa wywołania do destruktora, raz na wyjście jądra i raz na głównym wyjściu. Komunikat o błędzie jest związany z tym, że jeśli miejsca pamięci wskazywane przez d_state są zwolnione przy wyjściu jądra, nie mogą być już zwolnione przy wyjściu głównym. W związku z tym Destruktor musi być inny dla uruchamiania hosta i urządzenia. Dokonuje tego komentowany Destruktor w powyższym kodzie.

2

Author: JackOLantern,
Warning: date(): Invalid date.timezone value 'Europe/Kyiv', we selected the timezone 'UTC' for now. in /var/www/agent_stack/data/www/doraprojects.net/template/agent.layouts/content.php on line 54
2017-04-10 17:19:37

Struktura tablic jest koszmarem w cuda. Będziesz musiał skopiować każdy wskaźnik do nowej struktury, której urządzenie może użyć. Może zamiast tego mógłbyś użyć tablicy struktur? Jeśli nie jedynym sposobem, jaki znalazłem, jest zaatakowanie go tak, jak ty, co nie jest w żaden sposób ładne.

Edytuj: ponieważ nie mogę dodawać komentarzy do górnego posta: Krok 9 jest zbędny, ponieważ możesz zmienić krok 8 i 9 na

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

-3

Author: martiert,
Warning: date(): Invalid date.timezone value 'Europe/Kyiv', we selected the timezone 'UTC' for now. in /var/www/agent_stack/data/www/doraprojects.net/template/agent.layouts/content.php on line 54
2012-02-16 11:31:26

score 23 · Accepted Answer

Edit: CUDA 6 wprowadza zunifikowaną pamięć, co znacznie ułatwia ten problem "głębokiej kopii". Zobacz ten post Po Więcej Szczegółów.

Nie zapominaj, że możesz przekazać struktury według wartości do jądra. Ten kod działa:

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

Oznacza to, że wystarczy skopiować tablicę do urządzenia, a nie 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;