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

Question

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

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

---

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

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

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

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

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

Answer 1

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

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

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

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

64 KBの定数制限は、CUDAコンパイル単位であるCUmoduleごとです。CUmoduleの概念は、CUDAランタイムの下に隠されていますが、CUDAドライバーAPIからアクセスできます。