为什么我把一个浮点数组从host拷到device又拷回来,里面内容就变了？
我把device pointer作为函数参数传到函数里边再拷就会出问题
请楼主提供出问题的详细代码段。
或者建议楼主自己先逐步调试一下，看看每一步是否都正确执行了。
善哉。
按照楼主的说法，您实际上做过如下操作：
（1）从host memory A复制到device memory B,
（2）以B为参数传递给一个kernel, 然后执行，
（3）从device memory B复制回host memory A。
然后您才发现B变化了。那么可能是哪里的原因呢？
（一）其实您应该上去怀疑下您的(2)操作，因为你2次复制间夹杂的(2)很惹眼，对吧？
在您的(2)里，您是否有更改（例如写入)B的内容的操作呢？不妨按照ICE版的说法，检查一下。
（二）如果(2)里没有更改过B, 那我建议您怀疑下您的(1)(3)部操作，他们都成功了吗？（例如，源地址和目标地址是否正确？传输大小是否正确？传输类型是否正确？）
但不管怎样，我相当同意ICE大的"请楼主提供出问题的详细代码段"的提议。这样的确可以快速的发现问题的所在.(如果您的代码方便提供的话).
当然，您也可以简单的用语言描述您的代码的行为，如果您不方便的话。
祝您调试愉快！
如果是双精度的话要确认编译时加上-arch=sm_30(GK104), sm_20(Fermi)或sm_13(GT200)。否则不会支持双精度，从而会出现lz说的现象。
显然不是。cudaMemcpy和你要复制的数据的类型是无关的，它又不知道你复制的数据有何意义。
对吧！
感谢横同学指出这一点，看来是我记错了。但我的建议还是如果是双精度的话，必需要加上那些选项。否则第二步kernel里会自动把它们降为单精度处理。
建议至少在代码最后加上下面一行看看有没有执行错误：
printf(“%s\n”, cudaGetErrorString(cudaGetLastError()));
另外一个可能原因是即使kernel里没有写这个数组，但有内存越界访问。这可以通过用cuda-memcheck来检查一下。使用方法是：
cuda-memcheck [exe_name] [exe_parameters]
善哉。感谢P同学参与讨论。让我们继续热烈的讨论下去：
P:“另外一个可能原因是即使kernel里没有写这个数组，但有内存越界访问。”
我不认为有越界的情况下，会导致楼主的情形出现。
楼主的流程是: 复制内容到设备，再执行kernel, 再复制回来，发现内容变化了。如果如同P同学的说法，那么此kernel会执行失败(假设是正常普通读写，先不考虑cudaBoundaryModeZero之类的)。然后后一步的复制回来实际上不会被执行。自然楼主观察到的情形不会出现。
嘟嘟。
然后后一步的复制回来实际上不会被执行。自然楼主观察到的情形不会出现。
GPU上内存越界访问不一定会导致kernel运行失败。否则我们也不用专门开发cuda-memcheck来探测内存越界了。
越界，可能会存在3种情况，（1）分配粒度的原因导致有一定的分配裕量，读写这里的确不会导致kernel挂掉。（2）分配的大小是分配粒度的整数倍，越界，挂。（3）有分配粒度导致的富余量，但已经超出了，也挂。
其中的确只有(2)(3)情况会导致kernel挂掉。
P同学认为的可能有(1)情况存在，但楼主的帖子对此点给予了否认。楼主如果真是kernel读写这些多余的地方，何来复制回去能看到他自己的缓冲区内的值改变呢！他又不是观察的多余分配的部分。
所以P同学的“越界会导致复制回去出错”的评论，在(1)情况下不会得到楼主那样的观察情形; 在(2)(3)情况下楼主根本就复制不回去。鉴于此，我认为此段开始部分的P同学的评论不是很符合楼主。P同学认为呢？
感谢横扫千军的深入探讨。
内存越界是有可能导致lz的问题的。关键是内存越界导致kernel失败并不是超出cudaMalloc的分配粒度。比如下面一个小例子，n=64时就会重现lz汇报的问题。n=2^17时才会报ULF。
include <stdio.h>
global void ker(double *a, int n)
{
a[threadIdx.x+n] += 0.2;
}
int main()
{
const int size = 32;
double a[size], d_a, d_b;
memset(a, 0, sizesizeof(double));
cudaMalloc(&d_a, sizesizeof(double));
cudaMalloc(&d_b, sizesizeof(double));
cudaMemset(d_b, 0, sizesizeof(double));
cudaMemset(d_a, 0, size*sizeof(double));
for (int n = 1; n < 2000000; n = 2) {
printf(“n = %d\n”, n);
ker<<<1, 32>>>(d_a, n);
cudaMemcpy(a, d_b, sizesizeof(double), cudaMemcpyDeviceToHost);
printf(“%s\n”, cudaGetErrorString(cudaGetLastError()));
for (int i = 0; i < size; i++)
printf(“%f “, a[i]);
printf(”\n”);
}
cudaFree(d_a);
cudaFree(d_b);
}
编译运行：nvcc -arch=sm_20 main.cu ; a.out
部分输出：
n = 32
no error
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000
n = 64
no error
0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000
…
n = 65536
no error
0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000
n = 131072
unspecified launch failure
0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000 0.200000
[quote=“system”, post: 11, topic: 1838]
感谢横扫千军的深入探讨。
内存越界是有可能导致lz的问题的。关键是内存越界导致kernel失败并不是超出cu …[/quote]
include <stdio.h>
global void ker(double *a, int n)
{
a[threadIdx.x+n] += 0.2;
}int main()
{
const int size = 32;
double a[size], d_a, d_b;
memset(a, 0, sizesizeof(double));
cudaMalloc(&d_a, sizesizeof(double));
cudaMalloc(&d_b, sizesizeof(double));
cudaMemset(d_b, 0, sizesizeof(double));
cudaMemset(d_a, 0, size*sizeof(double));
for (int n = 1; n < 2000000; n = 2) {
printf(“n = %d\n”, n);
ker<<<1, 32>>>(d_a, n);
cudaMemcpy(a, d_b, sizesizeof(double), cudaMemcpyDeviceToHost);
printf(“%s\n”, cudaGetErrorString(cudaGetLastError()));
for (int i = 0; i < size; i++)
printf(“%f “, a);
printf(”\n”);
}
cudaFree(d_a);
cudaFree(d_b);
}
感谢PengWang的积极回复，这里面，有一个问题，是否在发帖时丢失了部分代码？
因为文中kernel 修改的是d_a，而cudaMemcpy回拷的是d_b；而且没见到d_a和d_b相互有什么操作。这使我十分不解。
没有丢失代码。目的就是为了展示没有对d_b进行任何操作的情况下d_b的值也被d_a的内存越界访问修改了（lz的问题）。
感谢PengWang的解说，前面确实没有想到这一点。
但是，如果要实现您给出的情况，我觉得需要有至少2点：
1：d_a，d_b两段分配的空间一定要有确定的物理位置关系，这样才能保证d_a一直向后写，会改写d_b中的值。在cudaMalloc的时候是否存在这样的位置关系，我作为一个外人，不得而知。
2：如果要在n=64的时候恰好出现d_b被改写，需要特定的cudaMalloc的分配粒度。
此外可能还有写越界到什么程度才能被检测到，这也是一个外人无法得知的问题。
不过，此例子确实说明了在某种精巧的情况下，写越界改写了另外数组的值，却没有使kernel挂掉。这一点值得学习。
感谢Peng同学的代码。
（1）Peng同学利用了分配粒度(512B)和连续的分配造成的内存布局。Peng同学分配了2次连续的256B, 但实际上分配了2个512B，而且在地址上是连续的。这样有d_b在后头顶着，d_a写超了后，先会写到多分配的256B上，然后继续写到d_b上。实际上无需64, 512B/8B = 64B, 64-31 = 33, 从n=33左右，d_b已经开始被改变了.
（2）Peng同学利用了分配算法和默认的初始堆(heap)大小。cudaMalloc用来进行分配的堆，初始大小是1MB, 而Peng同学熟知此分配算法是从头开始扫描的，导致d_a实际上从堆的最低可用地址开始分配的，而之后实际上有1MB的可用空间。所以不会导致挂掉。具体说，实际到1MB(2^20) / 8B(2^3) 然后 - 32即成功超越默认的1MB空间，即2^17 - 32以后，即从131041左右开始，kernel即挂.
（3）关于返回值，win*s上的实现是应该返回cudaErrorUnknown的，即"unknown error"错误字符串，当超越了1MB后。
因为Peng同学的分配顺序，内存布局，堆初始大小等因素，才会这么巧。
我不认为楼主只分配了一个缓冲区就满足Peng同学精心设计的如上条件。Peng同学认为呢？
Well, 这么简单代码就能重现的一个可能性，我们要么可以认为它不可能，要么可以花几秒钟运行cuda-memcheck来排除。在lz提供的关于这个问题的非常有限信息下，我个人的经验不足以让我先验的认为它不可能。