cudaStream 奇怪的表现答案

【问题标题】：cudaStream strange performancecudaStream 奇怪的表现
【发布时间】：2016-06-16 06:44:12
【问题描述】：

我尝试用 cudaStream 开发一个 sobel 的例子。这是程序：

void SobelStream(void)
{

    cv::Mat imageGrayL2 = cv::imread("/home/xavier/Bureau/Image1.png",0);


    u_int8_t *u8_PtImageHost;
    u_int8_t *u8_PtImageDevice;

    u_int8_t *u8_ptDataOutHost;
    u_int8_t *u8_ptDataOutDevice;

    u_int8_t u8_Used[NB_STREAM];

    u8_ptDataOutHost = (u_int8_t *)malloc(WIDTH*HEIGHT*sizeof(u_int8_t));
    checkCudaErrors(cudaMalloc((void**)&u8_ptDataOutDevice,WIDTH*HEIGHT*sizeof(u_int8_t)));

    u8_PtImageHost = (u_int8_t *)malloc(WIDTH*HEIGHT*sizeof(u_int8_t));
    checkCudaErrors(cudaMalloc((void**)&u8_PtImageDevice,WIDTH*HEIGHT*sizeof(u_int8_t)));


    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned char>();
    checkCudaErrors(cudaMallocArray(&Array_PatchsMaxDevice, &channelDesc,WIDTH,HEIGHT ));
    checkCudaErrors(cudaBindTextureToArray(Image,Array_PatchsMaxDevice));


    dim3 threads(BLOC_X,BLOC_Y);
    dim3 blocks(ceil((float)WIDTH/BLOC_X),ceil((float)HEIGHT/BLOC_Y));

    ClearKernel<<<blocks,threads>>>(u8_ptDataOutDevice,WIDTH,HEIGHT);


    int blockh = HEIGHT/NB_STREAM;


    Stream = (cudaStream_t *) malloc(NB_STREAM * sizeof(cudaStream_t));

    for (int i = 0; i < NB_STREAM; i++)
    {
        checkCudaErrors(cudaStreamCreate(&(Stream[i])));
    }

//    for(int i=0;i<NB_STREAM;i++)
//    {
//        cudaSetDevice(0);
//        cudaStreamCreate(&Stream[i]);
//    }


    cudaEvent_t Start;
    cudaEvent_t Stop;
    cudaEventCreate(&Start);
    cudaEventCreate(&Stop);

    cudaEventRecord(Start, 0);


    //////////////////////////////////////////////////////////
    for(int i=0;i<NB_STREAM;i++)
    {
        if(i == 0)
        {
            int localHeight  = blockh;
            checkCudaErrors(cudaMemcpy2DToArrayAsync( Array_PatchsMaxDevice,
                                                      0,
                                                      0,
                                                      imageGrayL2.data,//u8_PtImageDevice,
                                                      WIDTH,
                                                      WIDTH,
                                                      blockh,
                                                      cudaMemcpyHostToDevice  ,
                                                      Stream[i]));

            dim3 threads(BLOC_X,BLOC_Y);
            dim3 blocks(ceil((float)WIDTH/BLOC_X),ceil((float)localHeight/BLOC_Y));
            SobelKernel<<<blocks,threads,0,Stream[i]>>>(u8_ptDataOutDevice,0,WIDTH,localHeight-1);
            checkCudaErrors(cudaGetLastError());

            u8_Used[i] = 1;

        }else{


            int ioffsetImage =  WIDTH*(HEIGHT/NB_STREAM  );
            int hoffset = HEIGHT/NB_STREAM *i;
            int hoffsetkernel = HEIGHT/NB_STREAM -1 + HEIGHT/NB_STREAM* (i-1);
            int localHeight  = min(HEIGHT - (blockh*i),blockh);

            //printf("hoffset: %d hoffsetkernel %d localHeight %d rest %d ioffsetImage %d \n",hoffset,hoffsetkernel,localHeight,HEIGHT - (blockh +1 +blockh*(i-1)),ioffsetImage*i/WIDTH);

            checkCudaErrors(cudaMemcpy2DToArrayAsync( Array_PatchsMaxDevice,
                                                      0,
                                                      hoffset,
                                                      &imageGrayL2.data[ioffsetImage*i],//&u8_PtImageDevice[ioffset*i],
                            WIDTH,
                            WIDTH,
                            localHeight,
                            cudaMemcpyHostToDevice  ,
                            Stream[i]));


            u8_Used[i] = 1;
            if(HEIGHT - (blockh +1 +blockh*(i-1))<=0)
            {
                break;
            }
        }
    }



    ///////////////////////////////////////////
    for(int i=0;i<NB_STREAM;i++)
    {
        if(i == 0)
        {
            int localHeight  = blockh;


            dim3 threads(BLOC_X,BLOC_Y);
            dim3 blocks(1,1);
            SobelKernel<<<blocks,threads,0,Stream[i]>>>(u8_ptDataOutDevice,0,WIDTH,localHeight-1);
            checkCudaErrors(cudaGetLastError());

            u8_Used[i] = 1;

        }else{


            int ioffsetImage =  WIDTH*(HEIGHT/NB_STREAM  );
            int hoffset = HEIGHT/NB_STREAM *i;
            int hoffsetkernel = HEIGHT/NB_STREAM -1 + HEIGHT/NB_STREAM* (i-1);
            int localHeight  = min(HEIGHT - (blockh*i),blockh);


            dim3 threads(BLOC_X,BLOC_Y);
            dim3 blocks(1,1);

            SobelKernel<<<blocks,threads,0,Stream[i]>>>(u8_ptDataOutDevice,hoffsetkernel,WIDTH,localHeight);
            checkCudaErrors(cudaGetLastError());

            u8_Used[i] = 1;
            if(HEIGHT - (blockh +1 +blockh*(i-1))<=0)
            {
                break;
            }
        }
    }


    ///////////////////////////////////////////////////////
    for(int i=0;i<NB_STREAM;i++)
    {
        if(i == 0)
        {
            int localHeight  = blockh;
            checkCudaErrors(cudaMemcpyAsync(u8_ptDataOutHost,u8_ptDataOutDevice,WIDTH*(localHeight-1)*sizeof(u_int8_t),cudaMemcpyDeviceToHost,Stream[i]));
            u8_Used[i] = 1;

        }else{

            int ioffsetImage =  WIDTH*(HEIGHT/NB_STREAM  );
            int hoffset = HEIGHT/NB_STREAM *i;
            int hoffsetkernel = HEIGHT/NB_STREAM -1 + HEIGHT/NB_STREAM* (i-1);
            int localHeight  = min(HEIGHT - (blockh*i),blockh);

            checkCudaErrors(cudaMemcpyAsync(&u8_ptDataOutHost[hoffsetkernel*WIDTH],&u8_ptDataOutDevice[hoffsetkernel*WIDTH],WIDTH*localHeight*sizeof(u_int8_t),cudaMemcpyDeviceToHost,Stream[i]));

            u8_Used[i] = 1;
            if(HEIGHT - (blockh +1 +blockh*(i-1))<=0)
            {
                break;
            }
        }
    }


    for(int i=0;i<NB_STREAM;i++)
    {
        cudaStreamSynchronize(Stream[i]);
    }

    cudaEventRecord(Stop, 0);

    cudaEventSynchronize(Start);
    cudaEventSynchronize(Stop);


    float dt_ms;
    cudaEventElapsedTime(&dt_ms, Start, Stop);

    printf("dt_ms %f \n",dt_ms);

}

