CUDA PTX 코드 및 레지스터 메모리와의 혼동

Question

:) 커널 리소스를 관리하는 동안 PTX를 살펴보기로 결정했지만 이해할 수없는 몇 가지 사항이 있습니다. 여기에 내가 쓴 아주 간단한 커널은 다음과 같습니다 그럼 사용하여 컴파일 된

__global__ 
void foo(float* out, float* in, uint32_t n) 
{ 
    uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; 
    uint32_t one = 5; 
    out[idx] = in[idx]+one; 
} 

: nvcc --ptxas-options=-v -keep main.cu을 나는 콘솔에서이 출력을 가지고 :

ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z3fooPfS_j' for 'sm_10' 
ptxas info : Used 2 registers, 36 bytes smem 

을 그리고 결과 PTX는 다음과 같습니다 :

.entry _Z3fooPfS_j (
      .param .u64 __cudaparm__Z3fooPfS_j_out, 
      .param .u64 __cudaparm__Z3fooPfS_j_in, 
      .param .u32 __cudaparm__Z3fooPfS_j_n) 
    { 
    .reg .u16 %rh<4>; 
    .reg .u32 %r<5>; 
    .reg .u64 %rd<8>; 
    .reg .f32 %f<5>; 
    .loc 15 17 0 
$LDWbegin__Z3fooPfS_j: 
    .loc 15 21 0 
    mov.u16  %rh1, %ctaid.x; 
    mov.u16  %rh2, %ntid.x; 
    mul.wide.u16 %r1, %rh1, %rh2; 
    cvt.u32.u16  %r2, %tid.x; 
    add.u32  %r3, %r2, %r1; 
    cvt.u64.u32  %rd1, %r3; 
    mul.wide.u32 %rd2, %r3, 4; 
    ld.param.u64 %rd3, [__cudaparm__Z3fooPfS_j_in]; 
    add.u64  %rd4, %rd3, %rd2; 
    ld.global.f32 %f1, [%rd4+0]; 
    mov.f32  %f2, 0f40a00000;  // 5 
    add.f32  %f3, %f1, %f2; 
    ld.param.u64 %rd5, [__cudaparm__Z3fooPfS_j_out]; 
    add.u64  %rd6, %rd5, %rd2; 
    st.global.f32 [%rd6+0], %f3; 
    .loc 15 22 0 
    exit; 
$LDWend__Z3fooPfS_j: 
    } // _Z3fooPfS_j 

이제는 이해할 수없는 몇 가지 사항이 있습니다.

• ptx 어셈블리에 따르면 4 + 5 + 8 + 5 = 22 레지스터가 사용됩니다. 그렇다면 컴파일 중에 왜 used 2 registers이라고 표시됩니까?
• 어셈블리를 살펴보면 threadId, blockId 등의 데이터 유형이 임을 깨달았습니다. 이것이 CUDA 사양에 정의되어 있습니까? 또는이 버전은 CUDA 드라이버 버전마다 다를 수 있습니다.
• 누군가 내게이 줄을 설명해 줄 수 있습니까? mul.wide.u16 %r1, %rh1, %rh2;? %r1이 u32 인 이유는 u32 대신 wide이 사용됩니까?
• 레지스터 이름은 어떻게 선택합니까? 내 꽃병에 나는 %r 부분을 이해하지만 h (null), d 부분을 이해하지 못합니다. 데이터 유형 길이에 따라 선택됩니까? 즉 : 16 비트의 경우 h, 32 비트의 경우 null, 64 비트의 경우 d?
• 커널의 마지막 2 줄을 out[idx] = in[idx];으로 바꾼다면 프로그램을 컴파일 할 때 3 개의 레지스터가 사용된다고 말합니다! 지금 더 많은 레지스터를 어떻게 사용할 수 있습니까?

내 테스트 커널이 배열 인덱스가 범위를 벗어 났는지 확인하지 않는다는 사실을 무시하십시오.

대단히 감사합니다.

Answer 1

PTX는 여러 GPU 아키텍처에서 이식 가능하도록 설계된 중간 언어입니다. 특정 아키텍처에 대해 컴파일러 구성 요소 PTXAS에 의해 최종 머신 코드 (SASS라고도 함)로 컴파일됩니다. nvcc 옵션 -Xptxas -v은 PTXAS가 기계어 코드에 사용 된 물리적 레지스터의 수를 포함하여 생성 된 기계 코드에 대한 다양한 통계를보고하게합니다. cuobjdump --dump-sass으로 분해하여 기계 코드를 검사 할 수 있습니다.

PTX 코드에서 사용되는 레지스터의 수는 가상 레지스터이기 때문에 아무런 의미가 없습니다. CUDA 컴파일러는 SSA 형식 (정적 단일 지정, http://en.wikipedia.org/wiki/Static_single_assignment_form 참조)으로 알려진 PTX 코드를 생성합니다. 이것은 기본적으로 작성된 각각의 새 결과에 새 레지스터가 지정되었음을 의미합니다.

지침 mul.wide은 PTX 사양에 설명되어 있으며 현재 버전 (3.1)은 http://docs.nvidia.com/cuda/parallel-thread-execution/index.html입니다. 예제 코드에서 접미사 .u16은 두 개의 부호없는 16 비트 수량을 곱하고 부호없는 32 비트 결과를 반환한다는 것을 의미합니다. 즉 소스 피연산자의 전체, 두 배 너비 곱을 계산합니다.

PTX의 가상 레지스터가 입력되었지만 유형과 관계없이 이름을 자유롭게 선택할 수 있습니다. CUDA 컴파일러는 내부 구현 산출물이므로 문서화되지 않은 특정 규칙을 따르는 것으로 나타났습니다. PTX 코드를 살펴보면 현재 생성 된 레지스터 이름이 유형 정보를 인코딩한다는 것을 알 수 있습니다. 이는 쉽게 디버깅 할 수 있습니다. p<num>은 술어에 사용되며 r<num>은 32 비트 정수, rd<num>은 64 비트 정수, f<num> 32 비트 플로트의 경우는 fd<num>, 64 비트 double의 경우는 fd<num>입니다. 이러한 가상 레지스터를 생성하는 PTX 코드의 .reg 지시문을 보면 쉽게 알 수 있습니다.