请问一个关于__launch_bounds__() 问题

加速计算 CUDA
1

我用的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函数上，请问怎么办？？

2

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

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

3

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

4

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

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

5

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.
}

6

"VS会提示有错？”

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

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

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

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

7

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

8

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

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

9

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

10

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

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

欢迎再归来。

11

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

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

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

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

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

12

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

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

13

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