将常量参数传递给CUDA内核的最快(或最简洁)方法 [英] Fastest (or most elegant) way of passing constant arguments to a CUDA kernel

问题描述

让我们说我想要一个需要做很多事情的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.

这篇关于将常量参数传递给CUDA内核的最快(或最简洁)方法的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！