Constant Memory

GPGPU CUDA

またまたCUDAの話です。

CUDAには

  1. Host Memory
    • Normal/Pinned/Mapped
  2. Device Memory
  3. Shared Memory
  4. Local Memory
  5. Texture Memory
  6. Constant Memory
  7. Registers

と、メモリ階層がたくさんあります。

このうちConstant MemoryはGPUからは読み出し専用ですがキャッシュがあるということで、定数やテーブルを格納するのに向いています。
というわけでConstant Memoryのメモ。

キャッシュ

8KBです。
とはいえ64KB中の8KBなので、キャッシュミス時のlatencyはほぼ無視できるかと。

アクセス速度

Constant Memoryの特徴の一つが、アクセスの速さです。CUDAのドキュメント(CUDA Programming Guide 2.3)曰く、「キャッシュ上にある限りはレジスタと同等」となっています。
しかしよくみると、これには一つの条件がついています。
「(half-warp:16 threads)が同じアドレスを読み出す限りは、レジスタからの読み込みと同等である」
そして、「読み込みコストは、half-warpのthreadsによって読み込まれるアドレスの数に線形に比例する」とも書かれています。
ようは、Constant Memoryを使うときはすべてのスレッド(少なくともhalf-warp内)で同じ場所を読めということです。
とはいえ、線形に比例するだけ(らしい)ので、キャッシュに当たる限りは最大でもレジスタより16倍遅いぐらいで済むということ・・・かな？

詳しくはCUDA Programming Guide 2.3の5.1.2.3を参照してください。

容量制限

Constant Memoryの泣き所の一つが、"64KBしか割り当てられない"という謎の仕様です。
おそらく何らかのハードウェアの制限だと思いますが、理由は不明です。

問題なのは、容量制限を回避したい時にどうすればいいか?というあたりです。

一個のkernelで同時に64KB以上を使いたいとき

無理です。Texture Memoryを使うか、Global Memoryを直接使ってください。

二個以上のkernelで合計64KB以上を使いたいとき

Constant MemoryはFile Scopeなので、別々のファイルの場合は使用できるようです。しかし残念なことに、一個のkernelでは64KB未満しか使用しなくても(参照しなくても)、同じファイル中の合計が64KBを超える場合はコンパイルできないようです。

一個のkernelで別々に64KB以上使いたいとき

一つのkernelで64KB以上使いたいが，一回のkernel呼び出しで同時に64KB以上使わない場合です．

方法1: kernel呼び出しごとにConstant Memoryをセットする

kernel呼び出す都度，Constant Memoryを書きかえる方法です．そのまんま．

方法2: kernelを複数(のファイル)に分割する

別のファイルにすることで，同じkernelを別々のCUDAのモジュールとして扱います．

この場合，プリプロセッサとテンプレートを駆使することで(比較的)すっきりしたコードになります．たとえば，

  • test1.h
class Test1;
class Test2;
template <class CLASS> class TTest1 {
	public:
		static TTest1* getInstance(){
			return instance;
		}
		float run();
		void setConst(float c);
		static TTest1 *instance;
};
  • test1.cu
#define TTest1_SPECIALIZE Test1
#include "test1.cu.inl"
  • test2.cu
#define TTest1_SPECIALIZE Test2
#include "test1.cu.inl"
  • test1.cu.inl
#include <cutil_inline.h>
#include "test1.h"
__constant__ float cmem1[1024];
template <>
TTest1<TTest1_SPECIALIZE>* TTest1<TTest1_SPECIALIZE>::instance = new TTest1<TTest1_SPECIALIZE>();
template <class CLASS>
__global__ void test1_run(float* a){
	float t = cmem1[0];
	a[0] = t;
}
template <>
float TTest1<TTest1_SPECIALIZE>::run(){
	float *a;
	cudaMalloc( (void**)&a, 1*sizeof(float));
	test1_run<TTest1_SPECIALIZE><<< 1, 1>>>(a);
	return *a;
}
template <>
void TTest1<TTest1_SPECIALIZE>::setConst(float c){
	cudaMemcpyToSymbol(cmem1, &c, 1*sizeof(float), 0, cudaMemcpyHostToDevice);
}

など．

まあ、これ全部、Fermiになったらがらりと変わっちゃう話なんですけどね。