rt，在编程指导中说cudaMemcpy是Asynchronous transfer的，对其进行调用后就把控制权返回CPU。有Asynchronous transfer性质的还有好几个函数，一直不懂这个到底是怎么回事。斑竹能不能给个具体例子说明一下。。谢谢
cudaMemcpy*Async()系列API,
以及<<<>>>,
他们在执行的时候，host端的代码在继续执行，所以叫“异步”。
你既然在手册中看到了，你应该看到手册下文的具体例子，这里就不给出了。
谢谢斑竹的回答，但是自己对于异步执行还不是很理解。比如在下面例子中：
#include <stdio.h>
#include <stdlib.h>
#include <helper_cuda.h>
struct PStruct
{
int PSize;
float *PMatrix;
};
__global__ void KernelStruct(PStruct pstruct)
{
int PSize=pstruct.PSize;
int Idx=threadIdx.x;
if (Idx<PSize)
{
pstruct.PMatrix[Idx]=pstruct.PMatrix[Idx]+1;
}
}
int main()
{
int i;
PStruct CPU_Struct,GPU_Struct;
//初始化CPU
int PSize=10;
CPU_Struct.PSize=PSize;
CPU_Struct.PMatrix=(float*)malloc(PSize*sizeof(float));
for (i=0;i<PSize;i++){
CPU_Struct.PMatrix[i]=i;
}
//初始化GPU
GPU_Struct.PSize=PSize;
checkCudaErrors(cudaMalloc((void**)&GPU_Struct.PMatrix,sizeof(float)*PSize));
checkCudaErrors(cudaMemcpy(GPU_Struct.PMatrix,CPU_Struct.PMatrix,sizeof(float)*PSize,cudaMemcpyHostToDevice));
//核函数
dim3 grim(PSize,1,1);
KernelStruct <<<1,grim>>>(GPU_Struct);
//拷贝到CPU
cudaMemcpy(CPU_Struct.PMatrix,GPU_Struct.PMatrix,sizeof(float)*PSize,cudaMemcpyDeviceToHost);
//打印
FILE *pFile=fopen("Example.txt","w");
for (i=0;i<PSize;i++){
fprintf(pFile,"PMatrix[%d]%f\n",i,CPU_Struct.PMatrix[i]);
}
//释放
cudaFree(GPU_Struct.PMatrix);
free(CPU_Struct.PMatrix);
return 0;
}
有几个问题不懂：
（1）：在编程指南中指出主机和设备的异步执行包含了这两种情况：内核发射和主机和设备内拷贝小于64KB的存储器块。那么在上面代码中
checkCudaErrors(cudaMemcpy(GPU_Struct.PMatrix,CPU_Struct.PMatrix,sizeof(float)*PSize,cudaMemcpyHostToDevice));
//核函数
dim3 grim(PSize,1,1);
KernelStruct <<<1,grim>>>(&GPU_Struct);
//拷贝到CPU
cudaMemcpy(CPU_Struct.PMatrix,GPU_Struct.PMatrix,sizeof(float)*PSize,cudaMemcpyDeviceToHost);
这几句都是满足异步执行的条件的，但是在核函数发射的时候，不是需要上面从CPU中拷贝过来的值么？但是这里是异步的，会不会在数据还没有拷贝完全就执行核函数呢？当然下面从GPU到CPU数据拷贝类似会出现此类问题。
（2）：在上面核函数出现了分支语句if，在看胡文美老师在08年讲的视频中说这样实际上所有的线程都是要执行if语句的，不知道在现在还会不会出现这种情况？还有for循环需要人为的展开还是编译器可以自动展开啊？
（3）：在看http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=6344帖子，讲的是在global调用device函数的情况，不知道版主给出的方法是一种参考还是在cuda中必须这样啊?能不能像在c中一样包含头文件这种呢？（4）：不好意思，在问一下我怎么能找到我原来发的帖子呢？谢谢
LZ您好：
1：这样写是安全的，以及即便是异步版本的memcpy也是安全的，在同一个stream里面，任务是串行执行的。异步函数只是说返回CPU是异步的，这些任务发布给GPU任务队列的时候，依然是顺序的，会保证前一项结束后一项才执行。
2：kernel函数里面的语句都是每个线程都要执行的（分支部分的除外），所以kernel中的if依然是每个线程都要执行的，以及您的写法也是常见写法。
关于for循环，编译器会根据自身的判断和一些约束条件自行判断是否需要展开，以及展开到什么程度。您可以手工建议编译器展开某些循环，但是无法完全指定。
3：以前版本的CUDA Toolkit并不支持__device__函数和__global__函数写在不同的文件中，但是现在CUDA 5.0开始支持这样写。在头文件中声明应该是可以的。
4：论坛似乎还未支持搜索功能，您可以尝试google搜索： 关键词 site:cudazone.nvidia.cn 这样的形式看看。
祝您编码顺利~
楼主您好，你要知道“异步”是谁和谁异步。
我们讨论的异步API都是cuda调用和host代码的。而不是cuda调用和cuda调用之间的。
所以：
（1）您的cudaMemcpy的“异步”行为， 只是说在它“完全”完成之前和host的，而不是和后面的kernel的。
（2）此“异步”行为对从device->host的复制回来无效的（手册有错误。很正常）。
以及，综合1,2来说，您依然可以将cudaMemcpy是看成完全同步的，这样是安全的（不用理会手册）。
（如果要详细说一下的话，对于小于64KB的host->device的传输所展现出来的“异步”行为，实际上可以无视的。也就是说，cudaMemcpy返回后，您可以认为传输已经完成了，可以立刻释放或者修改源缓冲区了。因为实际上cudaMemcpy“总是”异步的，因为他是分步的，我们假设他内部使用了一个64KB的page-locked的缓冲区，您的普通内存，将先复制到此缓冲区，然后再从此缓冲区复制到显存。
对于您的小于此缓冲区容量的复制来说，复制到此缓冲区后就立刻返回了，此时不保证实际上已经到了显卡里（从此缓冲区到显卡的PCI-E传输可能依然进行中），但逻辑上已经无碍，您可以立刻修改源内容了。
对于您的大于此缓冲区容量的复制来说，将进行总大小 / 此缓冲区容量 次的staging, pci-e transfer步骤，而最后一次的总大小 % 此缓冲区容量的数据将展现类似上段的“异步”行为。
以及，只有cudaMemcpy*Async()系列函数才是真“异步”的，您在他们返回后不能动您的数据的，知道您执行了一次同步（cudaDeviceSynchronize()之类的）。
以及，曾经的某个版本的显卡驱动曾经使用了双缓冲区，对于大于里面的默认缓冲区大小的数据，将分别进行 缓冲区1填充-> pci-e复制 & 缓冲区2填充 → pci-e复制 (&缓冲区1填充)的循环。这样可以有效的提高复制效率。但后来被去掉了。现在的机制和我上文描述的一致。
）
这是我对您的“异步”的回答。再次重复，简单的说，您可以将所有的cudaMemcpy()看成和host是同步的，无论复制的数据大小。手册的说法（host<->device在<64KB的时候是异步的）您可以无视。以这里的回复为准。
关于您的“胡文美老师”的问题，我建议您直接咨询她。我没有参加过她的课程，无法知道她当时讲话的内容和当时的上下文，因此无法评价。
关于我对其他楼主的帖子，仅承诺对当时的楼主有效，您看不懂将不再解释（本来就不是您的问题）。
关于您如何找到自己的老帖子的功能，论坛不提供查看自己的所有发帖的。我也不能直接列出我所有的发过的帖子的。这个抱歉。
感谢来访。请多看点书。
谢谢版主耐心的解答。。基本明白了。。
谢谢版主的解答。。版主学识确实渊博啊。。