이상한 결과

Question

내 커널은이 같은 PTX 버전이 있습니다

.version 2.2 
.target sm_20, texmode_independent 

.entry histogram(
     .param .u32 .ptr .global .align 4 histogram_param_0, 
     .param .u32 .ptr .global .align 4 histogram_param_1 
) 
{ 
     .reg .f32  %f<2>; 
     .reg .s32  %r<12>; 

_histogram: 
     mov.u32   %r1, %tid.x; 
     mov.u32   %r2, %envreg3; 
     add.s32   %r3, %r1, %r2; 
     mov.u32   %r4, %ctaid.x; 
     mov.u32   %r5, %ntid.x; 
     mad.lo.s32  %r6, %r4, %r5, %r3; 
     shl.b32   %r7, %r6, 2; 
     ld.param.u32 %r8, [histogram_param_0]; 
     add.s32   %r9, %r8, %r7; 
     ld.param.u32 %r10, [histogram_param_1]; 
     ld.global.f32 %f1, [%r9]; 
     add.s32   %r11, %r10, %r7; 
     st.global.f32 [%r11], %f1; 
     ret; 
} 

를 나는 내가 계산으로, 내 커널 만 13 지침 (안가있다 ret 명령 포함). 작업 항목 수를 5120으로 설정하면 작업 그룹 크기는 64입니다. 각 SM에는 32 개의 스칼라 프로세서가있는 16 개의 SM이 있으므로 위의 코드는 SM에서 10 번 실행됩니다. 예상대로 실행 된 명령어의 수는 10 * 13 = 130이어야합니다.하지만 프로파일 링 한 후 결과는 발행 된 명령어 = 130, 실행 된 명령어 = 100입니다. 1. 발행 된 명령어의 수가 실행 된 명령어의 개수와 다른 이유는 무엇입니까? 가지가 없으므로 그들은 평등하지 않습니까? 2. 왜 실행 된 명령의 수가 예상보다 작습니까? ptx 버전의 모든 지침을 최소한 실행해야합니까? 3. 캐시 누락 (L1 및 L2)은 발급 된 명령 수 및 실행 된 명령 수에 영향을 줍니까? 감사합니다.

Answer 1

PTX는 컴파일 된 코드의 중간 표현 일뿐입니다. GPU가 실제로 실행하는 것은 아닙니다. GPU가 실행하는 코드를 방출하는 추가 어셈블리 단계가 있습니다.이 작업은 컴파일 할 때 또는 드라이버에서 JIT 컴파일을 사용하여 수행 할 수 있습니다. 결과적으로, 귀하의 명령 수 및 귀하가 그로부터 추측 한 내용은 모두 유효하지 않습니다.

는 NVIDIA는 PTX는 GPU에서 실행중인 정확히 아니라는 것을 명심 페르미 카드에 대해 생성 된 어셈블러 출력을 분해하고, GPU

Answer 2

에 실제 기계 코드의 실행을 표시 할 수 있습니다 cuobjdump라는 도구를 제공. PTX는 단지 중간 표현 일뿐입니다. 실제 코드는 .cubin 파일에 있습니다. 그렇기 때문에 ptx 소스 코드를 기반으로 정확한 계산을하는 것은 아무런 의미가 없습니다.

cuobjdump --sass CUDA 4.0과 함께 제공되는 도구를 사용하면 .cubin 파일의 GPU 어셈블리 코드를 더 읽기 쉬운 것으로 추출 할 수 있습니다.