Инициализация вложенных массивов на устройстве CUDA

Question

Я хочу инициализировать массив структур на устройстве CUDA, где каждая структура содержит массив значений типа double:

typedef struct TimeSeries
{
    double* values;
} TimeSeries;

TimeSeries* allTimeSeries;

Я пытался использовать управляемую память CUDA для инициализации массива структур allTimesSeries,затем инициализируйте каждый массив пар values.

Вот мой минимальный рабочий пример:

#include <cuda.h>
#include <cuda_runtime_api.h>
#include <array>
#include <iostream>

#define cudaCheckErrors(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line)
{   
    if (code != cudaSuccess) 
    {
        fprintf(stderr, "%s %s %d\n", cudaGetErrorString(code), file, line);
        exit(code);
    }
}   

typedef struct TimeSeries
{
    double* values;
} TimeSeries;

void print_host(TimeSeries* all)
{
    printf("CPU %f %f\n", all[0].values[0], all[0].values[1]);
    printf("CPU %f\n", all[1].values[0]);
}

__global__ void print_device(TimeSeries* all)
{
    printf("GPU %f %f\n", all[0].values[0], all[0].values[1]);
    printf("GPU %f\n", all[1].values[0]);
}

int main()
{   
    TimeSeries* all;
    cudaCheckErrors(cudaMallocManaged(reinterpret_cast<void**>(&all),
            2*sizeof(TimeSeries)));

    cudaCheckErrors(cudaMallocManaged(
            reinterpret_cast<void**>(&(all[0].values)), 2*sizeof(double)));
    cudaCheckErrors(cudaMallocManaged(
            reinterpret_cast<void**>(&(all[1].values)), 1*sizeof(double)));

    all[0].values[0] = 4.8;
    all[0].values[1] = 3.3;
    all[1].values[0] = 0.4;

    print_host(all);
    print_device<<<1, 1>>>(all);

    cudaCheckErrors(cudaFree(all[0].values));
    cudaCheckErrors(cudaFree(all[1].values));
    cudaCheckErrors(cudaFree(all));

    return 0;
}

Компиляция с nvcc (CUDA 10) и работа на Tesla K80 дает:

CPU 4.800000 3.300000
CPU 0.400000

Command terminated

Как можно инициализировать эти данные без сбоя графического процессора? Я не против, если решение использует унифицированную (управляемую) память или нет.

Answer 1

Неизбежно, я обнаружил проблему вскоре после публикации моего вопроса. Мне нужно было позвонить cudaDeviceSynchronize() перед вызовом cudaFree():

print_host(all);
print_device<<<1, 1>>>(all);

cudaCheckErrors(cudaDeviceSynchronize());

cudaCheckErrors(cudaFree(all[0].values));
cudaCheckErrors(cudaFree(all[1].values));
cudaCheckErrors(cudaFree(all));

Я нашел проблему с помощью отладчика cuda-gdb .