Использовать данные, распределенные динамически в ядре CUDA на хосте

Question

Я пытаюсь создать класс контейнера на устройстве, которое управляет некоторой памятью. Эта память распределяется динамически и заполняется во время построения объекта в ядре. В соответствии с документацией, которая может быть выполнена с помощью простого нового [] в ядре (с использованием CUDA 8.0 с возможностью вычисления 5.0 в Visual Studio 2012). Впоследствии я хочу получить доступ к данным внутри контейнеров в главном коде (например, для тестирования, если все значения верны).

Минимальная версия DeviceContainer класса выглядит следующим образом:

class DeviceContainer 
{ 
public: 
    __device__ DeviceContainer(unsigned int size); 
    __host__ __device__ ~DeviceContainer(); 

    __host__ __device__ DeviceContainer(const DeviceContainer & other); 
    __host__ __device__ DeviceContainer & operator=(const DeviceContainer & other); 

    __host__ __device__ unsigned int getSize() const { return m_sizeData; } 
    __device__ int * getDataDevice() const { return mp_dev_data; } 
    __host__ int* getDataHost() const; 

private: 
    int * mp_dev_data; 
    unsigned int m_sizeData; 
}; 


__device__ DeviceContainer::DeviceContainer(unsigned int size) : 
     m_sizeData(size), mp_dev_data(nullptr) 
{ 
    mp_dev_data = new int[m_sizeData]; 

    for(unsigned int i = 0; i < m_sizeData; ++i) { 
     mp_dev_data[i] = i; 
    } 
} 


__host__ __device__ DeviceContainer::DeviceContainer(const DeviceContainer & other) : 
    m_sizeData(other.m_sizeData) 
{ 
#ifndef __CUDA_ARCH__ 
    cudaSafeCall(cudaMalloc((void**)&mp_dev_data, m_sizeData * sizeof(int))); 
    cudaSafeCall(cudaMemcpy(mp_dev_data, other.mp_dev_data, m_sizeData * sizeof(int), cudaMemcpyDeviceToDevice)); 
#else 
    mp_dev_data = new int[m_sizeData]; 
    memcpy(mp_dev_data, other.mp_dev_data, m_sizeData * sizeof(int)); 
#endif 
} 


__host__ __device__ DeviceContainer::~DeviceContainer() 
{ 
#ifndef __CUDA_ARCH__ 
    cudaSafeCall(cudaFree(mp_dev_data)); 
#else 
    delete[] mp_dev_data; 
#endif 
    mp_dev_data = nullptr; 
} 


__host__ __device__ DeviceContainer & DeviceContainer::operator=(const DeviceContainer & other) 
{ 
    m_sizeData = other.m_sizeData; 

#ifndef __CUDA_ARCH__ 
    cudaSafeCall(cudaMalloc((void**)&mp_dev_data, m_sizeData * sizeof(int))); 
    cudaSafeCall(cudaMemcpy(mp_dev_data, other.mp_dev_data, m_sizeData * sizeof(int), cudaMemcpyDeviceToDevice)); 
#else 
    mp_dev_data = new int[m_sizeData]; 
    memcpy(mp_dev_data, other.mp_dev_data, m_sizeData * sizeof(int)); 
#endif 

    return *this; 
} 


__host__ int* DeviceContainer::getDataHost() const 
{ 
    int * pDataHost = new int[m_sizeData]; 
    cudaSafeCall(cudaMemcpy(pDataHost, mp_dev_data, m_sizeData * sizeof(int), cudaMemcpyDeviceToHost)); 
    return pDataHost; 
} 

Он просто управляет массив mp_dev_data. Массив создается и заполняется последовательными значениями во время построения, что должно быть возможно только на устройстве. (Обратите внимание, что на самом деле размер контейнеров может отличаться друг от друга.)

Я думаю, мне нужно предоставить конструктор копирования и оператор присваивания, так как я не знаю другого способа заполнить массив в ядро. (См. Вопрос № 3 ниже.) Поскольку копирование и удаление также могут произойти на хосте, __CUDA_ARCH__ используется для определения того, для какого пути выполнения мы компилируем. На хосте cudaMemcpy и cudaFree используется устройство, которое мы можем использовать только memcpy и delete[].

Ядро для создания объекта достаточно прост:

__global__ void createContainer(DeviceContainer * pContainer, unsigned int numContainer, unsigned int containerSize) 
{ 
    unsigned int offset = blockIdx.x * blockDim.x + threadIdx.x; 

    if(offset < numContainer) 
    { 
     pContainer[offset] = DeviceContainer(containerSize); 
    } 
} 

Каждый поток в одномерной решетке, которая находится в диапазоне создает единый объект-контейнер.

Основная функция затем выделяет массивы для контейнера (90000 в данном случае) на устройстве и хост, вызывает ядро ​​и пытается использовать объекты:

