cuda - ví dụ tối thiểu, cách sử dụng đăng ký cao

Question

Hãy xem xét 3 hạt nhỏ, tầm thường này. Việc sử dụng đăng ký của họ là nhiều hơn cao hơn tôi mong đợi. Tại sao?

A:

__global__ void Kernel_A() 
{ 
//empty 
} 

PTX tương ứng:

ptxas info : Compiling entry function '_Z8Kernel_Av' for 'sm_20' 
ptxas info : Function properties for _Z8Kernel_Av 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 2 registers, 32 bytes cmem[0] 

B:

template<uchar effective_bank_width> 
__global__ void Kernel_B() 
{ 
//empty 
} 

template 
__global__ void Kernel_B<1>(); 

PTX tương ứng:

ptxas info : Compiling entry function '_Z8Kernel_BILh1EEvv' for 'sm_20' 
ptxas info : Function properties for _Z8Kernel_BILh1EEvv 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 2 registers, 32 bytes cmem[0] 

C:

template<uchar my_val> 
__global__ void Kernel_C 
     (uchar *const device_prt_in, 
     uchar *const device_prt_out) 
{ 
//empty 
} 

PTX tương ứng:

ptxas info : Compiling entry function '_Z35 Kernel_CILh1EEvPhS0_' for 'sm_20' 
ptxas info : Function properties for _Z35 Kernel_CILh1EEvPhS0_ 
    16 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 10 registers, 48 bytes cmem[0] 

---

Câu hỏi:

Tại sao trống hạt nhân A và B sử dụng 2 thanh ghi? CUDA luôn sử dụng một thanh ghi tiềm ẩn, nhưng tại sao lại sử dụng thêm 2 thanh ghi rõ ràng?

Hạt nhân C thậm chí còn bực bội hơn. 10 thanh ghi? Nhưng chỉ có 2 con trỏ. Điều này cho phép 2 * 2 = 4 thanh ghi cho con trỏ. Ngay cả khi có thêm 2 thanh ghi bí ẩn (được gợi ý bởi Kernel A và Kernel B), điều này sẽ cho tổng cộng 6. Vẫn còn ít hơn 10!

---

Trong trường hợp bạn quan tâm, đây là mã ptx cho Kernel A. Mã ptx cho Kernel B là giống hệt nhau, modulo các giá trị số nguyên và tên biến.

.visible .entry _Z8Kernel_Av( 
) 
{   
     .loc 5 19 1 
func_begin0: 
     .loc 5 19 0 

     .loc 5 19 1 

func_exec_begin0: 
     .loc 5 22 2 
     ret; 
tmp0: 
func_end0: 
} 

Và đối với Kernel C ...

.weak .entry _Z35Kernel_CILh1EEvPhS0_(
     .param .u64 _Z35Kernel_CILh1EEvPhS0__param_0, 
     .param .u64 _Z35Kernel_CILh1EEvPhS0__param_1 
) 
{ 
     .local .align 8 .b8  __local_depot2[16]; 
     .reg .b64  %SP; 
     .reg .b64  %SPL; 
     .reg .s64  %rd<3>; 


     .loc 5 38 1 
func_begin2: 
     .loc 5 38 0 

     .loc 5 38 1 

     mov.u64   %SPL, __local_depot2; 
     cvta.local.u64 %SP, %SPL; 
     ld.param.u64 %rd1, [_Z35Kernel_CILh1EEvPhS0__param_0]; 
     ld.param.u64 %rd2, [_Z35Kernel_CILh1EEvPhS0__param_1]; 
     st.u64 [%SP+0], %rd1; 
     st.u64 [%SP+8], %rd2; 
func_exec_begin2: 
     .loc 5 836 2 
tmp2: 
     ret; 
tmp3: 
func_end2: 
} 

