분명히 중복 작업

Question

나는 글로벌 메모리 매트릭스 out에 글로벌 메모리 매트릭스 in의 간단한 임무를 수행하는 다음과 같은 커널이 있습니다

__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; 

해체 된 코드의 상단 또는 상단에있는 의견은 저만의 것입니다.

자세히 알 수 있듯이 은 쓸모없는 것입니다. 작업은 댓글에 ???으로 표시됩니다. 본질적으로, 그들은 등록기를 스스로 움직이는 것이다. 그들이 쓸모없는 경우

1. , 나는 그들이 쓸데없이 계산 시간을 소비하고 있다고 생각 :

   나는 다음 다음 두 가지 질문이 있습니다. 분해 된 마이크로 코드를 제거하여 최적화 할 수 있습니까?

2. PTX 파일은 CUDA 코드로 인라인 될 수 있습니다. 그러나 PTX는 GPU 간의 이식성에 필요한 중간 언어입니다. 어떻게 든 최적화 된 분해 된 마이크로 코드를 "인라인 (inline)"시킬 수 있습니까?

대단히 감사합니다.

EDIT : 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 */ 

EDIT를 컴파일 SAME CODE : 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 */ 

Answer 1

대답 릴리스 모드에서 컴파일 SAME CODE 두 가지 질문에 대해서는 아니오입니다.

최종 바이너리 페이로드에서 명령어를 삭제하려고 시도하는 경우. 코드 섹션의 길이를 변경하고 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; 
     ..................................... 

당신은 당신이 릴리스 설정으로 컴파일 보여 주었다 코드 있습니까?