SIGMA integration

[pinecone69]

I'm having trouble getting SIGMA_SHADOW_TRANSLUCENCY working (OUT_SHADOW_TRANSLUCENCY seems the exact same as the noisy shadow input, like there was no denoise op). I got REBLUR_DIFFUSE_SPECULAR working correctly just as I expected, following this Vulkan integration example. My assumption is that integrating SIGMA would be analogous to integrating REBLUR; I assume resources are created in the same way, pipeline creation is the same using NRDWrapper::createPipelines(), and denoising commands can be correctly recorded using the same NRDWrapper::dispatch() function, each denoiser requires its settings to be set, etc...

Is there anything I should expect to be different when integrating SIGMA, compared to when integrating REBLUR? If so, what is it? I will hugely appreciate any help whatsoever

[nvmheyer]

Hi there, apologies for responding so late.
Can you elaborate what errors you are running into?

[pinecone69]

Hi. I'm using a different shadow denoiser now, but from what I can remember, when I viewed the denoised shadow output from SIGMA it looked exactly the same as the noisy shadow input. Also, when viewing my scene at certain positions and angles there appeared to be a very small rectangle in the top left of my screen that was a fraction of my screen size, and it had some weird square shapes in it. Sadly I don't have an image of what that looked like. I don't remember getting any validation messages from vulkan. I don't think there is more for me to elaborate on besides that. I wasn't sure if this was happening because my inputs (IN_PENUMBRA only when I tried the SIGMA_SHADOW denoiser, and additionally IN_TRANSLUCENCY when I tried SIGMA_SHADOW_TRANSLUCENCY) were incorrect or if I had to integrate SIGMA differently for some reason, so that's why I was asking about that.

[nvmheyer]

NRD can be a bit tricky to setup. If you would like to provide your code changes (pointing me to a branch of yours) , I can take a look to see if there's anything I can spot. I personally have not yet used the shadow denoiser yet (the focus of this sample was to 'get off the ground quickly'). But this issue might be the reason to actually get started with it.
There's the official NRDSample at https://github.com/NVIDIA-RTX/NRD-Sample which includes a lot of techniques at the price of being more complex.

[pinecone69]

I should have been more clear about this from the start but this issue has nothing to do with my branch of this repository. I basically followed this integration example to add NRD to my personal project, so I don't have a branch with the exact changes I made trying to add SIGMA, which is what I think you are talking about. I posted an issue here because it seemed like the most relevant place to ask my question about integrating NRD into vulkan.

However I could show all of the functions I copied from this example that I think are relevant to the issue, like pipeline creation, resource creation, denoising commands, etc... and I could explain how I remember I was using them. Would that be helpful?

[nvmheyer]

Yes, it can't hurt to take a look.

[pinecone69]

Here is some relevant code (modified a bit to add the parts where I was trying to use SIGMA), might be missing some things. Incase it may be relevant, note I didn't record REBLUR dispatches while recording SIGMA dispatches because it seemed REBLUR was clearing the non-noisy input images so SIGMA couldn't use them, resulting in some weird behavior if I remember correctly. Also it seemed to me that OUT_SHADOW_TRANSLUCENCY was being used as history because NRDDescs.h said it was an input for SIGMA_SHADOW so I also tried forcing all frames in flight to share that same image, to no avail.

Everything here is almost a direct copy from this repo, and REBLUR_DIFFUSE_SPECULAR was working fine for me with this setup, but not SIGMA_SHADOW

void initializeDenoisers() {
	auto& denoisers = ctx->denoisers;

	denoisers.libraryDesc = nrd::GetLibraryDesc();
	if (denoisers.libraryDesc->versionMajor != NRD_VERSION_MAJOR || denoisers.libraryDesc->versionMinor != NRD_VERSION_MINOR) {
		throw std::runtime_error("NRD version mismatch detected!");
	}

	assert(denoisers.libraryDesc->normalEncoding == nrd::NormalEncoding::R10_G10_B10_A2_UNORM);

	std::vector<nrd::DenoiserDesc> denoiserDescs = {
		{ nrd::Identifier(nrd::Denoiser::REBLUR_DIFFUSE_SPECULAR), nrd::Denoiser::REBLUR_DIFFUSE_SPECULAR },
		{ nrd::Identifier(nrd::Denoiser::SIGMA_SHADOW), nrd::Denoiser::SIGMA_SHADOW },
	};

	nrd::InstanceCreationDesc instanceCreation = {};
	instanceCreation.denoisers = denoiserDescs.data();
	instanceCreation.denoisersNum = denoiserDescs.size();

	if (nrd::CreateInstance(instanceCreation, denoisers.instance) != nrd::Result::SUCCESS) {
		throw std::runtime_error("Failed to create NRD instance!");
	}

	denoisers.instanceDesc = nrd::GetInstanceDesc(*denoisers.instance);

	createDenoiserResources();
	createDenoiserPipelines();
}


