【问题标题】:CUDA kernel do I need to put const in all pass-by-value parameters?CUDA 内核是否需要将 const 放入所有按值传递的参数中?
【发布时间】:2020-11-26 03:50:31
【问题描述】:
__global__
void diff2D(int top, int left, (const int windowSize or int windowSize), float* a1, float* a2)
{
    (const int mid or int mid) = windowsize / 2;
    ... // templateSize, windowSize are all const. 
} 

我的问题是我是否需要显式使用“const int windowSize”关键字来使编译器将mid 放入常量内存/缓存中?由于编码标准,我不想将“const”放在按值传递的原始类型前面。

常量参数被放入常量缓存中。但我不知道内核内部评估的常量是否也会放在那里。

【问题讨论】:

    标签: c++ cuda constants compiler-optimization gpu-constant-memory


    【解决方案1】:

    任何类型的所有内核参数都通过常量内存传递。见here。参数的类型或修饰与此功能无关。

    除非您使用适当的__constant__ 装饰/用法,否则不会自动将任何类型的其他常量放入常量内存中。

    内核参数之外,指定:

    const int mid = windowsize/2;
    

    不会导致任何东西被放置在常量内存中。

    考虑常量内存的一种方式是它在全局范围内。这对于内核参数的使用(所有线程都可以访问这些参数)以及__constant__ 声明(所有线程都可以访问这些参数)都是有意义的。

    当你声明一个普通的局部变量时:

    const int mid = windowsize/2;
    

    不在全局范围内,它具有逻辑线程本地范围。

    对此的逻辑扩展是您不能执行以下操作:

    __constant__ int mid = windowsize/2;
    

    在您的内核代码中,突然在__constant__ 内存中声明一个变量。这些声明必须在全局/文件范围内进行。

    【讨论】:

    • 所以如果我想把“mid”放入常量内存而不是寄存器,我必须把它作为参数传递,对吗?
    • @DukeLe 你想减少寄存器的使用吗?
    • 是的。我不知道将常用变量(如循环次数、过滤器高度/宽度等)放入常量内存(而不是寄存器)是否值得考虑。
    • 说“代替注册”是不明智的。在大多数情况下,GPU SASS 指令从寄存器中提取操作数。在你对它做任何有用的事情之前,内存中的数据会被移动到一个寄存器中。实际上,对于直接从常量内存中提取的操作数或直接从指令流中提取的常量操作数,实际上有一些例外,但是 IMO 尝试关注这一点是不明智的。您基本上无法控制在常量内存中开始生命的操作数最终是否会使用寄存器。很多时候,它会。
    猜你喜欢
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 2012-12-16
    • 2019-12-07
    • 2013-08-09
    • 1970-01-01
    • 2014-04-13
    • 1970-01-01
    相关资源
    最近更新 更多