7

cプログラミングでasmを使用する場合と同様に、コストのかかる実行を減らすために、CUDACコードでアセンブリコードを使用したいと思います。

出来ますか？

4

2 に答える 2

20

CUDA 4.0以降、インラインPTXはCUDAツールチェーンでサポートされています。ツールキットには、それを説明するドキュメントがあります：Using_Inline_PTX_Assembly_In_CUDA.pdf

以下は、CUDA4.0でのインラインPTXの使用を示すコードです。このコードは、CUDAの組み込み__clz（）関数の代わりとして使用するべきではないことに注意してください。私は、新しいインラインPTX機能の側面を調査するためにこのコードを記述しただけです。

__device__ __forceinline__ int my_clz (unsigned int x)
{
    int res;

    asm ("{\n"
         "        .reg .pred iszero, gezero;\n"
         "        .reg .u32 t1, t2;\n"
         "        mov.b32         t1, %1;\n"
         "        shr.u32         %0, t1, 16;\n"
         "        setp.eq.b32     iszero, %0, 0;\n"
         "        mov.b32         %0, 0;\n"
         "@iszero shl.b32         t1, t1, 16;\n"
         "@iszero or.b32          %0, %0, 16;\n"
         "        and.b32         t2, t1, 0xff000000;\n"
         "        setp.eq.b32     iszero, t2, 0;\n"
         "@iszero shl.b32         t1, t1, 8;\n"
         "@iszero or.b32          %0, %0, 8;\n"
         "        and.b32         t2, t1, 0xf0000000;\n"
         "        setp.eq.b32     iszero, t2, 0;\n"
         "@iszero shl.b32         t1, t1, 4;\n"
         "@iszero or.b32          %0, %0, 4;\n"
         "        and.b32         t2, t1, 0xc0000000;\n"
         "        setp.eq.b32     iszero, t2, 0;\n"
         "@iszero shl.b32         t1, t1, 2;\n"
         "@iszero or.b32          %0, %0, 2;\n"
         "        setp.ge.s32     gezero, t1, 0;\n"
         "        setp.eq.b32     iszero, t1, 0;\n"
         "@gezero or.b32          %0, %0, 1;\n"
         "@iszero add.u32         %0, %0, 1;\n\t"
         "}"
         : "=r"(res)
         : "r"(x));
    return res;
}
于 2011-08-15T23:46:58.533 に答える
4

いいえ、できません。C/C++ の asm 構造のようなものはありません。できることは、生成された PTX アセンブリを微調整してから、CUDA で使用することです。

例については、これを参照してください。

ただし、GPU の場合、アセンブリの最適化は必要ありません。メモリの結合や占有など、他の最適化を最初に行う必要があります。詳細については、CUDA ベスト プラクティス ガイドを参照してください。

于 2010-09-09T14:10:41.620 に答える