ばぐだし

CUDA エクササイズ３は配列の反転。

単一ブロック、複数ブロック、複数ブロック高速版があった。

単一ブロック

// Part 1 of 1: implement the kernel
__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
    int in = threadIdx.x;
    int out = blockDim.x - 1 - threadIdx.x;
    d_out[out] = d_in[in];
}

複数ブロック

// Part3: implement the kernel
__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
    int inOffset  = blockDim.x * blockIdx.x;
    int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
    int in  = inOffset + threadIdx.x;
    int out = outOffset + (blockDim.x - 1 - threadIdx.x);
    d_out[out] = d_in[in];
}

複数ブロック高速版

// Part 2 of 2: implement the fast kernel using shared memory
__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
    extern __shared__ int s_data[];

    int inOffset  = blockDim.x * blockIdx.x;
    int in  = inOffset + threadIdx.x;

    // Load one element per thread from device memory and store it
    // *in reversed order* into temporary shared memory
    s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];

    // Block until all threads in the block have written their data to shared mem
    __syncthreads();

    // write the data from shared memory in forward order,
    // but to the reversed block offset as before

    int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);

    int out = outOffset + threadIdx.x;
    d_out[out] = s_data[threadIdx.x];
}

不思議なのは高速版で、sharedメモリ（ブロック内でのみアクセス）を介することで高速化しているらしい。
普通に見たらsharedメモリ挟む分速度が下がりそうだけど、ブロック単位にメモリアクセスを分ければ高速になるのだろうか。ぐぬぬ。

まぁこれで一通りのエクササイズが終了しました。
次はCUBLASについて調べようと思います。

続いてエクササイズ２、カーネルの起動いきまーす。

ソース一部抜粋

// Part 3 of 5: implement the kernel
__global__ void myFirstKernel(int *d_a)
{
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    d_a[idx] = 1000*blockIdx.x + threadIdx.x;
}

    // Part 2 of 5: configure and launch kernel
    dim3 dimGrid(numBlocks);
    dim3 dimBlock(numThreadsPerBlock);
    myFirstKernel<<<dimGrid,dimBlock>>>(d_a);

    // block until the device has completed
    cudaThreadSynchronize();

※以下私なりに丸めたちょっと怪しい表現を使います。正しい表現は本家ページで！
CUDAでは「関数名<<<グリッド内のブロックの数、１ブロックあたりのスレッドの数>>>(引数) 」で並列カーネル(デバイス上で実行される機能)を呼び出す。
CUDAで書いた関数(この場合myFirstKernel)が1つのスレッドに当たる。1度に実行できるのは1つのグリッドだけ。従って、一度に実行されるスレッド数(myFirstKernel)は、グリッド内のブロックの数(numBlocks)×１ブロックあたりのスレッドの数(numThreadsPerBlock)となる。
ちなみに、グリッド内のブロックの数、１ブロックあたりのスレッドの数は、2次元での表現も可能(ブロックあたりスレッド数は3次元まで可能)で、画像屋さんにはとても助かる。

また、カーネル内では、以下の変数が自動的に定義されており、アクセスができる。
dim3 gridDim;
ブロックのグリッドの次元（最大で2次元）
dim3 blockDim;
スレッドのブロックの次元
dim3 blockIdx;
グリッド内のブロックのインデックス
dim3 threadIdx;
ブロック内のスレッドのインデックス
てなわけでスレッドのインデックスとしては、

int idx = blockDim.x * blockIdx.x + threadIdx.x;

が1次元の並列カーネル呼び出しでは定番となる。

以上、次もエクササイズ！

CUDAってなんなのさ

CUDAをかじろうかと思う。
CUDAってのはGPGPUプログラムのための開発言語、開発環境のことで、まぁ詳しくはこちらを参照ってことで。
GPGPUに関してはこれから出るOpenCLの方が便利そうですし普及しそうだけど、そのベースになるCUDAも勉強しといて損はないかなと。
CUDAは以前サンプルを読んだりもしていたのですが、かなり断片的にしか分かっていないので今回は改めてお勉強。
基本的に本家本元NVIDIAのCUDAのページがとても親切に教えてくれるので、それを追う形でいこうかと。

