请问一个关于__launch_bounds__() 问题

我用的VS2010, CUDA5.0, K20C显卡,compute35,sm35. 为了提供寄存器的使用率,我试着用__launch_bounds__()手动给各block分配寄存器。 我的代码如下
global void
launch_bounds(MAXTHREADSPERBLOCK, MINBLOCKSPERMP)
TryMove()

在加入__launch_bounds__()之前,代码没有任何问题,已经可以进行运算。但是加了这段specifier之后,TryMove函数就挂掉了。VS提示了很多错误, 首先__launch_bounds__的()中的两个int值会被提示没有定义过,然后TryMove函数中的所有已建结构都不可识别,cuda lib里的变量也不可识别。

我发现这个问题是由TryMove函数内调用的device函数引起的,如果我把被调用的device函数改成global函数,再给这个函数同样加入__launch_bounds__(),问题就没有了,但是我现在需要TryMove函数调用一个device函数,而__launch_bounds__()不能被加载到device函数上,请问怎么办??

楼主您好,这本来就是给__global__用的。并没有单独限制__device__的方式(因为无法对__device__的启动进行资源分配,您觉得呢?)

请不要使用任何手册没有说明的行为。:frowning:

这个我知道的,手册里时明确说了只针对global函数,但是就没有任何别的方法可以强制对device函数进行设定么?

(1)如果__device__被inline了,那么它是调用它的__global__的一部分,自然__global__的寄存器数目分配。
(2)如果__device__没有被inline, 那么此时block已经在SM上执行了,寄存器的分配情况已经固定了,自然此时的寄存器数目依然是__global__的分配情况。

综合1,2,我认为__device__无法单独设置寄存器数目,也没有必要设置。
供您参考。

Hi,首先,谢谢你的回复与帮助。
确实device没有被inline,我对它强制inline后问题就解决了。一开始我的想法和你建议我的第二条是一样的,但是貌似事实证明不太一样。因为实际情况是,为global分配2个block per MP 与device实际的1个block per MP冲突,所以建立时VS会提示有错。我的代码其实如下
device YYYYYY(){
}

global void _launch_bounds(512, 2) XXXX(){ set number of blocks per MP to 1 by force

YYYYYY(); <-----device function call, 1 block per MP default.
}

"VS会提示有错?”

我不认为代码:
device void YYYYYY()
{
}

global void launch_bounds(512, 2) XXXX()
{
YYYYYY();
}

有任何问题。它应该毫无疑问的被立刻成功编译。

你可以说明你的问题是什么。

当我没有对device做内联,而是只用__launch_bounds__()人为的将global函数的block数设为2,也就是每个thread有64个寄存器时,VS编译会提示出错,提示的意思就是说我在用一个64 registers per thread的function call调用一个70 registers per thread的函数,70就是CUDA默认分配给该函数的register数量。

“当我用forinline人为的将block数设为2,VS编译会提示出错”

__forceinline__并不能修改寄存器分配数目的,楼主您还是好好整理下您的问题吧。

以及,您可以考虑使用-maxrregcount=64, 这依然无法限制您的__device__的寄存器使用数目么?

楼主我建议您先尝试使用-maxrregcount来限制一下,看看是否有效果。

以及,您如果依然感兴趣,您可以说下您的“当我用forinline人为…"这个是什么问题。

欢迎再归来。

我猜测这样会解决问题,以及楼主说的”一开始我的想法和你建议我的第二条是一样的,但是貌似事实证明不太一样“应该是对的。

运行时候的资源分配的确,的确只会受__global__的考虑,而__device__只会使用调用它的__global__的寄存器的分配数目。

但的确可能您这种存在编译时刻__device__超过64个寄存器(例如您的70个), 而导致__global__限制为64个,从来无法调用,编译出错。

此时使用-maxrregcount=64应该可以搞定。

以及感谢楼主带来此例子(大__device__使用了70个寄存器,从而导致无法成功编译)。这种案例之前我没有经验,未能快速给出正确解决方案,表示歉意。

sorry 之前是说错了, 我的意思是说 当我没用forceinline 对device内联, 而是只对global函数加了
launch_bounds()时,VS会提示错误哦,提示内容的意思大概就是说我在用一个64 registers per thread的function call调用一个70 registers per thread的函数,这个调用不被允许。70就是CUDA默认分配给该函数的register数量。

打错了内容实在抱歉,看了一天代码头有点晕当时。

恩 我有考虑过-maxrregcount 但是如果加入这个compiler flag的话,貌似对全部的函数都进行了限制,为了实现最优化,我不想笼统的覆盖所有函数,所以内联会是个不错的方法。谢谢版主的这么多回复,真的给了我很大的帮助!