void createDenoiserResources() {
	auto& denoisers = ctx->denoisers;

	const nrd::InstanceDesc* instanceDesc = denoisers.instanceDesc;

	VkImageUsageFlags textureUsage = VK_IMAGE_USAGE_SAMPLED_BIT | VK_IMAGE_USAGE_STORAGE_BIT;

	createBuffer(instanceDesc->constantBufferMaxDataSize, VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT,
		VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, denoisers.constantBuffer, denoisers.constantBufferMemory);

	denoisers.constantBufferSize = instanceDesc->constantBufferMaxDataSize;

	auto transitionToGeneral = [&](VkImage image) {
		singleTimeImageMemoryBarrier(image, VK_IMAGE_LAYOUT_GENERAL);
	};

	denoisers.permanentPool.resize(instanceDesc->permanentPoolSize);
	denoisers.permanentPoolMemory.resize(instanceDesc->permanentPoolSize);
	denoisers.permanentPoolViews.resize(instanceDesc->permanentPoolSize);

	for (uint32_t i = 0; i < instanceDesc->permanentPoolSize; i++) {
		VkFormat format = NRDToVkFormat(instanceDesc->permanentPool[i].format);

		createImage(ctx->renderResolution.width, ctx->renderResolution.height, format, VK_IMAGE_TILING_OPTIMAL, textureUsage,
			VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, denoisers.permanentPool[i], denoisers.permanentPoolMemory[i]);

		transitionToGeneral(denoisers.permanentPool[i]);

		denoisers.permanentPoolViews[i] = createImageView(denoisers.permanentPool[i], format, { VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1 });
	}

	denoisers.transientPool.resize(instanceDesc->transientPoolSize);
	denoisers.transientPoolMemory.resize(instanceDesc->transientPoolSize);
	denoisers.transientPoolViews.resize(instanceDesc->transientPoolSize);

	for (uint32_t i = 0; i < instanceDesc->transientPoolSize; i++) {
		VkFormat format = NRDToVkFormat(instanceDesc->transientPool[i].format);

		createImage(ctx->renderResolution.width, ctx->renderResolution.height, format, VK_IMAGE_TILING_OPTIMAL, textureUsage,
			VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, denoisers.transientPool[i], denoisers.transientPoolMemory[i]);

		transitionToGeneral(denoisers.transientPool[i]);

		denoisers.transientPoolViews[i] =
			createImageView(denoisers.transientPool[i], format, { VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1 });
	}

	denoisers.samplers.resize(instanceDesc->samplersNum);

	for (uint32_t i = 0; i < instanceDesc->samplersNum; i++) {
		nrd::Sampler sampler = instanceDesc->samplers[i];
		bool nearest = sampler == nrd::Sampler::NEAREST_CLAMP;

		denoisers.samplers[i] = createBasicSampler(nearest ? VK_FILTER_NEAREST : VK_FILTER_LINEAR, VK_SAMPLER_ADDRESS_MODE_CLAMP_TO_EDGE, 0.0f);
	}

	uint32_t maxNumUserPoolTextures = static_cast<uint32_t>(nrd::ResourceType::MAX_NUM);

	denoisers.userPool.resize(MAX_NUM_FRAMES_IN_FLIGHT);
	denoisers.userPoolMemory.resize(MAX_NUM_FRAMES_IN_FLIGHT);
	denoisers.userPoolViews.resize(MAX_NUM_FRAMES_IN_FLIGHT);

	for (uint32_t i = 0; i < MAX_NUM_FRAMES_IN_FLIGHT; i++) {

		// Required: IN_MV, IN_VIEWZ, IN_NORMAL_ROUGHNESS
		// According to NRDDescs.h, REBLUR_DIFFUSE_SPECULAR requires inputs: IN_DIFF_RADIANCE_HITDIST, IN_SPEC_RADIANCE_HITDIST
		// According to NRDDescs.h, SIGMA_SHADOW_TRANSLUCENCY requires inputs: IN_PENUMBRA, IN_TRANSLUCENCY, OUT_SHADOW_TRANSLUCENCY

		denoisers.userPool[i].resize(maxNumUserPoolTextures);
		denoisers.userPoolMemory[i].resize(maxNumUserPoolTextures);
		denoisers.userPoolViews[i].resize(maxNumUserPoolTextures);

		auto makeUserPoolImage = [&](nrd::ResourceType type, nrd::Format format) {
			uint32_t imageIndex = static_cast<uint32_t>(type);
			createImage(ctx->renderResolution.width, ctx->renderResolution.height, NRDToVkFormat(format),
				VK_IMAGE_TILING_OPTIMAL, textureUsage, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT,
				denoisers.userPool[i][imageIndex], denoisers.userPoolMemory[i][imageIndex]);

			transitionToGeneral(denoisers.userPool[i][imageIndex]);

			denoisers.userPoolViews[i][imageIndex] =
				createImageView(denoisers.userPool[i][imageIndex], NRDToVkFormat(format), {VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1});
		};

		// IN_MV (3D, RGBA16f+)
		makeUserPoolImage(nrd::ResourceType::IN_MV, nrd::Format::RGBA16_SFLOAT);
		
		// IN_VIEWZ (R16f+)
		makeUserPoolImage(nrd::ResourceType::IN_VIEWZ, nrd::Format::R16_SFLOAT);

		// IN_NORMAL_ROUGHNESS
		makeUserPoolImage(nrd::ResourceType::IN_NORMAL_ROUGHNESS, nrd::Format::R10_G10_B10_A2_UNORM);

		// IN_DIFF_RADIANCE_HITDIST (RGBA16f+)
		makeUserPoolImage(nrd::ResourceType::IN_DIFF_RADIANCE_HITDIST, nrd::Format::RGBA16_SFLOAT);

		// IN_SPEC_RADIANCE_HITDIST (RGBA16f+)
		makeUserPoolImage(nrd::ResourceType::IN_SPEC_RADIANCE_HITDIST, nrd::Format::RGBA16_SFLOAT);

		// OUT_DIFF_RADIANCE_HITDIST (RGBA16f+)
		makeUserPoolImage(nrd::ResourceType::OUT_DIFF_RADIANCE_HITDIST, nrd::Format::RGBA16_SFLOAT);

		// OUT_SPEC_RADIANCE_HITDIST (RGBA16f+)
		makeUserPoolImage(nrd::ResourceType::OUT_SPEC_RADIANCE_HITDIST, nrd::Format::RGBA16_SFLOAT);

		// DEBUG IMAGE:
		makeUserPoolImage(nrd::ResourceType::OUT_VALIDATION, nrd::Format::RGBA8_UNORM);

		// IN_PENUMBRA (R16+)
		makeUserPoolImage(nrd::ResourceType::IN_PENUMBRA, nrd::Format::R16_UNORM);

		// OUT_SHADOW_TRANSLUCENCY (R8+)
		makeUserPoolImage(nrd::ResourceType::OUT_SHADOW_TRANSLUCENCY, nrd::Format::R8_UNORM);
	}
}


