Copiar una estructura que contiene punteros a dispositivo CUDA

Question

que estoy trabajando en un proyecto donde tengo mi dispositivo CUDA para hacer cálculos en una estructura que contiene punteros.

typedef struct StructA { 
    int* arr; 
} StructA; 

Cuando asignar memoria para la estructura y luego copiarlo en el dispositivo, sólo se copiará la estructura y no el contenido del puntero. En este momento estoy trabajando en esto asignando el puntero primero, luego establezco la estructura del host para usar ese nuevo puntero (que reside en la GPU). El ejemplo de código siguiente describe este enfoque utilizando la estructura de arriba:

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

Mi pregunta es: Es esta la manera de hacerlo?

Parece que una gran cantidad de trabajo, y les recuerdo que esta es una estructura muy simple. Si mi estructura contiene muchos punteros o estructuras con punteros, el código de asignación y copia será bastante extenso y confuso.

Answer 1

Editar: CUDA 6 introduce de memoria unificada, lo que hace que este problema de "copia profunda" sea mucho más fácil. Vea this post para más detalles.

---

No olvide que usted puede pase estructuras de valor a los granos. Este código funciona:

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

Si lo hace, significa que sólo tiene que copiar la matriz para el dispositivo, no la estructura:

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

estructura de matrices es una pesadilla en CUDA. Tendrá que copiar cada uno de los punteros a una nueva estructura que el dispositivo puede usar. Tal vez, en su lugar, podría utilizar una serie de estructuras? Si no, la única forma que he encontrado es atacarlo de la manera en que lo haces, lo que de ninguna manera es bonito.

EDIT: ya que no puedo dar comentarios sobre el puesto más alto: Paso 9 es redundante, ya que se puede cambiar el paso 8 y 9 en

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

Answer 3

Como ha señalado Mark Harris, las estructuras se pueden pasar por los valores de Núcleos CUDA. Sin embargo, debe dedicarse un poco de cuidado para configurar un destructor adecuado ya que se llama al destructor a la salida del kernel.

consideremos el siguiente ejemplo

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

con el destructor sin comentar (no prestar demasiada atención en lo que realmente hace el código). Si se ejecuta el código, recibirá el siguiente resultado

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

Hay, pues, dos llamadas al destructor, una vez a la salida del núcleo y una vez en la salida principal. El mensaje de error está relacionado con el hecho de que, si las ubicaciones de memoria señaladas por d_state se liberan en la salida del núcleo, ya no pueden liberarse en la salida principal. En consecuencia, el destructor debe ser diferente para las ejecuciones de host y dispositivo. Esto se logra mediante el destructor comentado en el código anterior.