我在kernel函数中静态声明了两个共享内存数组,如下:
shared int angle[128];
shared int minInd[128];
然后发现调试时,不管中的数值是多少,angle和mingIND的大小恒是64,真心奇怪。
我在kernel函数中静态声明了两个共享内存数组,如下:
shared int angle[128];
shared int minInd[128];
然后发现调试时,不管中的数值是多少,angle和mingIND的大小恒是64,真心奇怪。
angle和mingIND的大小恒是63
这个指的是 angle[100]的值是63么?
一般来说,您需要先对shared memory里面的元素赋值,然后才用的。
一般的步骤是:多个线程赋值——__syncthreads()——多个线程使用。
就是这两个数组的元素个数都恒是64,不管我怎么改变中的数字。
这是我设置断点时,鼠标悬停查看的,只显示0号元素到63号元素,但是我在watch中,可以查看angle[255]甚至angle[256],但到319就看不了了。需要说明的是,这次我声明angle的大小为256。
这也是让人纳闷的。
悬停查看不了完整的数组,本身设备分配的数组大小也比你声明的大,至于具体多大又跟资源有关?
这个是个常见现象。因为nsight菜单里的start cuda debugging会,无视你的VS的选项里的自动build设置。
所以源代码经常和你调试的目标代码不一致的。
下图则演示了这一点。[attach]2981[/attach]
在该图中,定义了__shared__ int q[555];
但是sizeof(q)的结果却报告了q的大小是400B, 也就是100个元素。
这是因为演示之前用[100]编译过,然后修改了,但却忘记重新build. 导致了不一致。
解决方案:
rebuild all, 并重新开始调试。
(注:本回复给出了最常见的场景,但并不排除其他版主/会员/NV原厂支持可能存在的其他解释)
楼主您好。刚才在回复您楼主位的帖子的时候,被您插楼了。导致未能及时发现在我回复的同时您有发贴了。对延迟的回复表示歉意。
(1)您为何只能看前64个元素。
解释:那是因为nsight默认如此。
解决方案:您可以修改nsight的max array elements expansion为更高数值(默认64)。
(2)为何您可以越界的看。
解释:
(a)您的理解是错误的。不是设备分配了更多(2.x的shared memory是128B的分配粒度的。你的256个float不会导致过多分配)
(b)您能看到的原因是因为CUDA C/C++语言不检查数组越界,您可以继续往下看,您下面的其他__shared__数组的内容被您一起观察到了。
感谢您对CUDAZone China长久以来的支持。
祝您调试愉快。
谢谢!你的回答圆满的解决了我的疑问。
您客气了。我和ICE能服务您感到了极大的荣幸。