2010-04-27

CUDAコンパイラのバグその2

GPGPU CUDA

SDKをCUDA 3.0にしたら動かなくなったのでなんでかなーと思った結果。

C++のnamespace中でtextureを宣言すると、カーネル起動やメモリ操作関連APIで"unspecified driver error"が発生する模様。

Ref. http://forums.nvidia.com/index.php?showtopic=163545&pid=1042787&mode=threaded&show=&st=0

2010-02-12

Instruction throughputの意味

GPGPU CUDA

まだまだ続くCUDAねた。

CUDAにはCUDA Visual Profilerという便利なものがありまして、メモリの使用帯域とか各カーネルごとの実行時間をお手軽に測定することができます。
で、こいつで測定できる項目の一つにinstruction throughputというのがあります。
直感的にはいわゆる実行効率なのかなーと思うわけですが，1.0以上あるとか*1、ただの比率だとか*2いろいろ書かれていてよくわからんわけです。
というわけで簡単に調べたメモ。

そして身も蓋もない話で申し訳ないのですが、だいたいのことはマニュアル(同梱のcudaprof.html)に書いてありました。
それによりますと

instruction throughput: Instruction throughput ratio.
(略) calculated using the "instructions" profiler counter.
(略) calculated based on the GPU clock speed.

というように、instructionsカウンタの値から計算されるわけですね。
ただしinstructionsカウンタの解釈が問題で、
「Profiler counters for a single multiprocessor」の項目にあるように

These counters are incremented by one per each warp.

となっています。つまりthread単位ではなくwarp単位で積算されているため、全体の実行命令数を求めるにはこの値を32倍する必要があります。
そしてもう一つ、

the profiler can only target one of the multiprocessors in the GPU

と書かれているように、このプロファイラはどれか一つのSMを測定対象としています。
そのため、この実行命令数は測定対象のSMに関する統計になっているとみられます。

さて、全体の実行命令数がわかれば後は簡単です。
Instruction throughputはratioとなっていますから、最大実行可能命令数 = GPU(SM)のクロック * 実行時間(GPU Timeカウンタ) * 8 (SP)として最大実行可能命令数を求め、実行効率 = 全体の実行命令数 / 最大実行可能命令数とすればInstruction throughputの値が再現できます。

例えばGPU Time = 19684.3 [usec]、instructions = 4813980、クロック = 1.35 GHzでは

全体の実行命令数 = 4813980*32 = 154047360
最大実行可能命令数 = 1.35*10^9*19684.3*10^-6*8 = 212590440

となるので、

Instruction throughput = 154047360 / 212590440 = 0.724620355

となります。これは測定値 Instruction throughput = 0.724618とほぼ一致します。

・・・こんなんでいいのかな？

*1:http://forums.nvidia.com/index.php?showtopic=96866

*2:http://www.prace-project.eu/hpc-training/training_pdfs/2653.pdf

2010-01-13

いろんなGeForceのdeviceQuery結果

GPGPU CUDA

手元にあるGeForceのdeviceQuery結果のメモ。

NVIDIA GeForce GT 240

ボードはhttp://www.palit.biz/main/vgapro.php?id=1284

Device 0: "GeForce GT 240"
  CUDA Driver Version:                           3.0
  CUDA Runtime Version:                          2.30
  CUDA Capability Major revision number:         1
  CUDA Capability Minor revision number:         2
  Total amount of global memory:                 536543232 bytes
  Number of multiprocessors:                     12
  Number of cores:                               96
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       16384 bytes
  Total number of registers available per block: 16384
  Warp size:                                     32
  Maximum number of threads per block:           512
  Maximum sizes of each dimension of a block:    512 x 512 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
  Maximum memory pitch:                          262144 bytes
  Texture alignment:                             256 bytes
  Clock rate:                                    1.40 GHz
  Concurrent copy and execution:                 Yes
  Run time limit on kernels:                     Yes
  Integrated:                                    No
  Support host page-locked memory mapping:       Yes
  Compute mode:                                  Default (multiple host threads
can use this device simultaneously)

