Я пытаюсь обновить в cuda текстуру, используемую в DirectX12. Я могу что-то пропустить, но у меня нет подсказки по этому поводу.
- в правом верхнем углу изображения есть область "все время черная".
- только тогда, когда у меня естьRGB, имеющий одинаковое значение для всех пикселей, я получаю ожидаемый результат (по модулю первой проблемы), если нет неожиданных артефактов, как если бы массив не имел ожидаемой структуры.
Что делатьМне не хватает?
Вот создание текстуры:
{
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
EDIT Здесь возникают две проблемы.
Первым является формат текстуры: R32G32B32float, но RTV (?) Ожидает R32G32B32A32float. Совпадение всего на R32G32B32A32float может решить странные массивы цветов. Другой способ - сопоставить RTV с R32G32B32float текстурой, но я не вижу, как.
Вторая проблема - работать с cudaExternalMemoryGetMappedBuffer вместо cudaExternalMemoryGetMappedMipmappedArray;однако, как использовать его с текстурой, описанной D3D12_RESOURCE_DESC textureDesc{};, а также с 1D-массивом 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);