c++ - 将常量参数传递给 CUDA 内核的最快(或最优雅)方式

Question

假设我想要一个需要做很多事情的 CUDA 内核，但是有一些对所有内核都是不变的圆顶参数。此参数作为输入传递给主程序，因此它们不能在 #DEFINE 中定义。

内核将运行多次(大约 65K)并且它需要这些参数(和一些其他输入)来进行计算。

我的问题是:将这些常量传递给内核的最快(或者最优雅)的方法是什么？

常量是 2 或 3 个元素长度的 float* 或 int* 数组。他们大约有 5~10 个。

---

玩具示例:2 个常量 const1 和 const2

__global__ void kernelToyExample(int inputdata, ?????){
        value=inputdata*const1[0]+const2[1]/const1[2];
}

是不是更好了

__global__ void kernelToyExample(int inputdata, float* const1, float* const2){
        value=inputdata*const1[0]+const2[1]/const1[2];
}

或

__global__ void kernelToyExample(int inputdata, float const1x, float const1y, float const1z, float const2x, float const2y){
        value=inputdata*const1x+const2y/const1z;
}

或者在一些全局只读内存中声明它们并让内核从那里读取？如果是这样，L1，L2，全局？哪个？

还有我不知道的更好的方法吗？

在 Tesla K40 上运行。

Answer 1

只需按值传递它们。编译器会自动将它们放在最佳位置，以促进缓存广播到每个 block 中的所有线程 - 计算能力 1.x 设备中的共享内存，或计算能力 >= 2.0 设备中的常量内存/常量缓存。

例如，如果您有一长串要传递给内核的参数，按值传递的结构是一种干净的方法:

struct arglist {
    float magicfloat_1;
    float magicfloat_2;
    //......
    float magicfloat_19;
    int magicint1;
    //......
};

__global__ void kernel(...., const arglist args)
{
    // you get the idea
}

[标准免责声明:用浏览器编写，不是真正的代码，买者自负]

如果你的 magicint 中的一个实际上只取了你事先知道的少量值中的一个，那么模板是一个非常强大的工具:

template<int magiconstant1>
__global__ void kernel(....)
{
    for(int i=0; i < magconstant1; ++i) {
       // .....
    }
}

template kernel<3>(....);
template kernel<4>(....);
template kernel<5>(....);

编译器足够聪明，可以识别 magconstant 在编译时让循环行程已知，并会自动为您展开循环。模板是一个 very powerful technique用于构建快速、灵活的代码库，如果您还没有习惯，最好先习惯它。