在CUDA 中如何在 device 函数中复制内存到CPU呢？能否实现？
device float compute_d( float GpuArray_d, int AttributeN_d, int z_d, int i_d_d)
{
float tempdist_d = 0;
for (int j = 0; j < AttributeN_d ; j++ ) {
tempdist_d = tempdist_d + (powf(GpuArray_d[i_d_d*AttributeN_d+j] - GpuArray_d[z_dAttributeN_d+j],2.0));
}
return tempdist_d;
}
global void filterData_D(float *GpuArray, int Dataset_Dimension_d, int AttributeN_d, float tempArrary_d, int i_d,clock_t time_d)
{
const int bid = blockIdx.x;
const int tid = threadIdx.x;
float tempdist = 0;
for( int x=0; x< Dataset_Dimension_d; x++) // Dataset_Dimension_d=50000 ，。
for ( int z = bid * THREAD_NUM + tid; z < Dataset_Dimension_d; z += BLOCK_NUM * THREAD_NUM)
{
tempdist = compute_d(GpuArray, AttributeN_d, z, i_d);//调用
tempArrary_d[z] = (sqrtf(tempdist));
tempdist = 0;
}
这个程序的外循环是本来是在CPU上执行的，我感觉加速不够（只有37倍左右),所以想把外循环也实现并行，所以移到global里，这里我需要把每次的(总共50000)tempArrary_d[z] 复制到CPU的大小为5000050000向量里，收到显存限制，不能等到全部结束后在复制（5000050000），必须是外循环执行一次，结果被复制一次到CPU，这个需要在__global__复制内存，我在__device__,device __host__中 cudaMemcpy试图复制，但未能成功，
希望大家出谋划策。
ubuntu+nvcc编译环境 Tesla C2050
LZ您好：
1：不能在device端代码中使用cudaMemcpy()这个runtime API函数的。您可以查看一下官方的手册以确定cudaMemcpy()的具体使用方式。
2：按照您的叙述，您可以不必把大循环也整到kernel里面，您之前的实现还有访存合并上的优化潜力。
3：以及您可以使用多个stream来实现copy和计算的互相掩盖。
大致如此，由于LZ并未解释代码的实现意图和其他细节，无法继续脑补并给出建议了，上述建议中如有脑补失败的地方请LZ无视。
感谢您深夜来访，祝您编码顺利~
谢谢斑竹的回复。
这个代码是一个过滤数据代码。
以上代码中去掉外循环即（for( int x=0; x< Dataset_Dimension_d; x++) // Dataset_Dimension_d=50000 ）部分代码，那就是我现在用代码。
程序功能是：Cpu 代码较长，只能文字说明
在CPU端 for(i=0 i<=50000 i++)
{ 定义，初始化等
调用kernel,
cudaMemcpy(device To host)
用Thrust对反馈指针数组进行排序，//Thrust必须在host段执行，
根据threshold 判断是否过滤
i++
}
GPU端代码请如顶楼所示， 能实现37倍加速，
我试图把CPU的循环也放在 kernel（顶楼所示)，20秒左右也能完成计算，能实现150的加速比，但没办法复制每次的外循环结果向量加以判断是否过滤，如果能复制这个数据到CPU速度会有很大的改善。
根据这个程序能具体说明《访存合并》的使用吗？
真希望斑竹能提供帮助。谢谢。
自己顶自己顶自己顶自己顶
您别着急，
ICE给您指出了主要问题，就是最直接的帮助了。
而不是在于替您实现，替您写代码，那是为你打工。
您看您这里：
tempdist_d = tempdist_d + (powf(GpuArray_d[i_d_d*AttributeN_d+j] - GpuArray_d[z_d*AttributeN_d+j],2.0));
这里的访问：
GpuArray_d[z_d*AttributeN_d+j]将会严重影响您的显存读取性能的。
您的z_d是当前线程的工作编号，这样临近的2个线程：
z_d和z_d + 1之间，访问将间隔AttributeN_d个元素，这将非常严重的影响性能。
（当然，如果要较真的话，您要是说AttributeN_d=0或者=1之类的，那我无话说。。）
随便对您的代码举个例子：
如果您的AttributeN_d是16的话，您实际上将是：
GPUArray[id * 16 + j]
这将会导致临近的64B里只有4B是有用的，直接将您的L2 cache效率降低到6.7%。
所以您应当考虑这里如何尝试合并，您这代码实际上是简单的CPU代码的组合，而从未考虑过GPU的特性，为GPU的特性改编过。
请您考虑这一点，并根据我和ICE的建议指出的代码位置，看看能否改变算法，将访问合并，从而提高访存效率。
(请注意我们只能给出分析和建议，而不能直接为您设计算法、 编写代码、给出实现）
感谢您的夜间来访。
谢谢你的建议，
LZ您好：
昨日临时不在，让您久等了。
根据您补充的新信息，大致说几点：
1：您3#示意代码描述中，每次外循环均要执行定义和初始化，如果这里面涉及有缓冲区的分配和初始化等，请您考虑尽量复用，并将分配和初始化部分拿到循环外，以便提升效能。
2：您表示因为要使用thrust所以必须将数据copy回host端。我虽然不是thrust用户，但是thrust的手册在一开始就表明支持两种矢量容器，host_vector和device_vector，其中后者是存储于device端的，不知您能否利用这一点。
3：您无法在device端使用host端专属的一些Runtime API函数的，正如2#中指出的那样。您可以参阅手册得知这些runtime API函数的具体使用方法，这里不再赘述。
4：关于合并访问的问题，横扫斑竹已经给出了详细解释，不再赘述。
5：以及，如果您确实需要每步都将数据复制回来再做处理，您也可以考虑下使用多个stream的方式，如同2#中建议过的那样。
大致如此，祝您编码顺利~
真是好斑竹啊，谢谢指点。