找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 125|回复: 2

关于stream中使用并发内核但结果却串行执行的问题

[复制链接]
发表于 2018-9-24 17:29:43 | 显示全部楼层 |阅读模式
GTC
我想要使用并发内核实现下图这样的效果,但是实际结果却如右图所示,请问这种包含有数据传输的情况下能否实现内核并发呢?。

GPU1.png GPU2.png //数据输入
for (int i = 0; i < NSTREAM; i++)
        {
                int Coffset = i * m_nClassNumber;
                int ioffset = i * iElem;
                int joffset = i*iElem;
                if (i == 2)
                {
                        ioffset = (i - 2) * iElem;
                        joffset = (i - 1)*iElem;
                }
                if (i == 3)
                {
                        ioffset = (i - 2) * iElem;
                        joffset = (i - 3)*iElem;
                }
                cudaMemcpyAsync(&cuda_Start_l[ioffset], &start_CL[ioffset], StartNum / 2 * sizeof(int), cudaMemcpyHostToDevice, streams);
                cudaMemcpyAsync(&cuda_Start_a[ioffset], &start_CA[ioffset], StartNum / 2 * sizeof(int), cudaMemcpyHostToDevice, streams);
                cudaMemcpyAsync(&cuda_Start_b[ioffset], &start_CB[ioffset], StartNum / 2 * sizeof(int), cudaMemcpyHostToDevice, streams);
                cudaMemcpyAsync(&cuda_End_l[joffset], &end_CL[joffset], EndNum / 2 * sizeof(int), cudaMemcpyHostToDevice, streams);
                cudaMemcpyAsync(&cuda_End_a[joffset], &end_CA[joffset], EndNum / 2 * sizeof(int), cudaMemcpyHostToDevice, streams);
                cudaMemcpyAsync(&cuda_End_b[joffset], &end_CB[joffset], EndNum / 2 * sizeof(int), cudaMemcpyHostToDevice, streams);
        }
       
//核函数
                ColorInter4 << < StartNum / 2, EndNum / 2, 0, streams[0] >> >(cuda_allkde, &cuda_Start_a[0], &cuda_Start_b[0], &cuda_Start_l[0],
                        &cuda_End_a[0], &cuda_End_b[0], &cuda_End_l[0], m_nClassNumber, cuda_KDE, totalA, totalB, totalL, cuda_Best_KDE, times, XMIN,
                        XMAX, YMIN, YMAX, ZMIN, ZMAX, InterpolationNum, Scount);
                ColorInter41 << < StartNum / 2, EndNum / 2, 0, streams[1] >> >(cuda_allkde, &cuda_Start_a[500], &cuda_Start_b[500], &cuda_Start_l[500],
                        &cuda_End_a[500], &cuda_End_b[500], &cuda_End_l[500], m_nClassNumber, cuda_KDE, totalA, totalB, totalL, cuda_Best_KDE, times, XMIN,
                        XMAX, YMIN, YMAX, ZMIN, ZMAX, InterpolationNum, Scount);
                ColorInter42 << < StartNum / 2, EndNum / 2, 0, streams[2] >> >(cuda_allkde, &cuda_Start_a[0], &cuda_Start_b[0], &cuda_Start_l[0],
                        &cuda_End_a[500], &cuda_End_b[500], &cuda_End_l[500], m_nClassNumber, cuda_KDE, totalA, totalB, totalL, cuda_Best_KDE, times, XMIN,
                        XMAX, YMIN, YMAX, ZMIN, ZMAX, InterpolationNum, Scount);
                ColorInter43 << < StartNum / 2, EndNum / 2, 0, streams[3] >> >(cuda_allkde, &cuda_Start_a[500], &cuda_Start_b[500], &cuda_Start_l[500],
                        &cuda_End_a[0], &cuda_End_b[0], &cuda_End_l[0], m_nClassNumber, cuda_KDE, totalA, totalB, totalL, cuda_Best_KDE, times, XMIN,
                        XMAX, YMIN, YMAX, ZMIN, ZMAX, InterpolationNum, Scount);


//数据输出
for (int i = 0; i < NSTREAM; i++)
        {
                int Coffset = i * m_nClassNumber;
                int ioffset = i * iElem;
                int joffset = i*iElem;
                if (i == 2)
                {
                        ioffset = (i - 2) * iElem;
                        joffset = (i - 1)*iElem;
                }
                if (i == 3)
                {
                        ioffset = (i - 2) * iElem;
                        joffset = (i - 3)*iElem;
                }
                cudaMemcpyAsync(&resultl[Coffset], &totalL[0], 5 * sizeof(int), cudaMemcpyDeviceToHost, streams);
                cudaMemcpyAsync(&resulta[Coffset], &totalA[0], 5 * sizeof(int), cudaMemcpyDeviceToHost, streams);
                cudaMemcpyAsync(&resultb[Coffset], &totalB[0], 5 * sizeof(int), cudaMemcpyDeviceToHost, streams);
                cudaMemcpyAsync(&host_allkde, &cuda_Best_KDE[0], sizeof(int), cudaMemcpyDeviceToHost, streams);
        }

回复

使用道具 举报

发表于 2018-9-24 21:02:30 | 显示全部楼层
Jetson TX2
主要有两方面可能会影响到因素:

(1)多个流中发布的kernel启动, 如果想能同时进行(例如SDK中的例子那样),必须每个kernel的启动规模足够小,能一堆kernel在你的卡上同时执行。较大的kernel,线程和blocks数量很多,则可能无法同时让多个这样的kernel们同时运行。

(2)cudaMemcpyAsync在特定的流中启动的时候(注意你这里的排版,你可能丢失了下标[ i ], 光能看到你写streams), 需要使用page-locked memory(例如cudaMallocHost出来的, 而不是普通malloc), 才可以。如果不慎使用了普通的可分页内存(page-able memory), 则会自动退化成普通的cudaMemcpy, 并将引入自动的隐式样同步效果, 可能会影响多个流中的kernel的并发(如果上面的第(1)点不构成限制条件,本身能够并发执行的话)。

节日快乐!
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-9-25 08:53:14 | 显示全部楼层
Tesla P100
万分感谢您的指导!也祝您节日快乐。
回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表