v3: plastic, rough dielectric, per-texel emission, homogeneous medium

ddbb236 verified 19 days ago

41.6 kB

	// pathtracer-diff: a differentiable Monte Carlo path tracer as one kernel.
	//
	// Forward: a megakernel unidirectional path tracer. Materials: Lambertian
	// diffuse, GGX conductor (VNDF, height-correlated Smith), smooth dielectric,
	// rough plastic (diffuse base + GGX coat, F0 = 0.04), and rough dielectric
	// (GGX transmission, BSDF-sampled only). Lights: emissive faces with
	// per-texel emission textures, and an importance-sampled equirectangular
	// environment map. An optional homogeneous participating medium (absorption
	// + isotropic scattering) fills the scene, with distances sampled from a
	// detached rate frozen at Scene construction. Estimators: MIS (balance
	// heuristic) by default, plus legacy NEE-only and BRDF-only modes.
	// Per-texel albedo with bilinear filtering; binned-SAH BVH with
	// near-child-first traversal. One thread owns one pixel and accumulates its
	// spp samples serially, so the image is bitwise deterministic.
	//
	// Backward: exact path replay. Sampling depends only on geometry, frozen
	// material parameters (roughness, ior), the detached environment CDF, the
	// detached medium sampling rate, and the counter-based Philox stream --
	// never on the differentiable parameters (albedo texels, emission texels,
	// environment texels, sigma_a, sigma_s). The replay re-traces identical
	// paths with identical draws and no stored path state. Every radiance term
	// is trAcc * (prod of per-bounce factors) * S * E where each factor is
	// AFFINE in its vertex's albedo texels (Schlick Fresnel is affine in F0)
	// and E is a linear emission or environment texel; medium transmittance
	// ratios fold into trAcc and the factors, and their sigma-derivatives are
	// the closed-form log-derivatives -(D_total + d_term) and
	// -(D_total + d_term) + N_scatter/sigma_s. Gradients scatter through the
	// exact bilinear-footprint adjoints. Interior terms only here; geometry
	// gradients live in the companion geometry kernels.

	#include <cuda_runtime.h>
	#include <curand_kernel.h>

	#include <cstdint>

	#include "pathtracer_launch.h"

	namespace {

	constexpr int kThreads = 128;
	constexpr int kMaxBounces = 16;
	constexpr int kMaxMats = 64;
	constexpr int kSharedTexels = 2048; // shared-staged albedo grad limit
	constexpr int kSharedEmiTexels = 512; // shared-staged emission grad limit
	constexpr int kStack = 64;
	constexpr float kPi = 3.14159265358979323846f;
	constexpr float kInvPi = 0.31830988618379067154f;
	constexpr float kInv4Pi = 0.07957747154594766788f;
	constexpr float kRayEps = 1e-4f;
	constexpr float kShadowEps = 1e-3f;
	constexpr float kEnvDist = 1e30f;

	enum Mode { kModeBrdf = 0, kModeNee = 1, kModeMis = 2 };
	enum MatType {
	kDiffuse = 0,
	kConductor = 1,
	kDielectric = 2, // smooth
	kPlastic = 3, // diffuse base + GGX coat, F0 = 0.04
	kRoughDielectric = 4, // GGX transmission, BSDF-sampled only
	};
	constexpr float kF0Coat = 0.04f;

	// ---------------------------------------------------------------- float3 ops
	__device__ __forceinline__ float3 f3(float x, float y, float z) {
	return make_float3(x, y, z);
	}
	__device__ __forceinline__ float3 operator+(float3 a, float3 b) {
	return f3(a.x + b.x, a.y + b.y, a.z + b.z);
	}
	__device__ __forceinline__ float3 operator-(float3 a, float3 b) {
	return f3(a.x - b.x, a.y - b.y, a.z - b.z);
	}
	__device__ __forceinline__ float3 operator*(float3 a, float s) {
	return f3(a.x * s, a.y * s, a.z * s);
	}
	__device__ __forceinline__ float dot(float3 a, float3 b) {
	return a.x * b.x + a.y * b.y + a.z * b.z;
	}
	__device__ __forceinline__ float3 cross(float3 a, float3 b) {
	return f3(a.y * b.z - a.z * b.y, a.z * b.x - a.x * b.z, a.x * b.y - a.y * b.x);
	}
	__device__ __forceinline__ float3 normalize(float3 a) {
	return a * rsqrtf(fmaxf(dot(a, a), 1e-30f));
	}
	__device__ __forceinline__ float3 neg(float3 a) { return f3(-a.x, -a.y, -a.z); }

	// Duff et al. 2017, "Building an Orthonormal Basis, Revisited".
	__device__ __forceinline__ void onb(float3 n, float3& t, float3& b) {
	float s = copysignf(1.0f, n.z);
	float a = -1.0f / (s + n.z);
	float c = n.x * n.y * a;
	t = f3(1.0f + s * n.x * n.x * a, s * c, -s * n.x);
	b = f3(c, s + n.y * n.y * a, -n.y);
	}

	// ---------------------------------------------------------------- scene view
	struct DevScene {
	const float* tris; // [F, 9]
	const int* mats; // [F]
	const float* uvs; // [F, 3, 2]
	const float* nf; // [N, 6]
	const int* ni; // [N, 3] internal (l, r, axis<<1) \| leaf (start, count, odd)
	int n_nodes;
	const int* lf; // [L] emissive faces
	const float* lcdf; // [L]
	int nl;
	float larea;
	const float* tex; // [T, 3] albedo texels
	const int* thdr; // [M, 3] (offset, W, H)
	int nt;
	const float* etex; // [Te, 3] emission texels
	const int* ehdr; // [M, 3] (offset, W, H)
	int net;
	const int* mtype; // [M]
	const float* mrough; // [M]
	const float* mior; // [M]
	int nm;
	const float* med_sa; // [3] live absorption, or null
	const float* med_ss; // [3] live scattering
	float med_sbar; // frozen (detached) sampling rate
	int has_med;
	const float* env; // [Eh*Ew, 3] or null
	int ew, eh;
	const float* ecdf_m;
	const float* ecdf_c;
	const float* epdf;
	};

	struct DevCam {
	float p[3], f[3], r[3], u[3];
	};

	// grad buffers bundled to keep call signatures sane
	struct GCtx {
	const float* gs; // per-pixel dLoss/dpixel / spp
	float* g_tex; // [T, 3] global
	float* s_gt; // shared staging for g_tex or null
	float* g_etex; // [Te, 3] global
	float* s_get; // shared staging for g_etex or null
	float* g_env; // [Eh*Ew, 3] global or null
	float* s_gm; // shared [6]: d sigma_a[3], d sigma_s[3]
	};

	// ------------------------------------------------------------- intersection
	__device__ __forceinline__ bool tri_hit(const float* v, float3 ro, float3 rd,
	float tmin, float tmax, float& t,
	float& bu, float& bv, float3& ng) {
	float3 v0 = f3(v[0], v[1], v[2]);
	float3 e1 = f3(v[3], v[4], v[5]) - v0;
	float3 e2 = f3(v[6], v[7], v[8]) - v0;
	float3 p = cross(rd, e2);
	float det = dot(e1, p);
	if (fabsf(det) < 1e-12f) return false;
	float inv = 1.0f / det;
	float3 s = ro - v0;
	float u = dot(s, p) * inv;
	if (u < -1e-6f \|\| u > 1.0f + 1e-6f) return false;
	float3 q = cross(s, e1);
	float w = dot(rd, q) * inv;
	if (w < -1e-6f \|\| u + w > 1.0f + 1e-6f) return false;
	float tt = dot(e2, q) * inv;
	if (tt < tmin \|\| tt > tmax) return false;
	t = tt;
	bu = u;
	bv = w;
	ng = cross(e1, e2);
	return true;
	}

	__device__ __forceinline__ bool slab(const float* b, const float ro[3],
	const float inv[3], float tmax) {
	float t0 = kRayEps, t1 = tmax;
	#pragma unroll
	for (int a = 0; a < 3; ++a) {
	float lo = (b[a] - ro[a]) * inv[a];
	float hi = (b[3 + a] - ro[a]) * inv[a];
	if (lo > hi) {
	float tmp = lo;
	lo = hi;
	hi = tmp;
	}
	t0 = fmaxf(t0, lo);
	t1 = fminf(t1, hi);
	}
	return t0 <= t1;
	}

	__device__ __forceinline__ void inv_dir(float3 rd, float inv[3]) {
	float d;
	d = rd.x; if (fabsf(d) < 1e-12f) d = copysignf(1e-12f, d); inv[0] = 1.0f / d;
	d = rd.y; if (fabsf(d) < 1e-12f) d = copysignf(1e-12f, d); inv[1] = 1.0f / d;
	d = rd.z; if (fabsf(d) < 1e-12f) d = copysignf(1e-12f, d); inv[2] = 1.0f / d;
	}

	__device__ int bvh_closest(const DevScene& sc, float3 ro, float3 rd,
	float tmin, float& tbest, float& bu, float& bv,
	float3& ngbest) {
	float roa[3] = {ro.x, ro.y, ro.z};
	float dira[3] = {rd.x, rd.y, rd.z};
	float inv[3];
	inv_dir(rd, inv);
	int stack[kStack];
	int sp = 0;
	stack[sp++] = 0;
	int best = -1;
	while (sp > 0) {
	int nid = stack[--sp];
	if (!slab(&sc.nf[nid * 6], roa, inv, tbest)) continue;
	const int* n = &sc.ni[nid * 3];
	if (n[2] & 1) {
	for (int f = n[0]; f < n[0] + n[1]; ++f) {
	float t, u, v;
	float3 ng;
	if (tri_hit(&sc.tris[f * 9], ro, rd, tmin, tbest, t, u, v, ng)) {
	tbest = t;
	bu = u;
	bv = v;
	ngbest = ng;
	best = f;
	}
	}
	} else if (sp + 2 <= kStack) {
	int axis = n[2] >> 1;
	int near = (dira[axis] >= 0.0f) ? n[0] : n[1];
	int far = (dira[axis] >= 0.0f) ? n[1] : n[0];
	stack[sp++] = far;
	stack[sp++] = near;
	}
	}
	return best;
	}

	__device__ bool bvh_occluded(const DevScene& sc, float3 ro, float3 rd,
	float tmax) {
	float roa[3] = {ro.x, ro.y, ro.z};
	float inv[3];
	inv_dir(rd, inv);
	int stack[kStack];
	int sp = 0;
	stack[sp++] = 0;
	while (sp > 0) {
	int nid = stack[--sp];
	if (!slab(&sc.nf[nid * 6], roa, inv, tmax)) continue;
	const int* n = &sc.ni[nid * 3];
	if (n[2] & 1) {
	for (int f = n[0]; f < n[0] + n[1]; ++f) {
	float t, u, v;
	float3 ng;
	if (tri_hit(&sc.tris[f * 9], ro, rd, kRayEps, tmax, t, u, v, ng))
	return true;
	}
	} else if (sp + 2 <= kStack) {
	stack[sp++] = n[0];
	stack[sp++] = n[1];
	}
	}
	return false;
	}

	__device__ __forceinline__ int cdf_pick(const float* cdf, int n, float r) {
	int lo = 0, hi = n - 1;
	while (lo < hi) {
	int mid = (lo + hi) >> 1;
	if (cdf[mid] < r) lo = mid + 1; else hi = mid;
	}
	return lo;
	}

	// Bilinear fetch with repeat wrap over a flat texel block.
	__device__ __forceinline__ void bilinear(const float* block, int off, int W,
	int H, float u, float v, float rgb[3],
	int idx4[4], float w4[4]) {
	float uu = u * (float)W - 0.5f;
	float vv = v * (float)H - 0.5f;
	float fx = uu - floorf(uu);
	float fy = vv - floorf(vv);
	int x0 = (int)floorf(uu), y0 = (int)floorf(vv);
	int x0w = ((x0 % W) + W) % W;
	int x1w = (((x0 + 1) % W) + W) % W;
	int y0w = ((y0 % H) + H) % H;
	int y1w = (((y0 + 1) % H) + H) % H;
	idx4[0] = off + y0w * W + x0w; w4[0] = (1.0f - fx) * (1.0f - fy);
	idx4[1] = off + y0w * W + x1w; w4[1] = fx * (1.0f - fy);
	idx4[2] = off + y1w * W + x0w; w4[2] = (1.0f - fx) * fy;
	idx4[3] = off + y1w * W + x1w; w4[3] = fx * fy;
	rgb[0] = rgb[1] = rgb[2] = 0.0f;
	#pragma unroll
	for (int i = 0; i < 4; ++i) {
	const float* t = &block[idx4[i] * 3];
	rgb[0] += w4[i] * t[0];
	rgb[1] += w4[i] * t[1];
	rgb[2] += w4[i] * t[2];
	}
	}

	__device__ __forceinline__ void face_uv(const DevScene& sc, int face,
	float bu, float bv, float& u,
	float& v) {
	const float* U = &sc.uvs[face * 6];
	float w0 = 1.0f - bu - bv;
	u = w0 * U[0] + bu * U[2] + bv * U[4];
	v = w0 * U[1] + bu * U[3] + bv * U[5];
	}

	__device__ __forceinline__ void sample_albedo(const DevScene& sc, int m,
	int face, float bu, float bv,
	float rgb[3], int idx4[4],
	float w4[4]) {
	float u, v;
	face_uv(sc, face, bu, bv, u, v);
	const int* h = &sc.thdr[m * 3];
	bilinear(sc.tex, h[0], h[1], h[2], u, v, rgb, idx4, w4);
	}

	__device__ __forceinline__ void sample_emission(const DevScene& sc, int m,
	int face, float bu, float bv,
	float rgb[3], int idx4[4],
	float w4[4]) {
	float u, v;
	face_uv(sc, face, bu, bv, u, v);
	const int* h = &sc.ehdr[m * 3];
	bilinear(sc.etex, h[0], h[1], h[2], u, v, rgb, idx4, w4);
	}

	// ----------------------------------------------------------- environment map
	__device__ __forceinline__ void dir_to_equirect(float3 d, float& u, float& v) {
	float phi = atan2f(d.z, d.x);
	float theta = acosf(fminf(fmaxf(d.y, -1.0f), 1.0f));
	u = (phi + kPi) / (2.0f * kPi);
	v = theta / kPi;
	}

	__device__ __forceinline__ float3 equirect_to_dir(float u, float v) {
	float phi = u * 2.0f * kPi - kPi;
	float theta = v * kPi;
	float st = sinf(theta);
	return f3(st * cosf(phi), cosf(theta), st * sinf(phi));
	}

	__device__ __forceinline__ float env_pdf(const DevScene& sc, float3 d) {
	float u, v;
	dir_to_equirect(d, u, v);
	int x = min(sc.ew - 1, (int)(u * sc.ew));
	int y = min(sc.eh - 1, (int)(v * sc.eh));
	float st = fmaxf(sinf((y + 0.5f) * kPi / sc.eh), 1e-4f);
	float p_img = sc.epdf[y * sc.ew + x];
	float p_tab = p_img * (float)(sc.ew * sc.eh) / (2.0f * kPi * kPi * st);
	return 0.5f * p_tab + 0.5f * (1.0f / (4.0f * kPi));
	}

	__device__ __forceinline__ void env_fetch(const DevScene& sc, float3 d,
	float rgb[3], int idx4[4],
	float w4[4]) {
	float u, v;
	dir_to_equirect(d, u, v);
	v = fminf(fmaxf(v, 0.5f / sc.eh), 1.0f - 0.5f / sc.eh);
	bilinear(sc.env, 0, sc.ew, sc.eh, u, v, rgb, idx4, w4);
	}

	// ------------------------------------------------------------------- GGX
	__device__ __forceinline__ float ggx_lambda(float a2, float cs) {
	cs = fabsf(cs);
	float c2 = cs * cs;
	float t2 = fmaxf(0.0f, 1.0f - c2) / fmaxf(c2, 1e-12f);
	return 0.5f * (-1.0f + sqrtf(1.0f + a2 * t2));
	}

	__device__ __forceinline__ float ggx_d(float a2, float ch) {
	float d = ch * ch * (a2 - 1.0f) + 1.0f;
	return a2 / fmaxf(kPi * d * d, 1e-20f);
	}

	__device__ __forceinline__ float3 ggx_sample_vndf(float3 wi, float alpha,
	float u1, float u2) {
	float3 vh = normalize(f3(alpha * wi.x, alpha * wi.y, wi.z));
	float lensq = vh.x * vh.x + vh.y * vh.y;
	float3 T1 = lensq > 1e-12f ? f3(-vh.y, vh.x, 0.0f) * rsqrtf(lensq)
	: f3(1.0f, 0.0f, 0.0f);
	float3 T2 = cross(vh, T1);
	float r = sqrtf(u1);
	float phi = 2.0f * kPi * u2;
	float t1 = r * cosf(phi);
	float t2 = r * sinf(phi);
	float s = 0.5f * (1.0f + vh.z);
	t2 = (1.0f - s) * sqrtf(fmaxf(0.0f, 1.0f - t1 * t1)) + s * t2;
	float3 nh = T1 * t1 + T2 * t2 +
	vh * sqrtf(fmaxf(0.0f, 1.0f - t1 * t1 - t2 * t2));
	return normalize(f3(alpha * nh.x, alpha * nh.y, fmaxf(1e-6f, nh.z)));
	}

	__device__ __forceinline__ float ggx_pdf(float a2, float3 wi, float3 wo) {
	float3 h = normalize(wi + wo);
	float ch = fmaxf(h.z, 1e-6f);
	float wih = fmaxf(dot(wi, h), 1e-6f);
	float g1 = 1.0f / (1.0f + ggx_lambda(a2, wi.z));
	return g1 * ggx_d(a2, ch) * wih / fmaxf(wi.z, 1e-6f) / (4.0f * wih);
	}

	// f_spec * cos_o for a GGX lobe with Fresnel factor Fc (already evaluated):
	// D * Fc * G2 / (4 wi.z). The caller divides by pdf or folds geometry.
	__device__ __forceinline__ float ggx_spec_cos(float a2, float3 wi, float3 wo,
	float& sgl_out) {
	float3 h = normalize(wi + wo);
	sgl_out = 0.0f;
	float m = fminf(fmaxf(1.0f - fmaxf(dot(wi, h), 0.0f), 0.0f), 1.0f);
	float m2 = m * m;
	sgl_out = m2 * m2 * m;
	float G2 = 1.0f / (1.0f + ggx_lambda(a2, wi.z) + ggx_lambda(a2, wo.z));
	return ggx_d(a2, fmaxf(h.z, 1e-6f)) * G2 / (4.0f * fmaxf(wi.z, 1e-6f));
	}

	__device__ __forceinline__ float fresnel_dielectric(float cos_i, float eta) {
	// eta = n2/n1
	float s2 = (1.0f - cos_i * cos_i) / (eta * eta);
	if (s2 >= 1.0f) return 1.0f;
	float cos_t = sqrtf(1.0f - s2);
	float rs = (cos_i - eta * cos_t) / (cos_i + eta * cos_t);
	float rp = (eta * cos_i - cos_t) / (eta * cos_i + cos_t);
	return 0.5f * (rs * rs + rp * rp);
	}

	__device__ __forceinline__ float schlick_s(float cos_h) {
	float m = fminf(fmaxf(1.0f - cos_h, 0.0f), 1.0f);
	float m2 = m * m;
	return m2 * m2 * m;
	}

	__device__ __forceinline__ float mis_w(float pa, float pb) {
	return pa / fmaxf(pa + pb, 1e-20f);
	}

	__device__ __forceinline__ bool nee_capable(int mt) {
	return mt == kDiffuse \|\| mt == kConductor \|\| mt == kPlastic;
	}

	// ------------------------------------------------------------- path tracing
	__device__ __forceinline__ void scatter_buf(float* shared_buf, float* global_buf,
	int idx, int c, float v) {
	if (shared_buf) atomicAdd(&shared_buf[idx * 3 + c], v);
	else atomicAdd(&global_buf[idx * 3 + c], v);
	}

	struct Factor {
	float val[3];
	float dc[3];
	int fi[4];
	float fw[4];
	};

	// One radiance term: trAcc * (prod of nk factors, tf last when non-null) *
	// S * E, with E from an emission footprint (e_env=false), an env footprint
	// (e_env=true), or plain (ei null). dsa/dss are the medium log-derivative
	// sums for this term (0 when no medium).
	template <bool GRAD>
	__device__ void add_term(const Factor* fs, int nk, const float T[3],
	const float* tf, float S, const float E[3],
	const int* ei, const float* ew, bool e_env,
	const float trAcc[3], float dsa, const float dss[3],
	float3& acc, const GCtx& g) {
	float P[3] = {T[0], T[1], T[2]};
	if (tf) {
	P[0] *= tf[0];
	P[1] *= tf[1];
	P[2] *= tf[2];
	}
	if (!GRAD) {
	acc.x += P[0] * S * E[0];
	acc.y += P[1] * S * E[1];
	acc.z += P[2] * S * E[2];
	return;
	}
	// E-side gradients
	if (ei) {
	for (int c = 0; c < 3; ++c) {
	float base = g.gs[c] * P[c] * S;
	#pragma unroll
	for (int i = 0; i < 4; ++i) {
	if (e_env) atomicAdd(&g.g_env[ei[i] * 3 + c], base * ew[i]);
	else scatter_buf(g.s_get, g.g_etex, ei[i], c, base * ew[i]);
	}
	}
	}
	// medium sigma gradients: g * term * dln
	if (g.s_gm && (dsa != 0.0f \|\| dss[0] != 0.0f \|\| dss[1] != 0.0f \|\|
	dss[2] != 0.0f)) {
	for (int c = 0; c < 3; ++c) {
	float term = g.gs[c] * P[c] * S * E[c];
	atomicAdd(&g.s_gm[c], term * dsa);
	atomicAdd(&g.s_gm[3 + c], term * dss[c]);
	}
	}
	// albedo texel gradients via exclusion products, scaled by trAcc
	if (nk <= 0) return;
	float suf[kMaxBounces + 1][3];
	suf[nk][0] = suf[nk][1] = suf[nk][2] = 1.0f;
	for (int i = nk - 1; i >= 0; --i)
	for (int c = 0; c < 3; ++c) suf[i][c] = suf[i + 1][c] * fs[i].val[c];
	float pref[3] = {1.0f, 1.0f, 1.0f};
	for (int j = 0; j < nk; ++j) {
	for (int c = 0; c < 3; ++c) {
	if (fs[j].dc[c] != 0.0f) {
	float base = g.gs[c] * E[c] * S * trAcc[c] * pref[c] *
	suf[j + 1][c] * fs[j].dc[c];
	#pragma unroll
	for (int i = 0; i < 4; ++i)
	scatter_buf(g.s_gt, g.g_tex, fs[j].fi[i], c, base * fs[j].fw[i]);
	}
	pref[c] *= fs[j].val[c];
	}
	}
	}

	template <bool GRAD>
	__device__ void trace_path(const DevScene& sc, float3 ro, float3 rd,
	curandStatePhilox4_32_10_t& st, int B, int mode,
	float3& acc, const GCtx& g) {
	float T[3] = {1.0f, 1.0f, 1.0f}; // includes bsdf factors AND trAcc
	float trAcc[3] = {1.0f, 1.0f, 1.0f};
	Factor fs[kMaxBounces];
	int nk = 0;
	float prev_pdf = 0.0f;
	bool prev_delta = true;
	bool has_env = sc.env != nullptr;
	float Dtot = 0.0f; // medium distance so far
	int Nsc = 0; // medium scatter events so far
	float sa[3] = {0, 0, 0}, ss[3] = {0, 0, 0}, stt[3] = {0, 0, 0};
	if (sc.has_med) {
	for (int c = 0; c < 3; ++c) {
	sa[c] = sc.med_sa[c];
	ss[c] = fmaxf(sc.med_ss[c], 1e-8f);
	stt[c] = sa[c] + ss[c];
	}
	}

	auto med_dsa = [&](float dterm) { return sc.has_med ? -(Dtot + dterm) : 0.0f; };

	for (int k = 0; k < B; ++k) {
	float dmed = kEnvDist;
	if (sc.has_med) {
	float u = curand_uniform(&st); // in (0, 1]
	dmed = -logf(fmaxf(u, 1e-12f)) / sc.med_sbar;
	}

	float tbest = kEnvDist;
	float bu = 0.0f, bv = 0.0f;
	float3 ng;
	int face = bvh_closest(sc, ro, rd, kRayEps, tbest, bu, bv, ng);

	// ----------------------------------------------------- medium vertex
	if (sc.has_med && dmed < tbest) {
	Dtot += dmed;
	float3 x = ro + rd * dmed;
	// event factor: sigma_s * exp(-sigma_t d) / (sbar * exp(-sbar d))
	float denom = sc.med_sbar * expf(-sc.med_sbar * dmed);
	if (nk >= kMaxBounces) break;
	Factor& f = fs[nk++];
	for (int c = 0; c < 3; ++c) {
	f.val[c] = ss[c] * expf(-stt[c] * dmed) / denom;
	f.dc[c] = 0.0f;
	T[c] *= f.val[c];
	}
	#pragma unroll
	for (int i = 0; i < 4; ++i) { f.fi[i] = 0; f.fw[i] = 0.0f; }
	Nsc += 1;

	// light NEE from the medium point (isotropic phase, no cosine at x)
	bool draws_light = (mode != kModeBrdf) && sc.nl > 0 &&
	(mode == kModeMis \|\| k + 1 < B);
	if (draws_light) {
	float r1 = curand_uniform(&st);
	float r2 = curand_uniform(&st);
	float r3 = curand_uniform(&st);
	int li = cdf_pick(sc.lcdf, sc.nl, r1);
	int lface = sc.lf[li];
	const float* lv = &sc.tris[lface * 9];
	float3 lv0 = f3(lv[0], lv[1], lv[2]);
	float3 le1 = f3(lv[3], lv[4], lv[5]) - lv0;
	float3 le2 = f3(lv[6], lv[7], lv[8]) - lv0;
	float su = sqrtf(r2);
	float b0 = 1.0f - su, b1 = r3 * su;
	float3 y = lv0 + le1 * b0 + le2 * b1;
	float3 ln = normalize(cross(le1, le2));
	float3 dvec = y - x;
	float d2 = fmaxf(dot(dvec, dvec), 1e-12f);
	float d = sqrtf(d2);
	float3 wo_w = dvec * (1.0f / d);
	if (dot(ln, dvec) > 0.0f) ln = neg(ln);
	float cy = -dot(ln, wo_w);
	if (cy > 1e-6f && d > kShadowEps * 2.0f &&
	!bvh_occluded(sc, x, wo_w, d - kShadowEps)) {
	int lm = sc.mats[lface];
	float E[3];
	int ei[4];
	float ewt[4];
	sample_emission(sc, lm, lface, b0, b1, E, ei, ewt);
	float S = kInv4Pi * cy / d2 * sc.larea;
	float trL[3];
	float trm = 1.0f;
	for (int c = 0; c < 3; ++c) trL[c] = expf(-stt[c] * d);
	(void)trm;
	float w = 1.0f;
	if (mode == kModeMis) {
	float p_lw = d2 / fmaxf(cy * sc.larea, 1e-12f);
	w = mis_w(p_lw, kInv4Pi);
	}
	// fold shadow transmittance into the term factor
	float tf[3] = {trL[0], trL[1], trL[2]};
	float dss_t[3];
	float dsa_t = med_dsa(d);
	for (int c = 0; c < 3; ++c) dss_t[c] = dsa_t + Nsc / ss[c];
	add_term<GRAD>(fs, nk, T, tf, S * w, E, ei, ewt, false, trAcc,
	dsa_t, dss_t, acc, g);
	}
	}
	// isotropic phase continuation
	float p1 = curand_uniform(&st);
	float p2 = curand_uniform(&st);
	float z = 1.0f - 2.0f * p1;
	float rxy = sqrtf(fmaxf(0.0f, 1.0f - z * z));
	float ph = 2.0f * kPi * p2;
	rd = f3(rxy * cosf(ph), z, rxy * sinf(ph));
	ro = x;
	prev_delta = false;
	prev_pdf = kInv4Pi;
	continue;
	}

	// --------------------------------------------------- surface / miss
	if (sc.has_med && face >= 0) {
	// survival ratio to the surface: exp(-sigma_t t) / exp(-sbar t)
	Dtot += tbest;
	float esb = expf(sc.med_sbar * tbest); // 1 / exp(-sbar t)
	for (int c = 0; c < 3; ++c) {
	float r = expf(-stt[c] * tbest) * esb;
	trAcc[c] *= r;
	T[c] *= r;
	}
	}

	if (face < 0) {
	if (has_env && !sc.has_med) {
	float E[3];
	int ei[4];
	float ewd[4];
	env_fetch(sc, rd, E, ei, ewd);
	float w = 1.0f;
	if (mode == kModeMis && !prev_delta)
	w = mis_w(prev_pdf, env_pdf(sc, rd));
	if (mode != kModeNee \|\| prev_delta) {
	float z3[3] = {0, 0, 0};
	add_term<GRAD>(fs, nk, T, nullptr, w, E, ei, ewd, true, trAcc,
	0.0f, z3, acc, g);
	}
	}
	break;
	}

	int m = sc.mats[face];
	int mt = sc.mtype[m];
	float3 x = ro + rd * tbest;
	float3 n = normalize(ng);
	bool backface = dot(n, rd) > 0.0f;
	if (backface) n = neg(n);

	float am[3];
	int idx4[4];
	float w4[4];
	sample_albedo(sc, m, face, bu, bv, am, idx4, w4);

	// emission gather
	{
	float E[3];
	int ei[4];
	float ewt[4];
	sample_emission(sc, m, face, bu, bv, E, ei, ewt);
	if (E[0] > 0.0f \|\| E[1] > 0.0f \|\| E[2] > 0.0f) {
	float w = 1.0f;
	bool add = true;
	if (mode == kModeNee) {
	add = (k == 0);
	} else if (mode == kModeMis && !prev_delta && sc.nl > 0) {
	float cy = fabsf(dot(normalize(ng), rd));
	float p_l = (tbest * tbest) / fmaxf(cy * sc.larea, 1e-12f);
	w = mis_w(prev_pdf, p_l);
	}
	if (add) {
	float dss_t[3];
	float dsa_t = med_dsa(0.0f);
	for (int c = 0; c < 3; ++c)
	dss_t[c] = sc.has_med ? dsa_t + Nsc / ss[c] : 0.0f;
	add_term<GRAD>(fs, nk, T, nullptr, w, E, ei, ewt, false, trAcc,
	dsa_t, dss_t, acc, g);
	}
	}
	}

	// ---------------- smooth dielectric: delta lobes
	if (mt == kDielectric) {
	float c1 = curand_uniform(&st);
	float c2 = curand_uniform(&st);
	float c3 = curand_uniform(&st);
	(void)c2;
	(void)c3;
	float eta_r = backface ? sc.mior[m] : 1.0f / sc.mior[m]; // n1/n2
	float ci = -dot(rd, n);
	float F = fresnel_dielectric(ci, 1.0f / eta_r);
	float3 nd;
	if (c1 < F) {
	nd = rd + n * (2.0f * ci);
	ro = x + n * kRayEps;
	} else {
	float s2 = eta_r * eta_r * (1.0f - ci * ci);
	float ct = sqrtf(fmaxf(0.0f, 1.0f - s2));
	nd = rd * eta_r + n * (eta_r * ci - ct);
	ro = x - n * kRayEps;
	}
	rd = normalize(nd);
	if (nk < kMaxBounces) {
	Factor& f = fs[nk++];
	for (int c = 0; c < 3; ++c) { f.val[c] = 1.0f; f.dc[c] = 0.0f; }
	#pragma unroll
	for (int i = 0; i < 4; ++i) { f.fi[i] = idx4[i]; f.fw[i] = w4[i]; }
	}
	prev_delta = true;
	prev_pdf = 0.0f;
	continue;
	}

	// local frame
	float3 tang, bit;
	onb(n, tang, bit);
	float3 wi_w = neg(rd);
	float3 wi = f3(dot(wi_w, tang), dot(wi_w, bit), dot(wi_w, n));
	wi.z = fmaxf(wi.z, 1e-6f);
	float alpha = fmaxf(sc.mrough[m], 0.01f);
	float a2 = alpha * alpha;

	// ---------------- rough dielectric: BSDF-sampled only
	if (mt == kRoughDielectric) {
	float c1 = curand_uniform(&st);
	float c2 = curand_uniform(&st);
	float c3 = curand_uniform(&st);
	float3 h = ggx_sample_vndf(wi, alpha, c1, c2);
	float cih = dot(wi, h);
	float eta_r = backface ? sc.mior[m] : 1.0f / sc.mior[m]; // n1/n2
	float F = fresnel_dielectric(fmaxf(cih, 1e-6f), 1.0f / eta_r);
	float3 wo;
	bool transmit = false;
	if (c3 < F) {
	wo = h * (2.0f * cih) - wi;
	if (wo.z <= 1e-6f) break;
	} else {
	float s2 = eta_r * eta_r * (1.0f - cih * cih);
	if (s2 >= 1.0f) { // numerical TIR guard
	wo = h * (2.0f * cih) - wi;
	if (wo.z <= 1e-6f) break;
	} else {
	float ct = sqrtf(1.0f - s2);
	wo = neg(wi) * eta_r + h * (eta_r * cih - ct);
	if (wo.z >= -1e-6f) break;
	transmit = true;
	}
	}
	float li = ggx_lambda(a2, wi.z);
	float lo = ggx_lambda(a2, wo.z);
	float gw = (1.0f + li) / (1.0f + li + lo); // G2/G1
	if (nk < kMaxBounces) {
	Factor& f = fs[nk++];
	for (int c = 0; c < 3; ++c) { f.val[c] = gw; f.dc[c] = 0.0f; }
	#pragma unroll
	for (int i = 0; i < 4; ++i) { f.fi[i] = idx4[i]; f.fw[i] = w4[i]; }
	T[0] *= gw;
	T[1] *= gw;
	T[2] *= gw;
	}
	rd = normalize(tang * wo.x + bit * wo.y + n * wo.z);
	ro = transmit ? x - n * kRayEps : x + n * kRayEps;
	prev_delta = true; // no NEE partner: gathers take w = 1
	prev_pdf = 0.0f;
	continue;
	}

	// ---------------- light NEE (diffuse / conductor / plastic)
	if ((mode != kModeBrdf) && sc.nl > 0 &&
	(mode == kModeMis \|\| k + 1 < B)) {
	float r1 = curand_uniform(&st);
	float r2 = curand_uniform(&st);
	float r3 = curand_uniform(&st);
	int li = cdf_pick(sc.lcdf, sc.nl, r1);
	int lface = sc.lf[li];
	const float* lv = &sc.tris[lface * 9];
	float3 lv0 = f3(lv[0], lv[1], lv[2]);
	float3 le1 = f3(lv[3], lv[4], lv[5]) - lv0;
	float3 le2 = f3(lv[6], lv[7], lv[8]) - lv0;
	float su = sqrtf(r2);
	float b0 = 1.0f - su, b1 = r3 * su;
	float3 y = lv0 + le1 * b0 + le2 * b1;
	float3 ln = normalize(cross(le1, le2));
	float3 dvec = y - x;
	float d2 = fmaxf(dot(dvec, dvec), 1e-12f);
	float d = sqrtf(d2);
	float3 wo_w = dvec * (1.0f / d);
	if (dot(ln, dvec) > 0.0f) ln = neg(ln);
	float cx = dot(n, wo_w);
	float cy = -dot(ln, wo_w);
	if (cx > 1e-6f && cy > 1e-6f && d > kShadowEps * 2.0f &&
	!bvh_occluded(sc, x + n * kRayEps, wo_w, d - kShadowEps)) {
	int lm = sc.mats[lface];
	float3 wo = f3(dot(wo_w, tang), dot(wo_w, bit), cx);
	float S_geo = cx * cy / d2 * sc.larea;
	float tf_val[3], tf_dc[3];
	float S = S_geo;
	float p_bw = 0.0f;
	if (mt == kDiffuse) {
	S = S_geo * kInvPi;
	for (int c = 0; c < 3; ++c) { tf_val[c] = am[c]; tf_dc[c] = 1.0f; }
	p_bw = fmaxf(wo.z, 0.0f) * kInvPi;
	} else if (mt == kConductor) {
	float sgl;
	float gg = ggx_spec_cos(a2, wi, wo, sgl) / fmaxf(wo.z, 1e-6f);
	// gg here = D G2 / (4 wi.z wo.z); contribution fcos = ggwo.z
	S = S_geo * gg;
	for (int c = 0; c < 3; ++c) {
	tf_val[c] = am[c] * (1.0f - sgl) + sgl;
	tf_dc[c] = 1.0f - sgl;
	}
	p_bw = ggx_pdf(a2, wi, wo);
	} else { // plastic
	float sgl;
	float speco = ggx_spec_cos(a2, wi, wo, sgl); // D G2 / (4 wi.z)
	float Fc = kF0Coat + (1.0f - kF0Coat) * sgl;
	float spec_f = Fc * speco / fmaxf(wo.z, 1e-6f);
	for (int c = 0; c < 3; ++c) {
	tf_val[c] = am[c] * kInvPi + spec_f;
	tf_dc[c] = kInvPi;
	}
	p_bw = 0.5f * fmaxf(wo.z, 0.0f) * kInvPi + 0.5f * ggx_pdf(a2, wi, wo);
	}
	float w = 1.0f;
	if (mode == kModeMis) {
	float p_lw = d2 / fmaxf(cy * sc.larea, 1e-12f);
	w = mis_w(p_lw, p_bw);
	}
	float E[3];
	int ei[4];
	float ewt[4];
	sample_emission(sc, lm, lface, b0, b1, E, ei, ewt);
	float dsa_t = med_dsa(d);
	float dss_t[3] = {0, 0, 0};
	if (sc.has_med) {
	for (int c = 0; c < 3; ++c) {
	tf_val[c] = expf(-stt[c] d);
	tf_dc[c] = expf(-stt[c] d);
	dss_t[c] = dsa_t + Nsc / ss[c];
	}
	}
	if (nk < kMaxBounces) {
	Factor& f = fs[nk];
	for (int c = 0; c < 3; ++c) { f.val[c] = tf_val[c]; f.dc[c] = tf_dc[c]; }
	#pragma unroll
	for (int i = 0; i < 4; ++i) { f.fi[i] = idx4[i]; f.fw[i] = w4[i]; }
	add_term<GRAD>(fs, nk + 1, T, f.val, S * w, E, ei, ewt, false,
	trAcc, dsa_t, dss_t, acc, g);
	}
	}
	}

	// ---------------- environment NEE (MIS mode only)
	if (mode == kModeMis && has_env) {
	float e1 = curand_uniform(&st);
	float e2 = curand_uniform(&st);
	float e3 = curand_uniform(&st);
	float3 wo_w;
	if (e1 < 0.5f) {
	int y = cdf_pick(sc.ecdf_m, sc.eh, e2);
	int xcol = cdf_pick(&sc.ecdf_c[y * sc.ew], sc.ew, e3);
	float u = (xcol + 0.5f) / sc.ew;
	float v = (y + 0.5f) / sc.eh;
	wo_w = equirect_to_dir(u, v);
	} else {
	float z = 1.0f - 2.0f * e2;
	float rxy = sqrtf(fmaxf(0.0f, 1.0f - z * z));
	float ph = 2.0f * kPi * e3;
	wo_w = f3(rxy * cosf(ph), z, rxy * sinf(ph));
	}
	float cx = dot(n, wo_w);
	if (cx > 1e-6f &&
	!bvh_occluded(sc, x + n * kRayEps, wo_w, kEnvDist)) {
	float p = env_pdf(sc, wo_w);
	float3 wo = f3(dot(wo_w, tang), dot(wo_w, bit), cx);
	float tf_val[3], tf_dc[3];
	float S;
	float p_bw;
	if (mt == kDiffuse) {
	S = cx * kInvPi / p;
	for (int c = 0; c < 3; ++c) { tf_val[c] = am[c]; tf_dc[c] = 1.0f; }
	p_bw = fmaxf(wo.z, 0.0f) * kInvPi;
	} else if (mt == kConductor) {
	float sgl;
	float speco = ggx_spec_cos(a2, wi, wo, sgl); // D G2 / (4 wi.z)
	S = speco / p;
	for (int c = 0; c < 3; ++c) {
	tf_val[c] = am[c] * (1.0f - sgl) + sgl;
	tf_dc[c] = 1.0f - sgl;
	}
	p_bw = ggx_pdf(a2, wi, wo);
	} else { // plastic
	float sgl;
	float speco = ggx_spec_cos(a2, wi, wo, sgl);
	float Fc = kF0Coat + (1.0f - kF0Coat) * sgl;
	S = cx / p;
	float spec_f = Fc * speco / fmaxf(cx, 1e-6f);
	for (int c = 0; c < 3; ++c) {
	tf_val[c] = am[c] * kInvPi + spec_f;
	tf_dc[c] = kInvPi;
	}
	p_bw = 0.5f * fmaxf(wo.z, 0.0f) * kInvPi + 0.5f * ggx_pdf(a2, wi, wo);
	}
	float w = mis_w(p, p_bw);
	float E[3];
	int ei[4];
	float ewt[4];
	env_fetch(sc, wo_w, E, ei, ewt);
	float z3[3] = {0, 0, 0};
	if (nk < kMaxBounces) {
	Factor& f = fs[nk];
	for (int c = 0; c < 3; ++c) { f.val[c] = tf_val[c]; f.dc[c] = tf_dc[c]; }
	#pragma unroll
	for (int i = 0; i < 4; ++i) { f.fi[i] = idx4[i]; f.fw[i] = w4[i]; }
	add_term<GRAD>(fs, nk + 1, T, f.val, S * w, E, ei, ewt, true,
	trAcc, 0.0f, z3, acc, g);
	}
	}
	}

	// ---------------- BSDF continuation
	float c1 = curand_uniform(&st);
	float c2 = curand_uniform(&st);
	float c3 = curand_uniform(&st);
	float3 wo;
	float fv[3], fdc[3];
	if (mt == kDiffuse) {
	float rr = sqrtf(c1);
	float phi = 2.0f * kPi * c2;
	wo = f3(rr * cosf(phi), rr * sinf(phi), sqrtf(fmaxf(0.0f, 1.0f - c1)));
	prev_pdf = fmaxf(wo.z, 1e-8f) * kInvPi;
	for (int c = 0; c < 3; ++c) { fv[c] = am[c]; fdc[c] = 1.0f; }
	} else if (mt == kConductor) {
	float3 h = ggx_sample_vndf(wi, alpha, c1, c2);
	float wih = dot(wi, h);
	wo = h * (2.0f * wih) - wi;
	if (wo.z <= 1e-6f) break;
	float sgl = schlick_s(fmaxf(wih, 0.0f));
	float li = ggx_lambda(a2, wi.z);
	float lo = ggx_lambda(a2, wo.z);
	float gw = (1.0f + li) / (1.0f + li + lo);
	prev_pdf = ggx_pdf(a2, wi, wo);
	for (int c = 0; c < 3; ++c) {
	fv[c] = (am[c] * (1.0f - sgl) + sgl) * gw;
	fdc[c] = (1.0f - sgl) * gw;
	}
	} else { // plastic: 50/50 lobe mixture
	if (c3 < 0.5f) {
	float rr = sqrtf(c1);
	float phi = 2.0f * kPi * c2;
	wo = f3(rr * cosf(phi), rr * sinf(phi),
	sqrtf(fmaxf(0.0f, 1.0f - c1)));
	} else {
	float3 h = ggx_sample_vndf(wi, alpha, c1, c2);
	wo = h * (2.0f * dot(wi, h)) - wi;
	}
	if (wo.z <= 1e-6f) break;
	float pdf = 0.5f * wo.z * kInvPi + 0.5f * ggx_pdf(a2, wi, wo);
	prev_pdf = fmaxf(pdf, 1e-10f);
	float sgl;
	float speco = ggx_spec_cos(a2, wi, wo, sgl);
	float Fc = kF0Coat + (1.0f - kF0Coat) * sgl;
	float spec_f = Fc * speco / fmaxf(wo.z, 1e-6f);
	float scale = wo.z / prev_pdf;
	for (int c = 0; c < 3; ++c) {
	fv[c] = (am[c] * kInvPi + spec_f) * scale;
	fdc[c] = kInvPi * scale;
	}
	}
	prev_delta = false;
	rd = normalize(tang * wo.x + bit * wo.y + n * wo.z);
	ro = x + n * kRayEps;
	if (nk < kMaxBounces) {
	Factor& f = fs[nk++];
	for (int c = 0; c < 3; ++c) { f.val[c] = fv[c]; f.dc[c] = fdc[c]; }
	#pragma unroll
	for (int i = 0; i < 4; ++i) { f.fi[i] = idx4[i]; f.fw[i] = w4[i]; }
	}
	T[0] *= fv[0];
	T[1] *= fv[1];
	T[2] *= fv[2];
	}
	}

	__device__ __forceinline__ float3 camera_ray(const DevCam& cam, int px, int py,
	int W, int H, float jx, float jy) {
	float nx = 2.0f * ((px + jx) / (float)W) - 1.0f;
	float ny = 1.0f - 2.0f * ((py + jy) / (float)H);
	float3 d = f3(cam.f[0] + nx * cam.r[0] + ny * cam.u[0],
	cam.f[1] + nx * cam.r[1] + ny * cam.u[1],
	cam.f[2] + nx * cam.r[2] + ny * cam.u[2]);
	return normalize(d);
	}

	__global__ void k_forward(DevScene sc, DevCam cam, int H, int W, int spp,
	int B, int mode, unsigned long long seed,
	float* img) {
	int pid = blockIdx.x * blockDim.x + threadIdx.x;
	if (pid >= H * W) return;
	int px = pid % W, py = pid / W;
	float3 ro = f3(cam.p[0], cam.p[1], cam.p[2]);
	float3 acc = f3(0.0f, 0.0f, 0.0f);
	GCtx g = {};
	for (int s = 0; s < spp; ++s) {
	curandStatePhilox4_32_10_t st;
	curand_init(seed, (unsigned long long)pid * spp + s, 0, &st);
	float jx = curand_uniform(&st);
	float jy = curand_uniform(&st);
	float3 rd = camera_ray(cam, px, py, W, H, jx, jy);
	trace_path<false>(sc, ro, rd, st, B, mode, acc, g);
	}
	float inv_spp = 1.0f / (float)spp;
	img[pid * 3 + 0] = acc.x * inv_spp;
	img[pid * 3 + 1] = acc.y * inv_spp;
	img[pid * 3 + 2] = acc.z * inv_spp;
	}

	__global__ void k_backward(DevScene sc, DevCam cam, int H, int W, int spp,
	int B, int mode, unsigned long long seed,
	const float* gimg, float* g_tex, float* g_etex,
	float* g_env, float* g_med) {
	extern __shared__ float smem[];
	bool staged_t = sc.nt <= kSharedTexels;
	bool staged_e = sc.net <= kSharedEmiTexels;
	float* s_gt = staged_t ? smem : nullptr;
	int off = staged_t ? sc.nt * 3 : 0;
	float* s_get = staged_e ? smem + off : nullptr;
	if (staged_e) off += sc.net * 3;
	float* s_gm = smem + off; // [6]
	int total = off + 6;
	for (int i = threadIdx.x; i < total; i += blockDim.x) smem[i] = 0.0f;
	__syncthreads();

	int pid = blockIdx.x * blockDim.x + threadIdx.x;
	if (pid < H * W) {
	int px = pid % W, py = pid / W;
	float inv_spp = 1.0f / (float)spp;
	float gs[3] = {gimg[pid * 3 + 0] * inv_spp, gimg[pid * 3 + 1] * inv_spp,
	gimg[pid * 3 + 2] * inv_spp};
	GCtx g;
	g.gs = gs;
	g.g_tex = g_tex;
	g.s_gt = s_gt;
	g.g_etex = g_etex;
	g.s_get = s_get;
	g.g_env = g_env;
	g.s_gm = sc.has_med ? s_gm : nullptr;
	float3 ro = f3(cam.p[0], cam.p[1], cam.p[2]);
	float3 acc = f3(0.0f, 0.0f, 0.0f);
	for (int s = 0; s < spp; ++s) {
	curandStatePhilox4_32_10_t st;
	curand_init(seed, (unsigned long long)pid * spp + s, 0, &st);
	float jx = curand_uniform(&st);
	float jy = curand_uniform(&st);
	float3 rd = camera_ray(cam, px, py, W, H, jx, jy);
	trace_path<true>(sc, ro, rd, st, B, mode, acc, g);
	}
	}

	__syncthreads();
	if (staged_t) {
	for (int i = threadIdx.x; i < sc.nt * 3; i += blockDim.x)
	if (s_gt[i] != 0.0f) atomicAdd(&g_tex[i], s_gt[i]);
	}
	if (staged_e) {
	for (int i = threadIdx.x; i < sc.net * 3; i += blockDim.x)
	if (s_get[i] != 0.0f) atomicAdd(&g_etex[i], s_get[i]);
	}
	if (sc.has_med) {
	for (int i = threadIdx.x; i < 6; i += blockDim.x)
	if (s_gm[i] != 0.0f) atomicAdd(&g_med[i], s_gm[i]);
	}
	}

	DevScene make_scene(const PtdSceneArgs& a) {
	DevScene sc;
	sc.tris = a.tris;
	sc.mats = a.mat_ids;
	sc.uvs = a.uvs;
	sc.nf = a.nodes_f;
	sc.ni = a.nodes_i;
	sc.n_nodes = a.n_nodes;
	sc.lf = a.light_faces;
	sc.lcdf = a.light_cdf;
	sc.nl = a.n_lights;
	sc.larea = a.total_light_area;
	sc.tex = a.tex;
	sc.thdr = a.tex_hdr;
	sc.nt = a.n_texels;
	sc.etex = a.emi_tex;
	sc.ehdr = a.emi_hdr;
	sc.net = a.n_emi_texels;
	sc.mtype = a.mat_type;
	sc.mrough = a.mat_rough;
	sc.mior = a.mat_ior;
	sc.nm = a.n_mats;
	sc.med_sa = a.med_sa;
	sc.med_ss = a.med_ss;
	sc.med_sbar = a.med_sbar;
	sc.has_med = a.has_med;
	sc.env = a.env;
	sc.ew = a.env_w;
	sc.eh = a.env_h;
	sc.ecdf_m = a.env_cdf_m;
	sc.ecdf_c = a.env_cdf_c;
	sc.epdf = a.env_pdf;
	return sc;
	}

	DevCam make_cam(const float* cam) {
	DevCam c;
	for (int i = 0; i < 3; ++i) {
	c.p[i] = cam[i];
	c.f[i] = cam[3 + i];
	c.r[i] = cam[6 + i];
	c.u[i] = cam[9 + i];
	}
	return c;
	}

	} // namespace

	extern "C" void ptd_forward_launch(const PtdSceneArgs* args, const float* cam,
	int H, int W, int spp, int max_bounces,
	int mode, long long seed, float* image,
	cudaStream_t stream) {
	DevScene sc = make_scene(*args);
	DevCam dc = make_cam(cam);
	int blocks = (H * W + kThreads - 1) / kThreads;
	k_forward<<<blocks, kThreads, 0, stream>>>(
	sc, dc, H, W, spp, max_bounces, mode, (unsigned long long)seed, image);
	}

	extern "C" void ptd_backward_launch(const PtdSceneArgs* args, const float* cam,
	int H, int W, int spp, int max_bounces,
	int mode, long long seed,
	const float* grad_image, float* grad_tex,
	float* grad_emi_tex, float* grad_env,
	float* grad_med, cudaStream_t stream) {
	DevScene sc = make_scene(*args);
	DevCam dc = make_cam(cam);
	int blocks = (H * W + kThreads - 1) / kThreads;
	size_t smem = 6 * sizeof(float);
	if (args->n_texels <= kSharedTexels)
	smem += (size_t)args->n_texels * 3 * sizeof(float);
	if (args->n_emi_texels <= kSharedEmiTexels)
	smem += (size_t)args->n_emi_texels * 3 * sizeof(float);
	k_backward<<<blocks, kThreads, smem, stream>>>(
	sc, dc, H, W, spp, max_bounces, mode, (unsigned long long)seed,
	grad_image, grad_tex, grad_emi_tex, grad_env, grad_med);
	}