호스트의 CUDA 커널에서 동적으로 할당 된 데이터 사용

일부 메모리를 관리하는 장치에서 컨테이너 클래스를 작성하려고합니다. 이 메모리는 동적으로 할당되며 커널에서 객체를 생성하는 동안 채워집니다. 커널에있는 간단한 new []로 수행 할 수있는 문서에 따르면 (Visual Studio 2012에서 계산 능력 5.0과 함께 CUDA 8.0 사용). 이후 호스트 코드의 컨테이너 안의 데이터에 액세스하려고합니다 (예 : 모든 값이 올바른지 테스트).호스트의 CUDA 커널에서 동적으로 할당 된 데이터 사용

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)에 대한 배열을 할당은 커널을 호출하고 객체 사용을 시도 : 배열 malloc을 사용하는 I 가지고

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

을 호스트에 대한 생성은 DeviceContainer에 대한 기본 생성자를 갖고 싶지 않기 때문에. 내부적으로 cudaMemcpy을 호출하는 getDataHost()을 통해 컨테이너 내부의 데이터에 액세스하려고합니다.

cudaSafeCall 및 cudaCheckError은 oder 함수가 반환 한 cudaError을 평가하는 간단한 매크로로 마지막 오류를 적극적으로 폴링합니다. 완성도를 위해서 : 내가 커널의 "지정되지 않은 발사 실패"를받을 여기에 제시된대로 실행하면

#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 문제가있다. Nsight 디버거는 (생성자 또는 할당 연산자에서) mp_dev_data = new int[m_sizeData]; 줄에서 나를 멈추고 전역 메모리에 몇 가지 액세스 위반을보고합니다. 위반 횟수는 4에서 11 사이에서 무작위로 나타나며 연속되지 않은 스레드에서 발생하지만 항상 그리드의 상단 가까이에 있습니다 (블록 85 및 86).

내가 numContainer 10 줄이면
는, 커널은 그러나, getDataHost()의 cudaMamcpy 잘못된 인수 오류와 함께 실패, 원활하게 실행 - mp_dev_data가 0이 아닌 경우에도.(나는 할당이 잘못되어 메모리가 다른 객체에 의해 이미 삭제 된 것으로 의심된다.) 을 적절한 메모리 관리로 구현하는 방법을 알고 싶지만, 필자의 경우에도 마찬가지이다. 사본을 복사 할 수 없으며 양도 할 수 없도록하기에 충분합니다. 그러나 커널에서 컨테이너 배열을 올바르게 채우는 방법을 모르겠습니다. 소멸자 mp_dev_data을 삭제하여 문제를 일으킬 것

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

같은 아마 뭔가. 내가 수동으로 관리하는 메모리 삭제가 다소 더러운 느낌이 필요합니다.

는 또한 커널 코드 대신 new 및 delete에 malloc 및 free을 사용하려고하지만 결과는 동일했다.

나는 내 질문을 더 짧은 방법으로 틀을 잡을 수 없다는 점에 유감이다.

TL : DR : 커널에서 동적으로 메모리를 할당하고 호스트 코드에서도 사용할 수있는 클래스를 구현하는 방법은 무엇입니까? 복사하거나 할당 할 수없는 객체를 가진 커널에서 배열을 초기화하려면 어떻게해야합니까?

도움을 주시면 감사하겠습니다. 고맙습니다.

출처

2017-02-02 Devon Cornwall

분명히 대답은 : 내가하려는 것은 다소 불가능합니다. 커널에 new 또는 malloc으로 할당 된 메모리는 전역 메모리에 저장되지 않고 호스트에서 액세스 할 수없는 특수 힙 메모리에 저장됩니다.

호스트의 모든 메모리에 액세스하는 유일한 방법은 전역 메모리에 힙의 모든 요소를 저장하고 힙의 모든 요소를 전역 메모리로 복사하는 커널을 작성하는 배열을 먼저 할당하는 것입니다.

제한된 힙 크기로 인해 액세스 위반이 발생했습니다 (cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)에 의해 변경 될 수 있음)

출처

2017-02-02 13:38:26

호스트의 CUDA 커널에서 동적으로 할당 된 데이터 사용

답변

관련 문제