CPU 到 GPU 内存传输 - cudaMemcpy() 与使用 Map() 的 Direct3D 动态资源

【问题标题】：CPU to GPU memory transfer - cudaMemcpy() vs Direct3D dynamic resource with Map()CPU 到 GPU 内存传输 - cudaMemcpy() 与使用 Map() 的 Direct3D 动态资源
【发布时间】：2014-12-27 18:14:22
【问题描述】：

我有一个将 RGB32 帧编码为 H.264 的实时视频流管道。我的目标是 NVIDIA 硬件，所以我打算使用 CUDA 执行从 RGB32 到 NV12 的色彩空间转换。我查找了执行类似任务的内核示例，一切似乎都很好。然而，由于很多人提到数据传输速度是 CPU 到 GPU 通信的最关键点，我想知道是否有人有经验将 RGB32 数据提供给 CUDA 内核的更好方法：

  • 使用cudaMemcpy()（至少this 主题表明cudaMemcpy() 的性能优于操作系统图形堆栈
  • 使用在 cuda 注册并通过 Map() 从用户空间代码更新的动态 Direct3D11 资源

如果有人有这方面的经验，那么我很高兴听到它，否则 - 基准测试是:)

【问题讨论】：

    标签： cuda nvidia direct3d11


    【解决方案1】：

    由于没有任何活动，我冒昧地进行了基准测试，所以我将把所有内容都留在这里，以便任何人都可以使用它们或对改进发表评论。

    我比较了 1000 次迭代的时间：

    • Map/memcpy 到映射内存/Unmap 在每次迭代中动态 Direct3D 11 纹理 - 每次调用 3 毫秒
    • Map/Unmap 在每次迭代中使用动态 Direct3D 11 纹理（以获得映射/取消映射开销的概念） - 每次调用 1.4 毫秒
    • UpdateSubresource 在每次迭代中使用默认的 Direct3D 11 纹理（根据我的阅读，如果每帧有多个更新，这应该比动态表面慢） - 每次调用 2.13 毫秒
    • cudaMemcpy 从分配有 new 的临时指针到每次迭代中分配的 cudaMalloc 设备内存指针 - 每次调用 1.3 毫秒
    • cudaMemcpyAsync 从分配有new 的临时指针到每次迭代中分配的cudaMalloc 设备内存指针和最后一次迭代后的cudaDeviceSynchronize - 每次调用1.25ms
    • cudaMemcpyAsync 从 cudaMalloc 分配的主机内存指针到 cudaMalloc 分配的设备内存指针在每次迭代和最后一次迭代后 cudaDeviceSynchronize - 0.250ms 每次调用

    基本上，我似乎应该坚持使用 Cuda，因为它比使用 Direct3D 11 表面将数据从系统内存传输到 GPU 内存更快。

    在 Map/Unmap 本身很少被调用的情况下，Map/Unmap 本身被调用时，似乎Map/Unmap 方法胜过默认表面和UpdateSubresource。

    我将在下面发布基准代码（也可以在 GitHub 上获得） - 我很高兴收到任何反馈，因为基准可能存在问题，这可能会影响结果，因为我是新手到 Direct3D 11 和 Cuda。

    // STL
#include <iostream>
#include <cstdlib>
#include <memory>
#include <vector>

// ATL
#include <atlbase.h>

// CUDA
#include "cuda.h"
#include "cuda_runtime_api.h"

#pragma comment(lib, "cudart.lib")

// DXGI
#include <dxgi.h>
#pragma comment(lib, "dxgi.lib")

// D3D11
#include <d3d11.h>
#pragma comment(lib, "d3d11.lib")


