| 查看: 1214 | 回复: 4 | ||
| 当前只显示满足指定条件的回帖，点击这里查看本话题的所有回帖 | ||
relonfbq金虫 (正式写手)
|
[求助]
请问 “0: copyout memcpyAsync FAILED: 74 (misaligned address)” 问题怎么解决
|
|
|
我今天运行Fortran+cuda的程序，出现以上提示 “0: copyout memcpyAsync (host= , dev= , Size= ,Stream=) FAILED: 74 (misaligned address)”的错误，这是怎么回事？ 这个程序原来在我自己的实验室编译可以运行，现在换了个地方，也可以编译生成可执行程序，但就是出现这样的问题，不知道怎么解决。 在自己的实验室用的是cuda5.5编译的，pgfortran编译的 现在唯一改了用了cuda7.0编译的 (没办法，只有cuda7.0，页希望改成cuda5.5)，有什么解决办吗？现在在这儿机器不能随意安装编译器，即是可以，我也不知怎么弄？大家遇到过类似问题嘛，都是怎么解决的？ |
回复此楼» 猜你喜欢
336求调剂
已经有4人回复
-
求生物学调剂
已经有8人回复
-
08生物与医药专硕初试346找调剂
已经有6人回复
-
生物学327，求调剂
已经有6人回复
-
0710生物学求调剂
已经有7人回复
-
385分 生物学（071000）求调剂
已经有4人回复
-
材料考研调剂
已经有4人回复
-
366求调剂一志愿东北大学
已经有3人回复
-
08工科275分求调剂
已经有13人回复
-
材料工程322分
已经有6人回复
relonfbq
金虫 (正式写手)
- 应助: 1 (幼儿园)
- 金币: 1294.3
- 散金: 1140
- 红花: 2
- 帖子: 301
- 在线: 99.7小时
- 虫号: 744333
- 注册: 2009-04-09
- 专业: 金属材料的微观结构
4楼2016-03-31 18:54:29
relonfbq
金虫 (正式写手)
- 应助: 1 (幼儿园)
- 金币: 1294.3
- 散金: 1140
- 红花: 2
- 帖子: 301
- 在线: 99.7小时
- 虫号: 744333
- 注册: 2009-04-09
- 专业: 金属材料的微观结构
|
这个错误这么久了，还是没有解决， ========= Program hit cudaErrorInvalidValue (error 11) due to "invalid argument" on CUDA API call to cudaMemcpyAsync. ========= Saved host backtrace up to driver entry point at error ========= Host Frame:/usr/lib64/nvidia/libcuda.so.1 [0x2f31b3] ========= Host Frame:/opt/pgi/linux86-64/2015/cuda/7.5/lib64/libcudart.so.7.5 (cudaMemcpyAsync + 0x1b3) [0x435d3] ========= Host Frame:/opt/pgi/linux86-64/15.10/lib/libcudafor.so (pgf90_dev_copyout_async + 0x38) [0x1a7b0] ========= Host Frame:/opt/pgi/linux86-64/15.10/lib/libcudafor.so (cudamemcpyasync3r8out_ + 0x1c) [0x2346c] ========= Host Frame:./fasMD.exe [0x13de1f] ========= Host Frame:./fasMD.exe [0x13b8f6] ========= Host Frame:./fasMD.exe [0x12a62f] ========= Host Frame:./fasMD.exe [0x51af] ========= Host Frame:./fasMD.exe [0x7a2e] ========= Host Frame:./fasMD.exe [0x3b34] ========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1ed5d] ========= Host Frame:./fasMD.exe [0x3a29] ========= ========= ERROR SUMMARY: 1 error |
2楼2016-03-30 21:42:29
relonfbq
金虫 (正式写手)
- 应助: 1 (幼儿园)
- 金币: 1294.3
- 散金: 1140
- 红花: 2
- 帖子: 301
- 在线: 99.7小时
- 虫号: 744333
- 注册: 2009-04-09
- 专业: 金属材料的微观结构
3楼2016-03-31 17:43:55
relonfbq
金虫 (正式写手)
- 应助: 1 (幼儿园)
- 金币: 1294.3
- 散金: 1140
- 红花: 2
- 帖子: 301
- 在线: 99.7小时
- 虫号: 744333
- 注册: 2009-04-09
- 专业: 金属材料的微观结构
|
再把问题系统说一遍，该错误调试发现，程序终止于语句：ERR = cudaMemcpyAsync(DEN,m_DEN,dm_NPRT) DEN, m_DEN,都有声明和定义内存，分别是device和host内存。cudaMemcpyAsync这个语句在该行之前也出现过，前面的也通过的，不知为什么。 另外这个程序可以在我原来的机器(Cuda5.5/K20)上运行，最近换了个地方干活，换了台机器(CUDA7.5/GK110B)，就出现这个问题了。 --------------------------------------------------------------------------------------------------------------------------------- 然后我就利用cuda-memcheck 发现了以下错误提示： ------- ...... ========= Invalid __shared__ write of size 16 ========= at 0x00000250 in nbl_cal_neighborelist_devkernel2c_ ========= by thread (0,0,0) in block (14,0,0) ========= Address 0x00000004 is misaligned ========= Saved host backtrace up to driver entry point at kernel launch time ========= Host Frame:/usr/lib64/nvidia/libcuda.so.1 (cuLaunchKernel + 0x2cd) [0x15865d] ========= Host Frame:/opt/pgi/linux86-64/2015/cuda/7.5/lib64/libcudart.so.7.5 [0x146ad] ========= Host Frame:/opt/pgi/linux86-64/2015/cuda/7.5/lib64/libcudart.so.7.5 (cudaLaunch + 0x143) [0x2ece3] ========= Host Frame:./fasMD.exe [0xe8260] ========= Host Frame:./fasMD.exe [0xe7946] ========= Host Frame:./fasMD.exe [0xe3315] ========= Host Frame:./fasMD.exe [0x5002] ========= Host Frame:./fasMD.exe [0x77fe] ========= Host Frame:./fasMD.exe [0x3b34] ========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1ed5d] ========= Host Frame:./fasMD.exe [0x3a29] ========= Checking A003 0: copyout MemcpyAsync (host=0x203707e00, dev=0x704b27e00, size=32000, stream=0) FAILED: 4(unspecified launch failure) ========= ERROR SUMMARY: 168 errors [xxxx@oums-sfgpu MD]$ cuda-memcheck ./fasMD.exe ipa.thfu 0 1 ------- 查到这的时候，有个朋友说可能不是cudaMemcpyAsync的问题，好像有道理，我去查了一下cal_neighborelist_devkernel2c这个kernel子程序。注释掉，程序可以运行。 但是我仔细检查了这个kernel，也不能找出问题的所在。 我通过屏蔽些地方，或者添加些demo语句。发现可能是里面的shared变量有问题。但是我还是找不到问题的所在。 下面我贴出这个kernel，希望高手作答，谢谢。 -------------------------------------------------------------------------------------------------------------- use constant implicit none ! !--- DUMMY VARIABLES integer, value::NBPC,IP0, NC,NCX,NCY,NCZ, PDX, PDY, PDZ, NPART, CFROM, CTO, mxNAPDEV, mxKVOIS, IA1th0 real(KINDDF), value::cra011,cra021,cra031,cra012,cra022,cra032,cra013,cra023,cra033 real(KINDDF), device::XP(3,NPART) integer, device::ITYP(NPART), NAC(NC), IA1th(NC) integer,device::KVOIS(mxNAPDEV) integer,device::INDI(mxKVOIS,mxNAPDEV) integer(1),device::mvl(mxKVOIS,mxNAPDEV) !--- Local variables !nonshared by threads real(KINDSF):  OS(3), SEP(3)integer::IB, IB0, IT, IA, IA0,IA00, JA, NN, I, J, K, ITY !variables share by all thread integer,shared::NB, IC, IS0, STARTA, NCXY, NCXYZ,IX0, IY0, IZ0,IC0, NACC0 integer,shared::NS,NACC, IAC, IACE, FROM, TO integer, shared, dimension(mp_NNC)::CID, IX,IY, IZ, OUT integer(1), shared, dimension(mp_NNC)::I_mvl real(KINDSF), shared, dimension(3,mp_NNC)::CXYZ real(KINDSF), shared, dimension(3,mp_BLOCKSIZE)::SPOS real(KINDSF), shared::RC2(mp_MXGROUP,mp_MXGROUP) integer, shared, dimension(mp_BLOCKSIZE)::JTY IB = (blockidx%y-1) * griddim%x + blockidx%x-1 NB = blockdim%x*blockdim%y IB0 = IB/NBPC IP0 = (IB-IB0*NBPC)*NB IB0 = IB0 + CFROM-1 if(IB0 .GE. CTO) return IT = (threadidx%y-1)*blockdim%x + threadidx%x if(IT .EQ. 1) then RC2(1:mp_MXGROUP,1:mp_MXGROUP) = dcm_RU2(1:mp_MXGROUP,1:mp_MXGROUP) NCXY = NCX*NCY NCXYZ = NCXY*NCZ IS0 = IB0/NCXYZ IC = IB0-IS0*NCXYZ IZ0 = IC/NCXY IY0 = (IC-IZ0*NCXY)/NCX IX0 = IC-IZ0*NCXY-IY0*NCX IZ0 = IZ0 + 1 IY0 = IY0 + 1 IX0 = IX0 + 1 IC = IB0 + 1 STARTA = IA1th0 ! IA1th(CFROM) NACC0 = NAC(IC) end if call syncthreads() if(NACC0 .LE. 0) return if(IT .LE. mp_NNC) then OUT(IT)= 0 IZ(IT) = IZ0+mp_NIZ(IT) IY(IT) = IY0+mp_NIY(IT) IX(IT) = IX0+mp_NIX(IT) I_mvl(IT) = 0 CXYZ(1:3,IT) = 0.0d0 If(PDX .AND. IT.GT.1) Then IF( IX(IT).GT.NCX )THEN IX(IT) = 1 CXYZ(1,IT) = CXYZ(1,IT) + cra011 CXYZ(2,IT) = CXYZ(2,IT) + cra021 CXYZ(3,IT) = CXYZ(3,IT) + cra031 I_mvl(IT) = ior( I_mvl(IT) , 1) ELSE IF (IX(IT).LT.1) THEN IX(IT) = NCX CXYZ(1,IT) = CXYZ(1,IT) - cra011 CXYZ(2,IT) = CXYZ(2,IT) - cra021 CXYZ(3,IT) = CXYZ(3,IT) - cra031 I_mvl(IT) = ior( I_mvl(IT) , 2) ENDIF End If If(PDY .AND. IT.GT.1) Then IF( IY(IT).GT.NCY )THEN IY(IT) = 1 CXYZ(1,IT) = CXYZ(1,IT) + cra012 CXYZ(2,IT) = CXYZ(2,IT) + cra022 CXYZ(3,IT) = CXYZ(3,IT) + cra032 I_mvl(IT) = ior( I_mvl(IT) , 4) ELSE IF (IY(IT).LT.1) THEN IY(IT) = NCY CXYZ(1,IT) = CXYZ(1,IT) - cra012 CXYZ(2,IT) = CXYZ(2,IT) - cra022 CXYZ(3,IT) = CXYZ(3,IT) - cra032 I_mvl(IT) = ior( I_mvl(IT) , 8) ENDIF End If If(PDZ .AND. IT.GT.1) Then IF( IZ(IT).GT.NCZ )THEN IZ(IT) = 1 CXYZ(1,IT) = CXYZ(1,IT) + cra013 CXYZ(2,IT) = CXYZ(2,IT) + cra023 CXYZ(3,IT) = CXYZ(3,IT) + cra033 I_mvl(IT) = ior( I_mvl(IT) , 16) ELSE IF (IZ(IT).LT.1) THEN IZ(IT) = NCZ CXYZ(1,IT) = CXYZ(1,IT) - cra013 CXYZ(2,IT) = CXYZ(2,IT) - cra023 CXYZ(3,IT) = CXYZ(3,IT) - cra033 I_mvl(IT) = ior( I_mvl(IT) , 32) ENDIF End If IF( IX(IT) .GT. NCX .OR. IX(IT) .LT. 1) OUT(IT) = 1 IF( IY(IT) .GT. NCY .OR. IY(IT) .LT. 1) OUT(IT) = 1 IF( IZ(IT) .GT. NCZ .OR. IZ(IT) .LT. 1) OUT(IT) = 1 IF(OUT(IT) .EQ. 0) then CID(IT) = NCXY*(IZ(IT)-1)+NCX*(IY(IT)-1)+IX(IT)+IS0*NCXYZ ELSE CID(IT) = 0 END IF END IF call syncthreads() IA = (IT-1) + IA1th(IC) + IP0 IA0 = IA - STARTA + 1 IA00 = IA - IA1th(IC)+1 if(IA00.LE.NACC0) then POS(1:3)=XP(1:3,IA) ITY = ITYP(IA) NN = 0 endif K=1 NACC = NAC(CID(K)) NS = (NACC-1)/NB+1 IAC = IA1th(CID(K)) IACE = IAC + NACC -1 DO J=1, NS, 1 FROM = min((J-1)*NB+IAC,IACE) TO = min(FROM+NB-1, IACE) if(IT+FROM-1<=IACE)then SPOS(1:3,IT) = XP(1:3, IT+FROM-1) + CXYZ(1:3,K) JTY(IT) = ITYP(IT+FROM-1) end if call syncthreads() IF(IA00.LE.NACC0) then DO I=FROM, TO JA = I-FROM+1 SEP(1:3) = POS(1:3) - SPOS(1:3,JA) if( SEP(1)*SEP(1)+SEP(2)*SEP(2)+SEP(3)*SEP(3) .LE. RC2(ITY,JTY(JA)) ) then if(I.NE.IA) then if(NN<mxKVOIS) then NN = NN + 1 INDI(NN,IA0) = I mvl(NN,IA0) = I_mvl(K) end if end if end if END DO END IF call syncthreads() END DO DO K=2, mp_NNC IF(OUT(K)) cycle NACC = NAC(CID(K)) NS = min((NACC-1)/NB+1, NACC) IAC = IA1th(CID(K)) IACE = IAC + NACC -1 call syncthreads() DO J=1, NS FROM = min((J-1)*NB+IAC,IACE) TO = min(FROM+NB-1, IACE) if(IT+FROM-1<=IACE)then SPOS(1:3,IT) = XP(1:3, IT+FROM-1) + CXYZ(1:3,K) JTY(IT) = ITYP(IT+FROM-1) endif call syncthreads() IF(IA00.LE.NACC0) then DO I=FROM, TO JA = I-FROM+1 SEP(1:3) = POS(1:3) - SPOS(1:3,JA) if( SEP(1)*SEP(1)+SEP(2)*SEP(2)+SEP(3)*SEP(3) .LE. RC2(ITY,JTY(JA)) ) then if(NN<mxKVOIS) then NN = NN + 1 INDI(NN, IA0) = I mvl(NN, IA0) = I_mvl(K) end if end if END DO END IF call syncthreads() END DO END DO IF(IA00 .LE. NACC0) then KVOIS(IA0) = NN END IF return |
5楼2016-04-03 01:01:17
