[attach]3026[/attach]
某程序 计算 1616 16网格情况, 比如计算5000步,可以计算完成。
计算3232 32网格,计算到2000步时候出现了如上图所示的错误提示。
请求建议如何添加检查的程序代码和位置以及如何检查。
自己推测可能是在算BICGS前出差。
PS:
(1)自己曾经添加过
if(cudaMallocHost((void**)&device_tem, dofsizeof(double))!=cudaSuccess){
printf(“error:paged host memory alloc failed!”);
算到32 32*32 1000多步时候报这个错。
然后我把cudaMallocHost全部注销掉,但是仍然出现了 附图的错误。
(2)同意程序,在另个电脑算3232 32算400步就不能算了
PPS:羽毛球去之前,开始计算程序,本想回来时候可以看到一个圆满结局,谁知负心的程序让我心碎,心痛,流血不止!!!
在线等大神回复!
LZ您好,手机看到您的帖子。
建议您还是先休息吧,身体才是革 命的本钱,不要干亏本的事情。
我和其他版主明天一定会认真研究和回复您的帖子的,请您放心。
另外我手机看不清楚您的截图,就先不下什么结论了。
您如果还不放心,请尽量详细地描述您的问题即可。
祝您晚安。
好的!
谢谢ice版大 每次都这么贴心的回复。
我听您的~ 先早点睡了。
要是我帖子里面有不明确的,麻烦明天指出,我一定按照您提的要求尽量详细描述。
谢谢~
也祝您晚安!
楼主这个截图很奇怪。如果楼主确定此图无PS, 那么,
我说下我根据您提供的已有信息进行一些推测:
楼主的异常是个managed exception, 是托管环境下会出现的。所以您可能是在C++/CLI下,或者您的代码是在调用一个现成的库的(而该库使用了.net ),或者您干脆是在C#里面通过P/Invoke调用封装好的基于cuda的dll的。这是对运行环境的分析。
然后下面分析下什么原因到导致这个问题。
您途中的异常表明一次对host memory(注意不是device memory)的非法访问被捕获了,而stack trace居然表明了是在C Runtime的初始化部分!而此时您的代码已经执行好久了。所以这个很不合理,那么最大的可能就是stack trace信息已经被破坏。
那么最大的可能是您的代码(或者您的代码调用的代码)在内存上进行了一些无效的操作:
例如:
(1)没有检测是否内存分配成功就黑上使用(特别是有忘记释放的之前分配过的内存,导致下次分配失败)。
(2)您对一些对局部的数组(stack上),访问越界,导致写入了无效的数据,冲垮了里面保存的返回地址,导致在一些您的函数结束的,程序跑飞了,飞到了mainCRTStartup()地址范围里的随机处,导致非法访问。
从您的配图看,可能是以上1和2的原因的概率很大,特别是2, 非常大。建议您先检查您的host code里对局部数组、变量使用是否问题。然后再考虑1这个因素。
这是我在根据您已有的信息下给出的建议。并进行了可能性排序。
然后如果您都不是,而是在正常的<<<>>>的GPU kernel执行期间出现此问题,那么以上分析全部作废!然后我表示我在已有的信息情况下无法提供更多建议,需要您提供更多信息。
夜安。祝好!
横扫斑竹深夜详细回帖,令人感动!
我没什么异议,不过就顶楼LZ的一些说法稍微写点。
1:LZ在顶楼的PS.1里面写曾经在cudaMallocHost()这里检查了返回值,并发现在1000多步的时候报错,然后把cudaMallocHost()注销掉了,跑出截图的错误。
以及另外一台电脑400多步就不能运行了。
请问,您每次运行一步都要进行一次cudaMallocHost()么?如果是的话,您是一直申请而没有释放么?如果还是的话,那么host端的paged memory耗尽的时候就会申请就会报错,如同横扫斑竹在第一条中指出的那样。而且那个400步就报错的机器也许本身host端内存就少。
另外,您说1000多步的时候就报错了,此时您没有修正该问题,而只是注释掉了cudaMallocHost()。如果您只是注释掉了错误检查,那么问题应该还是存在的。
如果您是注释掉了cudaMallocHost(),那么不申请空间,您的程序是如何运行的?
请LZ反馈一下这里的情况,方便的话,也可以贴出代码段。
2:稍微说一下检查Error的问题。
您可以先cudaDeviceSynchronize();一下,然后printf(“%s”,cudaGetErrorString(cudaGetLastError()));
您可以将这两行自行封装为Jeffrey_test_CUDA_Error()函数,以便使用。
因为如kernel 启动等,Error不是立即返回的,所以前面加上了同步。
如果您在程序某处加上此判断,并得到cudaSuccess的话,那说明此处之前都是OK的。
您可以用此方法检查下您的代码,看哪里开始报错的。
祝您好运~
首先向横扫版大表示抱歉,临时出了趟短差,刚刚才回来,没有及时回复见谅。
其次,感谢您细致的回复。
最后,回复下您提出的疑问。(部分疑问我没法直接回答,只能从测试的情况去如实叙述)
(1)我的确没有检查内存分配成功与否,就黑上使用。但是我每一步都给予初值,且printf看了值,所以关于分配内存的应该没有问题。
对于您提到没有释放的,您的推测非常准确,经过我仔细排查,的却有6个host端局部数组没有释放,现在已经改过来了。
(2)关于您提到局部数组(stack)访存越界的情况,您能具体说下嘛?
不太明白此处如何排查。
我又更新了下图片,可以看到错误一般位于BICGS里面。但是此处里面局部显存和局部内存我都释放了。
内核的话也运行那么多次,应该也没啥错误。
目前测试情况的描述:
(一)1616 16网格: 计算15000步,过程无错,结果正确
(二)3232 32网格:
[attach]3034[/attach]
上图的错误是在笔记本电脑 有一个局部内存忘记释放的结果。
我按照横版大说的释放后,笔记本电脑结果现在目前运行到2030(截止到发此回复时候)尚未报错。
但是奇怪的是 我的台式机电脑, 当1000步(每10步输出)时候 运行到1000步没有报错;但是当10000(每100步输出)就报错了。但是我觉得这个步长因素不会影响结果啊。我测试了3遍,都是如此,甚为诡异。
综3232 32:一是电脑不同,结果不同;二是步长这个看似不影响的貌似影响,待我进一步考证
(三)6464 64网格:运行27步,出错。
但是这里出错原因应该是 求解矩阵的值没有收敛导致(flag=1 代表未收敛),与3232 32错误不一样。
这里PS下:在未释放那个内存前,6464 64计算5步就错了,错误是求解不出来;释放内存后,结果就是目前这个。
[attach]3035[/attach]
请版大进一步指导!
没有能更进一步的解释了。我来简单说一下你的一个疑问:
如下代码可以轻易跑飞:
void lulu(int *p)
{
int buffer[10];
for (int i = 0; i < 50; i++) buffer[ i ][i] = p[i][ i ];
for (int i = 0; i < 50; i++) p[i][ i ] = buffer[50 - i];
return; //Release下,飞了!(debug下可能会检测到错误。)
}
因为这里有个内存布局:
(低地址) buffer[…], 可能的保存的RBP(EBP), 返回地址 (高地址)
而上述代码将返回地址覆盖了。导致了不可预测的返回地址。所以返回到了未知地址。
没有更多的我能进行建议的了。[/i][/i][/i]
横扫斑竹深夜详细回帖,令人感动!
我没什么异议,不过就顶楼LZ的一些说法稍微写点。
同样向ice版大道个歉,因个人问题回复您回复晚了,见谅!
谢谢您的解答和提问!
回答您的疑问:
(1)我的cudaMallocHost 是分配一个数组用来测试用的,所以注销掉不影响整个程序运行。其他主机端数组都是用 double * mm; 这类进行分配
(2)经过我仔细排查,途中出现一次page host报错,是因为那时候我忘了free page host时候,用来cudaFree, 应该用cudaFreeHost。目前已经排查。
(3)代码的却非常长,所以不便给出。这是我写一个最可能出错的一个函数的布置形式。
最有可能出错是 计算possion方程的子函数。
Possion(double *,………………)
{
……局部显存分配d_A d_b等
<<<<内核一>>>:形成d_b的值
<<<<内核二>>>:形成d_A的值
调用子函数 bicgs( **, **, **,); -------此步返回要求的X的值、
释放d_A d_b。
}
其中bicgs子函数里面有局部显存分配、内存分配,并且会调用 很多内核(例如点乘、矩阵向量相乘等)用于求解矩阵x的值。
另外,向向您进一步请教,您说的例如Jeffrey_test_CUDA_Error() 函数。
您看加载Possion什么位置好?我觉得只能加很多个了。
另外因为 调用的 bicgs里面也用到多个内核。
所以我估计如果错,bicgs里面概率也很大。
最后请您结合我回横扫版大的帖子一起给下一步指点!
顺祝春祺!!
system
2013 年3 月 30 日 16:24
10
恩呢! 谢谢进一步解读。
我再仔细测试看看。如果有进一步进展也向您随时汇报!
安~
system
2013 年3 月 30 日 16:28
11
既然楼主能确定这里概率较大,那么:
最有可能出错是 计算possion方程的子函数。
Possion(double *,………………)
{
……局部显存分配d_A d_b等 //在分配后,和使用之前,有无host上
<<<<内核一>>>:形成d_b的值 //的code对d_A和d_B的赋值初始化之类?如有建议排查
<<<<内核二>>>:形成d_A的值
调用子函数 bicgs( **, **, **,); -------此步返回要求的X的值、
释放d_A d_b。
}
//如果没有,则请无视此建议。
system
2013 年3 月 30 日 16:59
12
好的!
按照您的说的继续测试。
回复下您提的几点:
(1) 在在分配后,和使用之前possion里面没有使用host, 但是在bicgs子函数里面有使用了host,并且host也是在bicgs子函数里面进行释放的
(2)d_A d_b均初始化为0 cudaMemset,但是没有检验分配成功,只是初始化了
(3)您指的返回x的值是返回 主机端x的值吗? 还是? 此处d_x是调用的行参
如下: res_norm_ref = BICGS(n_possion, dof,d_A,d_col,d_x,d_b);
我用返回残差看是否计算正确。
system
2013 年3 月 31 日 02:50
13
LZ您好,根据您的反馈,(1)和(2)已经排查。
(3)的说法横扫斑竹也给出了建议。
我再对(3)和其他内容稍加补充。
按照您的注释,内核一和内核二生成了d_A,d_b,然后一个包含许多kernel和CPU代码的函数bicgs()进行计算。我觉得检查一下bicgs()里面的数据依赖性,在循环的时候有没有累积误差造成的问题,有没有越界和泄露什么的,或许有助于解决问题。
至于查错函数放置的位置,原则上你放在A位置,那么能检查A位置之前的GPU操作成功与否。您可以放置在若干个地方,然后逐渐缩小出问题的范围。不过考虑到你这个问题可能会在循环若干次之后才出现,实际排查还是比较麻烦的。
其他没有什么合适的建议了,祝您好运!
system
2013 年3 月 31 日 10:45
14
好的 我先按照您和横大的意见进行排查!
如果新的进展也将向二位汇报!
谢谢鼓励!
按照ice的建议,从头开始检查。
测试如下:
main()
{
函数1(); //// 里面有一些内核
cudaDeviceSynchronize();
printf(“%s”,cudaGetErrorString(cudaGetLastError()));
函数2();
函数3();
}
测试结果是:(测试的位置是计算量函数1后面的)
(1)cudaGetErrorString的 printf 出来的结果是 ”no error“
(2)采用断点调试,,在函数一 内部加入多个断点,过程中,没有显示任何异常。
但是一结束函数一 后的断点,就出现
“.exe 中的 0x750dc41f 处最可能的异常: Microsoft C++ 异常: 内存位置 0x0053eb70 处的 cudaError_enum。”
我想问的是:
(1)这两个测试的结果是否矛盾? 一个 no error 一个 cudaError_enum。
(2)整个函数一断点调试中 均不出现 cudaError_enum。 但函数一断点结束后会出现。感觉很奇怪。
.exe 中的 0x750dc41f 处最可能的异常: Microsoft C++ 异常: 内存位置 0x0053eb70 处的 cudaError_enum。
请无视任何中文版VS的“最可能的异常”的字样,那其实只是"First chance exception", 而且可能已经被CUDA的一些api函数内部捕获并处理并作为错误代码返回了。
所以,只要你听从ICE建议,总是检查返回值,那么你可以无视这个提示。
但是你这个位置的返回值检查,真的100%能排除所有情况么?不能。
然后我继续说下。
如果你确定此2行:
cudaDeviceSynchronize();
printf(“%s”,cudaGetErrorString(cudaGetLastError()));
执行完毕后是"no error", 那么只能证明你没有kernel执行的问题(异步错误)。
但这不能排除同步调用cudaMalloc()之类的一定是成功的。
因为有的时候,cudaMalloc()的失败不会继续让下文的AP调用继续返回错误。(但是kernel之类的持续会)
所以这大致能排除你的kernel问题,以及,不能排除cudaMalloc之类的无问题,你需要在“函数1"里面的每次cudaMalloc都检查下返回值。
建议您如下操作:
每个cudaMalloc(…)都改成:
cudaError_t r = cudaMalloc(…);
if (r != cudaSucess) __debugbreak(); //你需要#include <intrin.h>
或者建议您改成:
cudaError_t r = cudaMalloc(…);
assert(r == cudaSuccess); //需要#include <assert.h>
然后一旦中途cudaMalloc()之类的有问题,您会被立刻打断执行的, 然后弹出调试器对话框的。
建议尝试次建议。以及,您可以无视此建议。
那么楼主,实际上上述说法有点不恰当,需要这么说:
(1)如果您每次错误都是用的cudaGetLastError()检测的。那么不用怕漏掉对cudaMalloc的检测,它会反应到下一次。
(2)如果您每次都是用的返回值检测的错误, (包括异步的kernel用的cudaErrot_t r = cudaDeviceSynchronize(); 来实现返回值风格的取回错误),那么您可能会漏掉对cudaMalloc的之前的错误的检测。
l那么,楼主如果总是用的cudaGetLastError,那么无需考虑单独错误检测, 可以忽视上文。反之,则需要。
楼主上文实际用的是前者,所以我多虑了。很安全。可以不考虑我上个楼的建议。
特此说明。
LZ您好,经过今天的试验和讨论,我前面说的cudaGetLastError()的用法有些不很完善,实际情况要复杂很多,如果对您造成了一定误导,还望见谅。
实际情况大致是这样的:
cudaGetLastError()实际上有一个缓冲机制,缓冲了前面的函数返回值,但是这个缓冲区数量有限,而且似乎没有官方对其行为的说明,所以有些混乱。
经过我们测试和讨论,大致是这样的:
cudaGetLastError()能缓冲一次cuda错误,如果后面都是成功的,那么这次错误一直都在,除非使用cudaGetLastError()输出该错误,并同时清掉缓冲。如果后面也有错误,那么前面的会被冲掉。
同时,cudaMalloc()这种函数即使失败了,并不会导致后面的kernel ,cudaMalloc等挂掉,而只是返回一个错误值(我们还可以用这个错误值来错误处理)。如果是kernel或者异步操作挂掉了,那么后面永远是报错的。(这一点横扫斑竹也说过)
那么,继续看一下下面几种情况:
1:如果我们用了多个cudaGetLastError(),最后一个是no error,那说明前面无问题么?
不一定,前面可能有cudaMalloc()的问题,记录在了缓冲区中,但是被更前面的cudaGetLastError()输出了。
2:如果我们看到cudaGetLastError()的结果是kernel挂掉(一般是 unknow error),那么前面就没有其他问题了么?
不一定,即使前面有cudaMalloc()的问题,而且没有输出,此时都会被强制刷为kernel挂掉的错误。
3:如果我们看到cudaGetLastError()的结果是cudaMalloc()等函数的错误,那么前面就只有一次错误么?
不一定,前面可能多次此类错误,但缓冲区只保留了最后一个。
4:如果cudaGetLastError()报某个错误,一定是最近的那个函数的问题么?
不一定,也可以,失败——成功——成功…——cudaGetLastError(),此时是最前面那个失败的错误。所以你会看到最近那个函数的返回值直接检测是 no error,而cudaGetLastError()报另外一个错误,看上去非常矛盾。
总之,似乎情况有些复杂了,只在一些特殊情况下能快速识别,比如,只用了一次cudaGetLastError(),并得到no error,这说明前面都是好的;或者如果得到unkonw error,这说明前面基本上一定有个kernel挂掉了。其他详细调试查错,似乎离不开密集地检测返回值并仔细检查了,这确实是个体力活。
最后,谨祝您身体健康,体力充沛,以便在和BUG做斗争中取得最后的胜利!
谢谢ice版大给予这么多的测试和结论。
体力活没事,只要能找到bug,哪怕一行一行找我都愿意!
我根据2位版大的建议,决定用 只用一次 cudaGetLastError() 来试一试。
这里我有2个小小疑问,劳烦ice版大帮我明确下:
(1)如果只用一次cudaGetLastError() ,且检测位置放在cuda分配和HostToDevice传递之后, 是否可以检测到错误?
(2)“.exe 中的 0x750dc41f 处最可能的异常: Microsoft C++ 异常: 内存位置 0x0053eb70 处的 cudaError_enum。” 这个真的不用考虑吗?
(这里解释下,这里不是怀疑横大解释的,只是自己检测的的错误的恰巧是在这个位置得到的数的根源所在(通过一点点的断点发现的。))
3ks!顺祝版大马上的假期愉快!
(3)