cudaMemcpy的Asynchronous transfer

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的时候是异步的)您可以无视。以这里的回复为准。

关于您的“胡文美老师”的问题,我建议您直接咨询她。我没有参加过她的课程,无法知道她当时讲话的内容和当时的上下文,因此无法评价。

关于我对其他楼主的帖子,仅承诺对当时的楼主有效,您看不懂将不再解释(本来就不是您的问题)。

关于您如何找到自己的老帖子的功能,论坛不提供查看自己的所有发帖的。我也不能直接列出我所有的发过的帖子的。这个抱歉。

感谢来访。请多看点书。

谢谢版主耐心的解答。。基本明白了。。

谢谢版主的解答。。版主学识确实渊博啊。。