GPGPU
重力を計算するところ、 float dx = diff_x(sx[i], x); float dy = diff_x(sy[i], y); float dist_square = dx * dx + dy * dy; float d = dist_square * sqrt(dist_square); ax += dx / d; ay += dy / d;ここで、sqrtを使っているが、実は、 sqrt(x) = x / …
前回遅かったのは、グローバルな配列に何度もアクセスしていたためと思われる。共有メモリというものがあって、こちらは速いので、これを使うことを考える。 説明しやすくするために、4096個の質点に0〜4095のIDを振っておこう。共有メモリは、ブロックの中…
ネタがないので、浮動小数点数の問題をみてみる。 n個の質点が互いに重力による相互作用を及ぼしあう問題を考える。なるべく多くの質点を用いたいが、ここでは4096個の質点の計算をする。また、これらはトーラス面上にあるとする。1×1の正方形を用意し、上下…
前にも取り上げたが、コラッツの問題というのは、どの自然数も、奇数なら3倍して1を足す、偶数なら2で割る、という操作を繰り返すといつか1になるか、というものである。例えば、 7→22→11→34→17→52→26→13→40→20→10→5→16→8→4→2→1 となる。 数値実験では、今の…
ガウスの故事にならって、1000個ずつ区切って素数の個数を数える。1001〜2000に135個、2001〜3000に127個ある。nを1500、2500、…として、横軸を1/log n、縦軸を素数の個数としてプロットした。 素数の密度は、ほぼ1/log nになっていることがわかる。これを素…
GPUというとふつうはfloatの計算だが、あまり興味はない。やっぱり整数の計算をしたい。整数の計算速度はどんなもんだろう。適当に加乗除のコードで試してみた。 __global__ void int_add(int *data, int width, int height, int nRep) { unsigned int xInde…
コードは書けるようになったものの、あまりよくわかっていない。しかし、いちおう解説をしておこう。GPUは、元々画像処理をするものなので、画像をGPUに送る。データは1次元だが、 const unsigned int size_x = 1024; const unsigned int size_y = 1024;とい…
まず、このへんから開発ツールを落としてくる。http://www.nvidia.co.jp/object/cuda_get_jp.htmlそれから、libとincludeにパスを通す。 サンプルを参考にして、前にやった数値積分の問題をCUDAで計算する。http://d.hatena.ne.jp/inamori/20080312/p1フルソ…
GPUを使うと計算速度が桁違いに上がることがあることは前にも紹介した。 しかし、GPUを動かすためのシェーダプログラミングはやや独特で、それ以上にGPUを動かすための準備に膨大なコードが必要で、かつ見通しも悪い。そこでCUDAである。 http://www.nvidia.…
開平（1） 開平の方法はたぶん中学で習ったと思うが、割り算と似たようなやり方がある。あれをコード化するとこんな感じだろうか。 #include using namespace std;void sqrt(int a, int& b, int& c); int get_bit_num(int a);int main() { for(int i = 0; i …
10進変換（4） 前回、ある程度小さい数を10進数にして、それを自乗していくという方法を考えた。今回は、その自乗をGPUで行う。 多倍長整数の乗算は、多項式の乗算と同じようなものだから、1ピクセル（一つの項）ごとに大量の乗算・加算をする。だからGPUで…
10進変換（3） 前回、16ビットずつ掛け算していったが、2100000 = (250000)2だから、250000を計算してそれを自乗したほうが速いような気がする。さらに、250000も同様に。 実際に組んでみたら、倍くらいの時間がかかるようになった。 ただし、この掛け算はGP…
10進変換（2） 前回は1ビットずつ計算していたが、さすがにそれでは遅いので16ビットずつに区切って10進計算を行ったところ、0.1秒くらいになった。 しかし、これでもGPUは使えない。
タイトルを変えた。 10進変換（1） 計算は2進数で行い、表示するときに10進数に直す。 2100000を10進数にするのに、1を10万回2倍すると非常に簡単に実装できる。適当な桁で区切って（下のコードは4桁）、ビットにしたがって、2倍とインクリメントを繰り返す…
前に効率的なべき乗の計算というのをやったが、これは多項式の計算がしたかったからである。 例えば、 という計算をやらせる。 代数的に簡単に計算できるが、それはおいといて、次のようになる。今回はテクスチャを2枚用意するので、久しぶりにフルソースを…
なにも考えずに素直にGLSLでGPUに解かせてみた。 こういうデータを用意する。 const int nMax = 100000; const int nWidth = 512; const int nHeight = 256; const int nData = nWidth * nHeight * 4; float *data = new float[nData]; for(int i = 0; i dat…
以下の話は、NVIDIAのG8シリーズ固有の話である。ただ、推定実行時間でソートしてGPUに投げるというのは恐らくどのGPUでも通用するテクニックだと思われる。 スレッド群 次のようにテクスチャを用意し、 int bit = atoi(argv[1]); int mask = 1 const int nW…
for文 例えば、128*128のデータを用意して、それになんらかの値が入っている。シェーダプログラムのほうは、0ならそのまま抜けて、それ以外なら重い計算（1000万回の繰り返し計算）をするようになっている。 uniform sampler2DRect texUnit0; // datauniform…
配列 配列は次のように宣言して使う。 int a[10];関数に引き渡すには次のようにする。 void f(in int a[10]);f(a);型にinがつくと、入力、 outがつくと、出力、 inoutがつくと、入出力になる。 次の例は、 an = a0 + ... + an-1 s = a0 + ... + an を、a0を…
素数分布（2） 前にも書いたように、GPUは分岐に弱い。 ある程度まとまったスレッド（16 or 32）が同じ命令を実行するため、 分岐すると、違う命令を実行するスレッドは、 今実行している命令が終わるのを待っていなければならない。 だから、分岐しないよう…
演算の速度 こういうデータを用意して、 const int nWidth = 256; const int nHeight = 256; const int nData = nWidth * nHeight * 4; float *data = new float[nData]; int nRep = 16384; for(int i = 0; i data[i] = (float)((i & 1) + 1); }こういう演算…
素数分布（1） 手始めに、ガウスの故事にならって、素数分布を求めよう。 1000ずつ区切って、そのうちの素数の数を数える。 1001〜1000万を1000ずつ区切ることにする。 const int nWidth = 128; const int nHeight = 128; const int nData = nWidth * nHeigh…
整数の演算 intの演算はどうなのだろうか。 uniform samplerRECT texUnit0;void main(void) { vec4 color = texRECT(texUnit0, gl_FragCoord.xy); ivec4 icolor = ivec4(color); icolor /= 3; gl_FragColor = vec4(icolor); }これを実行すると、 0.000000 0.…
全面的に書き直し。 整数の精度 整数のデータを入力したかったが、よくわからなかった。 検索しても出てこないし。 仕方がないので、floatの配列をGPUに送って、キャストすることにした。 const int nWidth = 4; const int nHeight = 2; const int nData = n…
浮動小数点数計算（2） ピクセルごとに少ない計算量だとあまり速くならないが、 ピクセルごとにループを回すと速くなる。 前回の問題を、 を計算するように変える。 コードはほとんど変わらない。 シェーダープログラムのほうを次のようにする。 //shader_nv…
浮動小数点数計算（1） まず、 http://aaa.jspeed.jp/~ohshima/cgi-bin/fswiki/wiki.cgi?page=OpenGL%2BGLSL(Windows)%A4%C7GPGPU_HelloWorld で必要な準備をする。 そして、一番下のコードをコピペし、次のようにコンパイルする。 > cl /O2 /GX gpgpu1.cpp…
GPGPUというのは、General Purpose GPUの略で、 本来CGに使うGPUを汎用的に使おうという話。 GPUは使いにくい面はあるが、CPUより圧倒的に速い。 革命的に速くなる。 しかし、整数計算ではまだあまり使われていないのではないだろうか。 （私も仕事では浮動…
（いろいろ修正）ここ3日くらいは、GPGPUをやっていた。 GPGPUはGeneral Purpose GPUの略で、 GPUは本来グラフィックの処理をやらせるものを、 それ以外のことでも使おうということ。 なぜなら、GPUのほうがCPUより圧倒的に速いから。 多くのデータを並べて…
10進変換（4） 前回、ある程度小さい数を10進数にして、それを自乗していくという方法を考えた。今回は、その自乗をGPUで行う。 多倍長整数の乗算は、多項式の乗算と同じようなものだから、1ピクセル（一つの項）ごとに大量の乗算・加算をする。だからGPUで…
