cuda - 逆アセンブルされた CUDA マイクロコードの明らかに冗長な操作

Question

inグローバルメモリマトリックスへのグローバルメモリマトリックスの単純な割り当てを実行する次のカーネルがありますout。

__global__ void simple_copy(float *outdata, const float *indata){

    int x = blockIdx.x * TILE_DIM + threadIdx.x;
    int y = blockIdx.y * TILE_DIM + threadIdx.y;

    int width = gridDim.x * TILE_DIM;

    outdata[y*width + x] = indata[y*width + x];

}

によってダンプされた逆アセンブルされたマイクロコードを調べていcuobjdumpます:

Function : _Z11simple_copyPfPKf
/*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100]; 
/*0008*/     /*0x80001de218000000*/     MOV32I R0, 0x20;            R0 = TILE_DIM
/*0010*/     /*0x00001c8614000000*/     LDC R0, c [0x0] [R0];       R0 = c
/*0018*/     /*0x90009de218000000*/     MOV32I R2, 0x24;            R2 = 36
/*0020*/     /*0x00209c8614000000*/     LDC R2, c [0x0] [R2];       R2 = c

int x = blockIdx.x * TILE_DIM + threadIdx.x;
/*0028*/     /*0x9400dc042c000000*/     S2R R3, SR_CTAid_X;         R3 = BlockIdx.x
/*0030*/     /*0x0c00dde428000000*/     MOV R3, R3;                 R3 = R3 ???
/*0038*/     /*0x84011c042c000000*/     S2R R4, SR_Tid_X;           R3 = ThreadIdx.x
/*0040*/     /*0x10011de428000000*/     MOV R4, R4;                 R4 = R4 ???
/*0048*/     /*0x8030dca32008c000*/     IMAD R3, R3, 0x20, R4;      R3 = R3 * TILE_DIM + R4  (contains x)

int y = blockIdx.y * TILE_DIM + threadIdx.y;
/*0050*/     /*0x98011c042c000000*/     S2R R4, SR_CTAid_Y;
/*0058*/     /*0x10011de428000000*/     MOV R4, R4;
/*0060*/     /*0x88015c042c000000*/     S2R R5, SR_Tid_Y;
/*0068*/     /*0x14015de428000000*/     MOV R5, R5;
/*0070*/     /*0x80411ca3200ac000*/     IMAD R4, R4, 0x20, R5;      R4 ...                   (contains y)

int width = gridDim.x * TILE_DIM;
/*0078*/     /*0x50015de428004000*/     MOV R5, c [0x0] [0x14];     R5 = c
/*0080*/     /*0x80515ca35000c000*/     IMUL R5, R5, 0x20;          R5 = R5 * TILE_DIM       (contains width)   

y*width + x
/*0088*/     /*0x14419ca320060000*/     IMAD R6, R4, R5, R3;        R6 = R4 * R5 + R3        (contains y*width+x)

Loads indata[y*width + x]
/*0090*/     /*0x08619c036000c000*/     SHL R6, R6, 0x2;            
/*0098*/     /*0x18209c0348000000*/     IADD R2, R2, R6;            
/*00a0*/     /*0x08009de428000000*/     MOV R2, R2;                 R2 = R2 ???
/*00a8*/     /*0x00209c8580000000*/     LD R2, [R2];                Load from memory - R2 = 

Stores outdata[y*width + x]
/*00b0*/     /*0x1440dca320060000*/     IMAD R3, R4, R5, R3;        
/*00b8*/     /*0x0830dc036000c000*/     SHL R3, R3, 0x2;
/*00c0*/     /*0x0c001c0348000000*/     IADD R0, R0, R3;            R0 = R0 + R3
/*00c8*/     /*0x00001de428000000*/     MOV R0, R0;                 R0 = R0 ???
/*00d0*/     /*0x00009c8590000000*/     ST [R0], R2;                Store to memory

/*00d8*/     /*0x40001de740000000*/     BRA 0xf0;
/*00e0*/     /*0x00001de780000000*/     EXIT;
/*00e8*/     /*0x00001de780000000*/     EXIT;
/*00f0*/     /*0x00001de780000000*/     EXIT;
/*00f8*/     /*0x00001de780000000*/     EXIT;

逆アセンブルされたコードの上または横のコメントは私自身のものです。

ご覧のとおり、明らかに役に立たない???操作がいくつかあり、コメントででマークされています。本質的に、それらはレジスター自体への移動です。

次に、次の2つの質問があります。

それらが役に立たない場合、それらは無駄に計算時間を消費していると思います。逆アセンブルされたマイクロコードを削除して最適化できますか?
PTX ファイルは、CUDA コードでインライン化できます。ただし、PTX は、GPU 間の移植性に必要な中間言語にすぎません。最適化された逆アセンブルされたマイクロコードをどうにかして「インライン化」できますか?

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

編集: SM = 2.0 のリリースモードでコンパイルされた同じコード

