1

CUDA コードの逆アセンブルは、場合によっては唯一ではないにしても、コンパイラの動作とパフォーマンス メトリックを理解するのに非常に役立つツールのようです。

残念ながら、CUDA BINARY UTILITIES Application Note から入手できるドキュメントは、CUDA アセンブリ命令を解釈するために必要なすべてのツールをユーザーに提供していないか、少なくともそのドキュメントから必要なすべての情報を推測することはできません。「CUDA ハンドブック」ブックには、CUDA BINARY UTILITIES ガイド以上の情報はありません。たとえば、指示をどのように解釈すればよいですか

ISETP.LT.AND P0, PT, R3, RZ, PT;

と

PSETP.AND.AND P0, PT, !P0, PT, PT;

指示の前に何をし@P0ますか? P0述語レジスタが真の場合に実行がそのラベルにジャンプするように、命令ラベルですか? CUDA アセンブリ命令を解釈する一般的な方法はありますか?

どうもありがとうございました。

NJUFFAのコメントに従って編集

次の単純なカーネルをコンパイルしました

__global__ void test_kernel(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if ((tid > 5) & (tid < 10)) a[tid] = tid;
    else b[tid] = tid;
}

その結果、

/*0000*/        MOV R1, c[0x1][0x100];                 /* 0x2800440400005de4 */
/*0008*/        S2R R0, SR_CTAID.X;                    /* 0x2c00000094001c04 */
/*0010*/        S2R R2, SR_TID.X;                      /* 0x2c00000084009c04 */
/*0018*/        IMAD R2, R0, c[0x0][0x8], R2;          /* 0x2004400020009ca3 */
/*0020*/        IADD R0, R2, -0x6;                     /* 0x4800ffffe8201c03 */
/*0028*/        ISETP.LT.U32.AND P0, PT, R0, 0x4, PT;  /* 0x188ec0001001dc03 */
/*0030*/        I2F.F32.S32 R0, R2;                    /* 0x1800000009201e04 */
/*0038*/   @!P0 ISCADD R3, R2, c[0x0][0x24], 0x2;      /* 0x400040009020e043 */
/*0040*/    @P0 ISCADD R2, R2, c[0x0][0x20], 0x2;      /* 0x4000400080208043 */
/*0048*/   @!P0 ST [R3], R0;                           /* 0x9000000000302085 */
/*0050*/    @P0 ST [R2], R0;                           /* 0x9000000000200085 */
/*0058*/        EXIT ;                                 /* 0x8000000000001de7 */

コンパイラは条件((tid > 5) & (tid < 10))を((i < 4) & (i >= 0)), withi = tid - 6に再キャストしたため、関連する命令は現在

ISETP.LT.U32.AND P0, PT, R0, 0x4, PT;
4

1 に答える 1