请问大家，在CUDAProf中，我想用流异步传输，但是，结果是所有流的异步传输结束后才开始进行计算，为什么不是一个流的数据加载完毕后这个流就进行计算呢？我不知道为什么我的程序用流异步传输没有效果……大家可以帮帮忙指点一下吗？书上的例子都是：加载数据用一个for循环，计算用一个for循环，但是为什么不是只用一个for循环呢？ 谢谢指点。。。
附件是用CUDAProf中的GPU Time Width Plot截图，谢谢帮忙指点。。。
使用cudaprofiler的时候所有的函数都是同步的。
谢谢您！我还想问您，0.8GB/s的glob mem read throughput是不是很低呀？您看看我这样分析对吗：
一个warp中的线程各自需要处理32bit字：
如果，让整个warp中的线程每人读自己的32bit字，结果half-warp中的线程一次操作，就只是64Bytes段的合并访问，所以一个warp指令，两次64Bytes段的合并访问；
如果，让half-warp中的线程每人读2*32bit字，另外half-warp不读，就是128Bytes段的合并访问，所以一个warp指令，一次128Bytes段的合并访问，带宽利用率会提高；但是，它们读完数据还得放到shared中，然后再分别读到自己的寄存器中，还会有warp内的一次分支，这样好吗？
具体要和你的卡比较，但是０．８是相当低的，一般此时意味着全局存储器读写问题。
你的分析大部分是没错的。只是“它们读完数据还得放到shared中，然后再分别读到自己的寄存器中”，为什么这样做？
另：ｗａｒｐ内分支很影响性能
我的线程都是只负责处理一个32位字，所以，它本来从Global中读到自己的寄存器中就可以开始计算了，但是，它要是负责把别人的数据下载下来，就得先放在shared中，别人在从shared中读到它的寄存器中处理。
这样的warp内分支只是前一部分干活，后一部分的”支“什么也不做，这样会影响性能吗？
我的程序还有一个很大的问题，您能指点一下吗？我如果不用MAPPED | WRITE COMBINED，而是用Memcpy进行数据加载，这部分时间竟然占了总共GPU时间的75%，kernel用的时间只用25%，同时kernel 中又有从global读和写global的时间，这样，我的算法部分好像即使优化也没有很明显的提升。我本来想让算法的吞吐量至少达到1GB/s,但是发现数据传输的吞吐量还没有1GB/s，glob mem read throughout才只有0.8GB/s，想增加glob mem read throughout，我想，是不是就得让线程一次读更多的数据，但是，一个线程要处理一个32bit字，一次读的数据只能是32bit，要是让一个线程干更多的活，之前的数据加载就需要加载更多数据，计算单元就得等更长的时间，我想又影响全局的性能。。。我该怎么突破这个瓶颈呢？
我还不明白一个问题：为什么流这样用，不是写在一个for中呢？
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
MyKernel<<<100, 512, 0, stream[i]>>> (outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost,
stream[i]);
cudaThreadSynchronize();
[ 本帖最后由 bzz168 于 2010-5-17 22:39 编辑 ]
放一个循环里是可以的
cudaMemcpyAsync( indp, hp, size, cudaMemcpyHostToDevice, stream[ 0 ] );
for( unsigned int sid=0; sid<n_streams; ++sid )
{
kernel<<< grid_layout, block_layout, stream[ sid ] >>>( outdp+isize, indp+isize, … );
if( sid<n_streams-1 ){
cudaMemcpyAsync( indp+( i+1 )size, hp+( i+1 )size, size, cudaMemcpyHostToDevice, stream[ sid+1 ] );
}
cudaMemcpyAsync( hp+isize, outdp+isize, size, cudaMemcpyDeviceToHost, stream[ sid ] );
}
cudaMemcpyAsync( hp+( n_streams-1 )*size, dp+( n_streams-1 )*size, size, cudaMemcpyDeviceToHost, stream[ n_streams-1 ] );
可以说会，也可以说不会，说会是因为有一半线程空闲，说不会是因为分支的一半线程基本上不消耗时间。
另：放到一个循环更好，更容易理解。
数据传输时间一直是异构计算的问题之一，要想办法隐藏它的负作用。
[ 本帖最后由 yyfn风辰 于 2010-5-18 09:01 编辑 ]
之前：一个线程负责处理一个32位字，下载它的32位数据时，一个warp指令，前后 half-warp两次read global操作，只能是两次64Bytes的合并访问，结果是：（264Bytes）/（2读global时间）；
为了增大glob mem read throughout,我想用128Bytes的合并访问，所以：
之后：一个线程负责处理两个32位字，用64位的寄存器读global,之后再(64bitdata >> 32)转换，下载它的64位数据时，一个warp指令，前后half-warp两次都是128Bytes合并访问global，结果是：（2128Bytes）/(2读global时间）；所以glob mem read throughout应当增加一倍。
但是，实验的结果是之后和之前差不多，而且略微差了一点。后者中，cudaprof显示全部满足128Bytes合并访问,但是glob mem read throughout还是没有增加，运行时间和以前相比差不多，还稍微多了一点。
我是不是上面的分析有问题呢？
这个，还真不好解释，答案是你的分析是错的，但是怎么说，又不太好说清楚
我看什么文章可以有助于理解清楚呢？
引用：
cudaMemcpyAsync( indp, hp, size, cudaMemcpyHostToDevice, stream[ 0 ] );
for( unsigned int sid=0; sid<n_streams; ++sid )
{
kernel<<< grid_layout, block_layout, stream[ sid ] >>>( outdp+isize, indp+isize, … );
if( sid<n_streams-1 ){
cudaMemcpyAsync( indp+( i+1 )size, hp+( i+1 )size, size, cudaMemcpyHostToDevice, stream[ sid+1 ] );
}
cudaMemcpyAsync( hp+isize, outdp+isize, size, cudaMemcpyDeviceToHost, stream[ sid ] );
}
cudaMemcpyAsync( hp+( n_streams-1 )*size, dp+( n_streams-1 )*size, size, cudaMemcpyDeviceToHost, stream[ n_streams-1 ] );
请问，这样是不是因为，在store调用之前，有一个隐式的store fence，如果全部放在一个for中，就成了顺序的了。如果像您这样写，是不是就在计算这一流的数据时，进行下一流的数据的加载的，然后才是store fence，这样能实现两个流的并行？要是全部的分开，就可以使数据的加载的for循环中，只要一个流的数据加载完毕，下一个流就可以开始进行，从而实现多个流的并行呢？
在http://www.drdobbs.com/architecture-and-design/217500110介绍WC主机端内存中，有“It is unclear if and when a CUDA programmer needs to take any action (such as using a memory fence) to ensure that the WC memory is in-place and ready for use by the host or graphics processor(s). The Intel documentation states that “[a] ‘memory fence’ instruction should be used to properly ensure consistency between the data producer and data consumer.” The CUDA driver does use WC memory internally and must issue a store fence instruction whenever it sends a command to the GPU. For this reason, the NVIDIA documentation notes, “the application may not have to use store fences at all” (emphasis added). A rough rule of thumb that appears to work is to look to the CUDA commands prior to referencing WC memory and assume they issue a fence instruction. Otherwise, utilize your compiler intrinsic operations to issue a store fence instruction and guarantee that every preceding store is globally visible. This is compiler dependent. Linux compilers will probably understand the _mm_sfence intrinsic while Windows compilers will probably use _WriteBarrier. ”，我的理解是：如果用WC的内存，CUDA驱动会自动在异步设备向主机写数据的函数之前加入store fence，您能看看我的这种理解正确吗？只有WC方式才会这样吗？其它的时候可以全部放在一个for循环中，而不是您提供的那种方法吗？另外，流的个数一般怎样设定呢？我不知道为什么，无论要处理的数据量是多少，用4个流都没有用两个流的效果好，是不是更多的流只会增加CPU调用kernel函数、进行数据加载和指令加载的频率，而掩盖了异步所带来的性能提升了呢？我又担心我的流使用错了，所以没有显著的性能优化。。。
[ 本帖最后由 bzz168 于 2010-5-21 03:40 编辑 ]
这个还没有足够的资料支持。可以这样理解，就是GPU中取数单元取数的速度是有限的，因此取数越多，时间就越多。
没有隐式的fence