7

そのため、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 呼び出しは必要ありませんが、私は動的割り当てアプローチを好みます。

4

6 に答える 6

5

このようなものはおそらくうまくいくはずです。

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

やってみて。

于 2008-09-17T03:52:41.950 に答える
1

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

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

ただし、グローバル データにアクセスしようとするたびに「アドバイザリ: グローバル メモリ空間を想定して、ポインタが何を指しているのかわかりません」というメッセージが表示されるようになったため、コンパイラが混乱していると思います。幸いなことに、仮定はたまたま正しいのですが、警告が煩わしいです。

とにかく、記録のために-私は多くの例を見て、ポイントが「正しい！」と言う出力を得ることであるnvidia演習を実行しました。とはいえ、私はそれらすべてを見たわけではありません。動的なグローバル デバイス メモリ割り当てを行う sdk の例を誰かが知っている場合でも、知りたいです。

于 2008-09-18T16:55:53.043 に答える
1

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++' を忘れないでください。

于 2008-09-17T02:24:59.977 に答える
0

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

于 2008-09-18T02:27:59.993 に答える
0

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

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

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

于 2008-09-17T02:35:32.073 に答える
0

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

-

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

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

于 2010-08-18T11:21:28.347 に答える