为什么不可以同时进行cudaMemcpyAsync（主机到设备）和CUDA内核？

Question

我已经加载了一张尺寸为1080 x 1920的图像（8位，无符号字符）。为了测试目的，我使用一个for循环对同一张图片进行了4次处理，然后生成了它的时间线分析。

策略：我将图像分成了3个部分。我为处理整个图像创建了三个流。

我提供了下面的最小工作示例。很抱歉它需要使用OpenCV加载图片，但我不知道如何在没有使用OpenCV加载图片的情况下模拟相同的情况。

问题：时间线分析显示第一个流已经完成传输数据，但是它分配的内核仍未启动。第一个流分配的内核和第三个流的数据传输是并行的。那么，我的问题是为什么第一个流的内核处理没有与第二个流的数据传输并行开始？

GPU：NVIDIA Quadro K2000，兼容3.0

时间线分析：每个流都被分配了不同的颜色。

我的代码：

__global__ void multiStream_ColorTransformation_kernel(int numChannels, int iw, int ih, unsigned char *ptr_source, unsigned char *ptr_dst)
{
    // Calculate our pixel's location
    int x = (blockIdx.x * blockDim.x) + threadIdx.x;
    int y = (blockIdx.y * blockDim.y) + threadIdx.y;

    // Operate only if we are in the correct boundaries
    if (x >= 0 && x < iw && y >= 0 && y < ih / 3)
    {
        ptr_dst[numChannels*  (iw*y + x) + 0] = ptr_source[numChannels*  (iw*y + x) + 0];
        ptr_dst[numChannels*  (iw*y + x) + 1] = ptr_source[numChannels*  (iw*y + x) + 1];
        ptr_dst[numChannels*  (iw*y + x) + 2] = ptr_source[numChannels*  (iw*y + x) + 2];

    }
}

void callMultiStreamingCudaKernel(unsigned char *dev_src, unsigned char *dev_dst, int numChannels, int iw, int ih, cudaStream_t *ptr_stream)
{

    dim3 numOfBlocks((iw / 20), (ih / 20)); //DON'T multiply by 3 because we have 1/3 data of image
    dim3 numOfThreadsPerBlocks(20, 20);
    multiStream_ColorTransformation_kernel << <numOfBlocks, numOfThreadsPerBlocks, 0, *ptr_stream >> >(numChannels, iw, ih, dev_src, dev_dst);

    return;
}

int main()
{

    cudaStream_t stream_one;
    cudaStream_t stream_two;
    cudaStream_t stream_three;

    cudaStreamCreate(&stream_one);
    cudaStreamCreate(&stream_two);
    cudaStreamCreate(&stream_three);

    Mat image = imread("DijSDK_test_image.jpg", 1);
    //Mat image(1080, 1920, CV_8UC3, Scalar(0,0,255));
    size_t numBytes = image.rows * image.cols * 3;
    int numChannels = 3;

    int iw = image.rows;
    int ih = image.cols;
    size_t totalMemSize = numBytes * sizeof(unsigned char);
    size_t oneThirdMemSize = totalMemSize / 3;

    unsigned char *dev_src_1, *dev_src_2, *dev_src_3, *dev_dst_1, *dev_dst_2, *dev_dst_3, *h_src, *h_dst;


    //Allocate memomry at device for SOURCE and DESTINATION and get their pointers
    cudaMalloc((void**)&dev_src_1, (totalMemSize) / 3);
    cudaMalloc((void**)&dev_src_2, (totalMemSize) / 3);
    cudaMalloc((void**)&dev_src_3, (totalMemSize) / 3);
    cudaMalloc((void**)&dev_dst_1, (totalMemSize) / 3);
    cudaMalloc((void**)&dev_dst_2, (totalMemSize) / 3);
    cudaMalloc((void**)&dev_dst_3, (totalMemSize) / 3);

    //Get the processed image 
    Mat org_dijSDK_img(image.rows, image.cols, CV_8UC3, Scalar(0, 0, 255));
    h_dst = org_dijSDK_img.data;

    //while (1)
    for (int i = 0; i < 3; i++)
    {
        std::cout << "\nLoop: " << i;

        //copy new data of image to the host pointer
        h_src = image.data;

        //Copy the source image to the device i.e. GPU
        cudaMemcpyAsync(dev_src_1, h_src, (totalMemSize) / 3, cudaMemcpyHostToDevice, stream_one);
        //KERNEL--stream-1
        callMultiStreamingCudaKernel(dev_src_1, dev_dst_1, numChannels, iw, ih, &stream_one);


        //Copy the source image to the device i.e. GPU
        cudaMemcpyAsync(dev_src_2, h_src + oneThirdMemSize, (totalMemSize) / 3, cudaMemcpyHostToDevice, stream_two);
        //KERNEL--stream-2
        callMultiStreamingCudaKernel(dev_src_2, dev_dst_2, numChannels, iw, ih, &stream_two);

        //Copy the source image to the device i.e. GPU
        cudaMemcpyAsync(dev_src_3, h_src + (2 * oneThirdMemSize), (totalMemSize) / 3, cudaMemcpyHostToDevice, stream_three);
        //KERNEL--stream-3
        callMultiStreamingCudaKernel(dev_src_3, dev_dst_3, numChannels, iw, ih, &stream_three);


        //RESULT copy: GPU to CPU
        cudaMemcpyAsync(h_dst, dev_dst_1, (totalMemSize) / 3, cudaMemcpyDeviceToHost, stream_one);
        cudaMemcpyAsync(h_dst + oneThirdMemSize, dev_dst_2, (totalMemSize) / 3, cudaMemcpyDeviceToHost, stream_two);
        cudaMemcpyAsync(h_dst + (2 * oneThirdMemSize), dev_dst_3, (totalMemSize) / 3, cudaMemcpyDeviceToHost, stream_three);

        // wait for results 
        cudaStreamSynchronize(stream_one);
        cudaStreamSynchronize(stream_two);
        cudaStreamSynchronize(stream_three);

        //Assign the processed data to the display image.
        org_dijSDK_img.data = h_dst;
        //DISPLAY PROCESSED IMAGE           
        imshow("Processed dijSDK image", org_dijSDK_img);
        waitKey(33);
    }

    cudaDeviceReset();
    return 0;
}

