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

Question

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 przydzielić pamięci dla struktury, a następnie skopiować go na urządzenie będzie kopiować tylko struct a nie zawartość wskaźnika. Obecnie pracuję nad tym, najpierw przydzielając wskaźnik, a następnie ustawiając strukturę hosta do korzystania z tego nowego wskaźnika (który znajduje się na GPU). Poniższy przykładowy kod opisuje to podejście za pomocą struct z góry:

#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 w ten sposób to zrobić?

Wygląda na to, że mam 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 przydzielania i kopiowania będzie dość obszerny i mylący.

Answer 1

EDIT: CUDA 6 wprowadza Unified Memory co sprawia, że ​​problem "głębokiej kopii" jest o wiele łatwiejszy. Aby uzyskać więcej informacji, patrz this post.

---

Nie zapominaj, że może struktury mijają wartość dla jąder. 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; 
} 

Może to oznacza, że ​​trzeba tylko skopiować tablicę do urządzenia, a nie struktura:

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; 

Answer 2

Struktura tablic jest koszmarem w cuda. Będziesz musiał skopiować każdy wskaźnik do nowej struktury, z której urządzenie może korzystać. Może zamiast tego możesz użyć tablicy struktur? Jeśli nie, jedynym sposobem, jaki znalazłem, jest zaatakowanie go tak, jak to robisz, co w żadnym wypadku nie jest piękne.

EDIT: ponieważ nie mogę dać komentarze na górnym stanowisku: Krok 9 jest zbędny, ponieważ można zmienić krok 8 i 9 do

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

Answer 3

Jak zauważył Mark Harris, struktury mogą być przekazywane przez wartości do Jądra CUDA. Należy jednak zadbać o skonfigurowanie właściwego destruktora, ponieważ destruktor jest wywoływany przy wyjściu z jądra.

Rozważmy 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 komentarzem destructor (nie płacą zbyt wiele uwagi na to, co rzeczywiście robi kod). Po uruchomieniu tego kodu, otrzymasz następujący komunikat

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

Istnieje wtedy dwa wywołania destruktora, raz przy wyjściu jądra i raz na głównym wyjściu. Komunikat o błędzie związany jest z tym, że jeśli lokalizacje pamięci wskazane przez d_state zostaną zwolnione przy wyjściu jądra, nie będą mogły zostać zwolnione przy głównym wyjściu. W związku z tym destruktor musi być inny dla hostów i urządzeń. Jest to realizowane przez komentowany destruktor w powyższym kodzie.