問題タブ [gpu-constant-memory]

シミュレーション パラメータを定数メモリに設定しようとしていますが、うまくいきません (CUDA.NET)。cudaMemcpyToSymbol 関数は cudaErrorInvalidSymbol を返します。cudaMemcpyToSymbol の最初のパラメーターは文字列です... シンボル名ですか? 実際、どのように解決できるかわかりません。どんな助けでも感謝します。

私の.cuファイルが含まれています

ワーキングソリューション

カーネルで使用されるメモリが一定であることを保証する最良の方法は何ですか。http://stackoverflow...r-pleasant-wayにも同様の質問があります。私は GTX580 を使用しており、2.0 機能のみをコンパイルしています。私のカーネルは次のようになります

ホストで次のコードを実行します。

別の方法は、追加することです

.cu ファイルに、カーネルからsrcポインターを削除して実行します

これら2つの方法は同等ですか、それとも最初の方法はグローバルメモリの代わりに定数メモリの使用を保証しませんか? サイズは動的に変化するため、私の場合、2 番目の方法は便利ではありません。

テキストファイルを読み込んで配列に保存したい。次に、配列をホストからデバイスに転送し、共有メモリに格納します。以下のコードを書きましたが、グローバルメモリを使用した場合に比べて実行時間が長くなりました。理由がわかりませんか？また、誰かが定数メモリを使用してこのコードを書くのを手伝ってくれると助かります。

「CUDACプログラミングガイド」によると、コンスタントメモリアクセスは、マルチプロセッサコンスタントキャッシュがヒットした場合にのみメリットがあります（セクション5.3.2.4）¹。そうしないと、合体したグローバルメモリ読み取りの場合よりも、ハーフワープに対してさらに多くのメモリ要求が発生する可能性があります。では、なぜ一定のメモリサイズが64KBに制限されているのでしょうか。

二度と聞かないためにもう一つ質問。私が理解している限り、Fermiアーキテクチャでは、テクスチャキャッシュはL2キャッシュと組み合わされています。テクスチャの使用はまだ意味がありますか、それともグローバルメモリの読み取りは同じ方法でキャッシュされますか？

¹定数メモリ（セクション5.3.2.4）

コンスタントメモリスペースはデバイスメモリに常駐し、セクションF.3.1およびF.4.1で説明したコンスタントキャッシュにキャッシュされます。

コンピューティング機能1.xのデバイスの場合、ワープに対する一定のメモリ要求は、最初に2つの要求に分割されます。1つはハーフワープごとに1つで、独立して発行されます。

次に、リクエストは最初のリクエストにある異なるメモリアドレスと同じ数の個別のリクエストに分割され、個別のリクエストの数に等しい係数でスループットが低下します。

結果として得られる要求は、キャッシュヒットの場合は一定のキャッシュのスループットで、それ以外の場合はデバイスメモリのスループットで処理されます。

ローカルメモリについて言及しているCUDAのドキュメントを読みました。（これは主に初期のドキュメントです。）device-propertiesは、local-memサイズ（スレッドごと）を報告します。「ローカル」メモリとはどういう意味ですか？「ローカル」メモリとは何ですか？「ローカル」メモリはどこにありますか？「ローカル」メモリにアクセスするにはどうすればよいですか？__device__記憶ですよね？

device-propertiesは、グローバル、共有、および一定のmemサイズも報告します。これらのステートメントは正しいですか： グローバルメモリは__device__メモリです。グリッドスコープとグリッド（カーネル）の存続期間があります。 コンスタントメモリは__device__ __constant__メモリです。グリッドスコープとグリッド（カーネル）の有効期間があります。 共有メモリは__device__ __shared__メモリです。単一のブロックスコープと（スレッドの）そのブロックの存続期間があります。

共有メモリはSMメモリだと思います。つまり、その単一のSMのみが直接アクセスできるメモリ。かなり限られたリソース。SMは一度にたくさんのブロックを割り当てられていませんか？これは、SMが異なるブロックの実行をインターリーブできる（またはインターリーブできない）ことを意味しますか？つまり、ブロック* A *スレッドをストールするまで実行します。次に、ブロック* B *スレッドを停止するまで実行します。次に、再びブロック* A *スレッドにスワップバックします。またはSMは、ブロック* A *のスレッドのセットを、ストールするまで実行しますか。次に、別のブロック* A *スレッドのセットがスワップインされます。このスワップは、ブロック* A *が使い果たされるまで続きます。それからそしてその時だけ、作業はブロック* Bで始まります*。共有メモリのためにお願いします。1つのSMが2つの異なるブロックからコードをスワップインしている場合、SMはどのようにして共有メモリチャンクをすばやくスワップイン/スワップアウトしますか？（後のsenerioは真であり、共有メモリスペースのスワップイン/スワップアウトはないと思います。Block* A *は完了するまで実行され、次にblock * B *が実行を開始します。注：block * A *は異なる場合があります。ブロックよりカーネル* B *。）