CUDAのしくみ

アーキテクチャはこんなん。

NVIDIAHPより
ここでは実装に重きを置いて書くことにするので、基本的なCUDAの仕組みはここの「CUDA ベーシックプログラミングガイド」で学んでくださいな。というか私も今はこのページをなぞるだけ。

環境構築

以前よりSDKのバージョンも上がってるし、動かないサンプルがあったので開発環境構築から。
グラボ：GeForce9600GT
OS：Vista 32bit
ここから
1. 最新CUDAドライバ(185.85)
2. CUDAツールキット一式(2.2)
3. CUDA SDKコードサンプル(2.21)
をダウンロード・インストール。

CUDAトレーニング コース

nVidiaのページの学習コースに乗っかる。第1巻： CUDAプログラミング入門 (日本語版)ってとこ。
・エクササイズ1: ホストとデバイス間でのコピー
ホストはCPUで、デバイスはGPU。ホスト-デバイスはGlobal Memoryを通じてデータをやりとりする。大容量だけど遅いから使用は最小限に抑える。

課題部分のソースを抜粋

    // Part 1 of 5: allocate device memory
    size_t memSize = dimA*sizeof(float);
    cudaMalloc((void**)&d_a, memSize);
    cudaMalloc((void**)&d_b, memSize);

    // Part 2 of 5: host to device memory copy
    cudaMemcpy(d_a, h_a, memSize, cudaMemcpyHostToDevice);

    // Part 3 of 5: device to device memory copy
    cudaMemcpy(d_b, d_a, memSize, cudaMemcpyDeviceToDevice);

    // Part 4 of 5: device to host copy
    cudaMemcpy(h_a, d_b, memSize, cudaMemcpyDeviceToHost);

    // Part 5 of 5: free device memory pointers d_a and d_b
    cudaFree(d_a);
    cudaFree(d_b);

とこんな風に、ホストとデバイスのメモリを意識して使い分ける。
おそらく、対象データをホストからデバイスにコピー→デバイス内で演算→結果データをデバイスからホストにコピーの流れになる。

今回は以上。次もエクササイズを追っていきます。

前回の記事で書きましたが、３Dモデルを描画すると後ろの物体が透けてしまいました。
よくよく見ると、どうやら物体の描画順序が遅いものが上書きされてしまっているようです。
で、いろいろ調べたのですが、Integration Of Memeplexesさんの記事で、

protected override void Draw(GameTime gameTime)
{
graphics.GraphicsDevice.Clear(
ClearOptions.Target | ClearOptions.DepthBuffer,    //色と深度バッファをクリア
Color.CornflowerBlue,   //色の初期値
0,    //深度バッファの初期値（いちばん近い所からはじめる）
0     //ステンシルバッファの初期値
);
graphics.GraphicsDevice.RenderState.DepthBufferFunction = CompareFunction.GreaterEqual;
blueTriangle.Draw(graphics.GraphicsDevice, basicEffect);
redTriangle.Draw(graphics.GraphicsDevice, basicEffect);
}
}

なるコードを発見。
GraphicsDevice.Clear関数は1つだけと思っていましたが、
複数の引数をとることができるようです。
そこで描画の基本部分をSpaceWarからコピペすることにしました。
といっても基本的には、

drawDepthBuffer = new DepthStencilBuffer(引数略);

graphics.GraphicsDevice.DepthStencilBuffer = drawDepthBuffer;

を追加して、

graphics.GraphicsDevice.Clear(ClearOptions.DepthBuffer, Color.CornflowerBlue, 1.0f, 0);

のようにClear()関数で引数を指定してやればよいでしょう。
以上で問題が解決して、無事に３Dモデルを表示できました。
これでやっと前に作ったもやしもんのオリゼー君をXNAで表示できます。


ただやすー