NVIDIA GeForce GTX 260

ボードはhttp://www.galaxytech.com/japan/Product_Details.asp?id=172

Device 0: "GeForce GTX 260"
  CUDA Driver Version:                           2.30
  CUDA Runtime Version:                          2.30
  CUDA Capability Major revision number:         1
  CUDA Capability Minor revision number:         3
  Total amount of global memory:                 939196416 bytes
  Number of multiprocessors:                     27
  Number of cores:                               216
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       16384 bytes
  Total number of registers available per block: 16384
  Warp size:                                     32
  Maximum number of threads per block:           512
  Maximum sizes of each dimension of a block:    512 x 512 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
  Maximum memory pitch:                          262144 bytes
  Texture alignment:                             256 bytes
  Clock rate:                                    1.35 GHz
  Concurrent copy and execution:                 Yes
  Run time limit on kernels:                     Yes
  Integrated:                                    No
  Support host page-locked memory mapping:       Yes
  Compute mode:                                  Default (multiple host threads
can use this device simultaneously)

NVIDIA GeForce 9600 GT

ボードはhttp://www.elsa-jp.co.jp/products/graphicsboard/gladiac_796_gt_sp/index.html

Device 0: "GeForce 9600 GT"
  CUDA Driver Version:                           3.0
  CUDA Runtime Version:                          2.30
  CUDA Capability Major revision number:         1
  CUDA Capability Minor revision number:         1
  Total amount of global memory:                 536543232 bytes
  Number of multiprocessors:                     8
  Number of cores:                               64
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       16384 bytes
  Total number of registers available per block: 8192
  Warp size:                                     32
  Maximum number of threads per block:           512
  Maximum sizes of each dimension of a block:    512 x 512 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
  Maximum memory pitch:                          262144 bytes
  Texture alignment:                             256 bytes
  Clock rate:                                    1.50 GHz
  Concurrent copy and execution:                 Yes
  Run time limit on kernels:                     Yes
  Integrated:                                    No
  Support host page-locked memory mapping:       No
  Compute mode:                                  Default (multiple host threads
can use this device simultaneously)

NVIDIA GeForce 9200M GS

Device 0: "GeForce 9200M GS"
  CUDA Driver Version:                           3.0
  CUDA Runtime Version:                          2.30
  CUDA Capability Major revision number:         1
  CUDA Capability Minor revision number:         1
  Total amount of global memory:                 536870912 bytes
  Number of multiprocessors:                     1
  Number of cores:                               8
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       16384 bytes
  Total number of registers available per block: 8192
  Warp size:                                     32
  Maximum number of threads per block:           512
  Maximum sizes of each dimension of a block:    512 x 512 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
  Maximum memory pitch:                          262144 bytes
  Texture alignment:                             256 bytes
  Clock rate:                                    1.30 GHz
  Concurrent copy and execution:                 No
  Run time limit on kernels:                     No
  Integrated:                                    No
  Support host page-locked memory mapping:       No
  Compute mode:                                  Default (multiple host threads
can use this device simultaneously)

NVIDIA GeForce 8600M GT

2010-04-27 追記

Device 0: "GeForce 8600M GT"
  CUDA Driver Version:                           3.0
  CUDA Runtime Version:                          3.0
  CUDA Capability Major revision number:         1
  CUDA Capability Minor revision number:         1
  Total amount of global memory:                 255393792 bytes
  Number of multiprocessors:                     4
  Number of cores:                               32
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       16384 bytes
  Total number of registers available per block: 8192
  Warp size:                                     32
  Maximum number of threads per block:           512
  Maximum sizes of each dimension of a block:    512 x 512 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             256 bytes
  Clock rate:                                    0.75 GHz
  Concurrent copy and execution:                 Yes
  Run time limit on kernels:                     No
  Integrated:                                    No
  Support host page-locked memory mapping:       No
  Compute mode:                                  Default (multiple host threads
can use this device simultaneously)