void createDenoiserPipelines() {
	auto& denoisers = ctx->denoisers;
	const nrd::InstanceDesc* instanceDesc = denoisers.instanceDesc;
	const nrd::LibraryDesc* libraryDesc = denoisers.libraryDesc;

	// These are the base binding indices for each type of binding
	const uint32_t constantBufferBindingOffset = libraryDesc->spirvBindingOffsets.constantBufferOffset;
	const uint32_t samplersBindingOffset = libraryDesc->spirvBindingOffsets.samplerOffset;
	const uint32_t texturesBindingOffset = libraryDesc->spirvBindingOffsets.textureOffset;
	const uint32_t storageTextureAndBufferOffset = libraryDesc->spirvBindingOffsets.storageTextureAndBufferOffset;

	// If samplers are in a separate descriptor set, create a descriptor set for
	// them now.
	// If NRD placed samplers in a separate descriptor set, create it now.
	bool samplersInSeparateSet = instanceDesc->constantBufferAndSamplersSpaceIndex != instanceDesc->resourcesSpaceIndex;
	assert(samplersInSeparateSet);

	{
		// Prepare immutable sampler descriptors
		std::vector<VkDescriptorSetLayoutBinding> setBindings(instanceDesc->samplersNum + 1);
		for (uint32_t s = 0; s < instanceDesc->samplersNum; ++s)
		{
			setBindings[s] = VkDescriptorSetLayoutBinding{ .binding = samplersBindingOffset + s,
														  .descriptorType = VK_DESCRIPTOR_TYPE_SAMPLER,
														  .descriptorCount = 1,
														  .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
				// We make use of immutable samplers as they will be shared among all the pipelines and
				// don't change over the lifetime of the pipelines
				.pImmutableSamplers = &denoisers.samplers[s] };
		}
		// Prepare constant buffer descriptor
		setBindings[instanceDesc->samplersNum] = VkDescriptorSetLayoutBinding{
			.binding = constantBufferBindingOffset,
			.descriptorType = VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER,
			.descriptorCount = 1,
			.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
		};

		const VkDescriptorSetLayoutCreateInfo samplerLayoutInfo{ .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
															.flags = 0,
															.bindingCount = (uint32_t)setBindings.size(),
															.pBindings = setBindings.data() };
		checkVkResult(vkCreateDescriptorSetLayout(ctx->device, &samplerLayoutInfo, nullptr, &denoisers.samplerConstBufferDescriptorSetLayout));

		const std::vector<VkDescriptorPoolSize> poolSizes = { {.type = VK_DESCRIPTOR_TYPE_SAMPLER, .descriptorCount = instanceDesc->samplersNum},
												  {.type = VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER, .descriptorCount = 1} };

		const VkDescriptorPoolCreateInfo poolInfo{ .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO,
												  .maxSets = 1,
												  .poolSizeCount = (uint32_t)poolSizes.size(),
												  .pPoolSizes = poolSizes.data() };
		checkVkResult(vkCreateDescriptorPool(ctx->device, &poolInfo, nullptr, &denoisers.samplerConstBufferDescriptorPool));

		const VkDescriptorSetAllocateInfo descriptorSetInfo{ .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO,
															.descriptorPool = denoisers.samplerConstBufferDescriptorPool,
															.descriptorSetCount = 1,
															.pSetLayouts = &denoisers.samplerConstBufferDescriptorSetLayout };

		checkVkResult(vkAllocateDescriptorSets(ctx->device, &descriptorSetInfo, &denoisers.samplerConstBufferDescriptorSet));
	}

	// Bind the constant buffer once and leave it there
	{
		VkDescriptorBufferInfo bufferInfo{};
		bufferInfo.buffer = denoisers.constantBuffer;
		bufferInfo.offset = 0;
		bufferInfo.range = VK_WHOLE_SIZE;

		VkWriteDescriptorSet constantBufferUpdate = { .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
													 .pNext = nullptr,
													 .dstSet = denoisers.samplerConstBufferDescriptorSet,
													 .dstBinding = constantBufferBindingOffset,
													 .descriptorCount = 1,
													 .descriptorType = VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER,
													 .pBufferInfo = &bufferInfo };
		vkUpdateDescriptorSets(ctx->device, 1, &constantBufferUpdate, 0, nullptr);
	}

	// Determine the maximum number of bindings a pipeline can have
	uint32_t maxNumTextureBindings = 0;
	for (uint32_t p = 0; p < instanceDesc->pipelinesNum; ++p)
	{
		const nrd::PipelineDesc& nrdPipelineDesc = instanceDesc->pipelines[p];

		uint32_t numResources = 0;
		for (uint32_t r = 0; r < nrdPipelineDesc.resourceRangesNum; ++r)
		{
			numResources += nrdPipelineDesc.resourceRanges[r].descriptorsNum;
		}
		maxNumTextureBindings = std::max(maxNumTextureBindings, numResources);
	}

	denoisers.pipelines.resize(instanceDesc->pipelinesNum);

#pragma omp parallel for
	for (int p = 0; p < static_cast<int>(instanceDesc->pipelinesNum); ++p)
	{
		logMessage("Compiling NRD pipeline %d\n", p);

		const nrd::PipelineDesc& pDesc = instanceDesc->pipelines[p];

		// Just reserve maximum number of bindings - we may not make use of them all for each pipeline
		// m textures (either for sampling or as storage)
		std::vector<VkDescriptorSetLayoutBinding> setBindings(maxNumTextureBindings);
		// On the main descriptor set, we make use of push descriptors which makes it so much easier to use and update.
		// But we can't use push descriptors for both; see VUID-VkPipelineLayoutCreateInfo-pSetLayouts-00293
		VkDescriptorSetLayoutCreateInfo setLayoutInfo{ .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
													  .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,
													  .pBindings = setBindings.data() };

		// Prepare image/texture descriptors
		for (uint32_t r = 0; r < pDesc.resourceRangesNum; ++r)
		{
			const nrd::ResourceRangeDesc& range = pDesc.resourceRanges[r];

			for (uint32_t b = 0; b < range.descriptorsNum; ++b)
			{
				VkDescriptorSetLayoutBinding binding{ .descriptorCount = 1, .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT };

				switch (range.descriptorType)
				{
				case nrd::DescriptorType::TEXTURE:
					binding.descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE;
					binding.binding = texturesBindingOffset + instanceDesc->resourcesBaseRegisterIndex + b;
					break;
				case nrd::DescriptorType::STORAGE_TEXTURE:
					binding.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE;
					binding.binding = storageTextureAndBufferOffset + instanceDesc->resourcesBaseRegisterIndex + b;
					break;
				default:
					assert(0);
				}
				setBindings[setLayoutInfo.bindingCount++] = binding;
			}
		}

		// Now let's build the layouts, descriptor sets and pipelines
		nrdPipeline& pipeline = denoisers.pipelines[p];

		checkVkResult(vkCreateDescriptorSetLayout(ctx->device, &setLayoutInfo, nullptr, &pipeline.resourceDescriptorSetLayout));

		pipeline.numBindings = setLayoutInfo.bindingCount;
		logMessage("Pipeline uses %d bindings\n", pipeline.numBindings);

		// NRD using these two set indexes is a hardcoded assumption that NRD promised not to break
		assert(instanceDesc->constantBufferAndSamplersSpaceIndex == 1 && instanceDesc->resourcesSpaceIndex == 0);

		// Each pipeline is accessing the global sampler+constant buffer set as well as a pipeline-specific set of textures
		const std::array<VkDescriptorSetLayout, 2> pipelineSetLayouts{ pipeline.resourceDescriptorSetLayout,
																	  denoisers.samplerConstBufferDescriptorSetLayout };
		const VkPipelineLayoutCreateInfo pipelineLayoutInfo{ .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
															.setLayoutCount = 2U,
															.pSetLayouts = pipelineSetLayouts.data() };

		checkVkResult(vkCreatePipelineLayout(ctx->device, &pipelineLayoutInfo, nullptr, &pipeline.pipelineLayout));

		VkShaderModule computeShaderModule = createShaderModule(pDesc.computeShaderSPIRV.bytecode, pDesc.computeShaderSPIRV.size);
		assert(computeShaderModule);

		const VkPipelineShaderStageCreateInfo stageCreateInfo{ .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
															  .stage = VK_SHADER_STAGE_COMPUTE_BIT,
															  .module = computeShaderModule,
															  .pName = "main" };

		const VkComputePipelineCreateInfo pipelineCreateInfo{ .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
															 .stage = stageCreateInfo,
															 .layout = pipeline.pipelineLayout };

		checkVkResult(vkCreateComputePipelines(ctx->device, VK_NULL_HANDLE, 1, &pipelineCreateInfo, nullptr, &pipeline.pipeline));

		vkDestroyShaderModule(ctx->device, computeShaderModule, nullptr);
	}
}

