忍者ブログ
CUDA+GPGPU、C++、C#などのプログラムについての備忘録がわり
[31] [30] [29] [28] [27] [26] [25] [24] [23] [22] [21
Posted by - 2019.10.17,Thu
×

[PR]上記の広告は3ヶ月以上新規記事投稿のないブログに表示されています。新しい記事を書く事で広告が消えます。

Posted by サンマヤ - 2008.12.09,Tue
CUDA関連でいろいろと見ていましたら、
かやの開発覚え書きさんというところで、
for文がおかしいという記述をみつけ、ちょっと試してみようと思いました。
内容はリンク先を参照していただくとして、
見事、現象の再現を確認(Linux x64)しましたw

んで、nvcc --ptxで、ptxを出してみたわけですが、これはビックリな結果に・・・
最初何がおきているのかよくわからなかったのですが、
少し計算してみると、unsignedの引き算に問題があることが分かりました。




まずは、問題のptxを引用します。
コメント部は私が書き足したものです。(g_dataは元記事のbufに相当)
また元記事でunsigned char*のところをfloat* にしています。
----以下、抜粋
$LBB1__Z10testKernelPfj:
    .loc    15    58    0
    ld.param.u64     %rd1, [__cudaparm__Z10testKernelPfj_g_data];    //  %rd1=g_data
    mov.s64     %rd2, %rd1;              // %rd2=%rd1(=g_data)copy_nの第1引数
    add.u64     %rd3, %rd1, 40;          //%rd3=%rd+40(=g_data+10) copy_nの第2引数
    ld.param.u32     %r1, [__cudaparm__Z10testKernelPfj_n];    //%r1=n
    cvt.u64.u32     %rd4, %r1;           // %rd4=n
    mul.lo.u64     %rd5, %rd4, 4;        // %rd5=4*n
    add.u64     %rd6, %rd1, %rd5;        // %rd6=%rd1+%rd5(=g_data+n)
    add.u64     %rd7, %rd6, 17179869180;    // !!問題の箇所 %rd7=%rd6+0x3FFFFFFFC %rd7=lastのつもりっぽい・・・
    setp.eq.u64     %p1, %rd7, %rd1;     // %p1 = (%rd1==%rd7 ? 1 : 0)
    @%p1 bra     $Lt_0_5;                // %p1がTrueならループを飛ばす
$Lt_0_7:
 //<loop> Loop body line 58, nesting depth: 1, estimated iterations: unknown
    .loc    15    52    0
    ld.global.f32     %f1, [%rd2+0];     //
    st.global.f32     [%rd3+0], %f1;     // *dest=*first
    .loc    15    50    0
    add.u64     %rd2, %rd2, 4;           // first++
    add.u64     %rd3, %rd3, 4;           // dest++
    setp.ne.u64     %p2, %rd7, %rd2;     //  %p2 =( first==%rd7 ? 1 : 0 )
    @%p2 bra     $Lt_0_7;                // %p2がTrueならループ
$Lt_0_5:
----ここまで
というわけで、!!問題の箇所としたところがコンパイラのバグっぽいです。
最初数字を見たときはなんだこりゃ、でしたが、
16進数に直すと0x3FFFFFFFC=0x400000000  - 4です。
ということは、インライン展開されたときに、
last =first + size = first + unsigned(n-1)
=first + n + unsgined(-1)
32bitの範囲で、0-1=0xFFFFFFFFで、これを64bitで4倍(floatなので4バイト)したのが
0x3FFFFFFFCというわけですね。。
32bitと64bitの使い分けができていないみたいです。
32ビットで符号拡張>64ビットにする>4倍する って感じでしょうか。
っていうか、引き算あるだろう・・・
元記事をよく見ると、引き算を使っているところでバグが出ているのが分かります。
最初は気づかなかったですがw
ってことで、sizeの宣言をunsigned int からintに変えてみると、問題のところは、
   sub.u64  %rd7, %rd6, 4
とちゃんと引き算してくれました。
結果もちゃんと戻ってきます。

今回、はじめてptxレベルで解析しましたが、
不可解な現象はコンパイラのバグに起因することがよくあるみたいです。
本当はデバッガから入るのが普通のやり方なんですが、
現象的にコンパイラのバグっぽかったのでいきなりptxからやってしまいました。
まあ、CUDAでは
・ループはなるべく使わない(これはGPUの仕組み上、パフォーマンスが落ちるため)
・ポインタの演算はあやしい(コンパイラがあやしい)
と考えてまず間違いないかと・・・
PR
Comments
無題
かやの覚え書きからきました。

なるほど、ptxで見るとやっぱ狂ってたんですね。もはやあきらめてまして、最近CUDAから離れてました。


逆にこちらへの記事の参照を張りたいのですが、よろしいでしょうか?
Posted by かや - 2008.12.10,Wed 18:11:03 / Edit
どうぞ
CUDAはむずかしいです。
私もしばらくやってませんでした。
C++とは別物と思った方がいいですね。
Posted by サンマヤ - 2008.12.10,Wed 21:39:04 / Edit
無題
ありがとうございます。
そうですね、基本的なコードさえ通るようになってからでないと、
仕事で有効利用出来なさそうです。
Posted by かや - 2008.12.11,Thu 00:41:54 / Edit
Post a Comment
Name :
Title :
E-mail :
URL :
Comments :
Pass :   Vodafone絵文字 i-mode絵文字 Ezweb絵文字
TrackBack URL
TrackBacks
カレンダー
09 2019/10 11
S M T W T F S
1 2 3 4 5
6 7 8 9 10 11 12
13 14 15 16 17 18 19
20 21 22 23 24 25 26
27 28 29 30 31
フリーエリア
最新コメント
[11/19 矢野 忠]
[02/25 山本義和]
[07/08 hirota]
[07/06 hirota]
[02/05 矢野 忠]
最新トラックバック
プロフィール
HN:
サンマヤ
性別:
非公開
職業:
プログラマ
趣味:
ゲーム
バーコード
ブログ内検索
カウンター
忍者アナライズ
Template by mavericyard*
Powered by "Samurai Factory"
忍者ブログ [PR]