本地对象参数至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）被加载到寄存器。内核的其他成员永远不会被加载。