今回は３Dモデル表示のための実装です。
前々回のXNA記事でお話したことをおさらいすると、３Dモデルの表示には以下の行列が必要になります。

• ワールド行列
• ビュー行列
• 射影変換行列

これらを定義できれば、とりあえず３Dモデルを出力することはできます。
ところで、３Dモデルを１つのクラスとしてみた場合、上記すべての行列を自分で生成・保持する必要はあるのでしょうか？
ワールド行列はモデルの位置姿勢を定義するものです。なのでモデルに依存します。
ビュー行列はカメラの位置姿勢ですね。なのでカメラに移るモデルに共有されなければなりません。
射影変換行列もビュー行列と同様です。
そんなわけでまずはビュー行列と射影変換行列を管理するCameraクラスを作ります。
これは各クラスから簡単に参照できるようにしたいので、staticオブジェクトとしてGameクラスが保持します。てかSpaceWarでそうしてたのでそうします。
次は3DモデルのクラスMover3Dを作ります。
こいつにはワールド行列と３Dモデルデータを保持させます。
そんなこんなで表示できました。


後側が透けて見えてしまっています。これはいかんですな。
ですがとりあえず表示とコントローラでの回転や移動はできました。
描画部分のソースはこちら。

//モデルの描画
foreach (ModelMesh mesh in _model.Meshes)
{
//エフェクトの指定
foreach (BasicEffect effect in mesh.Effects)
{
//ライティング
effect.EnableDefaultLighting();
//ワールド変換行列
effect.World =  _world;
//ビュー変換行列
effect.View = ShtGame.Camera.View;
//射影変換行列
effect.Projection = ShtGame.Camera.Projection;
}
mesh.Draw();
}

透けているのはeffectに与えるパラメータが不足しているからでしょうか。
それともGameDeviceのパラメータ関連？
次回はこれの解決方法を調べます。
あとModel、Mesh、Effectの関係がなんとなくしか分かっていないのと、
GraphicsDevice.RenderStateプロパティいじったらいろいろ出来そうなので、
これらについてもそのうち調べる予定です。
ま、我慢しきれなくなって製作に移っちゃうかもしれませんが。

前回に引き続き３Dモデルの描画について。



XNAではMatrixという構造体があり、その静的関数でMatrixを生成することが出来ます。

1個ずつ見ていきましょう。

・・・とまぁこんな感じです。

自信なしなのもありますが、前回書いたことが理解に役立ちました。



次回はいよいよプログラムを実装して実際に３Dモデルを表示してみます。

今回からいよいよ3Dモデルの描画に挑戦します。
早速ファイルを読み込んで表示…といきたいところですが、
今回は自分の復習のために３Dモデルがどのようにして描画されるかを書きたいと思います。
というのも、３次元の描画は行列を多用するので、仕組みが分かってないとわけわからなくなる恐れがあるからです。
３Dモデルの描画とは？
３Dモデルというのは、要は点の集まりです。点どうしが線を描いたり面(ポリゴン)を作ったりして3次元物体を表現するのです。
では、３Dモデルを画面に出力するってのはどういうことか？
それは3次元の点P（X,Y,Z）を画面（２次元）の点Q（x,y）に変換することです。
ではどうやって変換すればよいのか？
その答えは簡単。私たちの目やカメラの、物理的な仕組みをシミュレートすればよいのです。
３つの変換
３次元の点Pを２次元の点Qに変換するには以下の３つの変換を行います。
P -> （モデルビュー変換） -> （射影変換） -> （ビューポート変換） -> Q
モデルビュー変換は、３D空間中のモデルの位置・姿勢を表します。カメラの位置・姿勢も兼ねています。
射影変換は、いわば視界の定義です。ズームイン、ズームアウトや、アスペクト比などを設定します。
ビューポート変換は射影変換で出来た画像を画面に印刷する作業です。画面全体に表示することが多いです。
モデルビュー変換
モデルビュー変換は点P(X,Y,Z)に4次元目の値１を加えたベクトル[X,Y,Z,1]と４×４の行列の積により行います。
基本的なものとして以下のような座標変換行列があります。

