system
1
对于下面的函数
global void Test(int Length)
{
for(int i=0;i<COMPUTE_NUM_PER_THREAD;i++)
{
encode( Length);
}
}
调用方式:int length=10;Test<<<>>>(length);
global void Test()
{
int Length=10;
for(int i=0;i<COMPUTE_NUM_PER_THREAD;i++)
{
encode( Length);
}
}
调用方式:Test<<<>>>();
为什么调用这两个函数运行时间上会有较大差距(后者明显快于前者)?
system
2
后者为编译常量,编译器做了常量优化。不知道你的encode具体如何,我想跟这个参数是否常数关系非常大。
system
3
Length的值在encode中是不变的,我把__global__ void Test(int Length)
改成__global__ void Test(const int Length),时间没什么变化啊!
system
4
修饰符仅仅说明这个变量在函数内不被修改,但仍然是一个参数变量。
跟编译常量不是一个概念。
编译常量指的是编译时刻,函数已经没有参数。
system
6
楼主您好!看到您的发帖了。
我不认为您的(1)和(2)例子能有太大的差别。
在例子1中,Length作为参数传递给了kernel, kernel在使用的时候,在fermi和kepler上是通过constant cache访问的,可以直接传递给寄存器(1条指令), 也可以直接作为操作数(不占用任何额外指令, 如果被encode()被inline的话),延迟很低,带宽很大(所有线程的访问都是一致的)。
在例子2中,Length作为常数,此时可以直接依然是1条指令传递给寄存器,然后作为参数给您的encode()。或者,如果length被inline的话,甚至此常数可能可以直接作为立即数被编译到指令里(那就不占用额外指令了).
因为前者和后者只多了一次对constant cache的访问,但前文指出,这种是高效低延迟的,实际上可以认为是无影响(在fermi上). 加上从指令数看,实际上并无增加指令数。所以:
我不认为您的“后者明显快于前者”的现象有存在的可能。
对您有如下1条建议:
请给出您的全部测试代码,以便让论坛为您具体分析。
对您有如下1可选参考:
(如果您感兴趣,可以继续阅读,如果您不感兴趣,可以主要看看上文的说法和建议。)
下面将详细比较您的2个例子里面可能编译出来的指令的差异:
(1)例子有如下可能: mov Rx, c 或者 ldc Rx, c 或者 op …, c(作为操作数 )
(2)例子有如下可能: mov Rx, 10 或者 op …, 10 (作为立即数操作数).
您可以看到,最坏情况都是1条指令,最好情况都是不占用额外指令。所以在指令数上无差异。
然后您被告知:实际测试表明,无论c作为ldc,mov,还是普通指令的操作数和10作为立即数,实际中在测试表表现不出来。也就是您可以认为它和立即数一样的高效。所以这依然不是一个决定因素。
此可选参考解释了我做出了“ 我不认为您的(1)和(2)例子能有太大的差别。”的原因。
感谢您的来访,祝您工作愉快!