int main(int argc, char** argv)
{
    std::string sDeviceName("GeForce GTX 750 Ti");
    std::wstring sDeviceNameWide(sDeviceName.begin(), sDeviceName.end());
    const size_t nWidth = 1920, nHeight = 1080, nIterations = 1000;
#pragma region Direct3D 11
    CComPtr<IDXGIFactory1> pDXGIFactory1;
    ATLENSURE_SUCCEEDED(CreateDXGIFactory1(__uuidof(IDXGIFactory1), reinterpret_cast<void**>(&pDXGIFactory1)));
    ULONG nAdapterIndex = 0;
    CComPtr<IDXGIAdapter1> pDXGIAdapter1;
    DXGI_ADAPTER_DESC1 DXGIAdapterDescription1 = {};
    bool bD3D11AdapterFound = false;
    while (SUCCEEDED(pDXGIFactory1->EnumAdapters1(nAdapterIndex++, &pDXGIAdapter1)))
    {
        ATLENSURE_SUCCEEDED(pDXGIAdapter1->GetDesc1(&DXGIAdapterDescription1));
        std::wstring sDescription(DXGIAdapterDescription1.Description);
        if (sDescription.find(sDeviceNameWide) != std::string::npos)
        {
            bD3D11AdapterFound = true;
            break;
        }
    }
    if (bD3D11AdapterFound == false)
    {
        std::cout << "Direct3D 11 compatbile adapter named " << sDeviceName.c_str() << "was not found!" << std::endl;
        return EXIT_FAILURE;
    }
    const D3D_FEATURE_LEVEL RequestedFeatureLevels = D3D_FEATURE_LEVEL_11_0;
    D3D_FEATURE_LEVEL FeatureLevel;
    UINT nFlags = 0;
#ifdef _DEBUG
    nFlags |= D3D11_CREATE_DEVICE_DEBUG;
#endif
    CComPtr<ID3D11Device> pDevice;
    CComPtr<ID3D11DeviceContext> pDeviceContext;
    ATLENSURE_SUCCEEDED(D3D11CreateDevice(pDXGIAdapter1, D3D_DRIVER_TYPE_UNKNOWN, NULL, nFlags, &RequestedFeatureLevels, 1, D3D11_SDK_VERSION, &pDevice, &FeatureLevel, &pDeviceContext));
    std::unique_ptr<unsigned char[]> pFrame(new unsigned char[nWidth * nHeight * 3 / 2]);
    D3D11_TEXTURE2D_DESC TextureDescription = {};
    TextureDescription.Width = nWidth;
    TextureDescription.Height = nHeight;
    TextureDescription.Format = DXGI_FORMAT_NV12;
    TextureDescription.CPUAccessFlags = D3D11_CPU_ACCESS_WRITE;
    TextureDescription.Usage = D3D11_USAGE_DYNAMIC;
    TextureDescription.MipLevels = 1;
    TextureDescription.ArraySize = 1;
    TextureDescription.SampleDesc.Count = 1;
    TextureDescription.BindFlags = D3D11_BIND_DECODER;
    CComPtr<ID3D11Texture2D> pTexture;
    ATLENSURE_SUCCEEDED(pDevice->CreateTexture2D(&TextureDescription, NULL, &pTexture));
    CComQIPtr<ID3D11Resource> pResource(pTexture);
    D3D11_MAPPED_SUBRESOURCE MappedSubresource = {};
    {
        FILETIME StartFileTime = {};
        ::GetSystemTimeAsFileTime(&StartFileTime);
        for (size_t nIteration = 0; nIteration < nIterations; ++nIteration)
        {
            ATLENSURE_SUCCEEDED(pDeviceContext->Map(pResource, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedSubresource));
            _ASSERT(nWidth == MappedSubresource.RowPitch);
            {
                memcpy(MappedSubresource.pData, pFrame.get(), nWidth * nHeight * 3 / 2);
            }
            pDeviceContext->Unmap(pResource, 0);
        }
        FILETIME EndFileTime = {};
        ::GetSystemTimeAsFileTime(&EndFileTime);
        ULARGE_INTEGER StartTime = { StartFileTime.dwLowDateTime, StartFileTime.dwHighDateTime }, EndTime = { EndFileTime.dwLowDateTime, EndFileTime.dwHighDateTime };
        double fElapsedMiliseconds = static_cast<double>((EndTime.QuadPart - StartTime.QuadPart) / 10000.0f);
        std::cout << "Map/memcpy/Unmap total time: " << fElapsedMiliseconds << " ms, " << fElapsedMiliseconds / nIterations << " per call" << std::endl;
    }
    {
        FILETIME StartFileTime = {};
        ::GetSystemTimeAsFileTime(&StartFileTime);
        for (size_t nIteration = 0; nIteration < nIterations; ++nIteration)
        {
            ATLENSURE_SUCCEEDED(pDeviceContext->Map(pResource, 0, D3D11_MAP_WRITE_DISCARD, 0, &MappedSubresource));
            pDeviceContext->Unmap(pResource, 0);
        }
        FILETIME EndFileTime = {};
        ::GetSystemTimeAsFileTime(&EndFileTime);
        ULARGE_INTEGER StartTime = { StartFileTime.dwLowDateTime, StartFileTime.dwHighDateTime }, EndTime = { EndFileTime.dwLowDateTime, EndFileTime.dwHighDateTime };
        double fElapsedMiliseconds = static_cast<double>((EndTime.QuadPart - StartTime.QuadPart) / 10000.0f);
        std::cout << "Map/Unmap total time: " << fElapsedMiliseconds << " ms, " << fElapsedMiliseconds / nIterations << " per call" << std::endl;
    }
    TextureDescription.Usage = D3D11_USAGE_DEFAULT;
    TextureDescription.CPUAccessFlags = 0;
    pTexture.Release();
    ATLENSURE_SUCCEEDED(pDevice->CreateTexture2D(&TextureDescription, NULL, &pTexture));
    pResource = pTexture;
    {
        FILETIME StartFileTime = {};
        ::GetSystemTimeAsFileTime(&StartFileTime);
        for (size_t nIteration = 0; nIteration < nIterations; ++nIteration)
        {
            pDeviceContext->UpdateSubresource(pResource, 0, NULL, pFrame.get(), 1920, 0);
        }
        FILETIME EndFileTime = {};
        ::GetSystemTimeAsFileTime(&EndFileTime);
        ULARGE_INTEGER StartTime = { StartFileTime.dwLowDateTime, StartFileTime.dwHighDateTime }, EndTime = { EndFileTime.dwLowDateTime, EndFileTime.dwHighDateTime };
        double fElapsedMiliseconds = static_cast<double>((EndTime.QuadPart - StartTime.QuadPart) / 10000.0f);
        std::cout << "UpdateSubresource total time: " << fElapsedMiliseconds << " ms, " << fElapsedMiliseconds / nIterations << " per call" << std::endl;
    }
#pragma endregion
#pragma region Cuda
    int nCudaDeviceCount = 0;
    auto nCudaError = cudaGetDeviceCount(&nCudaDeviceCount);
    _ASSERT(nCudaError == CUDA_SUCCESS);
    std::vector<cudaDeviceProp> Devices;
    Devices.resize(nCudaDeviceCount);
    bool bCudaDeviceFound = false;
    int nCudaDevice = 0;
    for (; nCudaDevice < nCudaDeviceCount; ++nCudaDevice)
    {
        nCudaError = cudaGetDeviceProperties(&Devices[nCudaDevice], nCudaDevice);
        _ASSERT(nCudaError == CUDA_SUCCESS);
        if (Devices[nCudaDevice].name == sDeviceName)
        {
            bCudaDeviceFound = true;
            break;
        }
    }
    if (bCudaDeviceFound == false)
    {
        std::cout << "Cuda compatbile adapter named " << sDeviceName.c_str() << "was not found!" << std::endl;
        return EXIT_FAILURE;
    }
    nCudaError = cudaSetDevice(nCudaDevice);
    _ASSERT(nCudaError == CUDA_SUCCESS);
    void *pHostMemory = NULL, *pDeviceMemory = NULL;
    nCudaError = cudaMalloc(&pDeviceMemory, nWidth * nHeight * 3 / 2);
    _ASSERT(nCudaError == CUDA_SUCCESS);
    nCudaError = cudaMallocHost(&pHostMemory, nWidth * nHeight * 3 / 2);
    _ASSERT(nCudaError == CUDA_SUCCESS);
    {
        FILETIME StartFileTime = {};
        ::GetSystemTimeAsFileTime(&StartFileTime);
        for (size_t nIteration = 0; nIteration < nIterations; ++nIteration)
        {
            nCudaError = cudaMemcpy(pDeviceMemory, pFrame.get(), nWidth * nHeight * 3 / 2, cudaMemcpyHostToDevice);
            _ASSERT(nCudaError == CUDA_SUCCESS);
        }
        FILETIME EndFileTime = {};
        ::GetSystemTimeAsFileTime(&EndFileTime);
        ULARGE_INTEGER StartTime = { StartFileTime.dwLowDateTime, StartFileTime.dwHighDateTime }, EndTime = { EndFileTime.dwLowDateTime, EndFileTime.dwHighDateTime };
        double fElapsedMiliseconds = static_cast<double>((EndTime.QuadPart - StartTime.QuadPart) / 10000.0f);
        std::cout << "cudaMemcpy total time: " << fElapsedMiliseconds << " ms, " << fElapsedMiliseconds / nIterations << " per call" << std::endl;

    }
    {
        FILETIME StartFileTime = {};
        ::GetSystemTimeAsFileTime(&StartFileTime);
        for (size_t nIteration = 0; nIteration < nIterations; ++nIteration)
        {
            nCudaError = cudaMemcpyAsync(pDeviceMemory, pFrame.get(), nWidth * nHeight * 3 / 2, cudaMemcpyHostToDevice);
            _ASSERT(nCudaError == CUDA_SUCCESS);
        }
        cudaDeviceSynchronize();
        FILETIME EndFileTime = {};
        ::GetSystemTimeAsFileTime(&EndFileTime);
        ULARGE_INTEGER StartTime = { StartFileTime.dwLowDateTime, StartFileTime.dwHighDateTime }, EndTime = { EndFileTime.dwLowDateTime, EndFileTime.dwHighDateTime };
        double fElapsedMiliseconds = static_cast<double>((EndTime.QuadPart - StartTime.QuadPart) / 10000.0f);
        std::cout << "cudaMemcpyAsync total time: " << fElapsedMiliseconds << " ms, " << fElapsedMiliseconds / nIterations << " per call" << std::endl;
    }
    {
        FILETIME StartFileTime = {};
        ::GetSystemTimeAsFileTime(&StartFileTime);
        for (size_t nIteration = 0; nIteration < nIterations; ++nIteration)
        {
            nCudaError = cudaMemcpyAsync(pDeviceMemory, pHostMemory, nWidth * nHeight * 3 / 2, cudaMemcpyHostToDevice);
            _ASSERT(nCudaError == CUDA_SUCCESS);
        }
        cudaDeviceSynchronize();
        FILETIME EndFileTime = {};
        ::GetSystemTimeAsFileTime(&EndFileTime);
        ULARGE_INTEGER StartTime = { StartFileTime.dwLowDateTime, StartFileTime.dwHighDateTime }, EndTime = { EndFileTime.dwLowDateTime, EndFileTime.dwHighDateTime };
        double fElapsedMiliseconds = static_cast<double>((EndTime.QuadPart - StartTime.QuadPart) / 10000.0f);
        std::cout << "cudaMemcpyAsync with cudaMalloc'ed input memory total time: " << fElapsedMiliseconds << " ms, " << fElapsedMiliseconds / nIterations << " per call" << std::endl;
    }
    cudaFree(pDeviceMemory);
    cudaFree(pHostMemory);
#pragma endregion
    return EXIT_SUCCESS;
}

    【讨论】：

      猜你喜欢
      • 1970-01-01
      • 1970-01-01
      • 2015-07-02
      • 1970-01-01
      • 1970-01-01
      • 2018-10-22
      • 2017-10-15
      • 1970-01-01
      • 2012-07-07
      相关资源
      最近更新 更多
      热门标签
      Java Python linux javascript Mysql C# Docker 算法 前端 SpringBoot Redis Vue spring 设计模式 .net core .net kubernetes c++ 数据库 数据结构 大数据 js 机器学习 微服务 Android Go 程序员 面试 JVM ASP.net core 云原生 人工智能 后端 PHP git CSS golang k8s Nginx Django mybatis 深度学习 多线程 React 架构 devops 爬虫 云计算 Spring Boot LeetCode