CudaD3D12Update

CudaD3D12Update

关注
发信
关注(28)粉丝(399)

CUDA-DirectX 12 Texture2D(1D阵列)互操作

本文介绍了CUDA-DirectX 12 Texture2D(1D阵列)互操作的处理方法，对大家解决问题具有一定的参考价值，需要的朋友们下面随着小编来一起学习吧！

问题描述

我正在尝试在cuda中更新Directx12中使用的纹理.我可能会错过一些东西，但是对此我没有任何提示.

  1. 图像的右上角有一个一直黑"的区域.
  2. 只有当我的所有像素的R G B值都相同时，我才能得到预期的结果(对第一个问题求模)，如果没有，则我有意外的假象，就好像数组没有预期的结构一样.

我想念什么?

这是纹理的创建:

{
    TextureWidth = m_width;
    TextureHeight = m_height;
    auto nPixels = TextureWidth * TextureHeight * 3;
    auto pixelBufferSize = sizeof(float)* nPixels;

    D3D12_RESOURCE_DESC textureDesc{};
    textureDesc.MipLevels = 1;
    textureDesc.Format = DXGI_FORMAT_R32G32B32_FLOAT;
    textureDesc.Width = TextureWidth;
    textureDesc.Height = TextureHeight;
    textureDesc.Flags = D3D12_RESOURCE_FLAG_NONE;
    textureDesc.DepthOrArraySize = 1;
    textureDesc.SampleDesc.Count = 1;
    textureDesc.SampleDesc.Quality = 0;
    textureDesc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D;

    ThrowIfFailed(m_device->CreateCommittedResource(&CD3DX12_HEAP_PROPERTIES(D3D12_HEAP_TYPE_DEFAULT), D3D12_HEAP_FLAG_SHARED,
        &textureDesc, D3D12_RESOURCE_STATE_PIXEL_SHADER_RESOURCE, nullptr, IID_PPV_ARGS(&m_textureBuffer)));
    NAME_D3D12_OBJECT(m_textureBuffer);

    // Describe and create a SRV for the texture.
    {
        D3D12_SHADER_RESOURCE_VIEW_DESC srvDesc{};
        srvDesc.Shader4ComponentMapping = D3D12_DEFAULT_SHADER_4_COMPONENT_MAPPING;
        srvDesc.Format = textureDesc.Format;
        srvDesc.ViewDimension = D3D12_SRV_DIMENSION_TEXTURE2D;
        srvDesc.Texture2D.MipLevels = 1;
        m_device->CreateShaderResourceView(m_textureBuffer.Get(), &srvDesc, m_srvHeap->GetCPUDescriptorHandleForHeapStart());
        NAME_D3D12_OBJECT(m_srvHeap);
    }

    // Share m_textureBuffer with cuda
    {
        HANDLE sharedHandle{};
        WindowsSecurityAttributes windowsSecurityAttributes{};
        LPCWSTR name{};
        ThrowIfFailed(m_device->CreateSharedHandle(m_textureBuffer.Get(), &windowsSecurityAttributes, GENERIC_ALL, name, &sharedHandle));

        D3D12_RESOURCE_ALLOCATION_INFO d3d12ResourceAllocationInfo;
        d3d12ResourceAllocationInfo = m_device->GetResourceAllocationInfo(m_nodeMask, 1, &CD3DX12_RESOURCE_DESC::Buffer(pixelBufferSize));
        auto actualSize = d3d12ResourceAllocationInfo.SizeInBytes;

        cudaExternalMemoryHandleDesc externalMemoryHandleDesc;
        memset(&externalMemoryHandleDesc, 0, sizeof(externalMemoryHandleDesc));
        externalMemoryHandleDesc.type = cudaExternalMemoryHandleTypeD3D12Resource;
        externalMemoryHandleDesc.handle.win32.handle = sharedHandle;
        externalMemoryHandleDesc.size = actualSize;
        externalMemoryHandleDesc.flags = cudaExternalMemoryDedicated;

        checkCudaErrors(cudaImportExternalMemory(&m_externalMemory, &externalMemoryHandleDesc));

        cudaExternalMemoryBufferDesc externalMemoryBufferDesc;
        memset(&externalMemoryBufferDesc, 0, sizeof(externalMemoryBufferDesc));
        externalMemoryBufferDesc.offset = 0;
        externalMemoryBufferDesc.size = pixelBufferSize;
        externalMemoryBufferDesc.flags = 0;

        checkCudaErrors(cudaExternalMemoryGetMappedBuffer(&m_cudaDevVertptr, m_externalMemory, &externalMemoryBufferDesc));
        RunKernel(TextureWidth, TextureHeight, (float*)m_cudaDevVertptr, m_streamToRun, 1.0f);
        checkCudaErrors(cudaStreamSynchronize(m_streamToRun));
    }
}

