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

2022-01-10 00:00:00 cuda c++

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

Lets say I want a CUDA kernel that needs to do lots of stuff, but there are dome parameters that are constant to all the kernels. this arguments are passed to the main program as an input, so they can not be defined in a #DEFINE.

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

The kernel will run multiple times (around 65K) and it needs those parameters (and some other inputs) to do its maths.

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

My question is: whats the fastest (or else, the most elegant) way of passing these constants to the kernels?

常量是 2 或 3 个元素长度的 float* 或 int* 数组.它们将是其中的 5~10 个.

The constants are 2 or 3 element length float* or int* arrays. They will be around 5~10 of these.

玩具示例: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，全球?哪一个?

or maybe declare them in some global read only memory and let the kernels read from there? If so, L1, L2, global? Which one?

有没有更好的方法我不知道?

Is there a better way I don't know of?

在 Tesla K40 上运行.

Running on a Tesla K40.

推荐答案

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

Just pass them by value. The compiler will automagically put them in the optimal place to facilitate cached broadcast to all threads in each block - either shared memory in compute capability 1.x devices, or constant memory/constant cache in compute capability >= 2.0 devices.

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

For example, if you had a long list of arguments to pass to the kernel, a struct passed by value is a clean way to go:

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

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

[标准免责声明:用浏览器编写，不是真实代码，警告购买者]

[standard disclaimer: written in browser, not real code, caveat emptor]

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

If it turned out one of your magicint actually only took one of a small number of values which you know beforehand, then templating is an extremely powerful tool:

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

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

编译器足够聪明，可以识别 magconstant 在编译时知道循环行程，并会自动为您展开循环.模板是一种非常强大的技术，用于构建快速、灵活的代码库，如果您没有使用它，建议您习惯使用它'还没有这样做.

The compiler is smart enough to recognise magconstant makes the loop trip known at compile time and will automatically unroll the loop for you. Templating is a very powerful technique for building fast, flexible codebases and you would be well advised to accustom yourself with it if you haven't already done so.

相关文章