• なにもしない行列
• 移動
• 拡大縮小
• X軸回転・Y軸回転・Z軸回転
• 任意軸周りの回転

具体的な行列の値は下記参考サイト様を見てください。
後で説明しますが、これら行列の中身は知っておくべきですが、常に書けるようになる必要はありません。
上記の変換の他に、例えば座標(1,0,0)にZ軸周り90度で回転したいときは、
回転行列と移動行列の積を行えばOKです。
ところで、行列の合成は拡大縮小→回転→移動の順に行わないと、正しい位置に移動できないので注意が必要です。回転はモデル中心ではなく原点中心で行われるからです。
ところで、ここまでモデルの座標変換だけでカメラについては出てきませんでしたが、
カメラはモデルと相対関係（モデルが動くのとカメラが逆に動くのが同じ）にあるので、モデルビュー行列はそれらを一括で表現できるわけです。
といってもそれでは分かりづらいのでDirectXやXNAでは

• ワールド行列（モデルの配置）
• ビュー行列（カメラの配置）

の2つに分けて設定できるようになっています。
ついでに言っときますが、この行列は原点の移動も兼ねています。オブジェクトａから見たオブジェクトｂの位置姿勢が行列Bで表されるとき、原点から見たオブジェクトｂの位置姿勢（ワールド行列Ｂ’）はオブジェクトａのワールド行列Ａを用いて、
Ｂ’＝ＢＡ
で表すことができます。
これを使えば関節の動きなども表現できますね。
※DirectXはOpenGLとはかける順序が逆になるはずです。ご注意を。
射影変換
射影変換で設定する項目は以下のとおりです。

• 画角
• アスペクト比
• Nearプレーン、Farプレーン

画角はズームイン、ズームアウトを行うものです。画角が狭いとズームイン、画角が広いとズームアウトです。アスペクト比は、画面の縦横比です。描画する領域にあわせるとよいでしょう。NearプレーンとFarプレーンは手前と奥がどこまで見えるかを決めるものです。
ビューポート変換
ビューポート変換は印刷作業です。これまでの変換で出来た映像を画面に印刷します。
左右半分ずつの対戦画面は、ビューポート変換時に印刷対象として画面の半分を指定して出力しています。
このとき印刷する画面のアスペクト比が射影変換と同じでないと縦や横に伸びた映像になってしまうので注意が必要です。
ビューポート変換は他の2つの変換と違い、行列ではなくViewport構造体でパラメータを指定します。
XNAでの行列生成
モデルビュー変換、射影変換は行列を用いて行うものですが、いちいち自分で行列の値を1つずつ指定する必要はありません。
XNAには少ないパラメータで目的の行列を生成してくれるありがたい関数があるのです。
次回はXNAの行列生成関数について調べます。
※本記事を書くにあたりこちらとこちらを参考にさせていただきました。

今回はスプライト貼ってみます。あとそれをコントローラで動かせるようにします。
ここで役に立ったのがひにけにXNAさんの
SpriteBatchクラス その２

SpriteBatch.Beginメソッドにはブレンドモード、ソートモード、そしてステートモードの3つ引数を渡すメソッド以外に、引数を省略できる2つのオーバーライドがあります。引数を省略した場合、ブレンドモードはSpriteBlendMode.AlphaBlend、ソートモードはSpriteSortMode.Deferred、そしてステートモードはSaveStateMode.Noneとなります。
ここで重要なのはステートモードがSaveStateMode.Noneということで、SpriteBatch.Begin()、またはSpriteBatch.Begin(SpriteBlendMode blendMode)を呼んだ場合は、レンダーステートが変更されてしまうということです。変更されるレンダーステートは指定するSpriteBlendModeによって変わりますが、共通して以下のレンダーステートを変更します。

