需要做一个应用, kernel需要长期运行,但是kernel所需的数据不能一次性获得, 如何从主机源源不断地给kernel提供数据,并将处理结果返回呢?
都木有人…
为什么要求kernel要一直放在GPU上？
反复进行memcopy和启动kernel，每次启动的kernel处理对应的一部分数据不就可以了？
以及您还可以使用stream来使得计算和传输互相掩盖。
您不妨一试。
祝您编码愉快~
copy data to GPU → launch kernel → copy data to CPU → copy data to GPU → launch kernel…
是不是当kernel在运行的时候，GPU无法将数据拷贝到CPU中？
可以用stream来实现内核和内存拷贝并发，参考programming guide手册里面的介绍！注意，并发内核和内存拷贝，要求Host的memory必须是锁定页内存，也就是要用cudaHostAlloc分配标示为cudaHostAllocPortable的主存空间！
有点明白了，非常感谢
显然你没有明白，虽然我观察你一天了，希望你能明白. 
kernel长期运行，然后host不断传来/取走数据，这个说法，是不可以的。
其实具体的说，后半部分可以，但前半部分不行。
先说下可以的部分：
（1）host在kernel运行的过程中传递数据过来，这个可以(例如使用双缓冲区，一个kernel处理，一个host写入。然后通过某种方式通知运行中的kernel，然后kerne决定是否需要切换到下一个缓冲区).
（2）kernel在运行的时刻即将结果传回，这也可以（例如还是用host mapped memory, 然后加上对L2 cache write through的写操作(因为host mapped memory的结果写回，如果kernel还没结束，可能只是写回到L2中，不一定啥时候才能刷到主机内存里。)
然后回到前面，为何说你这个想法是不行的呢？因为正常的kernel，他的运行，是有线程的死亡以及诞生的。
假设你用的某卡，可以同时运行2万个线程(假设2048个线程每SM, 你有10个SM在该卡上)。那么你的一般正常的规模的kernel, 线程数要远大于2万。
我们假设说，你有个简单的线程，需要处理的缓冲区里是88万个独立数据，每个线程负责一个的话，那么你需要88万个线程。考虑到贵卡只能同时运行2万个，那么你的88万个线程的kernel, 实际上是：前一批2W个在GPU里运行，然后结束死掉，然后新的2W个又上来，如此反复，不断的老的死光，直到最后的2W个才能上去运行。
即：（1）老线程死了，新线程继续扑上来运行。（2）老的不死，新人一直在等待。
那么你想让你的kernel一直都在保持运行，却能反复上来一批数据，又一批数据，持续的处理，是可能的了。
因为你“处理完一批”意味着你的所有的线程都参与了一遍，然后再来新数据的时候，很多老线程已经死去了，无法复活参与新处理，必然下一批数据无法全体处理了。
所以，这个问题光从前面看，就已经不可能了。
那么如何才能复活这些老线程呢？
请参考这个意见“反复进行memcopy和启动kernel，每次启动的kernel处理对应的一批数据不就可以了”。
这个意见里的思路是正规的CUDA思路。因为如前文所说，线程在不断的死亡。只有重启kernel,才能让他们新生。
希望这么说楼主能真正的理解。
我的问题是这样的：比如说处理一张图片需要经过A–>B–>C三个步骤，传统的做法是将数据给A kernel处理，处理结束后返回CPU，然后启动B kernel，直到所有的步骤处理完毕，但是这中间有大量的CPU与GPU之间的数据传递，耗时较大。。。
我的思路是这样的：我先将三个步骤都以常驻kernel的形式启动起来，然后通过CPU不断地向GPU传递数据，直到所有的步骤都处理完毕再将结果取回，以此来节省数据在CPU与GPU之间的拷贝时间，不知道这样的处理思路有无问题？
cudaMemcpyToArray这个函数的最后一个参数设为cudaMemcpyDeviceToDevice，无需将中间数据传回host
1：如上文讨论，您的设想是不行的，请仔细研习8#的说法。
2：无需每一步都将数据copy回CPU端，除非CPU需要用到。您kernel A计算的结果本来就在GPU的显存里面，kernel B如需使用，只需要知道该结果保存数组的地址即可，而这一点，您可以简单通过将该地址作为kernel B的参数传给kernel B。
希望您能理解上述两点，并解决问题，祝您编码顺利~
感谢excellent_ld，但为了防止误导（因为翻页了），特意强调下，无需device to device copy ,直接启动下一步骤处理的kernel即可(诚如ICE所说)。
例如：
（1）kernel1<<<…>>>(…, input_buffer1, output_buffer2, …); //buffer 1,2都在device memory里。
（2）kernel2<<<…>>>>(…, output_buffer2(作为输入), output_buffer3); //buffer 3也在device memory里。
…(中间多个kernel, 但是他们都用上一个的输出作为输入，直接在显存里工作）。。。
(n) cudaMemcpy(…, cudaMemcpyDeviceToHost); //最后取回。
这样无需每步重复复制回来，复制过去。你说呢？
感谢横扫千军更正，恩恩，是的，这里我欠考虑了，表示歉意，确实直接就像ICE所说的即可！
客气客气，欢迎经常来论坛转转~
祝您新年快乐！
:2_31:哈哈，祝ICE新年快乐！
感些各位对新手的悉心指导，我辈自当发奋图强，同祝各位编码愉快！！！:lol