So-net無料ブログ作成

CUDAソースを読んでみる(dct8x8) [プログラミング]

意図が読めそうなソースをピックアップ。dct8x8にしました。とりあえず、VC++2008 Express Editionを入れてみた。Microsoftも他のIDEが軒並みタダで配られているから、配らざるを得なくなっちゃったみたいね。継続的にやっているのは良い事です。ちょっとしたソフトのコンパイルであればこれで十分(と思いたい)。



.cuってのが独自の拡張子だけれども、C言語とどこが違うのか分からない。んで、拡張子を登録していないので、デバッグとかブレークポイントをはれるのかな? どっちにしてもC言語の色分けをしてもらった方が見やすいので、プロジェクトのコンフィグを変更しておく。

メニューで「ツール>オプション>VC++プロジェクトの規定c/c++ファイル拡張子」に「;*.cu」を加える。他に「ツール>オプション>テキストエディタ>ファイル拡張子にcuを登録しておく。後者の方は必要ないかもしれない。そうすると次開く時に、きちんと色分けされたソースとして表示される。うむ、きちんとcuファイルでもブレークで止まりました。問題無し。

dctってのはDiscrete Cosine Transformの略で、JPEGとかMPEG1/2の基本となるピクチャ(何ピクチャかは忘れたけど)は8x8ピクセルのブロックごとに処理される。ジグザグに読んでフーリエとかなんとかで処理するんだったな、たぶん。浮動小数点な処理なので、GPUに処理させるには打ってつけな訳で。んで、8x8っていうのも数学の行列に当てはめやすいので、処理をそのまま当てはめる事ができる。

浮動小数点の計算だと、結果が微妙にずれるらしく、アルゴリズムが一緒でも全く同じ物はできない。だから、H.264だと配列に入れるのが整数みたいで、計算する環境によってばらつきは出ないらしい。先に読んだ通り、GPUは浮動小数点の計算に特化しているので、如何せん整数演算が弱い。というか、整数の計算はCPUでやってくれっていうスタンスだったと思う。幸か不幸か、3Dの計算は浮動小数点を扱うものばっかだったらしく、それまでに特別な整数演算用のレジスタとかがあんまりなかったのかもしれない。でも、GPGPUを標榜する限り、今後は整数関連も力を入れて行かなくてはならないのは必至である。

cuファイルに書かれている関数などは、定義場所を参照しても見られない事がありますね。コンパイルした時の参照は上手くいっているみたいだけれども、右クリックで見られる宣言場所はcuだとダメな傾向高いです。いちいち文字列でプロジェクト全体を探さないといけないので、ちょっと面倒かもしれない。<<< >>>という表記があるのだけれど、これはcu独特の表記なんでしょうか。テンプレートで<>は使っているけれども、単純にテンプレートとは書き方が違うし、演算子の再定義なんだろうか。ちょっと調べておこう。たぶん、GPU側で実行しますよー的なものだとは思うけれども確認してない。

ちょっと思いついたんだけど、アセンブラで表記されている部分ってのは、CPUのSIMDなものだろうから、GPUに持って行っても同じ程度しか変わらないんじゃないかって、ちょっと思った。さらに整数の演算はGPUは得意とするところではないし、これからサポートされるかもしれない程度のアドバンテージしかない。各種SSEを使わず、GPUと比較するならば、ものすごく差は出るかもしれないけれども、最近のアプリケーションは結構本気で最適化している感じなので、それほど差が出ないんじゃないかとふと感じた。

とはいえ、足踏みをしている場合じゃないので、どんどん進む。進めればの話ではあるが…。CPUはマルチコアでマルチスレッドに頼る程度の将来性しかないけれども、GPUは、特にGPGPUの部分はどうなるか分からないところがある。ただ、汎用的なプログラミングではなく、CPUで遅いから部分的にGPUにやらせておこうか、程度の期待しか実はなかったりもするのだが。とはいえ、H.264は処理的に重くてGPUにやらせるとそれなりに速くなるのは、AMD Streamingで実証済みだからやって損をする事だけはないと思われる。やっぱ頑張ろうっと。同じぐらいでも併用したり、GPUもマルチコアなどになったら状況は変わってくるだろうし。


