我有以下内核执行全局内存矩阵 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;
反汇编代码顶部或旁边的注释是我自己的。
如您所见，有一些显然无用的操作，在注释中用???标记。本质上，它们是将寄存器移动到自身。
我有以下两个问题:
- 如果它们没用，我相信它们是在无用地消耗计算时间。我可以通过删除它们来优化反汇编的微码吗？
- PTX 文件可以内联到 CUDA 代码中。然而，PTX 只是一种跨 GPU 可移植性所需的中间语言。我能否以某种方式“内联”优化的反汇编微代码？
非常感谢您。
编辑:在 SM = 2.0 的 Release模式下编译的相同代码
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 的 Release模式下编译的相同代码
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 */
最佳答案
这两个问题的答案是否定的。
如果您尝试从最终二进制负载中删除指令。您将更改代码段的长度并破坏 ELF 和 fatbinary 文件。要解决这个问题，需要手工制作其格式不易记录的 header ，这听起来像是为了优化几条指令而做的大量工作。
并且不支持内联原生汇编程序，但我相信您已经知道了。
最后，我无法使用 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 微代码中的明显冗余操作，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/19125998/