通过看一些资料什么的,我对流的理解是:每个流都是异步的,分成好多流以后,前一个流还在计算着,下一个流已经开始计算了,就是可以实现一边传输一边计算的目的。这样应用的场合是不是要在一开始把这些流所需要的总共的显存空间分配好然后再分开多流计算,比如说:我要计算1亿个数据,我必须要先申请这一亿个数据的空间,即先cudaMalloc(1亿数据空间),然后再分开比如说5个流,每个流再进行计算,计算的时候kernel的参数是不同的,也就是说,比如本来总的是de_a<就是指向1亿空间的指针>,分成多个流以后,是不是要重新分配5个 1亿/5的数据空间呢?还是通过下标来确定,用个循环?我看SDK里面那个simpleStream里面的例子,里面有句:
for (int i = 0; i < nstreams; i++)
{
cudaMemcpyAsync(hAligned_a + i * n / nstreams, d_a + i * n / nstreams, nbytes / nstreams, cudaMemcpyDeviceToHost, streams[i]);
}
这个我看不懂,大侠们这句怎么理解呢?我上面理解的那些对吗?
LZ您好:
您的问题大致说明如下:
1:stream内部有各种操作,包括数据copy,也包括kernel计算。stream内部是顺序执行的,stream之间是不保证顺序的。
2:一般情况都可以满足kerenl计算和copy互相掩盖。在合适的硬件下,以及合适发布规模的kernel,可实现kernel并行计算。
3:假定您的可用显存的数量是N,那么您可以申请5个N/5的存储空间,给5个stream使用,这里N/5未必要大于等于“1亿/5的数据量”,您可以在同一个stream里面多次copy和启动kernel计算的。
4:您给出的代码,循环变量i辅助用于memcpy中计算指针位置,而最后那个参数streams则是表示您copy命令所在的stream。另外值得您注意的是,需要使用异步版的cudaMemcpy,如上文中使用的那样。
大致如此,祝您编码愉快~
流可以从host和从device端的角度考虑,以及还可以从你的角度考虑。
对于host, 它是你发布命令的序列(例如cudaMemcpyAsync和<<<>>>), 每个流都单独维护自己的序列。
对于device, 它主要是用来获得多个独立操作的途径,gpu内部可能还有自己的多个独立的依赖链(这个数目可能<你的host上看到的流的数目, 也就是可能会将多个流组合到内部的少量的独立依赖序列)。
对于你,因为流提供了多个独立的命令依赖关系,从来使得你有可能充分利用设备的各个硬件组件(例如同时利用copy engine和SM, 例如能多个SM能上多个小kernel等等)。