CUDA - 커널이 예상보다 많은 레지스터를 사용합니까?

Question

나는 합계를 계산하는 커널을 가지고있다. 선언 된 변수의 수를 세는 커널을 살펴보면 커널 당 총 5 개의 레지스터가 있다고 가정합니다. 그러나 커널을 프로파일 링 할 때는 34 개의 레지스터가 사용됩니다. 내가 1024 스레드의 실행을 허용하기 위해 30 레지스터 내려야합니다.

무엇이 잘못되었는지 누가 볼 수 있습니까?

__global__ void sum_kernel(float* values, float bk_size, int start_idx, int end_idx, int resolution, float* avgs){ 

    // Allocate shared memory (assuming a maximum of 1024 threads). 
    __shared__ float sums[1024]; 

    // Boundary check. 
    if(blockIdx.x == 0){ 
     avgs[blockIdx.x] = values[start_idx]; 
     return; 
    } 
    else if(blockIdx.x == resolution-1) { 
     avgs[blockIdx.x] = values[start_idx+(end_idx-start_idx)-1]; 
     return; 
    } 
    else if(blockIdx.x > resolution -2){ 
     return; 
    } 

    // Iteration index calculation. 
    unsigned int idx_prev = floor((blockIdx.x + 0) * bk_size) + 1; 
    unsigned int from = idx_prev + threadIdx.x*(bk_size/blockDim.x); 
    unsigned int to = from + (bk_size/blockDim.x); 
    to = (to < (end_idx-start_idx))? to : (end_idx-start_idx); 

    // Partial average calculation using shared memory. 
    sums[threadIdx.x] = 0; 
    for (from; from < to; from++) 
    { 
     sums[threadIdx.x] += values[from+start_idx]; 
    } 

    __syncthreads(); 

    // Addition of partial sums. 
    if(threadIdx.x != 0) return; 
    from = 1; 
    for(from; from < 1024; from++) 
    { 
     sum += sums[from]; 
    } 
    avgs[blockIdx.x] = sum; 
} 

• 당 2 개 포인터 레지스터, 부호 INT 당 1 개 레지스터 상수 메모리에 저장된 인수를 가정.

Answer 1

선언 된 변수의 수로 사용 된 레지스터 수를 추정 할 수 없습니다. 컴파일러는 주소 계산을 위해 레지스터를 사용하거나 명시 적으로 등을 예를 들어

선언되지 않은 임시 변수를 저장하기 위해, 나는 다음을 가진, 즉 커널 함수의 첫 번째 부분,

__global__ void sum_kernel(float* values, float bk_size, int start_idx, int end_idx, int resolution, float* avgs){ 

    // Boundary check. 
    if(blockIdx.x == 0){ 
     avgs[blockIdx.x] = values[start_idx]; 
     return; 
    } 
    else if(blockIdx.x == resolution-1) { 
     avgs[blockIdx.x] = values[start_idx+(end_idx-start_idx)-1]; 
     return; 
    } 
    else if(blockIdx.x > resolution -2){ 
     return; 
    } 
} 

분해 한 결과

code for sm_20 
     Function : _Z10sum_kernelPffiiiS_ 
.headerflags @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)" 
/*0000*/  MOV R1, c[0x1][0x100];   /* 0x2800440400005de4 */ R1 = [0x1][0x100] 
/*0008*/  S2R R2, SR_CTAID.X;    /* 0x2c00000094009c04 */ R2 = BlockIdx.x 
/*0010*/  MOV R0, c[0x0][0x34];    /* 0x28004000d0001de4 */ R0 = [0x0][0x34] 
/*0018*/  ISETP.EQ.AND P0, PT, R2, RZ, PT; /* 0x190e0000fc21dc23 */ if (R2 == 0) 
/*0020*/ @P0 BRA 0x78;       /* 0x40000001400001e7 */ 
/*0028*/  MOV R0, c[0x0][0x30];    /* 0x28004000c0001de4 */ 
/*0030*/  IADD R0, R0, -0x1;    /* 0x4800fffffc001c03 */ 
/*0038*/  ISETP.NE.AND P0, PT, R2, R0, PT; /* 0x1a8e00000021dc23 */ 
/*0040*/ @P0 EXIT ;       /* 0x80000000000001e7 */ 
/*0048*/  MOV R0, c[0x0][0x2c];    /* 0x28004000b0001de4 */ 
/*0050*/  ISCADD R2, R2, c[0x0][0x34], 0x2; /* 0x40004000d0209c43 */ 
/*0058*/  ISCADD R0, R0, c[0x0][0x20], 0x2; /* 0x4000400080001c43 */ 
/*0060*/  LDU R0, [R0+-0x4];    /* 0x8bfffffff0001c85 */ 
/*0068*/  ST [R2], R0;      /* 0x9000000000201c85 */ 
/*0070*/  BRA 0x98;       /* 0x4000000080001de7 */ 
/*0078*/  MOV R2, c[0x0][0x28];    /* 0x28004000a0009de4 */ 
/*0080*/  ISCADD R2, R2, c[0x0][0x20], 0x2; /* 0x4000400080209c43 */ 
/*0088*/  LDU R2, [R2];      /* 0x8800000000209c85 */ R2 used for addressing and storing gmem data 
/*0090*/  ST [R0], R2;      /* 0x9000000000009c85 */ R0 used for addressing 
/*0098*/  EXIT ;       /* 0x8000000000001de7 */ 

위의 CUDA 코드 스 니펫에는 명시 적으로 선언 된 변수가 없습니다. 디스 어셈블 된 코드에서 볼 수 있듯이 컴파일러는 3 레지스터 (R0, R1 및 R2)를 사용했습니다. 이러한 레지스터는 기능상 상호 교환이 가능하며 상수, 메모리 주소 및 전역 메모리 값을 저장하는 데 사용됩니다.