实现了一个算法，在GPU上把一个长的一维数据进行了归约运算，这样最后的和就保存在这个一维数组的顶端，现在要把这个和取出来，采用cudaStatus = cudaMemcpy(&pSum, pDevArray, sizeof(double), cudaMemcpyDeviceToHost);
这种内存拷贝的方式，拷出第一个值就得到了这个规约后的和。
但是测试的时候发现效率太低了，主要是这个内存拷贝占了一半的时间。有什么好方法可以不用内存拷贝的方式来取GPU内存上的值啊？
LZ您好，DEVICE上的数据CPU是不能直接访问的，必须copy回去。
所以您说的不通过copy获取CPU显存上的数据的方式，我是不知道的。
祝您编码顺利。
以及，我再补充一下。不一定真的“占用了一半时间”。
（1）也可能是楼主的时间计算的问题，例如：
计时点1
bigmanvip_kernel<<<…>>>(…);
计时点2
cudaMemcpy(…);
计时点3
那么实际上2-3的时间是kernel+复制的时间，而不是仅仅是复制。（您可以在2位置插入cudaDeviceSynchronize()来规避可能的此项问题）。
以及，这个是猜测的，并根据猜测给出的建议。
（2）如果此猜测不成立，的确8B复制都占据了1半时间。。。那楼主的kernel可能的确太小，那么此时的工作建议在cpu上直接完成。除非是数据已经在gpu上了。
（3）如果在（2）的情况下，并且楼主决定在gpu上进行工作。您还可以考虑对host mapped memory进行写入结果，这个是常见的小数据量(例如4B,8B结果)返回结果的常见做法。然后看看时间是否满足您要求。
上述3点请着重考虑(1)。
cudaEventRecord( time1StartEvent, 0);
sum<<<m_cBlockPerGrid, m_cThreadPerBlock>>>(m_pDevSource, m_pDevAim);
cudaStatus = cudaThreadSynchronize();
cudaEventRecord( time1EndEvent, 0);
cudaEventSynchronize(time1EndEvent);
//把求和后的值复制到主机内存
cudaEventRecord( time2EventStart, 0);
cudaStatus = cudaMemcpy(&fSum, m_pDevSourceSumArray, sizeof(double), cudaMemcpyDeviceToHost);
cudaStatus = cudaMemcpy(&gSum, m_pDevAimSumArray, sizeof(double), cudaMemcpyDeviceToHost);
cudaEventRecord( time2EventEnd, 0);
cudaEventSynchronize(time2EventEnd);
float elapsedTime = 0;
cudaEventElapsedTime( &elapsedTime, time1StartEvent, time1EndEvent);
g_timeCuda1 = elapsedTime;
cudaEventElapsedTime( &elapsedTime, time2EventStart, time2EventEnd);
g_timeCuda2 = elapsedTime;
结果是g_timeCuda1耗时0.4ms，g_timeCuda2耗时0.2ms。
应该是你说的问题就是问题规模太小了，我做的是255个数的规约求和，直接用cpu就可以，或者我再试下你说的第三 种方法，实在感谢，这么晚还回帖。
善哉。感谢楼主的及时回复。
上文主要是怀疑您可能一并将kernel计入时间了，现在看则无此问题。
那么根据您的回帖，您无需尝试第三种方式了。
因为：这个时间已经超出了使用cudaEventSynchronize/cudaEventElapsedTime所能给出的最小时间间隔。
手册给出的cudaEventElapsedTime的时间分辨率（例如某卡能支持到0.5us级别的时间分辨率)是无问题的，但是，前提是，我们需要精心构造的测试代码，以及无干扰的host code执行环境(例如:一个非常快且极度空闲的CPU), 否则，例如说, host OS的线程调度/host cpu上的其他任务，以及其他因素(例如在同步中的各种上下文切换），对于您的这测试代码，都会严重干扰的。
所以您无需尝试上文第三种建议了，在您的具体回帖给出的情况下。
此时我建议您使用nvidia visual profiler来确定极其小的时间。这个时候的结果才是有意义的。(因为他们取得时间数据的操作方式不同）。
以及，如果您不想使用nvvp进行profiling, 您还可以选择如下建议：
使用cupti构建您自己的简单测试工具。
具体如何构建一个使用cupti的简单工具来测试您的代码已经超出了我的知识和能力范围。这个只是建议。以及依然推荐您使用现成的nvvp来测试。
以及，简单的概括说：时间太短，测不准了。实际上可能非常短（例如几个~几十个us), 但是众多因素导致您无法简单的测试出这个几个/几十个us的时间。（您依然可以用5#和6#建议的其他方式测试出）。
以及，这个只是可能的说法。也许贵卡/贵机器上就是将近1个ms也不是没可能的。
以及，您如果使用host mapped memory, 直接将结果写回host memory. 那么您的时间基本上就是kernel时间了。
这个做法也可以考虑。这样您实际上规避了第二次的那个测试，因为规避了，里面的众多干扰因素也不会存在，所以增加的时间也不存在。所以您可能会得到几乎无增加的时间（这也是我为何说少量数据推荐直接回写到host memory的原因）。
但是这增加了您的kernel的编写复杂程度。所以请自己考虑是无视这几个/几十个us甚至1个ms呢？还是要折腾自己的kernel下。
横扫斑竹所言甚是，使用host mapped memory虽然也是将GPU 显存里面的数据copy回去（是自动copy回去的），但是时间开销和cudaMemcpy相比，要小很多，适合小数据量返回结果。
以及LZ的数据量确实有些小，以及横扫斑竹给出了详细的分析并确定问题直至解决问题的方法。
相较之下，我在2#的回答太过就事论事和简略了，请LZ仔细参阅横扫斑竹的各楼回答。愿与LZ共勉。