void recordDenoisingCommands(const nrd::Identifier* denoisers, uint32_t denoisersNum, VkCommandBuffer cmd, uint32_t raytracingCommandIndex) {
	const nrd::DispatchDesc* dispatchDescs = nullptr;
	uint32_t dispatchDescsNum = 0;
	nrd::GetComputeDispatches(*ctx->denoisers.instance, reinterpret_cast<const nrd::Identifier*>(denoisers), denoisersNum,
		dispatchDescs, dispatchDescsNum);

	for (uint32_t d = 0; d < dispatchDescsNum; ++d)
	{
		const nrd::DispatchDesc& dDesc = dispatchDescs[d];

		recordDenoisingDispatch(cmd, dDesc, raytracingCommandIndex);
	}

}


void recordDenoisingDispatch(VkCommandBuffer cmd, const nrd::DispatchDesc& dispatchDesc, uint32_t raytracingCommandIndex) {
	auto& denoisers = ctx->denoisers;
	const nrd::LibraryDesc* lDesc = nrd::GetLibraryDesc();
	const nrd::InstanceDesc* iDesc = nrd::GetInstanceDesc(*denoisers.instance);
	const nrd::PipelineDesc& pDesc = iDesc->pipelines[dispatchDesc.pipelineIndex];

	// These are the base binding indices for each type of textures
	const uint32_t texturesBindingOffset = lDesc->spirvBindingOffsets.textureOffset;
	const uint32_t storageTextureAndBufferOffset = lDesc->spirvBindingOffsets.storageTextureAndBufferOffset;

	nrdPipeline& pipeline = denoisers.pipelines[dispatchDesc.pipelineIndex];

	std::vector<VkWriteDescriptorSet>  descriptorUpdates(pipeline.numBindings + iDesc->samplersNum,
		VkWriteDescriptorSet{ VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET, nullptr });
	std::vector<VkDescriptorImageInfo> descriptorImageInfos(pipeline.numBindings + iDesc->samplersNum);

	std::vector<VkImageMemoryBarrier> imageBarriers;

	auto transitionToShaderRead = [](VkImage image) {
		return makeImageMemoryBarrier(image, VK_ACCESS_SHADER_WRITE_BIT, VK_ACCESS_SHADER_READ_BIT,
			VK_IMAGE_LAYOUT_GENERAL, VK_IMAGE_LAYOUT_GENERAL);
	};

	auto transitionToShaderWrite = [](VkImage image) {
		return makeImageMemoryBarrier(image, VK_ACCESS_SHADER_READ_BIT, VK_ACCESS_SHADER_WRITE_BIT,
			VK_IMAGE_LAYOUT_GENERAL, VK_IMAGE_LAYOUT_GENERAL);
	};

	// Determine texture descriptor set updates and corresponding image layout transitions
	uint32_t numResourceUpdates = 0;  // Count and index the updates
	for (uint32_t r = 0; r < pDesc.resourceRangesNum; ++r)
	{
		const nrd::ResourceRangeDesc& resourceRange = pDesc.resourceRanges[r];
		const bool                    isStorage = resourceRange.descriptorType == nrd::DescriptorType::STORAGE_TEXTURE;

		uint32_t rangeBaseBindingIndex = isStorage ? storageTextureAndBufferOffset : texturesBindingOffset;

		for (uint32_t d = 0; d < resourceRange.descriptorsNum; ++d)
		{
			const nrd::ResourceDesc& nrdResource = dispatchDesc.resources[numResourceUpdates];

			VkWriteDescriptorSet& update = descriptorUpdates[numResourceUpdates];
			update.dstBinding = rangeBaseBindingIndex + d;
			update.descriptorCount = 1;
			update.descriptorType = NRDToVkDescriptorType(nrdResource.descriptorType);
			update.pImageInfo = &descriptorImageInfos[numResourceUpdates];

			assert(nrdResource.descriptorType == resourceRange.descriptorType);

			VkImage image = VK_NULL_HANDLE;
			VkImageView view = VK_NULL_HANDLE;
			if (nrdResource.type == nrd::ResourceType::TRANSIENT_POOL)
			{
				image = denoisers.transientPool[nrdResource.indexInPool];
				view = denoisers.transientPoolViews[nrdResource.indexInPool];
			}
			else if (nrdResource.type == nrd::ResourceType::PERMANENT_POOL)
			{
				image = denoisers.permanentPool[nrdResource.indexInPool];
				view = denoisers.permanentPoolViews[nrdResource.indexInPool];
			}
			else
			{
				image = denoisers.userPool[raytracingCommandIndex][(uint32_t)nrdResource.type];
				view = denoisers.userPoolViews[raytracingCommandIndex][(uint32_t)nrdResource.type];
			}

			assert(image && view);

			// We assume, images bound to storage bindings will be written to, while images bound to
			// texture bindings will be read from.
			// This is a rather simple scheme. If it turns out, these barriers cost too much performance,
			// we might want to be more clever about it by caching transitions between pipelines.
			isStorage ? transitionToShaderWrite(image) : transitionToShaderRead(image);

			VkDescriptorImageInfo descriptorImageInfo = { .sampler = nullptr, .imageView = view, .imageLayout = VK_IMAGE_LAYOUT_GENERAL };

			descriptorImageInfos[numResourceUpdates] = descriptorImageInfo;
			++numResourceUpdates;
		}
	}

	// Transition all resources into their appropriate state
	vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 0,
		nullptr, 0, nullptr, (uint32_t)imageBarriers.size(), imageBarriers.data());

	const bool samplersInSeparateSet = iDesc->constantBufferAndSamplersSpaceIndex != iDesc->resourcesSpaceIndex;
	assert(samplersInSeparateSet);

	if (pDesc.hasConstantData)
	{
		if (!dispatchDesc.constantBufferDataMatchesPreviousDispatch)
		{
			{
				VkBufferMemoryBarrier barrier = { VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER,
												 nullptr,
												 VK_ACCESS_SHADER_READ_BIT,
												 VK_ACCESS_TRANSFER_WRITE_BIT,
												 VK_QUEUE_FAMILY_IGNORED,
												 VK_QUEUE_FAMILY_IGNORED,
												 denoisers.constantBuffer,
												 0,
												 VK_WHOLE_SIZE };

				vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, 0, 0,
					nullptr, 1, &barrier, 0, nullptr);
			}

			vkCmdUpdateBuffer(cmd, denoisers.constantBuffer, 0, dispatchDesc.constantBufferDataSize, dispatchDesc.constantBufferData);

			{
				VkBufferMemoryBarrier barrier = { VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER,
												 nullptr,
												 VK_ACCESS_TRANSFER_WRITE_BIT,
												 VK_ACCESS_SHADER_READ_BIT,
												 VK_QUEUE_FAMILY_IGNORED,
												 VK_QUEUE_FAMILY_IGNORED,
												 denoisers.constantBuffer,
												 0,
												 VK_WHOLE_SIZE };

				vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 0,
					nullptr, 1, &barrier, 0, nullptr);
			}
		}
	}

	vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline.pipeline);

	// Bind the global set with the immutable samplers and constant buffer
	vkCmdBindDescriptorSets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline.pipelineLayout,
		iDesc->constantBufferAndSamplersSpaceIndex, 1, &denoisers.samplerConstBufferDescriptorSet, 0, nullptr);

	// Update the texture descriptors. Notice how push descriptors don't require us to make sure the
	// descriptors are not in use anymore.
	ctx->funcs.vkCmdPushDescriptorSetKHR(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline.pipelineLayout,
		iDesc->resourcesSpaceIndex, numResourceUpdates, descriptorUpdates.data());

	// Go!
	vkCmdDispatch(cmd, dispatchDesc.gridWidth, dispatchDesc.gridHeight, 1);
}


