テクスチャメモリ

108ページの tex1Dfetch() の説明がちょっと分かりにくい。次のように解釈した：

tex1Dfetch() は関数のように見えるが、実は関数ではなく、コンパイラに組み込まれている機能である。その実装では、tex1Dfetch() でアクセスするテクスチャが（リンク時や実行時でなく）コンパイル時に判明している―テクスチャ参照の宣言がコンパイラに見えている―必要がある。このため、tex1Dfetch() を用いてテクスチャにアクセスする blend_kernel() に例えば引数としてテクスチャ参照へのポインタを渡すことで入力バッファと出力バッファを切り替えることはできない。そこで今回のコードでは、テクスチャ参照2つをファイルスコープで（＝関数の外で）宣言し、どちらのバッファを使うかをフラグ dstOut で指定することにした。

ただ、テクスチャ参照は左辺値のようだし（sizeof できるしポインタも取れる）、tex1Dfetch() も関数のようだ（ポインタを取ろうとすると overloaded function と言われる）。本の nvcc と自分が使っている nvcc で版が違うのかも知れない。

$ nvcc -V
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2010 NVIDIA Corporation
Built on Thu_Nov_11_15:26:50_PST_2010
Cuda compilation tools, release 3.2, V0.2.1221

anim_gpu() 内で dstOut を volatile としてる理由は謎。

あと、試しにテクスチャ参照をブロック内で宣言してみたら、どうも勝手に static 扱いになってるみたいだった。

ドット積（内積）を求める

CUDA BY EXAMPLE 5.3.1節。 N次元ベクトル2つの内積を、256スレッドからなる複数のブロックを使って求める。 ブロック数は32個とするが、Nが小さい場合には必要な数だけのブロックを起動する。 各ブロックは、担当の部分の要素の積の和を求めて結果をCPUに返す。 よって32個以下の値がCPUに返ってくる。この総和（が内積である）をCPUは求めて表示する。

各スレッドの仕事は以下のようになっている。

• Nがちょうど256*32である場合、256*32個のスレッドそれぞれは1つのベクトル要素対を担当し、その積を求める。
• N=256*32*2である場合、ひとつのスレッドは、あるベクトル要素対と、そこから256*32個後ろのベクトル要素対の2つを担当し、それぞれの積の総和を求める。一般にN=256*32*Mの場合、各スレッドはM個のベクトル要素対の積の総和を求める。
• Nが256*32よりも小さい場合、スレッド数が足りる（つまりN以上となる）最小の数だけのブロックを起動する。

Nがぴったり256の倍数でない場合の処理は常套句の while (tid < N) で行われている。

各スレッドが自分の担当するところの積の和を求めたら、ブロック内で縮約して総和を求め、CPUに返す（カーネルの仮引数c、mainから見ると(dev_)partial_c）。CPUは返された32個（あるいはそれ以下）の値の総和を求める。

68ページにある最小倍数回りの分かりにくい部分は、こう読むといい：「データ要素がN個なので、ドット積を計算するのに必要なスレッドはN個である。スレッド数が256個*32ブロックでは多すぎる場合には、threadsPerBlock (=256) の倍数であってかつN以上であるような最も小さい個数のブロックを用いればよい。これはベクトルの和でも取り上げた。今回の場合は、(N+threadsPerBlock-1))/threadsPerBlock を計算することで、そのようなブロック数を求める。」

ビジー状態を保つはよく分からない。

グリッド、ブロック、スレッド

「CUDA BY EXAMPLE」の5.2.1節にある「長いベクトルの和」の項が分かりにくかったので覚えとして。

51ページ： この変数はすべてのブロックに対する定数であり、ブロックの各次元のスレッド数を含んでいます。 原文にあたってないけど、この部分は意味は多分： 「この変数は全てのブロックについて同じ（一定＝constant）であり、ブロックを構成する3次元に並んだスレッドたちの各次元のスレッド数をメンバとして持っています。」

まとめて書かれていないようだけど、どうやらグリッドとブロックとスレッドの関係はこんな感じらしい：

• グリッドは、2次元にならんだブロックたちである。変数 gridDim（41-42ページ）は、その各次元（x方向およびy方向）のブロック数を持っている。gridDim.x がx方向のブロック数、gridDim.y がy方向のブロック数である。
• ブロックは、3次元にならんだスレッドたちである。変数 blockDim は、その各次元（x, y, z方向）のスレッド数を持っている。blockDim.x がx方向のスレッド数、などとなる。

ついでに小さいところ。 50ページ：複数のブロックとスレッドを使う場合、インデックスを計算する方法は、2次元のインデックス空間を線形空間に変換するための標準的な手法に似てきます。 「線形空間」は「1次元 (linear) 空間」と読んだほうが分かりやすい。

GLUT を使う設定

CUDA BY EXAMPLE 4.2.2節でジュリア集合を計算する例が出てくる。GLUT を入れてなかったので少し手間だった。以下、作業内容。

1. MacPorts で mesa を入れる：sudo port install mesa。これで GLUT が入る（らしい）。
2. MacPorts はどうやってか知らないけど適切にライブラリ経路を設定してくれるようで、/opt/local/lib に入ったライブラリがちゃんと使える。よって .bashrc での DYLD_LIBRARY_PATH の設定は元のまま：export DYLD_LIBRARY_PATH=/usr/local/cuda/lib。

