AMD OpenCL 计算出错但在Nvidia 上正常

  • 3 replies
  • 156 views
AMD OpenCL 计算出错但在Nvidia 上正常
« 于: 七月 11, 2019, 08:05:57 pm »
最近把一个CPU计算代码翻译为OpenCL进行运行,已经调试并验证通过(使用GTX1060).
计算过程是一个按迭代过程,计算结果以残差(后一迭代步与前一迭代步直接的结果差值)方式呈现,残差变小的过程称为收敛过程;
计算环境VS2015,全部fp64,Kernel函数依次在同一命令队列中调用,基于event等待进行约束;
现象是:
1.使用CUDA 9.2(OpenCL1.2)在GTX1060上的运行,各迭代步结果与CPU串行运行基本一致(在小数点12后有细微偏差);
2.A卡具有Debug与Release的区别(仅改变VS的Debug与Release)
a)使用Debug直接运行,第3个迭代步就发散(计算结果nan),前面2步结果不正确且表现出随机性(多次运行结果均不一致);
b)使用Release运行,能够计算且不发散,但各迭代步的值和CPU/GTX1060有很大差别;
c)使得Debug调试,若对包含调用多个clEnqueueNDRangeKernel()的函数step over(逐过程)运行,结果出错,但是,如果进入该函数后,逐个运行clEnqueueNDRangeKernel()函数,结果又正确一致;
d)尝试更改AMD驱动版本,更换AMD显卡(2片r9 390,1片r9 280x),添加OpenCL编译选项(如-cl-std=CL1.2 -cl-opt-disable)等均没效果;
e)本身未启用乱序执行,怀疑有不按照流程执行,于是使用了clSetEventCallback设置回调函数监控时间,发现流程正确。
综上,尤其是那个逐个监控就正确,不看时就不正确,感觉是这个“薛定谔事件”极为不科学.

Re: AMD OpenCL 计算出错但在Nvidia 上正常
« 回复 #1 于: 七月 11, 2019, 08:36:19 pm »
补充
程序源码,附件code.zip,我使用的是Cmake配置VS工程;
在main.c中,默认启用的是OpenCL计算,如下
程序代码: [选择]
//#define USING_UNATOMIC  1
#define USING_OPENCL    1 //启用OpenCL计算
要改为CPU版本的话是
程序代码: [选择]
//#define USING_UNATOMIC  1 //启用CPU计算
//#define USING_OPENCL    1

程序测试案例,附件test_case.zip

以CPU为对照,N卡计算结果和CPU一样,而A卡不正确。
希望有[名词6]来指点下。

Re: AMD OpenCL 计算出错但在Nvidia 上正常
« 回复 #2 于: 七月 11, 2019, 10:30:20 pm »
没人回复偶 :'(
不过问题解决了 ;D

mark下,在main.c的Solve_OpenCL()函数中,
我个人的猜测是调用clEnqueueNDRangeKernel时,command_queue和kernel的相同组合不同同时存在,
于是解决方案可以是:
1)注释//s1.6一段的最后,添加clWaitForEvents(1, &event_s1_6);,注释//s2_4一段的最后添加clWaitForEvents(1, &event_s2_4);,其他类似,这样可使得多次调用如kernel_Flux_Jameson_INTERIOR_WithoutAF这样的kernel函数不同时出现;
2)将kernel_Flux_Jameson_INTERIOR_WithoutAF这样用到的kernel函数改为函数数组,每次调用用不同的;

当然,这个原因是我猜测的,毕竟clEnqueueNDRangeKernel参数中除了event_wait_list不一样外其他都一样,具体以什么区分还要再测试;
不过,不知道这是不是AMD特有的OpenCL的实现的问题呢?
反正在khronos的官方说明(https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clEnqueueNDRangeKernel.html)中没提到个这个样

这个问题从4月份困扰到现在了(虽然中间也没特意思考),在打算弃坑时居然突然调好了

Re: AMD OpenCL 计算出错但在Nvidia 上正常
« 回复 #3 于: 七月 12, 2019, 12:45:40 pm »
没人回复偶 :'(
不过问题解决了 ;D

mark下,在main.c的Solve_OpenCL()函数中,
我个人的猜测是调用clEnqueueNDRangeKernel时,command_queue和kernel的相同组合不同同时存在,
于是解决方案可以是:
1)注释//s1.6一段的最后,添加clWaitForEvents(1, &event_s1_6);,注释//s2_4一段的最后添加clWaitForEvents(1, &event_s2_4);,其他类似,这样可使得多次调用如kernel_Flux_Jameson_INTERIOR_WithoutAF这样的kernel函数不同时出现;
2)将kernel_Flux_Jameson_INTERIOR_WithoutAF这样用到的kernel函数改为函数数组,每次调用用不同的;

当然,这个原因是我猜测的,毕竟clEnqueueNDRangeKernel参数中除了event_wait_list不一样外其他都一样,具体以什么区分还要再测试;
不过,不知道这是不是AMD特有的OpenCL的实现的问题呢?
反正在khronos的官方说明(https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clEnqueueNDRangeKernel.html)中没提到个这个样

这个问题从4月份困扰到现在了(虽然中间也没特意思考),在打算弃坑时居然突然调好了

根据你的描述看,你应当对你的kernel对象的启动,使用一个host上的锁,保护从设定参数---EnqueueNDRangeKernel之间的一段。KHR的规范明确说明了,这里必须原子性的完成,AMD的实现应当是符合规范的。请检查是否是如此。

当然你建立多个kernel对象,然后分别的使用,应当也是可以的(这样问题的前提就不存在了)。

本回复只根据你的文字描述给出(因为代码太太太长了,光main的那个文件,就超过了三千行了)