初步接触CUDA编程，需要做kernel函数的并行，在手册上看到计算能力2.0以上的卡是支持最多16个kernel函数并行，用Telsa C2050的卡，测试是能够支持多个kernel函数并行的。用SDK中自带的例Concurrent可以并行，但是自己编写的不能并行，是通过stream编写的，还请各位大神们能不吝赐教，万分感激！！！:loveliness:
代码在那里？难道你让我们猜？好吧，我就来猜：
一、用了0号流。
二、不适当的同步。
三、流中中间的某些代码有隐藏等待
是应该用多流实现的，我刚写了一个，仿照SDK中的concurrent是可以实现并发的。
楼主还是贴代码吧，以后也得注意，用代码表达清楚。
建议仔细读读SDK里concurrent kernel的例子，看看与你的程序有什么不同的地方。另一个原因是concurrent kernel只有当每个kernel都比较小时才有效果。如果每个kernel都能用满GPU，则concurrent kernel自然没有用了。
呵呵，楼上说道要害了，估计楼主的问题应该是kernel规模太大了。
我的问题确实是kernel规模很大，我也试过例子程序，是可以并发执行的，<<<128,512>>>这个规模以后基本上就不能并行了。可是现在是要做到实时，数据量很大，kernel函数不能并行的话就不能实现，请问下这个除了并行还有其他的优化空间没。。。
我把代码贴下，欢迎大家批评指教~~~3Q!!!
for(j=0;j<SINGNAL_DIR_COUNT;j++)
{
//调用内核函数，进行err[i]与d_refer_signal_slip2点乘,将结果放入d_temp矩阵
for (i=0;i<FREQUENCY_NUM;i++)
DotProduct<<<grid1,block1,0,stream[i]>>>(SINGNAL_COUNT, d_target_signal+iSINGNAL_COUNTSINGNAL_DIR_COUNT+jSINGNAL_COUNT, d_refer_signal_slip2+iSINGNAL_COUNTSLIP_STEP2, d_refer_signal_slip2+iSINGNAL_COUNTSLIP_STEP2);
//调用内核函数，d_refer_signal_slip2每一列进行抽取个数据,将结果放入d_extract矩阵中
for (i=0;i<FREQUENCY_NUM;i++)
ExtractFilter<<<grid2,block2,0,stream[i]>>>(d_refer_signal_slip2+iSINGNAL_COUNTSLIP_STEP2, SINGNAL_COUNT, d_coeff+iCOEFF_SIZE, d_extract+iFFT_NUMSLIP_STEP2);
//对矩阵进行fft操作
for (i=0;i<FREQUENCY_NUM;i++)
cufftSafeCall(cufftExecC2C(plan[i], (cufftComplex )(d_extract+iFFT_NUMSLIP_STEP2), (cufftComplex )(d_extract+iFFT_NUMSLIP_STEP2), CUFFT_FORWARD));
//fftshift
for (i=0;i<FREQUENCY_NUM;i++)
fftshift_kernel<<<grid3,block3,0,stream[i]>>>(d_extract+iFFT_NUMSLIP_STEP2, FFT_NUM, (FFT_NUM/(2*BLOCK_SIZE_1));
//求模
for (i=0;i<FREQUENCY_NUM;i++)
Complex2Float<<<grid4,block4,0,stream[i]>>>(d_extract+iFFT_NUMSLIP_STEP2, FFT_NUM, SLIP_STEP2, d_final+iFFT_NUMSLIP_STEP2SINGNAL_DIR_COUNT+jFFT_NUM*SLIP_STEP2 );
}
这个是信号处理中做“距离-多普勒处理”的代码，规模比较大，很耗时，是用流写得，不能实现并行。。。不知道有什么更好的方式来优化。。。。求指导。。。。:loveliness:
你现在用流可以但是没多大意义，当kernel规模大的时候，kernel不能并发，只不过是见缝插针，效果非常小。所以，不必要用流，只能从其它方面优化：比如合并访问，bank冲突等技巧。
When a kernel is invoked or launched,it is executed as grid of parallel threads. Each CUDA thread grid typically is comprised of thousands to millions of lightweight GPU threads per kernel invocation.