Копирование структуры, содержащей указатели на устройство CUDA

Я работаю над проектом, где мне нужно, чтобы мое устройство CUDA выполняло вычисления в структуре, содержащей указатели.

typedef struct StructA {
    int* arr;
} StructA;

Когда я выделяю память для структуры и затем копирую ее на устройство, она копирует только структуру, а не содержимое указателя. Прямо сейчас я работаю над этим, сначала выделив указатель, а затем установил структуру узла, чтобы использовать этот новый указатель (который находится на графическом процессоре). Следующий пример кода описывает этот подход, используя структуру сверху:

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

Мой вопрос: Это способ сделать это?

Кажется, что очень много работы, и я напоминаю вам, что это очень простая структура. Если в моей структуре было много указателей или структур с указателями, код для размещения и копирования будет довольно обширным и запутанным.

Ответ 1

Изменить: CUDA 6 представляет унифицированную память, что значительно облегчает эту проблему с "глубокой копией". Подробнее см. этот пост.

Не забывайте, что вы можете передавать структуры по значению ядрам. Этот код работает:

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

Это означает, что вам нужно скопировать массив только на устройство, а не на структуру:

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;

Ответ 2

Как отметил Марк Харрис, структуры могут передаваться значениями в ядра CUDA. Однако следует позаботиться о создании надлежащего деструктора, поскольку деструктор вызывается при выходе из ядра.

Рассмотрим следующий пример

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

с uncommented destructor (не обращайте слишком много внимания на то, что на самом деле делает код). Если вы запустите этот код, вы получите следующий вывод

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

После этого два вызова деструктора, один раз на выходе ядра и один раз на основном выходе. Сообщение об ошибке связано с тем, что, если ячейки памяти, на которые указывает d_state, освобождаются при выходе ядра, они больше не могут быть освобождены на основном выходе. Соответственно, деструктор должен отличаться для исполнения хоста и устройства. Это выполняется комментарием деструктора в приведенном выше коде.

Ответ 3

Структура массивов - это кошмар в cuda. Вам придется скопировать каждый указатель на новую структуру, которую устройство может использовать. Может быть, вы могли бы использовать массив структур? Если бы не единственный способ, который я нашел, это атаковать его так, как вы делаете, что никоим образом не симпатично.

EDIT: так как я не могу комментировать верхнюю запись: Шаг 9 избыточен, так как вы можете изменить шаги 8 и 9 в

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