Thank you for your patient answers so many times , I finally find the reason which is my CUDA code in backward pass has some bug.Change the int i=0 in the first for loop to int i=index, then it pass the gradcheck with epsilon=1e-3 and 1e-4, but fail in 1e-5 or less, which i think is the problem of precision.(The gradcheck is fail with epsilon range from 1e-2 to 1e-7 before i correct the code)
int batchsize = blockIdx.x;
int index = threadIdx.x;
int stride = blockDim.x;
.......
//Wrong: for (int i=0;i<n;i+=stride)
//Correct:
for (int i=index;i<n;i+=stride){
	........
			
	for (int k=0;k<c;k++)
	{
		atomicAdd(grad_xyz1 + i + k*n, g*(xyz1[i + k*n]-xyz2[id + k*m]));
		atomicAdd(grad_xyz2 + id + k*m, -(g*(xyz1[i + k*n]-xyz2[id + k*m])));
	}

}




 albanD:

That might be the issue. In my experience, single precision floats are hard to get to work with gradcheck.





 albanD:

Can you add a simple float/double template to your function by any chance?


Thank you for your advice !  . I will add a template in my code.