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

假设我想要一个需要做很多事情的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,全球?哪一个?

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

在特斯拉K40上运行.

解决方法

只需按价值传递它们.编译器将自动将它们放在最佳位置,以便于缓存广播到每个块中的所有线程 – 计算能力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
}

[标准免责声明:用浏览器编写,不是真正的代码,需要注意事项]

如果事实证明你的魔法之一实际上只采用了一些你事先知道的值,那么模板是一个非常强大的工具:

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,如果你还没有这样做,你最好习惯使用它.

相关文章

文章浏览阅读315次。之前用C语言编过链表,这几天突然想用C+...
文章浏览阅读219次。碰到问题就要记录下来,防止遗忘吧。文章...
文章浏览阅读1.8k次,点赞11次,收藏37次。因为自己对决策树...
文章浏览阅读492次。C++ 设计模式之策略模式
文章浏览阅读683次。我也算是个C++的小白,对于C++中的谓语我...
文章浏览阅读225次。又看了一遍操作符的东西,感觉之前对操作...