cuda - 将 Directx 12 纹理导入 cuda 表面:列好,行拧紧
问题描述
我正在尝试从 cuda 编写 D3D12 纹理。为此,我将 D3D 纹理映射到 mipmap,然后是 cudaArray。我可以看到缺少许多列的预期结果。我错过了什么?
完整来源:https ://github.com/mprevot/CudaD3D12Update
D2D12部分:
TextureChannels = 4;
TextureWidth = m_width;
TextureHeight = m_height;
const auto textureSurface = TextureWidth * TextureHeight;
const auto texturePixels = textureSurface * TextureChannels;
const auto textureSizeBytes = sizeof(float)* texturePixels;
D3D12_RESOURCE_DESC d3dTexDesc{};
d3dTexDesc.MipLevels = 1;
d3dTexDesc.Format = TextureChannels == 4 ? DXGI_FORMAT_R32G32B32A32_FLOAT : DXGI_FORMAT_R32G32B32_FLOAT;
d3dTexDesc.Width = TextureWidth;
d3dTexDesc.Height = TextureHeight;
d3dTexDesc.Flags = D3D12_RESOURCE_FLAG_NONE;
d3dTexDesc.DepthOrArraySize = 1;
d3dTexDesc.SampleDesc.Count = 1;
d3dTexDesc.SampleDesc.Quality = 0;
d3dTexDesc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D;
ThrowIfFailed(m_device->CreateCommittedResource(&CD3DX12_HEAP_PROPERTIES(D3D12_HEAP_TYPE_DEFAULT), D3D12_HEAP_FLAG_SHARED,
&d3dTexDesc, D3D12_RESOURCE_STATE_PIXEL_SHADER_RESOURCE, nullptr, IID_PPV_ARGS(&TextureArray)));
NAME_D3D12_OBJECT(TextureArray);
导入/映射部分:
HANDLE sharedHandle{};
WindowsSecurityAttributes secAttr{};
LPCWSTR name{};
ThrowIfFailed(m_device->CreateSharedHandle(TextureArray.Get(), &secAttr, GENERIC_ALL, name, &sharedHandle));
const auto allocInfo = m_device->GetResourceAllocationInfo(m_nodeMask, 1, &d3dTexDesc);
cudaExternalMemoryHandleDesc cuExtmemHandleDesc{};
cuExtmemHandleDesc.type = cudaExternalMemoryHandleTypeD3D12Resource;
cuExtmemHandleDesc.handle.win32.handle = sharedHandle;
cuExtmemHandleDesc.size = allocInfo.SizeInBytes;
cuExtmemHandleDesc.flags = cudaExternalMemoryDedicated;
CheckCudaErrors(cudaImportExternalMemory(&m_externalMemory, &cuExtmemHandleDesc));
cudaExternalMemoryMipmappedArrayDesc cuExtmemMipDesc{};
cuExtmemMipDesc.extent = make_cudaExtent(TextureWidth, TextureHeight, 0);
cuExtmemMipDesc.formatDesc = cudaCreateChannelDesc<float4>();
cuExtmemMipDesc.numLevels = 1;
CheckCudaErrors(cudaExternalMemoryGetMappedMipmappedArray(&cuMipArray, m_externalMemory, &cuExtmemMipDesc));
CheckCudaErrors(cudaGetMipmappedArrayLevel(&cuArray, cuMipArray, 0));
cudaResourceDesc cuResDesc{};
cuResDesc.resType = cudaResourceTypeArray;
cuResDesc.res.array.array = cuArray;
checkCudaErrors(cudaCreateSurfaceObject(&cuSurface, &cuResDesc));
库达部分:
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)
pixel = make_float4(costyMany * 0.3, costyMany * 1, costyMany * 0.4, 1);
else if (y == height - 1)
pixel = make_float4(costyMany * 0.6, costyMany * 0.7, costyMany * 1, 1);
else if (x % 2 == 0)
{
if (x > width / 2)
pixel = make_float4(0.1, 0.5, costx * 1, 1);
else
pixel = make_float4(costx * 1, 0.1, 0.2, 1);
}
else if (x > width - margin - 1 | x <= margin)
pixel = make_float4(costxMany, costxMany * 0.9, costxMany * 0.6, 1);
else
pixel = make_float4(costx * 0.3, costx * 0.4, costx * 0.6, 1);
surf2Dwrite(pixel, surf, x, y);
}
void RunKernel(size_t textureW, size_t textureH, cudaSurfaceObject_t surfaceObject, cudaStream_t streamToRun, float animTime)
{
auto unit = 16;
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");
}
解决方案
所以,最后是 cuda 部分不正确:
surf2Dwrite
'sx
必须乘以 16。
surf2Dwrite(pixel, surf, x * 16, y);
推荐阅读
- c++ - 从 C++ 插入到 MYSQL 数据库(准备好的语句,MYSQL C API)失败,未知句柄(0)
- react-native - 我可以在导航 v6 的堆栈导航器中使用堆栈导航器吗?
- angular - 角度路线:pathMatch:“前缀”与“完整”?
- javascript - 如何在 Axios 中获取对象项?
- visual-studio - 使用 Windows 机器将无线部署到 iPhone
- kubernetes - prometheus在linux环境中将metric存储在哪个目录?
- python - Tkinter 树视图滚动条未正确显示
- javascript - Vue 是的验证
- python - Python数据框删除行 - 循环未完成
- css - 如何创建一个有 5 个块的 CSS 网格(第 1 行 2 个,第 2 行 3 个)