下面是更新此纹理的cuda代码:

int iDivUp(int a, int b) { return a % b != 0 ? a / b + 1 : a / b; }

__global__ void TextureKernel(float *pixels, unsigned int width, unsigned int height, float time)
{
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    if (y < height && x < width)
    {
        auto pos = (y * width + x) * 3;
        auto sint = __sinf(time) * 0.1f + 0.10f;
        auto sintAlt = (x / 32) % 2 == 0 ? 1.0f : sint;
        pixels[pos + 0] = sintAlt; //RED
        pixels[pos + 1] = 0; // (x + y) % 2 == 0 ? 1.0f : __sinf(time) * 0.25f + 0.75f; //GREEN
        pixels[pos + 2] = 0; // (x + y) % 2 == 0 ? 1.0f : 0.0f;                       //BLUE
        //pixels[pos + 0] = __sinf(time + 0.) * 0.5f + 0.5f;
        //pixels[pos + 1] = __sinf(time * 0.09) * 0.5f + 0.5f;
        //pixels[pos + 2] = __sinf(time + 2) * 0.5f + 0.5f;
    }
}

void RunKernel(size_t meshWidth, size_t meshHeight, float *texture_dev, cudaStream_t streamToRun, float animTime)
{
    //dim3 block(16, 16, 1);
    //dim3 grid(meshWidth / 16, meshHeight / 16, 1);
    auto unit = 32;
    dim3 threads(unit, unit);
    dim3 grid(iDivUp(meshWidth, unit), iDivUp(meshHeight, unit));
    TextureKernel <<<grid, threads, 0, streamToRun >>>(texture_dev, meshWidth, meshHeight, animTime);
    getLastCudaError("TextureKernel execution failed.\n");
}

以及通过以下代码得到的结果图像的一部分:

以及完整的回购(如果需要):

https://github.com/mprevot/CudaD3D12Update

编辑这里出现两个问题.

第一个是纹理的格式:R32G32B32float，但是RTV(?)实际期望的是R32G32B32A32float.匹配R32G32B32A32float处的所有内容可以解决奇怪的颜色数组.另一种方法是将RTV匹配到R32G32B32float纹理，但是我不知道如何.

第二个问题是使用cudaExternalMemoryGetMappedBuffer而不是cudaExternalMemoryGetMappedMipmappedArray；但是，如何将其与D3D12_RESOURCE_DESC textureDesc{};描述的纹理以及一维cuda数组float*结合使用尚不清楚.

我尝试使用以下代码(用于一维mipmap数组)，但没有成功(cudaErrorInvalidValue).

auto textureSurface = TextureWidth * TextureHeight;
auto texturePixels = textureSurface * TextureChannels;
cudaExternalMemoryMipmappedArrayDesc cuTexDesc{};
cuTexDesc.numLevels = 1;
cuTexDesc.extent = make_cudaExtent(texturePixels, 0, 0);
cuTexDesc.formatDesc = cudaCreateChannelDesc<float>();
auto result = cudaMallocMipmappedArray(&cuMipArray[0], &cuTexDesc.formatDesc, cuTexDesc.extent, cuTexDesc.numLevels);
解决方案

正确的做法是将纹理导入为外部存储器，然后导入为mipmap数组，然后使用该数组创建cuda曲面，然后对其进行修改在cuda内核中.

导入和映射是通过以下方式完成的:

cudaExternalMemoryMipmappedArrayDesc cuExtmemMipDesc{};
cuExtmemMipDesc.extent = make_cudaExtent(texDesc.Width, texDesc.Height, 0);
cuExtmemMipDesc.formatDesc = cudaCreateChannelDesc<float4>();
cuExtmemMipDesc.numLevels = 1;
cuExtmemMipDesc.flags = cudaArraySurfaceLoadStore;

cudaMipmappedArray_t cuMipArray{};
CheckCudaErrors(cudaExternalMemoryGetMappedMipmappedArray(&cuMipArray, m_externalMemory, &cuExtmemMipDesc));

cudaArray_t cuArray{};
CheckCudaErrors(cudaGetMipmappedArrayLevel(&cuArray, cuMipArray, 0));

cudaResourceDesc cuResDesc{};
cuResDesc.resType = cudaResourceTypeArray;
cuResDesc.res.array.array = cuArray;
checkCudaErrors(cudaCreateSurfaceObject(&cuSurface, &cuResDesc));
// where cudaSurfaceObject_t cuSurface{};

