CUDA: копия структуры, содержащей двойной указатель [duplicate]

Как отметил Марк Харрис, структуры могут передаваться значениями в ядра 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, освобождаются при выходе ядра, они больше не могут быть освобождены на основном выходе. Соответственно, деструктор должен отличаться для исполнения хоста и устройства. Это выполняется комментарием деструктора в приведенном выше коде.

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

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

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