CUDA グローバル (C と同様) デバイスメモリに割り当てられた動的配列

https://stackoverflow.com/questions/79121

cuda
nvidia

09-06-2019
|

質問

そこで、Nvidia の CUDA アーキテクチャを利用するコードを作成しようとしています。デバイスとの間でのコピーが全体的なパフォーマンスに悪影響を及ぼしていることに気づき、現在、大量のデータをデバイスに移動しようとしています。

このデータは多くの機能で使用されるため、グローバルにしたいと考えています。はい、ポインターを渡すことはできますが、この場合にグローバルを操作する方法を知りたいと思っています。

したがって、デバイスに割り当てられた配列にアクセスするデバイス関数があります。

理想的には、次のようなことができます。

__device__ float* global_data;

main()
{
  cudaMalloc(global_data);
  kernel1<<<blah>>>(blah); //access global data
  kernel2<<<blah>>>(blah); //access global data again
}

ただし、動的配列の作成方法がわかりません。次のように配列を宣言することで回避策を考え出しました。

__device__ float global_data[REALLY_LARGE_NUMBER];

それには cudaMalloc 呼び出しは必要ありませんが、私は動的割り当てアプローチを好みます。

解決

おそらくこのようなものが機能するはずです。

#include <algorithm>

#define NDEBUG
#define CUT_CHECK_ERROR(errorMessage) do {                                 \
        cudaThreadSynchronize();                                           \
         cudaError_t err = cudaGetLastError();                             \
         if( cudaSuccess != err) {                                         \
                     fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n",    \
                                             errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
                     exit(EXIT_FAILURE);                                                  \
                 } } while (0)


__device__ float *devPtr;

__global__
void kernel1(float *some_neat_data)
{
    devPtr = some_neat_data;
}

__global__
void kernel2(void)
{
    devPtr[threadIdx.x] *= .3f;
}


int main(int argc, char *argv[])
{
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    CUT_CHECK_ERROR("kernel1");

    kernel2<<<1,128>>>();

    CUT_CHECK_ERROR("kernel2");

    return 0;
}

やってみて。

他のヒント

NVIDIA が提供する豊富なドキュメントに集中して時間をかけてください。

プログラミングガイドより:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

これは、メモリを割り当てる方法の簡単な例です。さて、カーネルでは、次のように float へのポインタを受け入れる必要があります。

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x]++;
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= 0.3f;
}

したがって、次のようにして呼び出すことができます。

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

kernel1<<<1,128>>>(devPtr);
kernel2<<<1,128>>>(devPtr);

このデータは多数の機能で使用されているため、グローバルにしたいと思います。

グローバルを使用する正当な理由はほとんどありません。これは間違いなく一つではありません。この例を拡張して「devPtr」をグローバルスコープに移動する練習として残しておきます。

編集：

さて、根本的な問題は次のとおりです。カーネルはデバイスメモリにのみアクセスでき、カーネルが使用できる唯一のグローバルスコープポインタは GPU のものです。CPU からカーネルを呼び出すと、カーネルが実行される前に、バックグラウンドでポインタとプリミティブが GPU レジスタや共有メモリにコピーされます。

したがって、私が提案できる最も近いものは次のとおりです。目標を達成するには、cudaMemcpyToSymbol() を使用してください。しかし、背景としては、別のアプローチが正しいことである可能性があることを考慮してください。

#include <algorithm>

__constant__ float devPtr[1024];

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1];
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= devPtr[2];
}


int main(int argc, char *argv[])
{
    float some_data[256];
    for (int i = 0; i < sizeof(some_data) / sizeof(some_data[0]); i++)
    {
        some_data[i] = i * 2;
    }
    cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr) ));
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    kernel2<<<1,128>>>(otherDevPtr);

    return 0;
}

この例では「--host-compilation=c++」を忘れないでください。

私は一時ポインタを割り当て、それを kernel1 に似た単純なグローバル関数に渡すという解決策を試してみました。

良いニュースは、それが機能するということです:)

ただし、「アドバイザリー:グローバルデータにアクセスしようとするたびに、グローバルメモリ空間を想定して、どのポインタが指しているのかわかりません。」幸いなことに、その仮定はたまたま正しいのですが、警告は煩わしいものです。

とにかく、記録のために、私は多くの例を見て、「正しい!」という出力を得ることが重要な nvidia の演習を実行しました。ただし、見ていない全てそのうちの。動的グローバルデバイスメモリ割り当てを行う SDK の例を知っている人がいたら、ぜひ知りたいです。

えー、まさに devPtr をグローバルスコープに移動するという問題が私の問題でした。

私はまさにそれを行う実装を持っており、2つのカーネルが渡されたデータへのポインターを持っています。私はそれらのポインターを明示的に渡したくありません。

私はドキュメントをかなり注意深く読み、nvidia フォーラムにアクセスしました (Google で 1 時間ほど検索しました) が、実際に実行できるグローバル動的デバイス配列の実装は見つかりませんでした (コンパイルして実行するものをいくつか試しました)その後、新しくて興味深い方法で失敗します)。

SDK に含まれるサンプルを確認してください。これらのサンプルプロジェクトの多くは、例から学ぶのに適した方法です。

このデータは多くの機能で使用されるため、グローバルにしたいと考えています。

グローバルを使用する正当な理由はほとんどありません。これは間違いなく一つではありません。この例を拡張して、「devptr」をグローバルな範囲に移動することを含む運動として残します。

カーネルが配列で構成される大きな const 構造体で動作する場合はどうなるでしょうか?いわゆる定数メモリの使用は、サイズが非常に限られているため、オプションではありません。そうすると、グローバルメモリに置く必要があります...?

ライセンス： CC-BY-SA と帰属

所属していません StackOverflow

CUDA グローバル (C と同様) デバイス メモリに割り当てられた動的配列

CUDA グローバル (C と同様) デバイスメモリに割り当てられた動的配列