dct8x8での、CUDA2の処理の大まかな流れ。
処理内容 関数とか
CPU側のデータを確保。
送るデータの変換。(dctではintからfloat)
GPU側のメモリの確保 cudaMallocPitch()
CPU側からGPU側への処理データのコピー cudaMemcpy2D()
実行パラメータを渡してGPUで実行(DCT,idct,量子化) <<< >>>()って何?
スレッドの同期 cudaThreadSynchronize()
演算結果をCPUのメモリにコピー cudaMemcpy2D()
結果データの変換。(floatからbyteに)
GPUのメモリの解放 cudaFree()



GPUのメモリにコピーして、演算して、CPUのメモリに返す、というプロセスは、この他でも大して違いはないでしょうね。飛び抜けて難しい概念とかはないみたいですね。アセンブラは一文字もないし。フーリエ変換をコンピュータにやらせるためにコードに落とすということの方が、私にとって難しい事項です。いかにパフォーマンスを良く処理を選定してGPUに任せるのかが焦点になっていると思われます。という点で、cutXxxxTimerとかはパフォーマンスを測る術として多用する事になるかもなとも思われます。

ちなみにdct8x8のときもGoldっていう方が出てなくって、例に習って表示するようにさせたら、以下の結果が出てきていた。

Processing time (Gold 1)    : 25.852741 ms
Processing time (Gold 2)    : 3.865470 ms
Processing time (CUDA 1)    : 1.970371 ms
Processing time (CUDA 2)    : 0.369682 ms
Processing time (CUDA short): 0.453289 ms


やはり浮動小数点計算はそれなりに速いが、純粋にDCTのみの動作時間である事をお忘れなく。

ええと、一部だけじゃなくて、一連の処理を測ってみようと思う。そうすれば、データの転送などの時間もある程度、勘案する事ができると思う。サンプルソースだから手を抜いているとは考えたくないので、それなりに最適化していると信じよう。

CUDA sample DCT/IDCT implementation
===================================
Loading test image: barbara.bmp... [512 x 512]... Success
Running Gold 1 (CPU) version...
all time (gold1)    : 317.126343 ms
Success
Running Gold 2 (CPU) version...
all time (gold2)    : 82.745499 ms
Success
Running CUDA 1 (GPU) version...
all time (CUDA1)    : 80.248291 ms
Success
Running CUDA 2 (GPU) version...
all time (CUDA2)    : 24.764336 ms
Success
Running CUDA short (GPU) version...
all time (CUDA short)    : 7.980774 ms
Success


結果は以上。どこがどう時間がかかっているか、処理のボトルネックが分からないので何とも言えないけれども、最適化によって全体のパフォーマンスは上げる事ができるってことでしょう。でもSSEとマルチコアとマルチスレッドを駆使して最適化したもの以上のパフォーマンスが出るかどうかは分からない。

Core i7とかだとメモリコントローラがCPUにあるし、そこを結ぶバスも太い。それを考えると、無駄にスレッドを使っていたりしたら、無駄に速くなる事は想像するに難くない。とはいえ、グラフィックのバスも必要以上に太いし、何かをやるためには必要なだけメモリの量はあるし、メモリ自体も普通のDDRを使っているのではないのが普通なので、CPU側のメモリよりかは早くなっているはずだ。

んーやっぱりGPGPUも使いようによっては速くなるっていうスタンスが正しいのかもしれない。CPUと同時に使って違う事させて高速化するっていう方法がスマートでしょう。どっちか片方となると、どっちもそれなりに処理が頭打ちになりそうな気がする。その点では、GPUとCPUの違いを気にしないでプログラミングできると言われる、MacのOpenCLは良いやり方なのかもしれない。


タグ:GPGPU
nice!(0)  コメント(0) 
共通テーマ:パソコン・インターネット

nice! 0

コメント 0

コメントを書く

お名前:[必須]
URL:
コメント:
画像認証:
下の画像に表示されている文字を入力してください。