不同stream 是并行执行的吗

加速计算 CUDA
1

在一个循环里调用kernel时，传入不同的stream ，那么这些不同stream 是并行执行的吗

2

LZ您好：

不同stream之间是不保证先后顺序的，但不保证各个stream的kernel在同一时刻都在GPU上跑（可能在也可能不在）。

以及请您提供您的具体写法以及具体的对“并行执行”的定义，以便进一步讨论。

大致如此，祝您好运~

3 
for(int i = 0 ;i < n; i++)
{
   ......
   for(int iband = 0 ;iband < nBands; iband++){
   
   cudaStream_t stream;
   cudaStreamCreate( &stream );

   ors_byte* inbuf = pImgData->getBandBuf(iband);
   memset(outBuf,0,wid*hei/*npoolSize*/ * sizeof(BYTE));

   orthCorrection( dev_inBuf,dev_outBuf,inbuf, outBuf, nWid ,nHei,wid, hei, dem_X0, dem_Y0,dem_DX,dem_DY,  dem_row,dem_col,p_dem_data,        m_Boundary[n] ,bac, M,Xs,Ys,Zs,x0,y0,f,stream );

   outWriter->WriteBandRect( iband, 0, 0, wid, hei, outBuf);
   cudaStreamDestroy( stream );
   }
}

void orthCorrection(unsigned char* dev_inBuf,unsigned char* dev_outBuf, unsigned char* m_inBuf, unsigned char* m_outBuf, int nWid ,int nHei,int wid,int hei,
   double dem_X0,double dem_Y0,double dem_DX,double dem_DY, int dem_row,int dem_col,float* p_dem_data,Boundary m_Boundary ,
   rMatrix bac, double M,         double Xs,double Ys ,double Zs,double x0,double y0,double f,cudaStream_t stream)
{

   cudaError_t err;
   
   long n_inSize = nWid*nHei*sizeof(unsigned char);
   long n_outSize = wid*hei*sizeof(unsigned char);

   err = cudaMemcpy( dev_inBuf, m_inBuf , n_inSize, cudaMemcpyHostToDevice );

   err = cudaMemcpy( dev_outBuf ,m_outBuf, n_outSize, cudaMemcpyHostToDevice);
   
   float* dev_dem_data;
   long dem_data_size = dem_row*dem_col*sizeof(float);

   err = cudaMalloc((void**)&dev_dem_data, dem_data_size);
   err = cudaMemcpy( dev_dem_data, p_dem_data,dem_data_size,cudaMemcpyHostToDevice );
   
   kernel_orthCorrection<<<BLOCK_NUM,THREAD_NUM,0,stream>>>( dev_inBuf, dev_outBuf, nWid ,nHei,wid,hei,
   dem_X0,dem_Y0,dem_DX,dem_DY, dem_row,dem_col,dev_dem_data,        m_Boundary ,bac,M, Xs,Ys ,Zs,x0, y0, f);
   
   err = cudaMemcpy( m_outBuf ,dev_outBuf , n_outSize, cudaMemcpyDeviceToHost );

   cudaFree( dev_dem_data );
}

代码是这样的：
有两层循环，在第二层循环里用stream ，结果时间比不用还长一些。orthoCorrection 里调用了kernel

4

LZ您好：

stream一般不是您这样使用的。

一般可以用stream分离多个逻辑依赖链，使不同逻辑依赖链/stream里面的计算和传输可以互相掩盖。
以及不用每次循环都创建一个stream然后再销毁掉，创建和销毁stream是有开销的。
以及不必启动多少次kernel就创建多少次stream。
以及单纯调用kernel不含数据传输的话，除了线程规模很小的kernel可以同时在GPU上跑以外，其他情况都是一个kernel结束另外一个kernel才开始执行的。而LZ自己用函数封装了kernel调用，因此无法继续脑补了。

以及，LZ仍未解释1#中“并行执行”所想要表达的含义。

大致如此，请LZ重新参阅Programming guide有关stream的内容。

祝您好运~

5

我想说的“并行执行” 的意思是 给每个kernel 传不同stream，这些有不同stream的kernel是否在同时执行，

比如
{
kernel1<<<1,1,0,stream1>>>

kernel2<<<1,1,0,stream2>>>
}
kernel1 是否和 kernel2 在同时执行 。

看了斑竹的2#的意思，应该是kernel1 和 kernel2 有可能同时 ，也有可能没有同时执行。那就是说上

面这种使用方法有可能和使用默认stream的效率相当，或者比默认的快？

6

如果kernel1已经可以把GPU SM 资源都用上了，kernel2就要等一等。kernel2就要等有空闲的SM时才能开始执行。

7

LZ您好：

因为您在我4#回复之后，编辑了3#，重新提供了您自己封装函数的信息，这里做一定的建议：

1：您在封装的函数中有cudaMalloc()和cudaFree()，这使得您每次调用该封装函数的时候，都会进行申请和释放缓冲区。一般来说这个是不必要的，请您尽量一次性申请，然后多次复用。

2：您使用了cudaMemcpy()的同步版本，这将无法实现计算和传输互相掩盖，即便他们在不同的stream里面。

以及，请补充新内容的时候开新的回帖，如果实在要编辑原帖的话，请注明哪些部分是添加/改动的，以利于看帖人理清逻辑。

祝您好运~

8

以及，并不保证使用/不使用某种特性就一定能直接加速您的程序的，您需要根据实际情况，合理地设计一个实现，才能发挥GPU的效能。

希望您能明白这一点。

9

LZ您好：

4#和6#都已经为您解释了这一点，您也可以参考一下CUDA Samples里面的相关内容。

祝您好运~

10

恩，理解了kernel 运行的一些基本原理，谢谢 ice 斑竹 和 zehuanwang ！