GLUT のヘッダファイルとライブラリが nvcc から見えるように、コンパイルはこう↓する。

nvcc -I/opt/local/include -m 64 -o julia_cpu julia_cpu.cu -L/opt/local/lib -lglut -lGL

MacPorts で入れた libglut などが64ビットだったので -m 64 が必要になった。

これで動いた。終了するには画像ウィンドウの方で Esc を押す。

julia_gpu の方も同じやり方で動いた。4章終わり。

追記：Mac OS X には標準で OpenGL や GLUT が入っているようだけど、それを使うような設定が分からなかったので結局 mesa にした。うまくやればできそう。参考：GLUTによる「手抜き」OpenGL入門

CUDA BY EXAMPLE も始めてみる

和訳の方。1章2章はさくさく読み進めて3章。

../common/book.h は SDK ではなくこの本のサイトで提供されているファイル（.zip に入ってる）だった。.zip ファイルを展開して chapter03/ の下に行き、nvcc hello_world.cu したら a.out ができて、ちゃんと動いた。

楽するために簡単な Makefile を作った。

.SUFFIXES: .cu
.cu:
＜タブ＞nvcc -o $@ $<

これで make hello_world とかするとコンパイルできる。

Mac で CUDA を始めてみる

いつまで続けられるか分からないけど、Mac で CUDA してみようと思い立った。

マシンは最近買った MacBook Air 13インチ、2.13GHz Core 2 Duo、メモリ4GB。ディスプレイのチップセットは NVIDIA GeForce 320M、VRAM 256MB。Xcode は3.2.5がすでに入れてある。

とりあえず NVIDIA の Developer Zone から CUDA Toolkit 3.2 の Mac 用をもらってきて、Mac Getting Started Guide の通りにインストールとコンパイルをしてみたら、deviceQuery はちゃんと動いたようだ。

bandwidthTest の結果はこんな感じ：

[bandwidthTest]
./bandwidthTest Starting...

Running on...

 Device 0: GeForce 320M
 Quick Mode

 Host to Device Bandwidth, 1 Device(s), Paged memory
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     1261.8

 Device to Host Bandwidth, 1 Device(s), Paged memory
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     1181.8

 Device to Device Bandwidth, 1 Device(s)
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     9999.2


[bandwidthTest] - Test results:
PASSED

まずはここまで。

rd2 が動かなくなった

ひさしぶりに RT を使った RD ファイルに rd2 (rdrt2) をかけたらこんなエラーメッセージが出て失敗した。

$ rdrt2 -r rd/rd2html-ext-lib --out-code=utf8 --native-inline --ref-extension aaa.rd > aaa.html
/opt/local/lib/ruby/site_ruby/1.8/rd/search-file.rb:7:in `+': can't convert nil into String (TypeError)
        from /opt/local/lib/ruby/site_ruby/1.8/rd/search-file.rb:7:in `search_file'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/search-file.rb:6:in `each'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/search-file.rb:6:in `search_file'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/search-file.rb:5:in `each'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/search-file.rb:5:in `search_file'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/output-format-visitor.rb:23:in `apply_to_Include'
        from (eval):2:in `visit_Include'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/block-element.rb:65:in `accept'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/visitor.rb:20:in `visit_children'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/element.rb:45:in `each_child'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/element.rb:44:in `each'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/element.rb:44:in `each_child'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/visitor.rb:19:in `visit_children'
        from (eval):2:in `visit_DocumentElement'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/element.rb:144:in `accept'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/tree.rb:78:in `accept'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/visitor.rb:14:in `visit'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/rd2html-lib.rb:62:in `visit'
        from /opt/local/lib/ruby/site_ruby/1.8/rd/rd2html-ext-lib.rb:42:in `visit'
        from /opt/local/bin/rd2:241
$ 

調べてみたら、どうも rd2 が optparse の古いあるいは隠し機能を使ってるみたいだった。こんな風に変えてみた。

--- rd2.orig    2008-01-14 16:40:04.000000000 +0900
+++ rd2 2010-06-23 18:25:52.000000000 +0900
@@ -127,8 +127,14 @@
   
   # accept "PART:FILTER" and "PART"
   q.on("--with-part=PART",
-       /(\w+)(?:\s*:\s*(\w+))?/,
-       "include PART with Filter") do |src, part, filter|
+       String,
+       "include PART with Filter") do |i|
+    if /(\w+)(?:\s*:\s*(\w+))?/ =~ i
+      part   = $1
+      filter = $2
+    else
+      raise OptionParser::InvalidArgument, "--with-part=#{i}"
+    end
     with_part.push([part, filter || part])
     unless include_path.index(RD::RDTree.tmp_dir)
       include_path.push(RD::RDTree.tmp_dir)

optparse も Ruby もあまり詳しくないからもしかしたら適切な直し方じゃないかも知れないけど、一応動いたみたい。

ん、Firefox で見ると右の方が切れて見えるなぁ。デザイン変えた方がいいのかな。コピペするのは大丈夫みたいだけど。