我在执行我的程序时有一个非常奇怪的表现。我决定分析我的例子，我明白了：

我不明白似乎每个流都在互相等待。有人可以帮我解决这个问题吗？

【问题讨论】：

问题在哪里？有什么意外？

标签： cuda

【解决方案1】：

首先，以后请提供完整的代码。我还在处理您的交叉发布 here 以填写一些细节，例如内核大小。

您有两个问题需要解决：

首先，任何时候您希望使用cudaMemcpyAsync，您很可能希望使用固定主机分配。如果您使用创建的分配，例如使用malloc，就异步并发执行而言，您不会从cudaMemcpyAsync 获得预期的行为。 programming guide：

如果复制涉及到主机内存，它必须是页面锁定的。

因此，对您的代码进行的第一个更改是将其转换为：

u8_PtImageHost   = (u_int8_t *)malloc(WIDTH*HEIGHT*sizeof(u_int8_t));
u8_ptDataOutHost = (u_int8_t *)malloc(WIDTH*HEIGHT*sizeof(u_int8_t));

到这里：

checkCudaErrors(cudaHostAlloc(&u8_PtImageHost, WIDTH*HEIGHT*sizeof(u_int8_t), cudaHostAllocDefault));
checkCudaErrors(cudaHostAlloc(&u8_ptDataOutHost, WIDTH*HEIGHT*sizeof(u_int8_t), cudaHostAllocDefault));

根据我的测试，仅通过该更改，您的执行时间就会从大约 21 毫秒下降到 7 毫秒。这样做的原因是，如果不进行更改，我们将不会有任何重叠：

随着更改，复制活动可以相互重叠（H->D 和 D->H）并与内核执行重叠：

您面临并发内核执行的第二个问题是您的内核太大（太多块/线程）：

#define WIDTH   6400
#define HEIGHT  4800
#define NB_STREAM 10

#define BLOC_X 32
#define BLOC_Y 32

    dim3 threads(BLOC_X,BLOC_Y);
    dim3 blocks(ceil((float)WIDTH/BLOC_X),ceil((float)HEIGHT/BLOC_Y));

我建议，如果这些是您需要运行的内核大小，那么尝试争取内核重叠可能没有多大好处 - 每个内核都在启动足够的块来“填充”GPU，所以您已经暴露了足够的并行性以保持 GPU 忙碌。但是，如果您迫切希望见证内核并发性，您可以让您的内核使用较小数量的块，同时让每个内核花费更多时间执行。我们可以通过启动 1 个块来做到这一点，并让每个块中的线程执行图像过滤。

【讨论】：