2

PTX を別のファイルに書き込む場合、次の方法でカーネル パラメーターをレジスターに読み込むことができます。

.reg .u32 test;
ld.param.u32 test, [test_param];

ただし、インライン PTX を使用する場合、Using Inline PTX Assembly in CUDA (バージョン 01) アプリケーション ノートでは、パラメーターの読み込みが別の操作に密接にリンクしている構文について説明しています。次の例を示します。

asm("add.s32 %0, %1, %2;" : "=r"(i) : "r"(j), "r"(k));

生成するもの:

ld.s32 r1, [j];
ld.s32 r2, [k];
add.s32 r3, r1, r2;
st.s32 [i], r3;

多くの場合、2 つの操作を分離する必要があります。たとえば、パラメーターをループ外のレジスターに格納し、ループ内でレジスターを再利用して変更したい場合があります。これを行うために私が見つけた唯一の方法は、追加の mov 命令を使用して、パラメーターを暗黙的にロードされたレジスターから、後で使用できる別のレジスターに移動することです。

別のファイルの PTX からインライン PTX に移動するときに、この追加の mov 命令を回避する方法はありますか?

4

1 に答える 1

4

もし私があなたなら、これらの mov 操作についてあまり心配しないでしょう。

PTX は最終的なアセンブリ コードではないことに注意してください。PTX は、カーネルの起動前にさらに CUBIN にコンパイルされます。特に、この最後のステップでは、レジスタの割り当てが実行され、不要なmov操作がすべて削除されます。

特に、 から移動し%r1てから%r2まったく使用しない場合、アルゴリズムは同じハードウェア レジスタにおよび%r1を割り当て、移動を削除する可能性があります。%r1%r2

于 2012-04-03T05:48:50.797 に答える