CUDA部分如下所示:

int iDivUp(int a, int b) { return a % b != 0 ? a / b + 1 : a / b; }

__global__ void UpdateSurface(cudaSurfaceObject_t surf, unsigned int width, unsigned int height, float time)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (y >= height | x >= width) return;

    auto xVar = (float)x / (float)width;
    auto yVar = (float)y / (float)height;
    auto cost = __cosf(time) * 0.5f + 0.5f;
    auto costx = __cosf(time) * 0.5f + xVar;
    auto costy = __cosf(time) * 0.5f + yVar;
    auto costxx = (__cosf(time) * 0.5f + 0.5f) * width;
    auto costyy = (__cosf(time) * 0.5f + 0.5f) * height;
    auto costxMany = __cosf(y * time) * 0.5f + yVar;
    auto costyMany = __cosf((float)x/100 * time) * 0.5f + xVar;
    auto margin = 1;

    float4 pixel{};
    if (y == 0) // paint the first row
        pixel = make_float4(costyMany * 0.3, costyMany * 1, costyMany * 0.4, 1);
    else if (y == height - 1) // paint the last row
        pixel = make_float4(costyMany * 0.6, costyMany * 0.7, costyMany * 1, 1);
    else if (x % 5 == 0) // paint a column of 1 pixel wide every 5 pixels
    {
        if (x > width / 2) // a certain color for the right half
            pixel = make_float4(0.1, 0.5, costx * 1, 1);
        else // another color for the left half
            pixel = make_float4(costx * 1, 0.1, 0.2, 1);
    }
    else if (x > width - margin - 1 | x <= margin) // first and last columns
        pixel = make_float4(costxMany, costxMany * 0.9, costxMany * 0.6, 1);
    else // all the rest of the texture
        pixel = make_float4(costx * 0.3, costx * 0.4, costx * 0.6, 1);
    surf2Dwrite(pixel, surf, x * 16, y);
}

void RunKernel(size_t textureW, size_t textureH, cudaSurfaceObject_t surfaceObject, cudaStream_t streamToRun, float animTime)
{
    auto unit = 10;
    dim3 threads(unit, unit);
    dim3 grid(iDivUp(textureW, unit), iDivUp(textureH, unit));
    UpdateSurface <<<grid, threads, 0, streamToRun >>> (surfaceObject, textureW, textureH, animTime);
    getLastCudaError("UpdateSurface execution failed.\n");
}

