我的主机端有一块较大的数据，但内核函数只访问其中一部分，而且比较有规律，是能够预先知道位置的。
但是这些数据比较零散的分布在大数据集合里面，每一小部分可能只有数百字节，内核可能会访问数以万计的这些较小零散分布的数据块。但这些小块总和只占总数据量非常小的一个比例。
现在有如下2种方式：
1.主机端拼接
在主机端申请一块大内存，将这些数据拼接起来，然后cudamemcpyasync一起传输入GPU。
2.GPU端拼接
反复调用cudamemcpyasync该函数，每次只传输少量数据，在GPU端拼接起来。
我是很想用方式1，但是另外需要大内存。主程序有很多线程这样做的话，需要内存太大了，不好弄了。
不知道方式2是否存在问题？数据传输函数每次调用只传输几百字节，反复调用数万次的话，是否合适？
LZ您好，答案是“不合适，会慢死的”。
调用cudaruntime API的开销要比传输几百字节大得多，反复调用数万次，效率十分低下。
一般都建议减少内存传输次数。
此外，您还需要评估一下，您的计算是否真的需要GPU来实现。如果先用CPU拼接，拼接之后再copy给GPU，然后处理，然后再copy回来CPU做后处理，这样流程的时间和CPU直接搞定哪个更快？
以及，您既然说“小块总和只占总数据量的非常小的一个比例”，那又为何在主机端拼接需要“申请一大块内存”？
大致如上，供LZ参考。祝您编码愉快~
感谢您的答复！
我想主机有多快GPU卡，每个卡有多个流，每个流有多个预分配的计算任务队列，这样的乘积效应，会需要很多的内存需求。累积起来会很大了。
比如2卡8流4任务=64，那么事先需要64个这种内存，如果某个因子再上调，可能会需要更多。
另外看书中所说对某流的状态查寻函数效率很低，不知道一般都是怎样判断某个流的某个函数是否执行结束了？然后好处理输出结果，然后立即分配下一个任务？一般应该如何做才好？
我同意二楼的说法，以及楼主位的（1）做法。（在host memory上拼接，一次传输。而不是多次传输，device端拼接）。
因为CPU对内存的访问的总带宽是一定的。你同时为多个卡上的任务进行拼接，不如考虑下集中为一个卡拼接完，传过去，然后再拼接下一个。你说呢。你同时拼接多个，平均每个的时间，在CPU上，也不会比集中精力拼接完一个的时间快。
对于，对“对某流的状态查寻函数效率很低”，如果你指的是cudaStreamQuery(), 那么从日常应用中来说，我没有感觉到这个函数的“效率很低”。此外，对于cudaStreamSynchronize()或者cudaEventSynchronize()，我也没有感觉到这2者“效率很低”。我习惯于，在流中发布活，然后直接cudaStreamSynchronize()。
这是对你的1楼和上述楼层的一些问题的看法。仅供参考。
书中所说，同步本身是轮询而非中断方式进行的，会耗费大量CPU资源。
那么我的程序异步方式进行轮询，通过Sleep(1)这种方式进行轮询，是不是效果会更好？
这时我的程序只需一个线程控制所有的流，线程数也少了。
哦？请问你看到是什么书？
（1）同步在你的host process的执行的代码里，可以表现为轮询，也可以表现为阻塞在某个内核对象上。后者，当显卡在完成某个任务后，通过中断等机制，激活显卡驱动中的某部分负责的代码，然后驱动程序根据情况，进一步设置某些内核同步对象，然后唤醒你的在对某个同步对象上阻塞等待进程。
所以同步不是一定在反复polling导致"耗费大量CPU"。
如果想进一步了解详情，本论坛有10年-12年我的朋友的帖子，讨论同步都有实际哪几种情况。
此外，简单的说一下，在Windows上，
有反复循环轮询的，有带有Sleep(0)的反复循环轮询的(0表示立刻放弃时间片), 以及有WaitForSingleObject或其等效的阻塞等待通知的。
一般同步的时候会自动为你选择其中一种方式的，你也可以用cudaSetDeviceFlags()来明确指定一种等待方式。
（2）自己手工的cudaStreamQuery()，以及应用Sleep(1)的循环，是否效果好。要看你如何定义“效果”。从你的上下文看，你的效果可能指定是“线程对CPU的占用量“，也可能是指的“逻辑上的方便、简洁”。前者，不一定（看前文）。后者，也不一定，要看你的实现和cudaStreamSynchronize()的实际实现是那个简洁，以及要看你的逻辑结构上的实际需要。
此外，需要提醒的是，就算是表面上的直接用win32 api“阻塞”在某个同步对象上，实际上也不是立刻阻塞。现在的Windows OS都比较复杂，API是内部很多层次以及考虑到很多情况的复杂实现，这里的看上的“阻塞”，也可能是OS或者OS提供的用户态API内部在spin一个短暂的时间，然后才阻塞等侯（以应付快速已经就绪或者即将就绪的同步对象，以及可能需要长期等待的）。
综合的说，除非你的逻辑需要，否则我不建议你手工循环Query.
我自己在实际应用中，我一般喜欢用cudaDeviceScheduleBlockingSync(也就是以前的cudaDeviceBlockingSync)。这个一般情况下对我来说很小的kernel, 会立刻完成的，以及较长的，需要长期等待的，都适用。
以上仅供参考。可以采纳也可以无视。
简单的说，上去cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync|你的其他标志); 然后黑上cudaStreamSynchronize(你的流);即可。
除非你的流程/逻辑上不允许你这样做。
1.张舒那本书的P80倒数第3行。
2.一般来说每个GPU开多少个流比较合适呢？
太多的话，每个流能够获得的显存就小了，内核规模也就小了；太少的话，数据传输隐藏效果可能会差些。
3.每个流的任务队列多少比较合适？
如您所说，1的话，比较简单，但是同步的时候该流会有段时间空闲没事做。感觉上似乎不太妥当。
当然，这个空闲的时间片可能其它流会跟上填满GPU。但感觉上还是多几个任务在那排队更好。
如您介绍的，好像只需考虑1个任务就可以了，没必要弄得很复杂？
（1）我没有看过这本书，无法对此书“同步会耗费大量CPU“作出任何评价。建议重新参考我上面2楼对同步和CPU占用的讨论。
（2）开多少个流，需要根据你的实际需求。此外，流和显存无关系的。流只是任务指派的逻辑关系还分。你原来需要100个任务，在1个流里。你分成10个流，每个10个任务，还是100个。
（3）我没有表达过每个流只能给1个任务的说法。请不要乱扣帽子。
此外，在你给予每个流1个任务然后同步，会导致该流空闲神马的说法我看不懂：只要让GPU满负荷工作即可。
正因为你用1个流，但因为你总是需要准备数据，执行，然后CPU上后续处理神马的，导致容易无法始终让GPU满负荷。所以你才需要多个，每个的逻辑都简单，依然是这样，但是却有其他流的及时填上去的其他工作让GPU不闲着。
你的流可以看成是一堆发号施令的小领导，而GPU则是干活的手下员工。既然现在有一堆小领导了，哪怕每个小领导的都是漫不经心的，但是却能让GPU作为员工去100%的忙碌干活了。此时，你去在乎小领导们是否全心全意，实际上已经和贵司的单位时间内能出产的产品无关了。
此外，你实际上最后的回复已经是另外一个问题了。我建议你重新开一个主题，以便增加我的工作量。
EDIT: 您也可以继续在本帖回复。只是建议。
我建议只传输必要的部分，因为数据量一大，内存拷贝还是比较耗时的！
请仔细看楼主的初始问题的描述。人家即使是一开始假设的“拼接”，也是只传输部分数据，然后在device上组合。而不是传输全部到device上，再选取有用的部分。