有没有人遇到这样的问题：0: copyin MemcpyAsync (dev=0x704b27e00, host=0x203707e00, size=32000

[bqfuthu]

如题，该错误调试发现，程序终止于语句：ERR = cudaMemcpyAsync(DEN,m_DEN,dm_NPRT)
DEN, m_DEN,都有声明和定义内存，分别是devic和host内存。cudaMemcpyAsync这个语句在该行之前也出现过，前面的也通过的，不知为什么。
另外这个程序可以在我原来的机器(Cuda5.5/K20)上运行，最近换了个地方干活，换了台机器(CUDA7.5/GK110B)，就出现这个问题了。
---------------------------------------------------------------------------------------------------------------------------------

然后我就利用cuda-memcheck 发现了以下错误提示：
-------
========= Invalid __shared__ write of size 16
========= at 0x00000388 in nbl_cal_neighborelist_devkernel2c_
========= by thread (0,0,0) in block (14,0,0)
========= Address 0x00000074 is misaligned
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usrb64/nvidiabcuda.so.1 (cuLaunchKernel + 0x2cd) [0x15865d]
========= Host Frame:/opt/pginux86-64/2015/cuda/7.5b64bcudart.so.7.5 [0x146ad]
========= Host Frame:/opt/pginux86-64/2015/cuda/7.5b64bcudart.so.7.5 (cudaLaunch + 0x143) [0x2ece3]
========= Host Frame:./fasMD.exe [0xe8200]
========= Host Frame:./fasMD.exe [0xe78e0]
========= Host Frame:./fasMD.exe [0xe3315]
========= Host Frame:./fasMD.exe [0x5002]
========= Host Frame:./fasMD.exe [0x77fe]
========= Host Frame:./fasMD.exe [0x3b34]
========= Host Frame:b64bc.so.6 (__libc_start_main + 0xfd) [0x1ed5d]
========= Host Frame:./fasMD.exe [0x3a29]
-------
查到这的时候，有个朋友说可能不是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)::POS(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