1. Tại sao nó đầu tiên khai báo một biến cục bộ nhớ (.local)?
2. Tại sao hai con trỏ (được cho là đối số hàm) được lưu trữ trong sổ đăng ký? Không có một không gian đặc biệt cho họ?
3. Có lẽ hai con trỏ đối số hàm thuộc về thanh ghi - giải thích hai dòng .reg .b64. Nhưng dòng .reg .s64 là gì? Tại sao nó lại ở đó?

---

Nó trở nên tệ hơn vẫn:

D:

template<uchar my_val> 
__global__ void Kernel_D 
     (uchar * device_prt_in, 
     uchar *const device_prt_out) 
{ 
    device_prt_in = device_prt_in + blockIdx.x*blockDim.x + threadIdx.x; 
} 

cho

ptxas info : Used 6 registers, 48 bytes cmem[0] 

Vì vậy, thao tác lập luận (một con trỏ) giảm 10-6 sổ đăng ký?

Answer 1

Điểm đầu tiên cần thực hiện là nếu bạn lo lắng về sổ đăng ký, đừng nhìn vào mã PTX, vì nó sẽ không cho bạn biết bất cứ điều gì. PTX sử dụng dạng gán đơn tĩnh và mã được trình biên dịch phát ra không bao gồm bất kỳ "trang trí" nào được yêu cầu để tạo một điểm nhập mã máy có thể chạy được.

Với điều đó ra khỏi con đường, chúng ta hãy nhìn vào hạt nhân A:

$ nvcc -arch=sm_20 -m64 -cubin -Xptxas=-v null.cu 
ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z8Kernel_Av' for 'sm_20' 
ptxas info : Function properties for _Z8Kernel_Av 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 2 registers, 32 bytes cmem[0] 

$ cuobjdump -sass null.cubin 

    code for sm_20 
     Function : _Z8Kernel_Av 
    /*0000*/  /*0x00005de428004404*/  MOV R1, c [0x1] [0x100]; 
    /*0008*/  /*0x00001de780000000*/  EXIT; 
     ............................. 

Có hai thanh ghi của bạn. Các hạt trống không tạo ra hướng dẫn bằng không.

Ngoài ra, tôi không thể sao chép những gì bạn đã hiển thị. Nếu tôi nhìn vào hạt nhân C của bạn như được đăng, tôi nhận được điều này (trình biên dịch phát hành CUDA 5):

$ nvcc -arch=sm_20 -m64 -cubin -Xptxas=-v null.cu 
ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z8Kernel_CILh1EEvPhS0_' for 'sm_20' 
ptxas info : Function properties for _Z8Kernel_CILh1EEvPhS0_ 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 2 registers, 48 bytes cmem[0] 


$ cuobjdump -sass null.cubin 

code for sm_20 
    Function : _Z8Kernel_CILh1EEvPhS0_ 
/*0000*/  /*0x00005de428004404*/  MOV R1, c [0x1] [0x100]; 
/*0008*/  /*0x00001de780000000*/  EXIT; 
    ........................................ 

ie. giống hệt 2 mã đăng ký cho hai nhân đầu tiên.

và tương tự cho Kernel D:

$ nvcc -arch=sm_20 -m64 -cubin -Xptxas=-v null.cu 
ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z8Kernel_DILh1EEvPhS0_' for 'sm_20' 
ptxas info : Function properties for _Z8Kernel_DILh1EEvPhS0_ 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 2 registers, 48 bytes cmem[0] 

$ cuobjdump -sass null.cubin 
code for sm_20 
    Function : _Z8Kernel_DILh1EEvPhS0_ 
/*0000*/  /*0x00005de428004404*/  MOV R1, c [0x1] [0x100]; 
/*0008*/  /*0x00001de780000000*/  EXIT; 
    ........................................ 

Một lần nữa, 2 thanh ghi.

Đối với hồ sơ, phiên bản nvcc Tôi đang sử dụng là:

$ nvcc --version 
nvcc: NVIDIA (R) Cuda compiler driver 
Copyright (c) 2005-2012 NVIDIA Corporation 
Built on Fri_Sep_28_16:10:16_PDT_2012 
Cuda compilation tools, release 5.0, V0.2.1221