メモリを管理するデバイスにコンテナクラスを構築しようとしています。 このメモリは動的に割り当てられ、カーネルのオブジェクト構築時に埋められます。 カーネルの単純な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）を割り当てる：
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));
}
Iはアレイのmallocを使用する必要が私は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）。mp_dev_dataが0でなくても - 私は10にnumContainerを減らす場合、カーネルがスムーズに実行、しかし、
getDataHost()でcudaMamcpyは無効な引数エラーで失敗します。（私は割り当てが間違っていると思われ、メモリは既に別のオブジェクトによって削除されていると思われます）DeviceContainerを適切なメモリ管理で正しく実装する方法を知りたいそれをコピー不可能かつ割り当て不可能にするのに十分である。しかし、カーネル内のコンテナ配列を正しく埋める方法がわかりません。たぶん、デストラクタでmp_dev_dataを削除して問題につながるDeviceContainer dc(5); memcpy(&pContainer[offset], &dc, sizeof(DeviceContainer));のようなもの。むしろ汚れていると感じるメモリ削除を手動で管理する必要があります。
newと
deleteに
mallocと
freeを使用しようとしましたが、結果は同じでした。
短い質問で私の質問に答えることができなかったことは申し訳なく思っています。
TL; DR：カーネルに動的にメモリを割り当て、ホストコードでも使用できるクラスを実装する方法は？カーネル内のオブジェクトをコピーまたは割り当てできない配列を初期化するにはどうすればよいですか？
何か助けていただければ幸いです。ありがとうございました。