本地對象參數至CUDA內核

Question

我從this answer明白，參數給CUDA內核經由恆定存儲器傳遞（計算能力2.0或更高）和副本，如果更改它們被存儲爲本地副本在任一寄存器或在堆棧上。如果參數是一個對象，並且只有一些成員被內核修改，會發生什麼？整個對象是否必須存儲在本地，還是僅爲修改後的成員創建副本？

Answer 1

這是一個有趣的問題我以前沒有考慮，答案似乎是僅該結構的使用的部件被加載到寄存器（至少憑經驗根據恰好一個例子）。

考慮以下人爲的例子：

struct parameters 
{ 
    float w,x,y,z; 
    int a,b,c,d; 
}; 

__global__ 
void kernel(float *in, float *out, parameters p) 
{ 
    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    float val_in = in[tid]; 

    p.b += 10; 
    p.w *= 2.0f; 
    p.z /= 5.0f; 

    out[tid] = (p.b>0) ? (p.w*val_in) : (p.z*val_in); 
} 

如果編譯器只加載所需的參數，我們應該只看到p 3對32位參數負荷進行註冊。編譯器發射的有效PTX（Cuda的5.0版編譯器sm_30）看起來是這樣的：

// 
// Generated by NVIDIA NVVM Compiler 
// Compiler built on Sat Sep 22 02:35:14 2012 (1348274114) 
// Cuda compilation tools, release 5.0, V0.2.1221 
// 

.version 3.1 
.target sm_30 
.address_size 64 

    .file 1 "/tmp/tmpxft_00000b1a_00000000-9_parameters.cpp3.i" 
    .file 2 "/home/talonmies/parameters.cu" 
    .file 3 "/opt/cuda-5.0/bin/../include/device_functions.h" 

.visible .entry _Z6kernelPfS_10parameters(
    .param .u64 _Z6kernelPfS_10parameters_param_0, 
    .param .u64 _Z6kernelPfS_10parameters_param_1, 
    .param .align 4 .b8 _Z6kernelPfS_10parameters_param_2[32] 
) 
{ 
    .reg .pred %p<2>; 
    .reg .s32 %r<9>; 
    .reg .f32 %f<8>; 
    .reg .s64 %rd<8>; 


    ld.param.u64 %rd1, [_Z6kernelPfS_10parameters_param_0]; 
    ld.param.u64 %rd2, [_Z6kernelPfS_10parameters_param_1]; 
    ld.param.f32 %f1, [_Z6kernelPfS_10parameters_param_2+12]; 
    ld.param.f32 %f2, [_Z6kernelPfS_10parameters_param_2]; 
    ld.param.u32 %r1, [_Z6kernelPfS_10parameters_param_2+20]; 
    cvta.to.global.u64 %rd3, %rd2; 

///home/talonmies/parameters.cu:11  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    .loc 2 11 1 
    mov.u32  %r2, %ntid.x; 
    mov.u32  %r3, %ctaid.x; 
    mov.u32  %r4, %tid.x; 
    mad.lo.s32 %r5, %r2, %r3, %r4; 
    cvta.to.global.u64 %rd4, %rd1; 

///home/talonmies/parameters.cu:12  float val_in = in[tid]; 
    .loc 2 12 1 
    mul.wide.u32 %rd5, %r5, 4; 
    add.s64  %rd6, %rd4, %rd5; 

///home/talonmies/parameters.cu:14  p.b += 10; 
    .loc 2 14 1 
    add.s32  %r6, %r1, 10; 

///home/talonmies/parameters.cu:15  p.w *= 2.0f; 
    .loc 2 15 1 
    add.f32  %f3, %f2, %f2; 

///opt/cuda-5.0/bin/../include/device_functions.h:2399 return a/b; 
    .loc 3 2399 3 
    div.rn.f32 %f4, %f1, 0f40A00000; 

///home/talonmies/parameters.cu:18  out[tid] = (p.b>0) ? (p.w*val_in) : (p.z*val_in); 
    .loc 2 18 1 
    setp.gt.s32  %p1, %r6, 0; 
    selp.f32 %f5, %f3, %f4, %p1; 

///home/talonmies/parameters.cu:12  float val_in = in[tid]; 
    .loc 2 12 1 
    ld.global.f32 %f6, [%rd6]; 

///home/talonmies/parameters.cu:18  out[tid] = (p.b>0) ? (p.w*val_in) : (p.z*val_in); 
    .loc 2 18 1 
    mul.f32  %f7, %f5, %f6; 
    add.s64  %rd7, %rd3, %rd5; 
    st.global.f32 [%rd7], %f7; 

///home/talonmies/parameters.cu:19 } 
    .loc 2 19 2 
    ret; 
} 

你可以看到，只有_Z6kernelPfS_10parameters_param_2（這是p.w）_Z6kernelPfS_10parameters_param_2+12（這是p.z）和_Z6kernelPfS_10parameters_param_2+20（這是p.b）被加載到寄存器。內核的其他成員永遠不會被加載。