スプライトのみの描画を行う場合はいいですが、３Dオブジェクトも描画する場合はレンダーステートに気をつけなければなりません。
んまぁこれぐらいで、案外すんなりできました。
今までのDirectXを使ったプログラミングではここまでくるのにも結構苦労しましたが、
さすがXNA、すぐにできますね。
C#も使いやすく、アマチュアのゲーム作りには打ってつけじゃないでしょうか。
(製作物を使う側の制約はおいといて
できたのはこんなかんじ。

出演はOtomaniaさんより、はちゅねミクでした。

階層化GameComponentを作り始めるとソースが肥大化して分かりにくくなると思うので、しばらく放置しておきます。
その間はGame.Componentsプロパティを使用します。…あれ？こうやって使うのか？
とりあえず単純なゲームなら階層化は必要ないでしょう。
さて、今回はデバッグ用コンポーネントを作ります。
何を作るにしても、FPS等のデバッグ情報は表示しておきたいですからね。
まず画面に文字を表示するために一苦労あります。
SpriteFontというXMLファイルを作り、それをフォントとして読み込む必要があるのです。
フォントが読み込めたら、SpriteBatch.DrawString()関数で任意の場所に文字が書けます。
SpriteFontの仕様や文字の書き方については、XNA Game Studioメモさんの記事を参考にさせてもらいました。
次はFPSの計測です。
えーこれについては箱○で始めた（備忘録）さんのソースコードを丸パクリさせてもらいました。
パクってばかりですねすいませんorz
箱○で始めたさんありがとうございます。
てなわけで完成したソースコード公開です。自己責任でご自由にお使いください。

public class DebugMessage : DrawableGameComponent
{
private double _fps;
private double _updateInterval = 1.0;
private double _timeSinceLastUpdate = 0.0;
private double _framecount = 0.0;
private double _elapsed;
private SpriteFont _font;
private SpriteBatch _spriteBatch;
private Vector2 _position = Vector2.Zero;
public DebugMessage(Game game, SpriteBatch spriteBatch)
: base(game)
{
// TODO: Construct any child components here
_spriteBatch = spriteBatch;
}
public override void Initialize()
{
// TODO: Add your initialization code here
base.Initialize();
}
protected override void LoadContent()
{
base.LoadContent();
_font = this.Game.Content.Load("DebugFont");
}
public override void Draw(GameTime gameTime)
{
_elapsed = gameTime.ElapsedRealTime.TotalSeconds;
_framecount++;
_timeSinceLastUpdate += _elapsed;
if (_timeSinceLastUpdate > _updateInterval)
{
_fps = _framecount / _timeSinceLastUpdate;
_timeSinceLastUpdate -= _updateInterval;
_framecount = 0;
// FPSをコンソールに出力
System.Diagnostics.Debug.WriteLine("FPS: " + _fps.ToString("00.000") + " RT: " + gameTime.ElapsedRealTime.TotalSeconds);
}
base.Draw(gameTime);
// FPSを画面に出力
_spriteBatch.DrawString(_font, "FPS: " + _fps.ToString("00.000"), _position, Color.Black);
}
}

ちなみに実行するとこんな感じ。(クリックで拡大)

とりあえずサンプルのSpaceWarのソースを見てみると、Game.Componentsプロパティは一切使ってないご様子。
SpaceWarでは、自前で階層化したオブジェクト管理クラスSceneItemを用いて更新と描画を行っていました。そのクラスが継承していたのがGameComponentsCollectionでもGameComponentでもなく、Listだったのが不思議。
更新と描画をするなら何でもGameComponentにすればいいわけではないのか。
そういえばXNAGSEのGameComponentってもともとWindowsのコントロールみたいにドラッグ＆ドロップで手軽にゲームが作成できるようにっていう概念がありましたよね。あれって今なくなってるのかな？
Game.Componentsプロパティってその名残で、とりあえず初めて作る人はここに追加すれば動きますよ的なものなのかも。
とりあえずSpaceWarにならって自分で階層化可能なオブジェクト管理クラスを作ることにします。
Game.ComponentsプロパティはFPSの表示などデバッグ的な部分のみに留めましょう。