<commit_before>#include "svgf/svgf_pt.h"

#include "kernel/context.cuh"
#include "kernel/light.cuh"
#include "kernel/material.cuh"
#include "kernel/intersect.cuh"
#include "kernel/bvh.cuh"
#include "kernel/compaction.h"
#include "kernel/pt_common.h"

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

#include "cuda/helper_math.h"
#include "cuda/cudautil.h"
#include "cuda/cudamemory.h"

#include "aten4idaten.h"

inline __device__ float4 sampleTex(
	cudaSurfaceObject_t tex, 
	int ix, int iy,
	int width, int height)
{
	ix = clamp(ix, 0, width - 1);
	iy = clamp(iy, 0, height - 1);

	float4 ret;
	surf2Dread(&ret, tex, ix * sizeof(float4), iy);

	return ret;
}

__global__ void estimateVariance(
	cudaSurfaceObject_t dst,
	cudaSurfaceObject_t* aovs,
	int width, int height)
{
	int ix = blockIdx.x * blockDim.x + threadIdx.x;
	int iy = blockIdx.y * blockDim.y + threadIdx.y;

	if (ix >= width && iy >= height) {
		return;
	}

	float4 centerMoment;
	surf2Dread(
		&centerMoment,
		aovs[idaten::SVGFPathTracing::AOVType::moments],
		ix * sizeof(float4), iy);

	int frame = (int)centerMoment.w;

	centerMoment /= centerMoment.w;

	// UvZ.
	float var = centerMoment.x - centerMoment.y * centerMoment.y;

	if (frame < 4) {
		// ώZt[S or DisoccludedĂ.
		// 7x7birateral filterŋPxvZ.

		static const int radius = 3;
		static const float sigmaN = 0.005f;
		static const float sigmaD = 0.005f;
		static const float sigmaS = 8;

		float4 centerNormal;
		surf2Dread(
			&centerNormal,
			aovs[idaten::SVGFPathTracing::AOVType::normal],
			ix * sizeof(float4), iy);

		float4 vdepth;
		surf2Dread(
			&vdepth,
			aovs[idaten::SVGFPathTracing::AOVType::depth_meshid],
			ix * sizeof(float4), iy);
		float centerDepth = vdepth.x;

		float4 sum = make_float4(0, 0, 0, 0);
		float weight = 0.0f;

		for (int v = -radius; v <= radius; v++)
		{
			for (int u = -radius; u <= radius; u++)
			{
				auto moment = sampleTex(aovs[idaten::SVGFPathTracing::AOVType::moments], ix + u, iy + v, width, height);
				auto tmpNml = sampleTex(aovs[idaten::SVGFPathTracing::AOVType::normal], ix + u, iy + v, width, height);
				auto tmpDepth = sampleTex(aovs[idaten::SVGFPathTracing::AOVType::depth_meshid], ix + u, iy + v, width, height);

				moment /= moment.w;

				float4 sampleNml = tmpNml;

				float n = 1 - dot(sampleNml, centerNormal);
				float Wn = exp(-0.5f * n * n / (sigmaN * sigmaN));

				float sampleDepth = tmpDepth.x;

				float d = 1 - min(centerDepth, sampleDepth) / max(centerDepth, sampleDepth);
				float Wd = exp(-0.5f * d * d / (sigmaD * sigmaD));

				float Ws = exp(-0.5f * (u * u + v * v) / (sigmaS * sigmaS));

				float W = Ws * Wn * Wd;
				sum += moment * W;
				weight += W;
			}
		}

		if (weight > 0.0f) {
			sum /= weight;
		}

		var = sum.x - sum.y * sum.y;
	}

	// TODO
	// U̓}CiXɂȂȂEEEE
	var = abs(var);

	surf2Dwrite(
		make_float4(var, var, var, var),
		aovs[idaten::SVGFPathTracing::AOVType::var],
		ix * sizeof(float4), iy,
		cudaBoundaryModeTrap);

	surf2Dwrite(
		make_float4(var, var, var, var),
		dst,
		ix * sizeof(float4), iy,
		cudaBoundaryModeTrap);
}

namespace idaten
{
	void SVGFPathTracing::onEstimateVariance(
		cudaSurfaceObject_t outputSurf,
		int width, int height)
	{
		dim3 block(BLOCK_SIZE, BLOCK_SIZE);
		dim3 grid(
			(width + block.x - 1) / block.x,
			(height + block.y - 1) / block.y);

		auto& curaov = getCurAovs();

		estimateVariance << <grid, block >> > (
		//estimateVariance << <1, 1 >> > (
			outputSurf,
			curaov.ptr(),
			width, height);

		checkCudaKernel(estimateVariance);
	}
}
<commit_msg>Check if meshid is same.<commit_after>#include "svgf/svgf_pt.h"

#include "kernel/context.cuh"
#include "kernel/light.cuh"
#include "kernel/material.cuh"
#include "kernel/intersect.cuh"
#include "kernel/bvh.cuh"
#include "kernel/compaction.h"
#include "kernel/pt_common.h"

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

