在一个循环里调用kernel时,传入不同的stream ,那么这些不同stream 是并行执行的吗
LZ您好:
不同stream之间是不保证先后顺序的,但不保证各个stream的kernel在同一时刻都在GPU上跑(可能在也可能不在)。
以及请您提供您的具体写法以及具体的对“并行执行”的定义,以便进一步讨论。
大致如此,祝您好运~
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
LZ您好:
stream一般不是您这样使用的。
一般可以用stream分离多个逻辑依赖链,使不同逻辑依赖链/stream里面的计算和传输可以互相掩盖。
以及不用每次循环都创建一个stream然后再销毁掉,创建和销毁stream是有开销的。
以及不必启动多少次kernel就创建多少次stream。
以及单纯调用kernel不含数据传输的话,除了线程规模很小的kernel可以同时在GPU上跑以外,其他情况都是一个kernel结束另外一个kernel才开始执行的。而LZ自己用函数封装了kernel调用,因此无法继续脑补了。
以及,LZ仍未解释1#中“并行执行”所想要表达的含义。
大致如此,请LZ重新参阅Programming guide有关stream的内容。
祝您好运~
我想说的“并行执行” 的意思是 给每个kernel 传不同stream,这些有不同stream的kernel是否在同时执行,
比如
{
kernel1<<<1,1,0,stream1>>>
kernel2<<<1,1,0,stream2>>>
}
kernel1 是否和 kernel2 在同时执行 。
看了斑竹的2#的意思,应该是kernel1 和 kernel2 有可能同时 ,也有可能没有同时执行。那就是说上
面这种使用方法有可能和使用默认stream的效率相当,或者比默认的快?
如果kernel1已经可以把GPU SM 资源都用上了,kernel2就要等一等。kernel2就要等有空闲的SM时才能开始执行。
LZ您好:
因为您在我4#回复之后,编辑了3#,重新提供了您自己封装函数的信息,这里做一定的建议:
1:您在封装的函数中有cudaMalloc()和cudaFree(),这使得您每次调用该封装函数的时候,都会进行申请和释放缓冲区。一般来说这个是不必要的,请您尽量一次性申请,然后多次复用。
2:您使用了cudaMemcpy()的同步版本,这将无法实现计算和传输互相掩盖,即便他们在不同的stream里面。
以及,请补充新内容的时候开新的回帖,如果实在要编辑原帖的话,请注明哪些部分是添加/改动的,以利于看帖人理清逻辑。
祝您好运~
以及,并不保证使用/不使用某种特性就一定能直接加速您的程序的,您需要根据实际情况,合理地设计一个实现,才能发挥GPU的效能。
希望您能明白这一点。
LZ您好:
4#和6#都已经为您解释了这一点,您也可以参考一下CUDA Samples里面的相关内容。
祝您好运~
恩,理解了kernel 运行的一些基本原理,谢谢 ice 斑竹 和 zehuanwang !