[attach]3026[/attach]
某程序 计算 161616网格情况， 比如计算5000步，可以计算完成。
计算323232网格，计算到2000步时候出现了如上图所示的错误提示。
请求建议如何添加检查的程序代码和位置以及如何检查。
自己推测可能是在算BICGS前出差。
PS：
（1）自己曾经添加过
if(cudaMallocHost((void**)&device_tem, dofsizeof(double))!=cudaSuccess){
printf(“error:paged host memory alloc failed!”);
算到3232*32 1000多步时候报这个错。
然后我把cudaMallocHost全部注销掉，但是仍然出现了 附图的错误。
（2）同意程序，在另个电脑算323232算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里面。但是此处里面局部显存和局部内存我都释放了。
内核的话也运行那么多次，应该也没啥错误。
目前测试情况的描述：
（一）161616网格： 计算15000步，过程无错，结果正确
（二）323232网格：
[attach]3034[/attach]
上图的错误是在笔记本电脑 有一个局部内存忘记释放的结果。
我按照横版大说的释放后，笔记本电脑结果现在目前运行到2030（截止到发此回复时候）尚未报错。
但是奇怪的是 我的台式机电脑， 当1000步（每10步输出）时候 运行到1000步没有报错；但是当10000（每100步输出）就报错了。但是我觉得这个步长因素不会影响结果啊。我测试了3遍，都是如此，甚为诡异。
综323232：一是电脑不同，结果不同；二是步长这个看似不影响的貌似影响，待我进一步考证
（三）646464网格：运行27步，出错。
但是这里出错原因应该是 求解矩阵的值没有收敛导致（flag=1 代表未收敛），与323232错误不一样。
这里PS下：在未释放那个内存前，646464计算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]
同样向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里面概率也很大。
最后请您结合我回横扫版大的帖子一起给下一步指点！
顺祝春祺！！
恩呢！ 谢谢进一步解读。
我再仔细测试看看。如果有进一步进展也向您随时汇报!
安~
既然楼主能确定这里概率较大，那么：
最有可能出错是 计算possion方程的子函数。
Possion（double *,………………）
{
……局部显存分配d_A d_b等 //在分配后，和使用之前，有无host上
<<<<内核一>>>：形成d_b的值 //的code对d_A和d_B的赋值初始化之类？如有建议排查
<<<<内核二>>>：形成d_A的值
调用子函数 bicgs( **, **, **,); -------此步返回要求的X的值、
释放d_A d_b。
}
//如果没有，则请无视此建议。
好的！
按照您的说的继续测试。
回复下您提的几点：
（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);
我用返回残差看是否计算正确。
LZ您好，根据您的反馈，（1）和（2）已经排查。
（3）的说法横扫斑竹也给出了建议。
我再对（3）和其他内容稍加补充。
按照您的注释，内核一和内核二生成了d_A,d_b，然后一个包含许多kernel和CPU代码的函数bicgs()进行计算。我觉得检查一下bicgs()里面的数据依赖性，在循环的时候有没有累积误差造成的问题，有没有越界和泄露什么的，或许有助于解决问题。
至于查错函数放置的位置，原则上你放在A位置，那么能检查A位置之前的GPU操作成功与否。您可以放置在若干个地方，然后逐渐缩小出问题的范围。不过考虑到你这个问题可能会在循环若干次之后才出现，实际排查还是比较麻烦的。
其他没有什么合适的建议了，祝您好运！
好的 我先按照您和横大的意见进行排查！
如果新的进展也将向二位汇报！
谢谢鼓励！
按照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）