void main() 
{ 
    const unsigned int numContainer = 90000; 
    const unsigned int containerSize = 5; 

    DeviceContainer * pDevContainer; 
    cudaSafeCall(cudaMalloc((void**)&pDevContainer, numContainer * sizeof(DeviceContainer))); 

    dim3 blockSize(1024, 1, 1); 
    dim3 gridSize((numContainer + blockSize.x - 1)/blockSize.x , 1, 1); 

    createContainer<<<gridSize, blockSize>>>(pDevContainer, numContainer, containerSize); 
    cudaCheckError(); 

    DeviceContainer * pHostContainer = (DeviceContainer *)malloc(numContainer * sizeof(DeviceContainer)); 
    cudaSafeCall(cudaMemcpy(pHostContainer, pDevContainer, numContainer * sizeof(DeviceContainer), cudaMemcpyDeviceToHost)); 

    for(unsigned int i = 0; i < numContainer; ++i) 
    { 
     const DeviceContainer & dc = pHostContainer[i]; 

     int * pData = dc.getDataHost(); 
     for(unsigned int j = 0; j < dc.getSize(); ++j) 
     { 
     std::cout << pData[j]; 
     } 
     std::cout << std::endl; 
     delete[] pData; 
    } 

    free(pHostContainer); 
    cudaSafeCall(cudaFree(pDevContainer)); 
} 

я должен использовать malloc для массива создание на хосте, поскольку я не хочу иметь конструктор по умолчанию для DeviceContainer. Я пытаюсь получить доступ к данным внутри контейнера через getDataHost(), который внутренне просто вызывает cudaMemcpy.

cudaSafeCall и cudaCheckError - это простые макросы, которые оценивают cudaError, возвращаемые функцией oder, активно опроса последней ошибки. Для полноты картины:

#define cudaSafeCall(error) __cudaSafeCall(error, __FILE__, __LINE__) 
#define cudaCheckError() __cudaCheckError(__FILE__, __LINE__) 

inline void __cudaSafeCall(cudaError error, const char *file, const int line) 
{ 
    if (error != cudaSuccess) 
    { 
     std::cerr << "cudaSafeCall() returned:" << std::endl; 
     std::cerr << "\tFile: " << file << ",\nLine: " << line << " - CudaError " << error << ":" << std::endl; 
     std::cerr << "\t" << cudaGetErrorString(error) << std::endl; 

     system("PAUSE"); 
     exit(-1); 
    } 
} 


inline void __cudaCheckError(const char *file, const int line) 
{ 
    cudaError error = cudaDeviceSynchronize(); 
    if (error != cudaSuccess) 
    { 
     std::cerr << "cudaCheckError() returned:" << std::endl; 
     std::cerr << "\tFile: " << file << ",\tLine: " << line << " - CudaError " << error << ":" << std::endl; 
     std::cerr << "\t" << cudaGetErrorString(error) << std::endl; 

     system("PAUSE"); 
     exit(-1); 
    } 
} 

У меня есть 3 проблемы с этим кодом:

1. Если она выполняется, как представлено здесь, я получаю в «энное сбой запуска» ядра. Отладчик Nsight останавливает меня на линии mp_dev_data = new int[m_sizeData]; (либо в конструкторе, либо в операторе присваивания) и сообщает о нарушении прав доступа в глобальной памяти. Количество нарушений, по-видимому, является случайным между 4 и 11, и они встречаются в непоследовательных потоках, но всегда рядом с верхним концом сетки (блоки 85 и 86).

2. Если я уменьшить numContainer до 10, ядро ​​проходит гладко, однако, cudaMamcpy в getDataHost() завершается с ошибкой недопустимого аргумента - даже если mp_dev_data не 0.(Я подозреваю, что назначение неисправно, и память уже удалена другим объектом.)

3. Хотя я хотел бы знать, как правильно реализовать DeviceContainer с надлежащим управлением памятью, в моем случае это также было бы достаточным для того, чтобы сделать его не скопированным и не подлежащим передаче. Однако я не знаю, как правильно заполнить контейнер-массив в ядре. Может быть что-то вроде

   DeviceContainer dc(5); memcpy(&pContainer[offset], &dc, sizeof(DeviceContainer));

   Что приведет к проблемам с удалением mp_dev_data в деструкторе. Мне нужно вручную управлять удалением памяти, которое кажется довольно грязным.

Я также попытался использовать malloc и free в коде ядра вместо new и delete, но результаты были одинаковыми.

Прошу прощения за то, что я не смог задать свой вопрос короче.

TL; DR: Как реализовать класс, который динамически распределяет память в ядре и может также использоваться в коде хоста? Как я могу инициализировать массив в ядре с объектами, которые нельзя скопировать или назначить?

Любая помощь приветствуется. Спасибо.

Answer 1

Видимо, ответ: то, что я пытаюсь сделать, более или менее невозможно. Память, выделенная new или malloc в ядре, не помещается в глобальную память, а скорее в специальную память кучи, недоступную для хоста.

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

Нарушение прав доступа связано с ограниченным размером кучи (который может быть изменен на cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size).