我更新了git repo以反映这些更改( https://github.com/mprevot/CudaD3D12Update)

I'm trying to update in cuda a texture used in directx12. I may miss something but I have no tip about it.

  1. there is an "all the time black" area in the top right area of the image.
  2. only when I have R G B having the same value for all pixels, I get the expected result (modulo the first problem), if not I have unexpected artefacts, as if the array was not having the expected structure.

What do I miss ?

Here is the creation of the texture:

{
    TextureWidth = m_width;
    TextureHeight = m_height;
    auto nPixels = TextureWidth * TextureHeight * 3;
    auto pixelBufferSize = sizeof(float)* nPixels;

    D3D12_RESOURCE_DESC textureDesc{};
    textureDesc.MipLevels = 1;
    textureDesc.Format = DXGI_FORMAT_R32G32B32_FLOAT;
    textureDesc.Width = TextureWidth;
    textureDesc.Height = TextureHeight;
    textureDesc.Flags = D3D12_RESOURCE_FLAG_NONE;
    textureDesc.DepthOrArraySize = 1;
    textureDesc.SampleDesc.Count = 1;
    textureDesc.SampleDesc.Quality = 0;
    textureDesc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D;

    ThrowIfFailed(m_device->CreateCommittedResource(&CD3DX12_HEAP_PROPERTIES(D3D12_HEAP_TYPE_DEFAULT), D3D12_HEAP_FLAG_SHARED,
        &textureDesc, D3D12_RESOURCE_STATE_PIXEL_SHADER_RESOURCE, nullptr, IID_PPV_ARGS(&m_textureBuffer)));
    NAME_D3D12_OBJECT(m_textureBuffer);

    // Describe and create a SRV for the texture.
    {
        D3D12_SHADER_RESOURCE_VIEW_DESC srvDesc{};
        srvDesc.Shader4ComponentMapping = D3D12_DEFAULT_SHADER_4_COMPONENT_MAPPING;
        srvDesc.Format = textureDesc.Format;
        srvDesc.ViewDimension = D3D12_SRV_DIMENSION_TEXTURE2D;
        srvDesc.Texture2D.MipLevels = 1;
        m_device->CreateShaderResourceView(m_textureBuffer.Get(), &srvDesc, m_srvHeap->GetCPUDescriptorHandleForHeapStart());
        NAME_D3D12_OBJECT(m_srvHeap);
    }

    // Share m_textureBuffer with cuda
    {
        HANDLE sharedHandle{};
        WindowsSecurityAttributes windowsSecurityAttributes{};
        LPCWSTR name{};
        ThrowIfFailed(m_device->CreateSharedHandle(m_textureBuffer.Get(), &windowsSecurityAttributes, GENERIC_ALL, name, &sharedHandle));

        D3D12_RESOURCE_ALLOCATION_INFO d3d12ResourceAllocationInfo;
        d3d12ResourceAllocationInfo = m_device->GetResourceAllocationInfo(m_nodeMask, 1, &CD3DX12_RESOURCE_DESC::Buffer(pixelBufferSize));
        auto actualSize = d3d12ResourceAllocationInfo.SizeInBytes;

        cudaExternalMemoryHandleDesc externalMemoryHandleDesc;
        memset(&externalMemoryHandleDesc, 0, sizeof(externalMemoryHandleDesc));
        externalMemoryHandleDesc.type = cudaExternalMemoryHandleTypeD3D12Resource;
        externalMemoryHandleDesc.handle.win32.handle = sharedHandle;
        externalMemoryHandleDesc.size = actualSize;
        externalMemoryHandleDesc.flags = cudaExternalMemoryDedicated;

        checkCudaErrors(cudaImportExternalMemory(&m_externalMemory, &externalMemoryHandleDesc));

        cudaExternalMemoryBufferDesc externalMemoryBufferDesc;
        memset(&externalMemoryBufferDesc, 0, sizeof(externalMemoryBufferDesc));
        externalMemoryBufferDesc.offset = 0;
        externalMemoryBufferDesc.size = pixelBufferSize;
        externalMemoryBufferDesc.flags = 0;

        checkCudaErrors(cudaExternalMemoryGetMappedBuffer(&m_cudaDevVertptr, m_externalMemory, &externalMemoryBufferDesc));
        RunKernel(TextureWidth, TextureHeight, (float*)m_cudaDevVertptr, m_streamToRun, 1.0f);
        checkCudaErrors(cudaStreamSynchronize(m_streamToRun));
    }
}

And here the cuda code for updating this texture:

int iDivUp(int a, int b) { return a % b != 0 ? a / b + 1 : a / b; }

__global__ void TextureKernel(float *pixels, unsigned int width, unsigned int height, float time)
{
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    if (y < height && x < width)
    {
        auto pos = (y * width + x) * 3;
        auto sint = __sinf(time) * 0.1f + 0.10f;
        auto sintAlt = (x / 32) % 2 == 0 ? 1.0f : sint;
        pixels[pos + 0] = sintAlt; //RED
        pixels[pos + 1] = 0; // (x + y) % 2 == 0 ? 1.0f : __sinf(time) * 0.25f + 0.75f; //GREEN
        pixels[pos + 2] = 0; // (x + y) % 2 == 0 ? 1.0f : 0.0f;                       //BLUE
        //pixels[pos + 0] = __sinf(time + 0.) * 0.5f + 0.5f;
        //pixels[pos + 1] = __sinf(time * 0.09) * 0.5f + 0.5f;
        //pixels[pos + 2] = __sinf(time + 2) * 0.5f + 0.5f;
    }
}

void RunKernel(size_t meshWidth, size_t meshHeight, float *texture_dev, cudaStream_t streamToRun, float animTime)
{
    //dim3 block(16, 16, 1);
    //dim3 grid(meshWidth / 16, meshHeight / 16, 1);
    auto unit = 32;
    dim3 threads(unit, unit);
    dim3 grid(iDivUp(meshWidth, unit), iDivUp(meshHeight, unit));
    TextureKernel <<<grid, threads, 0, streamToRun >>>(texture_dev, meshWidth, meshHeight, animTime);
    getLastCudaError("TextureKernel execution failed.\n");
}

And an extract of the resulting image I get with this code:

And the full repo if needed:

https://github.com/mprevot/CudaD3D12Update

EDITTwo problems occur here.

The first is the format of texture: R32G32B32float, but the RTV (?) is expecting actually R32G32B32A32float. Matching everything at R32G32B32A32float can solve the weird colors arrays. The other way is to match the RTV to a R32G32B32float texture, but I don't see how.

The second problem is to work with cudaExternalMemoryGetMappedBuffer instead of cudaExternalMemoryGetMappedMipmappedArray; however how to use it with the texture described by D3D12_RESOURCE_DESC textureDesc{}; as well as a 1D cuda array float* is no clear yet.

I tried with the following code (for a 1D mipmap array), without success (cudaErrorInvalidValue).

auto textureSurface = TextureWidth * TextureHeight;
auto texturePixels = textureSurface * TextureChannels;
cudaExternalMemoryMipmappedArrayDesc cuTexDesc{};
cuTexDesc.numLevels = 1;
cuTexDesc.extent = make_cudaExtent(texturePixels, 0, 0);
cuTexDesc.formatDesc = cudaCreateChannelDesc<float>();
auto result = cudaMallocMipmappedArray(&cuMipArray[0], &cuTexDesc.formatDesc, cuTexDesc.extent, cuTexDesc.numLevels);
解决方案

The right thing to do is to import the texture as external memory, then as mipmap array, then use this array to create a cuda surface, and then modify this surface in the cuda kernel.

The import and mapping is done this way:

cudaExternalMemoryMipmappedArrayDesc cuExtmemMipDesc{};
cuExtmemMipDesc.extent = make_cudaExtent(texDesc.Width, texDesc.Height, 0);
cuExtmemMipDesc.formatDesc = cudaCreateChannelDesc<float4>();
cuExtmemMipDesc.numLevels = 1;
cuExtmemMipDesc.flags = cudaArraySurfaceLoadStore;

cudaMipmappedArray_t cuMipArray{};
CheckCudaErrors(cudaExternalMemoryGetMappedMipmappedArray(&cuMipArray, m_externalMemory, &cuExtmemMipDesc));

cudaArray_t cuArray{};
CheckCudaErrors(cudaGetMipmappedArrayLevel(&cuArray, cuMipArray, 0));

cudaResourceDesc cuResDesc{};
cuResDesc.resType = cudaResourceTypeArray;
cuResDesc.res.array.array = cuArray;
checkCudaErrors(cudaCreateSurfaceObject(&cuSurface, &cuResDesc));
// where cudaSurfaceObject_t cuSurface{};

the cuda part looks like this:

int iDivUp(int a, int b) { return a % b != 0 ? a / b + 1 : a / b; }

__global__ void UpdateSurface(cudaSurfaceObject_t surf, unsigned int width, unsigned int height, float time)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (y >= height | x >= width) return;

    auto xVar = (float)x / (float)width;
    auto yVar = (float)y / (float)height;
    auto cost = __cosf(time) * 0.5f + 0.5f;
    auto costx = __cosf(time) * 0.5f + xVar;
    auto costy = __cosf(time) * 0.5f + yVar;
    auto costxx = (__cosf(time) * 0.5f + 0.5f) * width;
    auto costyy = (__cosf(time) * 0.5f + 0.5f) * height;
    auto costxMany = __cosf(y * time) * 0.5f + yVar;
    auto costyMany = __cosf((float)x/100 * time) * 0.5f + xVar;
    auto margin = 1;

    float4 pixel{};
    if (y == 0) // paint the first row
        pixel = make_float4(costyMany * 0.3, costyMany * 1, costyMany * 0.4, 1);
    else if (y == height - 1) // paint the last row
        pixel = make_float4(costyMany * 0.6, costyMany * 0.7, costyMany * 1, 1);
    else if (x % 5 == 0) // paint a column of 1 pixel wide every 5 pixels
    {
        if (x > width / 2) // a certain color for the right half
            pixel = make_float4(0.1, 0.5, costx * 1, 1);
        else // another color for the left half
            pixel = make_float4(costx * 1, 0.1, 0.2, 1);
    }
    else if (x > width - margin - 1 | x <= margin) // first and last columns
        pixel = make_float4(costxMany, costxMany * 0.9, costxMany * 0.6, 1);
    else // all the rest of the texture
        pixel = make_float4(costx * 0.3, costx * 0.4, costx * 0.6, 1);
    surf2Dwrite(pixel, surf, x * 16, y);
}

void RunKernel(size_t textureW, size_t textureH, cudaSurfaceObject_t surfaceObject, cudaStream_t streamToRun, float animTime)
{
    auto unit = 10;
    dim3 threads(unit, unit);
    dim3 grid(iDivUp(textureW, unit), iDivUp(textureH, unit));
    UpdateSurface <<<grid, threads, 0, streamToRun >>> (surfaceObject, textureW, textureH, animTime);
    getLastCudaError("UpdateSurface execution failed.\n");
}

I updated the git repo to reflect those changes (https://github.com/mprevot/CudaD3D12Update)

这篇关于CUDA-DirectX 12 Texture2D(1D阵列)互操作的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持！

09-04 23:08
Powered by LMLPHP ©2024 CudaD3D12Update 0.051487