2010-01-02

Texture Memoryとテンプレート

GPGPU CUDA

あけましておめでとうございます。

またまたCUDAネタです。

Texture Memoryは型変換や補間などけっこう便利なのですが、同一ファイル内でグローバル変数として宣言しなければならないなど、使用が非常に面倒です。
そして、どうにかして楽に使う方法はないかと探しているときに遭遇した不思議な現象がこのお話です。

初めはテクスチャの宣言を固定長配列としておこない、後から参照する方法を考えていましたが早々に断念。テクスチャを使用した際のptxコードを見ればわかるとおり、どのテクスチャを使用するかはコンパイル段階で決定されていないといけないようです。

次に考えたのは、型テンプレートを使用する方法。texture宣言をstaticメンバで含む型をテンプレート化するもので、

test.cu

...
template<int n>
struct Tex{
　static texture<float4, 2, cudaReadModeElementType> tex;
};
texture<float4, 2, cudaReadModeElementType> Tex<0>::tex;
...
texture<float4, 2, cudaReadModeElementType> Tex<10>::tex;
...

のような.cuコードを使い、使用する場所でTex<0-10>::texと指定します。

ところがこれをnvccにかけると、次のようなエラーが。

error C2720: 'Tex<n>::tex' : 'static ' ストレージ クラスの指定子が識別子に対して誤って指定されています。

このエラー、発生場所をみるとptxの作成やGPU側オブジェクトコードの作成は問題なく済んだあと、HOST側コードのコンパイル段階で発生している。
MSDNでみると、Compiler Error C2720 | Microsoft Docsとなっている。元の.cuではstaticはつけていないのに、はて？と思ってnvccによって生成された.cu.cppをみてびっくり。

test.cu.cpp

...
template<int n> 
struct Tex { 
static   texture< float4, 2, cudaReadModeElementType>  tex; 
}; 
static   texture< float4, 2, cudaReadModeElementType>  Tex< 0> ::tex; 
...

なんかついてる・・・！

ついでに微妙に変なスペースもはいってる。

ここから先は想像ですが、nvccが.cu.cppを生成する際にtexture宣言された変数に手当たり次第static修飾子をつけているのではないだろうか？
通常のグローバル変数の場合はそれで問題ないのだが、staticメンバとして宣言された場合は、static修飾子は型のほうにすでにあるため文法エラーになる。しかしnvccがそれを認識せずにstatic修飾子をつけるため、エラーになっている。
そしてこの.cu.cppを手動で編集しstatic修飾子を取り除くと、問題なくコンパイルされました。ようはnvccの手抜き実装が原因？

ちなみにテクスチャにアクセスする時に使うtex1Dなどの関数はtexture_fetch_functions.hで定義されているので、Texture Memoryを使う場合は確認しておいたほうがいいです。

2010-01-04 追記
staticメンバのテンプレート特殊化の宣言を.cuではなく別の.cppファイルに移すことで、特殊化の部分のみnvccを迂回してコンパイルすればとりあえず回避できるようです。コンパイラはCUDA 2.3 win32 xp, VC++ 2008 Expressです。

test.cu

...
template<int n>
struct Tex{
　static texture<float4, 2, cudaReadModeElementType> tex;
};
...

test.cpp

texture<float4, 2, cudaReadModeElementType> Tex<0>::tex;
...
texture<float4, 2, cudaReadModeElementType> Tex<10>::tex;

2010-01-04 追記その2
Texクラスのテンプレート引数を増やすと、nvcc内の__text_varマクロの展開で落ちる模様。がんばれば処理できそうだけど、大人しくテンプレートを使わず普通にテクスチャ宣言したほうがよさそう。

2009-12-22

Constant Memory

GPGPU CUDA

またまたCUDAの話です。

CUDAには

Host Memory
- Normal/Pinned/Mapped
Device Memory
Shared Memory
Local Memory
Texture Memory
Constant Memory
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になったらがらりと変わっちゃう話なんですけどね。