#include "cuda/helper_math.h"
#include "cuda/cudautil.h"
#include "cuda/cudamemory.h"

#include "aten4idaten.h"

inline __device__ float4 sampleTex(
	cudaSurfaceObject_t tex, 
	int ix, int iy,
	int width, int height)
{
	ix = clamp(ix, 0, width - 1);
	iy = clamp(iy, 0, height - 1);

	float4 ret;
	surf2Dread(&ret, tex, ix * sizeof(float4), iy);

	return ret;
}

__global__ void estimateVariance(
	cudaSurfaceObject_t dst,
	cudaSurfaceObject_t* aovs,
	int width, int height)
{
	int ix = blockIdx.x * blockDim.x + threadIdx.x;
	int iy = blockIdx.y * blockDim.y + threadIdx.y;

	if (ix >= width && iy >= height) {
		return;
	}

	float4 centerMoment;
	surf2Dread(
		&centerMoment,
		aovs[idaten::SVGFPathTracing::AOVType::moments],
		ix * sizeof(float4), iy);

	int frame = (int)centerMoment.w;

	centerMoment /= centerMoment.w;

	// UvZ.
	float var = centerMoment.x - centerMoment.y * centerMoment.y;

	if (frame < 4) {
		// ώZt[S or DisoccludedĂ.
		// 7x7birateral filterŋPxvZ.

		static const int radius = 3;
		static const float sigmaN = 0.005f;
		static const float sigmaD = 0.005f;
		static const float sigmaS = 8;

		float4 centerNormal;
		surf2Dread(
			&centerNormal,
			aovs[idaten::SVGFPathTracing::AOVType::normal],
			ix * sizeof(float4), iy);

		float4 depth_meshid;
		surf2Dread(
			&depth_meshid,
			aovs[idaten::SVGFPathTracing::AOVType::depth_meshid],
			ix * sizeof(float4), iy);

		float centerDepth = depth_meshid.x;
		int centerMeshId = (int)depth_meshid.y;

		float4 sum = make_float4(0, 0, 0, 0);
		float weight = 0.0f;

		for (int v = -radius; v <= radius; v++)
		{
			for (int u = -radius; u <= radius; u++)
			{
				auto moment = sampleTex(aovs[idaten::SVGFPathTracing::AOVType::moments], ix + u, iy + v, width, height);
				auto sampleNml = sampleTex(aovs[idaten::SVGFPathTracing::AOVType::normal], ix + u, iy + v, width, height);
				auto sampleDepthMeshId = sampleTex(aovs[idaten::SVGFPathTracing::AOVType::depth_meshid], ix + u, iy + v, width, height);

				moment /= moment.w;

				float sampleDepth = sampleDepthMeshId.x;
				int sampleMeshId = (int)sampleDepthMeshId.y;

				float n = 1 - dot(sampleNml, centerNormal);
				float Wn = exp(-0.5f * n * n / (sigmaN * sigmaN));

				float d = 1 - min(centerDepth, sampleDepth) / max(centerDepth, sampleDepth);
				float Wd = exp(-0.5f * d * d / (sigmaD * sigmaD));

				float Ws = exp(-0.5f * (u * u + v * v) / (sigmaS * sigmaS));

				float Wm = centerMeshId == sampleMeshId ? 1.0f : 0.0f;

				float W = Ws * Wn * Wd * Wm;
				sum += moment * W;
				weight += W;
			}
		}

		if (weight > 0.0f) {
			sum /= weight;
		}

		var = sum.x - sum.y * sum.y;
	}

	// TODO
	// U̓}CiXɂȂȂEEEE
	var = abs(var);

	surf2Dwrite(
		make_float4(var, var, var, var),
		aovs[idaten::SVGFPathTracing::AOVType::var],
		ix * sizeof(float4), iy,
		cudaBoundaryModeTrap);

	surf2Dwrite(
		make_float4(var, var, var, var),
		dst,
		ix * sizeof(float4), iy,
		cudaBoundaryModeTrap);
}

namespace idaten
{
	void SVGFPathTracing::onEstimateVariance(
		cudaSurfaceObject_t outputSurf,
		int width, int height)
	{
		dim3 block(BLOCK_SIZE, BLOCK_SIZE);
		dim3 grid(
			(width + block.x - 1) / block.x,
			(height + block.y - 1) / block.y);

		auto& curaov = getCurAovs();

		estimateVariance << <grid, block >> > (
		//estimateVariance << <1, 1 >> > (
			outputSurf,
			curaov.ptr(),
			width, height);

		checkCudaKernel(estimateVariance);
	}
}

{'fraction_non_alphanumeric': 0.14848360093316865, 'fraction_numerical': 0.017428296967201867, 'mean_word_length': 3.552154903185509, 'pattern_counts': {'":': 0, '<': 25, '<?xml version=': 0, '>': 25, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '6604857', 'n_tokens_mistral': 3360, 'n_tokens_neox': 2968, 'n_words': 668}