Twitter

最近のトラックバック

無料ブログはココログ

« 歯が口内炎が | トップページ | Ubuntuどうしよう »

2012年1月21日 (土)

CUDA上の定数

CUDAでプログラムを作るときに,double型を使っていないにもかかわらず

ptxas /tmp/tmpxft_000008ef_00000000-8_fdirdet_dev.ptx, line 95; warning : Double is not supported. Demoting to float

という警告が消えませんでした.
nvccに渡したオプションは,macでは

-m64 -O2 -lcutil_x86_64 --compiler-options -fpermissive

という感じです.

しょうがないのでptxを出力させたところ,こんな結果となりました(関連部分を抜粋)

ld.global.f32 %f1, [%r12+0];
cvt.f64.f32 %fd1, %f1;
mov.f64 %fd2, 0dbfe0f5c28f5c28f6; // -0.53
mul.f64 %fd3, %fd1, %fd2;
cvt.rn.f32.f64 %f3, %fd3;

-0.53の部分は,ソースファイルの中で#defineした定数です.
何故か-0.53とかけ算するために,一度doubleに変換し,その後floatに戻しているように見えます(多分)
たったこれだけの計算でいろんなレジスタを使っていて,いかにも遅そうです.

ちなみに95行目のコードは

mul.f64 %fd3, %fd1, %fd2;

です.サポートしてないという割には64bitっぽい命令ですが,恐らくcubinに変換されるときに,float型の計算になるのではないのでしょうか.

問題は,-0.53がdouble型となっていることが推測されるので,ソースコードを

-0.53f

と変更してみました.
すると,警告はなくなり,周辺のコードはすっきりしました.
やってることも一目瞭然でいい感じです.

ld.global.f32 %f1, [%r12+0];
mov.f32 %f3, 0fbf07ae14; // -0.53
mul.f32 %f4, %f1, %f3;

そういえば,Cの実数定数の型はdoubleだった気がします.
意外と律儀に型を守ってくれるみたいですね.その方が嬉しいです.

オプションでアーキテクチャの指定をすると,この警告がなくなるというのは読んだことがあるので,
これでは今まで効率の悪いcubinが出力されていた可能性がありますね.

ただ,速度は全く上がりませんでしたw
computeprofレベルで見れば変わっているかもしれませんね.

ところで,アーキテクチャの指定をしていなくても,64bitからの変換や,64bitレジスタの使用には警告が出ないと言うことは,どんなに古いCUDAの動くGPUでも64bitの値を保持する仕組みはあったということでしょうか?

いずれにしても,警告が0になって一安心です.
これで堂々と論文に貼れますw

ちなみに,この警告を直している間にほかのバグを二つ発見しました・・・危なかった・・・

« 歯が口内炎が | トップページ | Ubuntuどうしよう »

技術関連」カテゴリの記事

卒研」カテゴリの記事

CUDA」カテゴリの記事

コメント

この記事へのコメントは終了しました。

トラックバック


この記事へのトラックバック一覧です: CUDA上の定数:

« 歯が口内炎が | トップページ | Ubuntuどうしよう »