xptxasを使用してレジスタの使用状況を確認すると、次のようになります。

現在、cmem[x]を明確に説明しているドキュメントはあるのでしょうか。コンスタントメモリを複数のバンクに分割することのポイントは何ですか、合計でいくつのバンクがありますか、そして0、2、14、16以外の他のバンクは何に使用されますか？

補足として、@ njuffa（特にあなたに感謝します）は以前にnvidiaのフォーラムで銀行0,2,14,16とは何かを説明しました：

使用される定数メモリは、定数プログラムの「変数」（バンク1）と、コンパイラによって生成された定数（バンク14）に分割されます。

cmem [0]：カーネル引数

cmem [2]：ユーザー定義の定数オブジェクト

cmem [16]：コンパイラによって生成された定数（ソースコードのリテラル定数に対応するものもあります）

手書きのカーネルの各 CUDA スレッドのリソース使用量を理解しようとしています。

kernel.cuファイルをファイルkernel.oにコンパイルしましたnvcc -arch=sm_20 -ptxas-options=-v

そして、次の出力を得ました（パススルーc++filt）：

上記の出力を見ると、次のように言うのは正しいですか?

• 各 CUDA スレッドは 46 個のレジスタを使用していますか?
• ローカルメモリに流出するレジスタはありませんか?

出力の理解にも問題があります。

• 私のカーネルはたくさんの__device__関数を呼び出しています。__global__関数と__device__関数のスタック フレームのメモリの合計は 72 バイトですか?

• 0 byte spill storesとはどう違いますか0 bytes spill loads

• の情報cmem(これは定数メモリであると想定しています) が異なる数値で 2 回繰り返されるのはなぜですか? カーネル内では、定数メモリを使用していません。それは、コンパイラが内部で GPU に一定のメモリを使用するように指示するということですか?

_{この質問は次のセクションに「続きます」: ptxas の詳細出力の解釈、パート II}

CUDA 定数メモリ領域を 2 次元で割り当ててコピーする方法はありますか? cudaMemcpyToSymnbolof __constant__ はオプションではないようです。

定数メモリ (グローバル変数) に配列があり、関数呼び出し cudaGetSymbolAddress によってその配列への参照を取得しました。グローバル変数を使用するのではなく、この参照を使用して定数データをフェッチすると、カーネルの実行が遅くなります。これの理由は何ですか？

これはサンプル コードで、ワープ内のすべてのスレッドが同じ場所を同時にロードします。コメント化されたコードは、定数メモリに直接アクセスすることによるものです

定数メモリ キャッシュを使用しない理由の説明（talonmiesによる）

その理由は、定数キャッシュの欠如です。キャッシュされたアクセスは、定数状態空間にあると明示的にマークされた変数に対してコンパイラが特定の PTX 命令 (ld.const) を発行した場合にのみ発生します。そして、コンパイラがこれを行うことを認識する方法は、変数が宣言されたときです。これは、__constant__コード生成に影響を与える静的なコンパイル時の属性です。実行時に同じプロセスが発生することはありません。

グローバル メモリにポインターを渡し、そのポインターが定数状態空間にあるとコンパイラが判断できない場合、定数キャッシュを介してそのメモリにアクセスするための正しい PTX が生成されません。その結果、アクセスが遅くなります。

未回答の問題

配列gが変数として宣言されている場合でも__device__、それを参照するとコードが遅くなるのはなぜですか。PTXグローバルメモリをレジスタにロードするためのコードを見ると、次のようになります。

• ld.global.s32レジスタに 4 バイトをロードする2 つの命令が使用されます。(参照を使用するコード内)
• ld.global.v2.s328 バイトを 2 つのレジスタにロードする1 つの命令が使用されます (グローバル変数を使用するコード内)。

違いは何ですか。ドキュメントの参照をいただければ幸いです。

問題

コンスタントメモリを使用してサンプルCUDAコードを1つ用意しました。これはcuda4.2で正常に実行できますが 、CUDA 5を使用してコンパイルすると、 「無効なデバイスシンボル」が表示されます。ここにサンプルコードを添付しました。

コード

WINDOWS：CUDA 5.0 Production Releaseを試しましたが、グラフィックカードはGTX590
です。ご協力いただければ幸い です。