// Before each frame, denoiser settings are set:

void setDenoiserSettings() {
	float FPS = 1.0f / ctx->deltaTime;
	float accumulationPeriodInSeconds = 1.0f;

	nrd::ReblurSettings reblurSettings = {};
	reblurSettings.maxAccumulatedFrameNum = uint32_t(accumulationPeriodInSeconds * FPS);
	reblurSettings.maxFastAccumulatedFrameNum = reblurSettings.maxAccumulatedFrameNum;
	reblurSettings.checkerboardMode = 
		ctx->userSettings.nrdCheckerboardEnabled == 1 ? nrd::CheckerboardMode::BLACK : nrd::CheckerboardMode::OFF;

	nrd::SetDenoiserSettings(*ctx->denoisers.instance, nrd::Identifier(nrd::Denoiser::REBLUR_DIFFUSE_SPECULAR), &reblurSettings);

	nrd::SigmaSettings sigmaSettings = {};
	sigmaSettings.lightDirection = ctx->camera.shadow;

	nrd::SetDenoiserSettings(*ctx->denoisers.instance, nrd::Identifier(nrd::Denoiser::SIGMA_SHADOW), &sigmaSettings);
}

// Before each frame, denoiser common settings are set:

void setDenoiserCommonSettings(uint32_t raytracingCommandIndex) {
	nrd::CommonSettings settings = {};

	static glm::mat4x4 prevView;
	static glm::mat4x4 prevProj;

	// Write to previous matrices
	memcpy(&settings.viewToClipMatrixPrev, &prevProj, sizeof(nrd::CommonSettings::viewToClipMatrixPrev));
	memcpy(&settings.worldToViewMatrixPrev, &prevView, sizeof(nrd::CommonSettings::worldToViewMatrixPrev));

	// Write to new matrices
	uniformBuffer* u = reinterpret_cast<uniformBuffer*>(ctx->uniformBuffersMapped[raytracingCommandIndex]);

	mat4x4 proj = u->proj;
	proj[1][1] *= -1.f;

	mat4x4 view = u->view;

	// Remember
	prevProj = proj;
	prevView = view;

	memcpy(&settings.viewToClipMatrix, &proj, sizeof(nrd::CommonSettings::viewToClipMatrix));
	memcpy(&settings.worldToViewMatrix, &view, sizeof(nrd::CommonSettings::worldToViewMatrix));

	settings.frameIndex = u->frame;

	settings.accumulationMode = settings.frameIndex == 0 ?
		nrd::AccumulationMode::CLEAR_AND_RESTART : nrd::AccumulationMode::CONTINUE;

	// 3D motion vectors
	settings.motionVectorScale[0] = 1.f;
	settings.motionVectorScale[1] = 1.f;
	settings.motionVectorScale[2] = 1.f;

	// Jitter
	settings.cameraJitterPrev[0] = 0.0f;
	settings.cameraJitterPrev[1] = 0.0f;
	settings.cameraJitter[0] = 0.0f;
	settings.cameraJitter[1] = 0.0f;

	// Dynamic resolution scaling support
	uint32_t width = ctx->renderResolution.width;
	uint32_t height = ctx->renderResolution.height;
	settings.resourceSize[0] = width;
	settings.resourceSize[1] = height;
	settings.resourceSizePrev[0] = width;
	settings.resourceSizePrev[1] = height;
	settings.rectSize[0] = width;
	settings.rectSize[1] = height;
	settings.rectSizePrev[0] = width;
	settings.rectSizePrev[1] = height;

	// Recommended to true
	settings.isMotionVectorInWorldSpace = true;

	settings.disocclusionThreshold = 0.01f;

	// Validation
#if ENABLE_DENOISER_VALIDATION
	settings.enableValidation = true;
#endif

	nrd::SetCommonSettings(*ctx->denoisers.instance, settings);
}


// This is how denoising commands are submitted in the same command buffer as the raytracing commands:

{
	// (Submit raytracing commands)

	// (Make denoiser inputs visible to denoising commands)

	std::vector<nrd::Identifier> denoisers = { 
		nrd::Identifier(nrd::Denoiser::SIGMA_SHADOW)
	};
	recordDenoisingCommands(denoisers.data(), denoisers.size(), cmd, raytracingCommandBufferIndex);

	// (Make writes to denoiser outputs visible to compositing)
	
	// (Compositing commands)

}

// Raytracing commands create denoiser input images. I remember packing penumbra like this:

{
	// I'm pretty sure I remember trying both radians and degrees because it wasn't clear to me which I should use
	penumbraImage[ivec2(LaunchID.xy)] = SIGMA_FrontEnd_PackPenumbra(distToOccluder, tan(sunAngularRadiusDegrees));
}

// Compositing commands read denoiser ouput images. I remember unpacking like this:

{

	float shadow = SIGMA_Backend_UnpackShadow(outShadowImage[gl_globalInvocationID.xy];
	vec3 finalColor = baseColor * shadow;

}