Function : _Z11simple_copyPfPKf
.headerflags    @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
/*0000*/        MOV R1, c[0x1][0x100];            /* 0x2800440400005de4 */
/*0008*/        S2R R0, SR_CTAID.Y;               /* 0x2c00000098001c04 */
/*0010*/        S2R R2, SR_TID.Y;                 /* 0x2c00000088009c04 */
/*0018*/        S2R R3, SR_CTAID.X;               /* 0x2c0000009400dc04 */
/*0020*/        S2R R4, SR_TID.X;                 /* 0x2c00000084011c04 */
/*0028*/        MOV R5, c[0x0][0x14];             /* 0x2800400050015de4 */
/*0030*/        ISCADD R2, R0, R2, 0x5;           /* 0x4000000008009ca3 */
/*0038*/        ISCADD R3, R3, R4, 0x5;           /* 0x400000001030dca3 */
/*0040*/        SHL R0, R5, 0x5;                  /* 0x6000c00014501c03 */
/*0048*/        IMAD R2, R0, R2, R3;              /* 0x2006000008009ca3 */
/*0050*/        ISCADD R0, R2, c[0x0][0x24], 0x2; /* 0x4000400090201c43 */
/*0058*/        ISCADD R2, R2, c[0x0][0x20], 0x2; /* 0x4000400080209c43 */
/*0060*/        LD R0, [R0];                      /* 0x8000000000001c85 */
/*0068*/        ST [R2], R0;                      /* 0x9000000000201c85 */
/*0070*/        EXIT ;                            /* 0x8000000000001de7 */

編集：SM = 2.1のリリースモードでコンパイルされた同じコード

Function : _Z11simple_copyPfPKf
.headerflags    @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
/*0000*/        MOV R1, c[0x1][0x100];            /* 0x2800440400005de4 */
/*0008*/        NOP;                              /* 0x4000000000001de4 */
/*0010*/        MOV R0, c[0x0][0x14];             /* 0x2800400050001de4 */
/*0018*/        S2R R2, SR_CTAID.Y;               /* 0x2c00000098009c04 */
/*0020*/        SHL R0, R0, 0x5;                  /* 0x6000c00014001c03 */
/*0028*/        S2R R3, SR_TID.Y;                 /* 0x2c0000008800dc04 */
/*0030*/        ISCADD R3, R2, R3, 0x5;           /* 0x400000000c20dca3 */
/*0038*/        S2R R4, SR_CTAID.X;               /* 0x2c00000094011c04 */
/*0040*/        S2R R5, SR_TID.X;                 /* 0x2c00000084015c04 */
/*0048*/        ISCADD R2, R4, R5, 0x5;           /* 0x4000000014409ca3 */
/*0050*/        IMAD R2, R0, R3, R2;              /* 0x200400000c009ca3 */
/*0058*/        ISCADD R0, R2, c[0x0][0x24], 0x2; /* 0x4000400090201c43 */
/*0060*/        ISCADD R2, R2, c[0x0][0x20], 0x2; /* 0x4000400080209c43 */
/*0068*/        LD R0, [R0];                      /* 0x8000000000001c85 */
/*0070*/        ST [R2], R0;                      /* 0x9000000000201c85 */
/*0078*/        EXIT ;                            /* 0x8000000000001de7 */

score 1 · Accepted Answer

両方の質問に対する答えはノーです。

最終的なバイナリペイロードから命令を削除しようとした場合。コードセクションの長さを変更し、ELF および fatbinary ファイルを分割します。これを修正するには、フォーマットが文書化されていないヘッダーを手作業で作成する必要があり、いくつかの命令を最適化するだけでも大変な作業のように思えます。

また、インラインネイティブアセンブラはサポートされていません。

最後に、CUDA 5.0 を使用して再現できません。

Fatbin elf code:
================
arch = sm_20
code version = [1,6]
producer = cuda
host = mac
compile_size = 32bit
identifier = pumpkinhead.cu

    code for sm_20
        Function : _Z11simple_copyPfPKf
    /*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
    /*0008*/     /*0x98001c042c000000*/     S2R R0, SR_CTAid_Y;
    /*0010*/     /*0x88009c042c000000*/     S2R R2, SR_Tid_Y;
    /*0018*/     /*0x9400dc042c000000*/     S2R R3, SR_CTAid_X;
    /*0020*/     /*0x84011c042c000000*/     S2R R4, SR_Tid_X;
    /*0028*/     /*0x08001ca340000000*/     ISCADD R0, R0, R2, 0x5;
    /*0030*/     /*0x10309ca340000000*/     ISCADD R2, R3, R4, 0x5;
    /*0038*/     /*0x50001ca350004000*/     IMUL R0, R0, c [0x0] [0x14];
    /*0040*/     /*0x08009ca340000000*/     ISCADD R2, R0, R2, 0x5;
    /*0048*/     /*0x90201c4340004000*/     ISCADD R0, R2, c [0x0] [0x24], 0x2;
    /*0050*/     /*0x80209c4340004000*/     ISCADD R2, R2, c [0x0] [0x20], 0x2;
    /*0058*/     /*0x00001c8580000000*/     LD R0, [R0];
    /*0060*/     /*0x00201c8590000000*/     ST [R2], R0;
    /*0068*/     /*0x00001de780000000*/     EXIT;
        .....................................

あなたが示したコードは、リリース設定でコンパイルされたものですか?

cuda - 逆アセンブルされた CUDA マイクロコードの明らかに冗長な操作

1 に答える 1

Related

Reference