CUDAコンパイラ、unsignedの引き算ができない？

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の仕組み上、パフォーマンスが落ちるため）
・ポインタの演算はあやしい（コンパイラがあやしい）
と考えてまず間違いないかと・・・