更新-1：如果我删除第一个流的内核调用，那么第二个内核和第三个流的H2D复制会以某种方式重叠（但不是完全重叠），如下所示。

更新-2：我甚至尝试使用10个流，情况仍然保持不变。只有在第十个流的H2D复制后，第一个流的内核处理才开始。

- skm

---

但是你正在展示的那个配置文件数据中，每个流中都发生了复制/执行重叠。 - talonmies

主机内存是否被固定？该文档说明必须对主机内存进行页面锁定才能实现重叠。 - Jez

如果您还没有这样做，我建议阅读CUDA C编程指南的3.2.5节，其中讨论了并发执行以及可能涉及的所有奇怪问题。特别是，对于CC小于等于3.0的设备，有一些额外的限制适用。 - reirab

OpenCV有自己的方法将Mat数据复制到GPU，可以查看gpu :: CudaMem。如果您不想使用这些方法，也许您需要深度复制数据或使用固定内存作为cv :: Mat或其他输入指针的输入。虽然可能完全错误，但我尚未将openCV与Cuda结合使用。 - Micka

你的原始无符号字符数据来自哪里？你不能提供一些内存区域，在那里它将首先被“生成”吗？通常相机API等会提供这样的功能。 - Micka

Answer 1

正如评论者所指出的，主机内存必须是页面锁定。

不需要通过cudaHostAlloc分配额外的主机内存，您可以在现有的OpenCV图像上使用cudaHostRegister。

cudaHostRegister(image.data, totalMemSize, cudaHostRegisterPortable)

- m.s.

---

实际上，我的图像不会是OpenCV图像。我将以 unsigned char形式获取图像数据，然后需要对其进行处理。我通过在主机上分配可固定的内存并使用 memcpy() 将我的图像数据复制到此固定内存中来解决了这个问题。我担心 memcpy() 是否是一种有效的方法。 - skm

@skm，您仍然可以使用cudaHostRegister，它不依赖于OpenCV图像。只需将您的unsigned char*指针作为第一个参数传递即可。 - m.s.

你能告诉我，是否也可以重叠内核执行吗？我尝试使用20个流来减小内核大小，但仍然无法重叠内核处理。 - skm

@skm 这取决于您的GPU和内核启动配置："为给定内核的线程块进行调度，如果所有前面内核的线程块都已经被调度，并且仍然有SM资源可用"，同时也请阅读此文档。 - m.s.