到 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 加载到寄存器的三个 32 位参数。编译器（用于 sm_30 的 Cuda 5.0 版本编译器）生成的 PTX 如下所示：

//
// 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）被加载到寄存器中。该结构的其他成员永远不会被内核加载。