Copiar una estructura que contiene punteros al dispositivo CUDA

Question

Copiar una estructura que contiene punteros al dispositivo CUDA

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 asigne memoria para la estructura y luego la copie en el dispositivo, solo copiará la estructura y no el contenido del puntero. En este momento estoy trabajando alrededor de esto asignando el puntero primero, luego configure la estructura host para usar ese nuevo puntero (que reside en la GPU). El siguiente ejemplo de código 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 mucho trabajo, y les recuerdo que esta es una estructura muy simple. Si mi estructura contiene muchos punteros o estructuras con punteros, el código para la asignación y copia será bastante extenso y confuso.

25

struct device pointers host cuda

Author: Guru Swaroop, 2012-02-16

Source

3 answers

Como señaló Mark Harris, las estructuras se pueden pasar por valores a núcleos CUDA. Sin embargo, se debe tener cierto cuidado para configurar un destructor adecuado, ya que el destructor se llama al salir del núcleo.

Considere 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 no comentado (no preste 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

Hay entonces dos llamadas al destructor, una vez en 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 apuntadas por d_state se liberan en la salida del núcleo, ya no se pueden liberar en la salida principal. En consecuencia, el destructor debe ser diferente para las ejecuciones de host y dispositivo. Esto es logrado por el destructor comentado en el código anterior.

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/ajaxhispano.com/template/agent.layouts/content.php on line 61
2017-04-10 17:19:37

Estructura de matrices es una pesadilla en cuda. Tendrá que copiar cada puntero a una nueva estructura que el dispositivo pueda usar. Tal vez en su lugar podría utilizar una matriz de estructuras? Si no la única manera que he encontrado es atacarlo como lo haces tú, lo cual no es de ninguna manera bonito.

EDITAR: dado que no puedo dar comentarios en la publicación principal: El paso 9 es redundante, ya que puede cambiar los pasos 8 y 9 a

// 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/ajaxhispano.com/template/agent.layouts/content.php on line 61
2012-02-16 11:31:26

score 23 · Accepted Answer

Edit: CUDA 6 introduce la Memoria unificada, lo que hace que este problema de "copia profunda" sea mucho más fácil. Ver este post 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 el array al 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;