CUDAで一定のメモリサイズが制限されるのはなぜですか？

Question

"CUDA C Programming Guide"によれば、マルチプロセッサ定数キャッシュ（セクション5.3.2.4）がヒットした場合のみ、定数メモリアクセスが有効です。さもなければ、合体された大域メモリ読み出しの場合よりも、半ワープに対するメモリ要求がさらに多くなる可能性がある。なぜ、一定のメモリサイズが64 KBに制限されているのですか？

さらにもう一度質問しないように質問します。私が理解する限り、Fermiアーキテクチャでは、テクスチャキャッシュはL2キャッシュと結合されています。テクスチャの使用はまだ有効か、グローバルメモリの読み込みは同じ方法でキャッシュされますか？

---

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

が一定のメモリ空間はデバイスメモリ内に常駐し、セクションF.3.1およびFに記載された一定のキャッシュ内にキャッシュされています4.1。

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

要求は、最初の要求に異なるメモリアドレスが存在する場合と同じように多くの個別の要求に分割され、個別の要求の数に等しい要因でスループットが低下します。

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

Answer 1

コンピューティング能力1.0-3.0デバイスの定数メモリサイズは64 KBです。キャッシュワーキングセットはわずか8KBです（CUDAプログラミングガイドv4.2の表F-2を参照）。

定数メモリは、ドライバ、コンパイラ、および変数__device__ __constant__で宣言された変数によって使用されます。ドライバは、定数メモリを使用してパラメータ、テクスチャバインディングなどを伝達します。コンパイラは多くの命令で定数を使用します（逆アセンブリを参照）。

定数メモリに配置された変数は、ホストランタイム関数cudaMemcpyToSymbol()およびcudaMemcpyFromSymbol()（CUDAプログラミングガイドv4.2のセクションB.2.2を参照）を使用して読み書きできます。一定のメモリはデバイスメモリにありますが、定数キャッシュを通してアクセスされます。

フェルミのテクスチャでは、定数L1およびI-Cacheは、すべてのSM内またはその周りのすべてのレベル1キャッシュです。すべてのレベル1キャッシュは、L2キャッシュを介してデバイスメモリにアクセスします。

64 KBの制限は、CUDAコンパイル単位であるCUmoduleあたりです。 CUmoduleの概念はCUDAランタイムの下に隠されていますが、CUDAドライバーAPIによってアクセス可能です。