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

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

26 голосов
/ 16 февраля 2012

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

typedef struct StructA {
    int* arr;
} StructA;

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

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

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

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

Ответы [ 3 ]

24 голосов
/ 17 февраля 2012

Редактировать: 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 голосов
/ 10 апреля 2017

Как отметил Марк Харрис, структуры могут передаваться по значениям ядрам 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));
}

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

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

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

0 голосов
/ 16 февраля 2012

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

РЕДАКТИРОВАТЬ: так как я не могу давать комментарии к верхнему посту: Шаг 9 излишний,Вы можете изменить шаги 8 и 9 на 

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