2009-12-21

CUDAコンパイラの最適化バグ?

GPGPU CUDA

狙ってやっているわけではないのですが、四半期ぶりの日記です。

ここ半年ばかりGPGPUを使った研究に取り組んでいて、なかなかすばらしい性能を出してくれています。
ところが、ある日コードを書いていると突然コンパイラが落ちるようになりました。

2>Compiling with CUDA Build Rule...
2>"C:\CUDA\bin\nvcc.exe" ...
...
2>nvopencc ERROR: C:\CUDA\bin/../open64/lib//be.exe returned non-zero status -1073741819

その時は適当にコードを変えていたら出なくなったのですが、その後再発し開発が進まなくなってしまいました。
というわけで簡単に調べたメモ。

まず計算自体はShared Memory上のデータとConstant Memory上のフィルタの畳み込み積分を計算するもので、コードにするとこんな感じになります(実際はもう少し複雑です)。

float out[2][2];
out[0][0] = out[0][1] = ... = 0;
#pragma unroll
  for(uint x = 0; x < SIZE; x++){
    #pragma unroll
    for(uint p = 0; p < 2; p++){
      out[v][p] += smem[threadIdx.x+x]*cmem[v][x][p];
    }
  }
}

問題が起きたのは、ここでoutの要素数やSIZEを増やした時。
CUDAではローカル変数はLocal Memoryというところに保存されるのですが、レジスタが空いている時はコンパイラが変数をレジスタに割り当ててくれます。
ただし配列の場合は、ループアンロールを行うなどしてアクセスする要素をコンパイル時に決定できるようにしないとレジスタ割当が行われません(ただの経験則です、念のため)。

そして問題のコンパイルエラーは配列の要素数を増やしたときもしくはSIZEを大きくしたときに発生しました。詳しくは解析していませんが、どうやら、配列への命令がある一定数(数千〜?)に達するとコンパイルできなくなるようです。

試行錯誤した結果、配列を二つに分割することで解決しました。先ほどの例でいえば

float out0[2], out1[2];
out0[0] = out1[0] = ... = 0;
#pragma unroll
  for(uint x = 0; x < SIZE; x++){
    out0[v] += smem[threadIdx.x+x]*cmem[v][x];
    out1[v] += smem[threadIdx.x+x]*cmem[v][x];
  }
}

と、一部の配列を手動でアンロールして別々の配列に割り当てることで解決。

おそらく配列をレジスタに割り当てるアルゴリズムがオーバーフローを起こしているのだと思いますが、よくわかりません。

2009-09-09

Logicool Qcam Pro 9000シリーズの光学系

webcam

最近また画像処理をやっているのでメモ。
Logicoolさんがデータを出してくれないので、ぐぐった結果のまとめです。
かなり適当な調べ方なので間違ってるかも。

Spec

CMOSイメージセンサ素子	非公開
素子解像度	1600x1200
素子サイズ	6mm x 4mm*1
ピッチ	2.8um (たぶん)
Focus	3.7mm
F値	2.0
対角画角	75度(公称)
水平画角	60度(推定)
垂直画角	45度(推定)
水平画角	102度(測定値1*2 )
水平画角	84度(測定値1)
垂直画角	62度(測定値2*3 )

~~体感では水平60度ってことはない気がしますが・・・自分が測ったほうがいいかな？~~
測ってみた。カメラから1m離れた地点で、だいたい横1.2m、縦0.86mの物体が撮影できます。
なので、水平画角は60度〜62度、垂直画角は45度〜47度程度のようです。

Note

製品名は同じで内部リビジョンがけっこう変わっているようで、CMOSセンサもバージョンによって違うかも。

References

*1:http://forums.quickcamteam.net/showthread.php?tid=265

*2:http://forums.quickcamteam.net/showthread.php?tid=107

*3:価格.com - ロジクール Qcam Communicate Deluxe QCAM-130X のクチコミ掲示板 (2ページ目)