pointers - Copiando una estructura que contiene punteros al dispositivo CUDA

device host

Question

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

typedef struct StructA {
    int* arr;
} StructA;

Cuando asigno memoria para la estructura y luego la copio en el dispositivo, solo 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 siguiente ejemplo de código describe este enfoque utilizando la estructura desde 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 muchísimo 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

Como señaló Mark Harris, las estructuras se pueden pasar por valores a los núcleos de 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.

Considera 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 comentario (no prestes demasiada atención a lo que el código realmente hace). Si ejecuta ese código, recibirá la siguiente salida

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

Entonces hay dos llamadas al destructor, una en la salida del kernel y otra en la salida principal. El mensaje de error está relacionado con el hecho de que, si las ubicaciones de memoria apuntadas 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.

Answer 2

struct of arrays 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.

EDITAR: ya que no puedo dar comentarios sobre la publicación principal: el Paso 9 es redundante, ya que puede cambiar los pasos 8 y 9 en

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

Answer 3

Editar: CUDA 6 presenta la memoria unificada, lo que hace que este problema de "copia profunda" sea mucho más fácil. Vea esta publicación para más detalles.

No olvide que puede pasar estructuras por valor a los núcleos. 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;
}

Hacerlo significa que solo tiene que copiar la matriz en el dispositivo, no en 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;