cuda - directx 12 texture2D（在一维数组中）互操作
cuda - directx 12 texture2D (in 1D array) interop
我正在尝试在 cuda 中更新 directx12 中使用的纹理。我可能会错过一些东西，但我没有任何提示。
- 图片右上角有一个"all the time black"区域。
- 只有当我的 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* 一起使用尚不清楚。
我尝试使用以下代码（对于 1D 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);
您假设具有三个浮点类型通道的 2D 纹理图像将具有简单的逐行线性内存布局。正如您的结果所证明的那样，这通常是不正确的。
纹理针对空间相干访问进行了优化。他们的内存布局旨在让 n 维纹理 space 中接近的东西在内存中保持接近。通过简单的行优先内存布局，对于多于一维的任何事物都无法实现这一点。特定纹理图像的确切内存布局通常不是您可以假设知道或依赖的东西。这将取决于您使用的 GPU（通常，数据将以某种方式存储，例如平铺或 Morton order，并在适当的位置填充以保持对齐）。
正如您自己注意到的，您想要做的是使用 cudaExternalMemoryGetMappedMipmappedArray() 将 CUDA 数组（数组是纹理图像的 CUDA 模拟）映射到来自 D3D12 的外部数据。此 CUDA 数组的格式必须与在 D3D12 中创建的纹理格式相匹配。然后，您应该能够使用 CUDA 运行时的纹理或表面函数 API 来访问此 CUDA 数组表示的纹理图像…
正确的做法是将纹理导入为外存，然后导入为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 存储库以反映这些更改 (https://github.com/mprevot/CudaD3D12Update)
我正在尝试在 cuda 中更新 directx12 中使用的纹理。我可能会错过一些东西，但我没有任何提示。
- 图片右上角有一个"all the time black"区域。
- 只有当我的 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* 一起使用尚不清楚。
我尝试使用以下代码（对于 1D 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);
您假设具有三个浮点类型通道的 2D 纹理图像将具有简单的逐行线性内存布局。正如您的结果所证明的那样，这通常是不正确的。
纹理针对空间相干访问进行了优化。他们的内存布局旨在让 n 维纹理 space 中接近的东西在内存中保持接近。通过简单的行优先内存布局，对于多于一维的任何事物都无法实现这一点。特定纹理图像的确切内存布局通常不是您可以假设知道或依赖的东西。这将取决于您使用的 GPU（通常，数据将以某种方式存储，例如平铺或 Morton order，并在适当的位置填充以保持对齐）。
正如您自己注意到的，您想要做的是使用 cudaExternalMemoryGetMappedMipmappedArray() 将 CUDA 数组（数组是纹理图像的 CUDA 模拟）映射到来自 D3D12 的外部数据。此 CUDA 数组的格式必须与在 D3D12 中创建的纹理格式相匹配。然后，您应该能够使用 CUDA 运行时的纹理或表面函数 API 来访问此 CUDA 数组表示的纹理图像…
正确的做法是将纹理导入为外存，然后导入为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 存储库以反映这些更改 (https://github.com/mprevot/CudaD3D12Update)