インラインPTX asm（）命令を使用する場合、 'volatile'とは何ですか？

Question

我々は例えば、私たちの一般C/C++ CUDAコードでインラインPTXアセンブリを書く：

__device__ __inline__ uint32_t bfind(uint32_t val) 
{ 
    uint32_t ret; 
    asm ("bfind.u32 %0, %1;" : "=r"(ret): "r"(val)); 
    return ret; 
} 

我々は例えば、asm後volatileキーワードを追加することができます。

__device__ __inline__ uint32_t bfind(uint32_t val) 
{ 
    uint32_t ret; 
    asm volatile ("bfind.u32 %0, %1;" : "=r"(ret): "r"(val)); 
    return ret; 
} 

CUDA documentation on inline PTX assembly言う：

出力オペランドを変更する以外は、asm()ステートメントに副作用がないものとします。 asmが削除またはPTXの世代の中に移動されていないことを確認するために、あなたはvolatileキーワード

私はそれが何を意味するのか理解していないを使用する必要があります。したがって、

• asm()はなぜ削除されるのですか？コンパイラがそれに気付かない場合、なぜそれが削除されているのでしょうか？
• PTXの生成中にasm()が移動した場合、なぜ問題になるのですか？それは最適化プロセスの一部ですね。
• 不揮発性で揮発性の命令にそれぞれ対応するとき、コンパイラの動作をより正確にどのように特徴付けていますか？

Answer 1

なぜ私のasm（）が削除されるのですか？コンパイラがそれに気づいた場合は、 は何の効果もありませんが、なぜ私はそれが削除されているはずですか？

コンパイラがインラインPTXは、スレッドローカルスコープで以外での状態を変更するには寄与しないことを検出した場合は、最適化として、それを削除すること自由に感じています。 一般的に言えば、それはまさにあなたが起こりたいものです。しかし、時にはそうではありません。あなたの意図とコンパイラの最適化戦略が、あなたが望むまたは期待する方法で交差するとは限りません。注意してくださいemptorとすべてのこと。

私のasm（）がPTXの生成中に動かされたのはなぜですか？ それは最適化プロセスの一部ですね。

これは問題ではなく、最適化プロセスの一部です。時にはそれを回避したいかもしれません。マイクロベンチマークを作成し、コンパイラがインラインPTXでコーディングした注意深く設計された命令の順序を変更すると決めたとします（古典的なケースでは、コードを間違った場所に移動してタイミングセクションやメモリトランザクションパターンの設計が壊れる）。結果はあなたの意図通りではありません。私はそれがかなりイライラするかもしれないと思います。

がそれぞれvolatile asm（）命令に直面したとき、どのようにコンパイラの動作をより正確に特徴づけることができますか？

標準的なCUDAカーネルコードと同様に、volatileは、コンパイラがコード解析によって最適化されることなく、出力で特定のインラインPTX操作を発行することを保証します。