Skip to content

Commit 8846466

Browse files
yozhijkyozhijk
committed
BDPT proto
1 parent e0d1915 commit 8846466

18 files changed

+2592
-112
lines changed

‎App/CL/camera.cl

Lines changed: 294 additions & 56 deletions
Original file line numberDiff line numberDiff line change
@@ -27,20 +27,21 @@ THE SOFTWARE.
2727
#include <../App/CL/sampling.cl>
2828
#include <../App/CL/utils.cl>
2929
#include <../App/CL/path.cl>
30+
#include <../App/CL/vertex.cl>
3031

3132
// Pinhole camera implementation.
3233
// This kernel is being used if aperture value = 0.
33-
KERNEL
34+
KERNEL
3435
void PerspectiveCamera_GeneratePaths(
3536
// Camera
3637
GLOBAL Camera const* restrict camera,
3738
// Image resolution
3839
int output_width,
3940
int output_height,
40-
// Pixel domain buffer
41-
GLOBAL int const* restrict pixel_idx,
42-
// Size of pixel domain buffer
43-
GLOBAL int const* restrict num_pixels,
41+
// Pixel domain buffer
42+
GLOBAL int const* restrict pixel_idx,
43+
// Size of pixel domain buffer
44+
GLOBAL int const* restrict num_pixels,
4445
// RNG seed value
4546
uint rng_seed,
4647
// Current frame
@@ -59,9 +60,9 @@ void PerspectiveCamera_GeneratePaths(
5960
// Check borders
6061
if (global_id < *num_pixels)
6162
{
62-
int idx = pixel_idx[global_id];
63-
int y = idx / output_width;
64-
int x = idx % output_width;
63+
int idx = pixel_idx[global_id];
64+
int y = idx / output_width;
65+
int x = idx % output_width;
6566

6667
// Get pointer to ray & path handles
6768
GLOBAL ray* my_ray = rays + global_id;
@@ -124,69 +125,69 @@ void PerspectiveCamera_GeneratePaths(
124125
// Physical camera implemenation.
125126
// This kernel is being used if aperture > 0.
126127
KERNEL void PerspectiveCameraDof_GeneratePaths(
127-
// Camera
128-
GLOBAL Camera const* restrict camera,
129-
// Image resolution
130-
int output_width,
131-
int output_height,
132-
// Pixel domain buffer
133-
GLOBAL int const* restrict pixel_idx,
134-
// Size of pixel domain buffer
135-
GLOBAL int const* restrict num_pixels,
136-
// RNG seed value
137-
uint rng_seed,
138-
// Current frame
139-
uint frame,
140-
// Rays to generate
141-
GLOBAL ray* restrict rays,
142-
// RNG data
143-
GLOBAL uint* restrict random,
144-
GLOBAL uint const* restrict sobolmat,
145-
// Path buffer
146-
GLOBAL Path* restrict paths
147-
)
128+
// Camera
129+
GLOBAL Camera const* restrict camera,
130+
// Image resolution
131+
int output_width,
132+
int output_height,
133+
// Pixel domain buffer
134+
GLOBAL int const* restrict pixel_idx,
135+
// Size of pixel domain buffer
136+
GLOBAL int const* restrict num_pixels,
137+
// RNG seed value
138+
uint rng_seed,
139+
// Current frame
140+
uint frame,
141+
// Rays to generate
142+
GLOBAL ray* restrict rays,
143+
// RNG data
144+
GLOBAL uint* restrict random,
145+
GLOBAL uint const* restrict sobolmat,
146+
// Path buffer
147+
GLOBAL Path* restrict paths
148+
)
148149
{
149-
int global_id = get_global_id(0);
150+
int global_id = get_global_id(0);
150151

151-
// Check borders
152-
if (global_id < *num_pixels)
153-
{
154-
int idx = pixel_idx[global_id];
155-
int y = idx / output_width;
156-
int x = idx % output_width;
152+
// Check borders
153+
if (global_id < *num_pixels)
154+
{
155+
int idx = pixel_idx[global_id];
156+
int y = idx / output_width;
157+
int x = idx % output_width;
157158

158-
// Get pointer to ray & path handles
159-
GLOBAL ray* my_ray = rays + global_id;
160-
GLOBAL Path* my_path = paths + y * output_width + x;
159+
// Get pointer to ray & path handles
160+
GLOBAL ray* my_ray = rays + global_id;
161+
GLOBAL Path* my_path = paths + y * output_width + x;
161162

162-
// Initialize sampler
163-
Sampler sampler;
163+
// Initialize sampler
164+
Sampler sampler;
164165
#if SAMPLER == SOBOL
165-
uint scramble = random[x + output_width * y] * 0x1fe3434f;
166+
uint scramble = random[x + output_width * y] * 0x1fe3434f;
166167

167-
if (frame & 0xF)
168-
{
169-
random[x + output_width * y] = WangHash(scramble);
170-
}
168+
if (frame & 0xF)
169+
{
170+
random[x + output_width * y] = WangHash(scramble);
171+
}
171172

172-
Sampler_Init(&sampler, frame, SAMPLE_DIM_CAMERA_OFFSET, scramble);
173+
Sampler_Init(&sampler, frame, SAMPLE_DIM_CAMERA_OFFSET, scramble);
173174
#elif SAMPLER == RANDOM
174-
uint scramble = x + output_width * y * rng_seed;
175-
Sampler_Init(&sampler, scramble);
175+
uint scramble = x + output_width * y * rng_seed;
176+
Sampler_Init(&sampler, scramble);
176177
#elif SAMPLER == CMJ
177-
uint rnd = random[x + output_width * y];
178-
uint scramble = rnd * 0x1fe3434f * ((frame + 133 * rnd) / (CMJ_DIM * CMJ_DIM));
179-
Sampler_Init(&sampler, frame % (CMJ_DIM * CMJ_DIM), SAMPLE_DIM_CAMERA_OFFSET, scramble);
178+
uint rnd = random[x + output_width * y];
179+
uint scramble = rnd * 0x1fe3434f * ((frame + 133 * rnd) / (CMJ_DIM * CMJ_DIM));
180+
Sampler_Init(&sampler, frame % (CMJ_DIM * CMJ_DIM), SAMPLE_DIM_CAMERA_OFFSET, scramble);
180181
#endif
181182

182183
// Generate pixel and lens samples
183184
float2 sample0 = Sampler_Sample2D(&sampler, SAMPLER_ARGS);
184185
float2 sample1 = Sampler_Sample2D(&sampler, SAMPLER_ARGS);
185186

186-
// Calculate [0..1] image plane sample
187-
float2 img_sample;
188-
img_sample.x = (float)x / output_width + sample0.x / output_width;
189-
img_sample.y = (float)y / output_height + sample0.y / output_height;
187+
// Calculate [0..1] image plane sample
188+
float2 img_sample;
189+
img_sample.x = (float)x / output_width + sample0.x / output_width;
190+
img_sample.y = (float)y / output_height + sample0.y / output_height;
190191

191192
// Transform into [-0.5, 0.5]
192193
float2 h_sample = img_sample - make_float2(0.5f, 0.5f);
@@ -221,4 +222,241 @@ KERNEL void PerspectiveCameraDof_GeneratePaths(
221222
}
222223
}
223224

225+
226+
KERNEL
227+
void PerspectiveCamera_GenerateVertices(
228+
// Camera
229+
GLOBAL Camera const* restrict camera,
230+
// Image resolution
231+
int output_width,
232+
int output_height,
233+
// Pixel domain buffer
234+
GLOBAL int const* restrict pixel_idx,
235+
// Size of pixel domain buffer
236+
GLOBAL int const* restrict num_pixels,
237+
// RNG seed value
238+
uint rng_seed,
239+
// Current frame
240+
uint frame,
241+
// RNG data
242+
GLOBAL uint* restrict random,
243+
GLOBAL uint const* restrict sobolmat,
244+
// Rays to generate
245+
GLOBAL ray* restrict rays,
246+
// Eye subpath vertices
247+
GLOBAL PathVertex* restrict eye_subpath,
248+
// Eye subpath length
249+
GLOBAL int* restrict eye_subpath_length,
250+
// Path buffer
251+
GLOBAL Path* restrict paths
252+
)
253+
254+
{
255+
int global_id = get_global_id(0);
256+
257+
// Check borders
258+
if (global_id < *num_pixels)
259+
{
260+
int idx = pixel_idx[global_id];
261+
int y = idx / output_width;
262+
int x = idx % output_width;
263+
264+
// Get pointer to ray & path handles
265+
GLOBAL ray* my_ray = rays + global_id;
266+
267+
GLOBAL PathVertex* my_vertex = eye_subpath + BDPT_MAX_SUBPATH_LEN * idx;
268+
GLOBAL int* my_count = eye_subpath_length + idx;
269+
GLOBAL Path* my_path = paths + idx;
270+
271+
// Initialize sampler
272+
Sampler sampler;
273+
#if SAMPLER == SOBOL
274+
uint scramble = random[x + output_width * y] * 0x1fe3434f;
275+
276+
if (frame & 0xF)
277+
{
278+
random[x + output_width * y] = WangHash(scramble);
279+
}
280+
281+
Sampler_Init(&sampler, frame, SAMPLE_DIM_CAMERA_OFFSET, scramble);
282+
#elif SAMPLER == RANDOM
283+
uint scramble = x + output_width * y * rng_seed;
284+
Sampler_Init(&sampler, scramble);
285+
#elif SAMPLER == CMJ
286+
uint rnd = random[x + output_width * y];
287+
uint scramble = rnd * 0x1fe3434f * ((frame + 133 * rnd) / (CMJ_DIM * CMJ_DIM));
288+
Sampler_Init(&sampler, frame % (CMJ_DIM * CMJ_DIM), SAMPLE_DIM_CAMERA_OFFSET, scramble);
289+
#endif
290+
291+
// Generate sample
292+
float2 sample0 = Sampler_Sample2D(&sampler, SAMPLER_ARGS);
293+
294+
// Calculate [0..1] image plane sample
295+
float2 img_sample;
296+
img_sample.x = (float)x / output_width + sample0.x / output_width;
297+
img_sample.y = (float)y / output_height + sample0.y / output_height;
298+
299+
// Transform into [-0.5, 0.5]
300+
float2 h_sample = img_sample - make_float2(0.5f, 0.5f);
301+
// Transform into [-dim/2, dim/2]
302+
float2 c_sample = h_sample * camera->dim;
303+
304+
// Calculate direction to image plane
305+
my_ray->d.xyz = normalize(camera->focal_length * camera->forward + c_sample.x * camera->right + c_sample.y * camera->up);
306+
// Origin == camera position + nearz * d
307+
my_ray->o.xyz = camera->p + camera->zcap.x * my_ray->d.xyz;
308+
// Max T value = zfar - znear since we moved origin to znear
309+
my_ray->o.w = camera->zcap.y - camera->zcap.x;
310+
// Generate random time from 0 to 1
311+
my_ray->d.w = sample0.x;
312+
// Set ray max
313+
my_ray->extra.x = 0xFFFFFFFF;
314+
my_ray->extra.y = 0xFFFFFFFF;
315+
Ray_SetExtra(my_ray, 1.f);
316+
317+
PathVertex v;
318+
PathVertex_Init(&v,
319+
camera->p,
320+
camera->forward,
321+
camera->forward,
322+
0.f,
323+
1.f,
324+
1.f,
325+
1.f,
326+
kCamera,
327+
-1);
328+
329+
*my_count = 1;
330+
*my_vertex = v;
331+
332+
// Initlize path data
333+
my_path->throughput = make_float3(1.f, 1.f, 1.f);
334+
my_path->volume = -1;
335+
my_path->flags = 0;
336+
my_path->active = 0xFF;
337+
}
338+
}
339+
340+
KERNEL
341+
void PerspectiveCameraDof_GenerateVertices(
342+
// Camera
343+
GLOBAL Camera const* restrict camera,
344+
// Image resolution
345+
int output_width,
346+
int output_height,
347+
// Pixel domain buffer
348+
GLOBAL int const* restrict pixel_idx,
349+
// Size of pixel domain buffer
350+
GLOBAL int const* restrict num_pixels,
351+
// RNG seed value
352+
uint rng_seed,
353+
// Current frame
354+
uint frame,
355+
// RNG data
356+
GLOBAL uint* restrict random,
357+
GLOBAL uint const* restrict sobolmat,
358+
// Rays to generate
359+
GLOBAL ray* restrict rays,
360+
// Eye subpath vertices
361+
GLOBAL PathVertex* restrict eye_subpath,
362+
// Eye subpath length
363+
GLOBAL int* restrict eye_subpath_length,
364+
// Path buffer
365+
GLOBAL Path* restrict paths
366+
)
367+
368+
{
369+
int global_id = get_global_id(0);
370+
371+
// Check borders
372+
if (global_id < *num_pixels)
373+
{
374+
int idx = pixel_idx[global_id];
375+
int y = idx / output_width;
376+
int x = idx % output_width;
377+
378+
// Get pointer to ray & path handles
379+
GLOBAL ray* my_ray = rays + global_id;
380+
GLOBAL PathVertex* my_vertex = eye_subpath + BDPT_MAX_SUBPATH_LEN * (y * output_width + x);
381+
GLOBAL int* my_count = eye_subpath_length + y * output_width + x;
382+
GLOBAL Path* my_path = paths + y * output_width + x;
383+
384+
// Initialize sampler
385+
Sampler sampler;
386+
#if SAMPLER == SOBOL
387+
uint scramble = random[x + output_width * y] * 0x1fe3434f;
388+
389+
if (frame & 0xF)
390+
{
391+
random[x + output_width * y] = WangHash(scramble);
392+
}
393+
394+
Sampler_Init(&sampler, frame, SAMPLE_DIM_CAMERA_OFFSET, scramble);
395+
#elif SAMPLER == RANDOM
396+
uint scramble = x + output_width * y * rng_seed;
397+
Sampler_Init(&sampler, scramble);
398+
#elif SAMPLER == CMJ
399+
uint rnd = random[x + output_width * y];
400+
uint scramble = rnd * 0x1fe3434f * ((frame + 133 * rnd) / (CMJ_DIM * CMJ_DIM));
401+
Sampler_Init(&sampler, frame % (CMJ_DIM * CMJ_DIM), SAMPLE_DIM_CAMERA_OFFSET, scramble);
402+
#endif
403+
404+
// Generate pixel and lens samples
405+
float2 sample0 = Sampler_Sample2D(&sampler, SAMPLER_ARGS);
406+
float2 sample1 = Sampler_Sample2D(&sampler, SAMPLER_ARGS);
407+
408+
// Calculate [0..1] image plane sample
409+
float2 img_sample;
410+
img_sample.x = (float)x / output_width + sample0.x / output_width;
411+
img_sample.y = (float)y / output_height + sample0.y / output_height;
412+
413+
// Transform into [-0.5, 0.5]
414+
float2 h_sample = img_sample - make_float2(0.5f, 0.5f);
415+
// Transform into [-dim/2, dim/2]
416+
float2 c_sample = h_sample * camera->dim;
417+
418+
// Generate sample on the lens
419+
float2 lens_sample = camera->aperture * Sample_MapToDiskConcentric(sample1);
420+
// Calculate position on focal plane
421+
float2 focal_plane_sample = c_sample * camera->focus_distance / camera->focal_length;
422+
// Calculate ray direction
423+
float2 camera_dir = focal_plane_sample - lens_sample;
424+
425+
// Calculate direction to image plane
426+
my_ray->d.xyz = normalize(camera->forward * camera->focus_distance + camera->right * camera_dir.x + camera->up * camera_dir.y);
427+
// Origin == camera position + nearz * d
428+
my_ray->o.xyz = camera->p + lens_sample.x * camera->right + lens_sample.y * camera->up;
429+
// Max T value = zfar - znear since we moved origin to znear
430+
my_ray->o.w = camera->zcap.y - camera->zcap.x;
431+
// Generate random time from 0 to 1
432+
my_ray->d.w = sample0.x;
433+
// Set ray max
434+
my_ray->extra.x = 0xFFFFFFFF;
435+
my_ray->extra.y = 0xFFFFFFFF;
436+
Ray_SetExtra(my_ray, 1.f);
437+
438+
PathVertex v;
439+
PathVertex_Init(&v,
440+
camera->p,
441+
camera->forward,
442+
camera->forward,
443+
0.f,
444+
1.f,
445+
1.f,
446+
1.f,
447+
kCamera,
448+
-1);
449+
450+
*my_count = 1;
451+
*my_vertex = v;
452+
453+
// Initlize path data
454+
my_path->throughput = make_float3(1.f, 1.f, 1.f);
455+
my_path->volume = -1;
456+
my_path->flags = 0;
457+
my_path->active = 0xFF;
458+
}
459+
}
460+
461+
224462
#endif // CAMERA_CL

0 commit comments

Comments
 (0)