@@ -284,7 +284,8 @@ __kernel void ShadeSurface(
284
284
// Indirect rays
285
285
__global ray * indirectrays ,
286
286
// Radiance
287
- __global float3 * output
287
+ __global float3 * output ,
288
+ __global int * numshadowrays
288
289
)
289
290
{
290
291
int globalid = get_global_id (0 );
@@ -307,6 +308,9 @@ __kernel void ShadeSurface(
307
308
// Only applied to active rays after compaction
308
309
if (globalid < * numhits )
309
310
{
311
+ if (globalid == 0 )
312
+ * numshadowrays = * numhits * 2 ;
313
+
310
314
// Fetch index
311
315
int hitidx = hitindices [globalid ];
312
316
int pixelidx = pixelindices [globalid ];
@@ -343,6 +347,10 @@ __kernel void ShadeSurface(
343
347
sample3 .y = SobolSampler_Sample1D (sampler -> seq , GetSampleDim (bounce , kIndirectV ), sampler -> s0 , sobolmat );
344
348
345
349
float sample4 = SobolSampler_Sample1D (sampler -> seq , GetSampleDim (bounce , kRR ), sampler -> s0 , sobolmat );
350
+
351
+ float2 sample5 ;
352
+ sample5 .x = SobolSampler_Sample1D (sampler -> seq , GetSampleDim (bounce , kIndirectU ), sampler -> s0 , sobolmat );
353
+ sample5 .y = SobolSampler_Sample1D (sampler -> seq , GetSampleDim (bounce , kIndirectV ), sampler -> s0 , sobolmat );
346
354
#else
347
355
// Prepare RNG
348
356
Rng rng ;
@@ -351,6 +359,7 @@ __kernel void ShadeSurface(
351
359
float2 sample1 = UniformSampler_Sample2D (& rng );
352
360
float2 sample2 = UniformSampler_Sample2D (& rng );
353
361
float2 sample3 = UniformSampler_Sample2D (& rng );
362
+ float2 sample5 = UniformSampler_Sample2D (& rng );
354
363
float sample4 = UniformSampler_Sample2D (& rng ).x ;
355
364
#endif
356
365
@@ -388,23 +397,25 @@ __kernel void ShadeSurface(
388
397
// Terminate if emissive
389
398
if (Bxdf_IsEmissive (& diffgeo ))
390
399
{
391
- if (ndotwi > 0.f )
400
+ if (bounce == 0 && ndotwi > 0.f )
392
401
{
393
402
// There can be two cases: first we hit after specular vertex or primary ray
394
403
// where MIS can't be applied. In this case we simply pass radiance as is
395
404
// because we were using BRDF based estimator at previous step
396
405
{
397
406
// In this case we hit after an application of MIS process at previous step.
398
407
// That means BRDF weight has been already applied.
399
- output [pixelidx ] += Path_GetThroughput ( path ) * Emissive_GetLe (& diffgeo , TEXTURE_ARGS );
408
+ output [pixelidx ] += Emissive_GetLe (& diffgeo , TEXTURE_ARGS );
400
409
}
401
410
}
402
411
403
412
Path_Kill (path );
404
- Ray_SetInactive (shadowrays + globalid );
413
+ Ray_SetInactive (shadowrays + 2 * globalid );
414
+ Ray_SetInactive (shadowrays + 2 * globalid + 1 );
405
415
Ray_SetInactive (indirectrays + globalid );
406
416
407
- lightsamples [globalid ] = 0.f ;
417
+ lightsamples [2 * globalid ] = 0.f ;
418
+ lightsamples [2 * globalid + 1 ] = 0.f ;
408
419
return ;
409
420
}
410
421
@@ -443,64 +454,140 @@ __kernel void ShadeSurface(
443
454
float3 wo ;
444
455
float bxdfweight = 1.f ;
445
456
float lightweight = 1.f ;
457
+ float3 throughput = Path_GetThroughput (path );
446
458
447
- int lightidx = num_lights > 0 ? Scene_SampleLight (& scene , sample0 .y , & selection_pdf ) : -1 ;
459
+ // Sample light
460
+ int light_idx = Scene_SampleLight (& scene , sample0 .y , & selection_pdf );
461
+ bool light_singular = Light_IsSingular (scene .lights + light_idx );
462
+ bool bxdf_singular = Bxdf_IsSingular (& diffgeo );
448
463
449
- float3 throughput = Path_GetThroughput (path );
464
+ if (light_singular && bxdf_singular )
465
+ {
466
+ Path_Kill (path );
467
+ Ray_SetInactive (shadowrays + 2 * globalid );
468
+ Ray_SetInactive (indirectrays + 2 * globalid );
469
+ Ray_SetInactive (shadowrays + 2 * globalid + 1 );
470
+ Ray_SetInactive (indirectrays + 2 * globalid + 1 );
450
471
451
- // Sample bxdf
452
- float3 bxdf = Bxdf_Sample (& diffgeo , wi , TEXTURE_ARGS , sample2 , & bxdfwo , & bxdfpdf );
472
+ lightsamples [2 * globalid ] = 0.f ;
473
+ lightsamples [2 * globalid + 1 ] = 0.f ;
474
+ return ;
475
+ }
453
476
454
- // If we have light to sample we can hopefully do mis
455
- if (lightidx > -1 )
477
+ /*if (!bxdf_singular)
456
478
{
457
479
// Sample light
458
- float3 le = Light_Sample (lightidx , & scene , & diffgeo , TEXTURE_ARGS , sample1 , & lightwo , & lightpdf );
480
+ float3 le = Light_Sample(light_idx , &scene, &diffgeo, TEXTURE_ARGS, sample1, &lightwo, &lightpdf);
459
481
lightbxdfpdf = Bxdf_GetPdf(&diffgeo, wi, normalize(lightwo), TEXTURE_ARGS);
460
482
lightweight = BalanceHeuristic(1, lightpdf, 1, lightbxdfpdf);
461
483
462
- // Sample BxDF
463
- bxdflightpdf = Light_GetPdf (lightidx , & scene , & diffgeo , bxdfwo , TEXTURE_ARGS );
464
- bxdfweight = BalanceHeuristic (1 , bxdfpdf , 1 , bxdflightpdf );
465
-
466
484
// Apply MIS to account for both
467
- if (NON_BLACK (le ) && lightpdf > 0.0f && ! Bxdf_IsSingular ( & diffgeo ) )
485
+ if (NON_BLACK(le) && lightpdf > 0.0f)
468
486
{
469
487
wo = lightwo;
470
488
float ndotwo = fabs(dot(diffgeo.n, normalize(wo)));
471
489
radiance = le * Bxdf_Evaluate(&diffgeo, wi, normalize(wo), TEXTURE_ARGS) * throughput * ndotwo * lightweight / lightpdf / selection_pdf;
490
+
491
+ if (NON_BLACK(radiance))
492
+ {
493
+ // Generate shadow ray
494
+ float shadow_ray_length = (1.f - 2.f * CRAZY_LOW_DISTANCE) * length(wo);
495
+ float3 shadow_ray_dir = normalize(wo);
496
+ float3 shadow_ray_o = diffgeo.p + CRAZY_LOW_DISTANCE * s * diffgeo.n;
497
+ int shadow_ray_mask = Bxdf_IsSingular(&diffgeo) ? 0xFFFFFFFF : 0x0000FFFF;
498
+
499
+ Ray_Init(shadowrays + 2 * globalid, shadow_ray_o, shadow_ray_dir, shadow_ray_length, 0.f, shadow_ray_mask);
500
+
501
+ // Apply the volume to shadow ray if needed
502
+ int volidx = Path_GetVolumeIdx(path);
503
+ if (volidx != -1)
504
+ {
505
+ radiance *= Volume_Transmittance(&volumes[volidx], &shadowrays[2 * globalid], shadow_ray_length);
506
+ radiance += Volume_Emission(&volumes[volidx], &shadowrays[2 * globalid], shadow_ray_length) * throughput;
507
+ }
508
+
509
+ // And write the light sample
510
+ lightsamples[2 * globalid] = REASONABLE_RADIANCE(radiance);
511
+ }
512
+ else
513
+ {
514
+ // Otherwise save some intersector cycles
515
+ Ray_SetInactive(shadowrays + 2 * globalid);
516
+ lightsamples[2 * globalid] = 0;
517
+ }
518
+ }
519
+ else
520
+ {
521
+ // Otherwise save some intersector cycles
522
+ Ray_SetInactive(shadowrays + 2 * globalid);
523
+ lightsamples[2 * globalid] = 0;
472
524
}
473
525
}
526
+ else
527
+ {*/
528
+ // Otherwise save some intersector cycles
529
+ Ray_SetInactive (shadowrays + 2 * globalid );
530
+ lightsamples [2 * globalid ] = 0 ;
531
+ //}
474
532
475
- // If we have some light here generate a shadow ray
476
- if (NON_BLACK (radiance ))
533
+ if (1 /*!light_singular*/ )
477
534
{
478
- // Generate shadow ray
479
- float shadow_ray_length = (1.f - 2.f * CRAZY_LOW_DISTANCE ) * length (wo );
480
- float3 shadow_ray_dir = normalize (wo );
481
- float3 shadow_ray_o = diffgeo .p + CRAZY_LOW_DISTANCE * s * diffgeo .n ;
482
- int shadow_ray_mask = Bxdf_IsSingular (& diffgeo ) ? 0xFFFFFFFF : 0x0000FFFF ;
535
+ // Sample bxdf
536
+ float3 bxdf = Bxdf_Sample (& diffgeo , wi , TEXTURE_ARGS , sample2 , & bxdfwo , & bxdfpdf );
537
+ bxdfwo *= CRAZY_HIGH_DISTANCE ;
483
538
484
- Ray_Init (shadowrays + globalid , shadow_ray_o , shadow_ray_dir , shadow_ray_length , 0.f , shadow_ray_mask );
539
+ bxdflightpdf = Light_GetPdf (light_idx , & scene , & diffgeo , bxdfwo , TEXTURE_ARGS );
540
+ bxdfweight = BalanceHeuristic (1 , bxdfpdf , 1 , bxdflightpdf );
485
541
486
- // Apply the volume to shadow ray if needed
487
- int volidx = Path_GetVolumeIdx (path );
488
- if (volidx != -1 )
542
+ // Apply MIS to account for both
543
+ if (NON_BLACK (bxdf ) && bxdfpdf > 0.0f )
489
544
{
490
- radiance *= Volume_Transmittance ( & volumes [ volidx ], & shadowrays [ globalid ], shadow_ray_length ) ;
491
- radiance += Volume_Emission ( & volumes [ volidx ], & shadowrays [ globalid ], shadow_ray_length ) * throughput ;
492
- }
545
+ wo = bxdfwo ;
546
+ float ndotwo = fabs ( dot ( diffgeo . n , normalize ( wo ))) ;
547
+ radiance = Light_GetLe ( light_idx , & scene , & diffgeo , & wo , TEXTURE_ARGS ) * bxdf * throughput * ndotwo * bxdfweight / bxdfpdf / selection_pdf ;
493
548
494
- // And write the light sample
495
- lightsamples [globalid ] = REASONABLE_RADIANCE (radiance );
549
+ if (NON_BLACK (radiance ))
550
+ {
551
+ // Generate shadow ray
552
+ float shadow_ray_length = CRAZY_HIGH_DISTANCE ;
553
+ float3 shadow_ray_dir = normalize (wo );
554
+ float3 shadow_ray_o = diffgeo .p + CRAZY_LOW_DISTANCE * s * diffgeo .n ;
555
+ int shadow_ray_mask = Bxdf_IsSingular (& diffgeo ) ? 0xFFFFFFFF : 0x0000FFFF ;
556
+
557
+ Ray_Init (shadowrays + 2 * globalid + 1 , shadow_ray_o , shadow_ray_dir , shadow_ray_length , 0.f , shadow_ray_mask );
558
+
559
+ // Apply the volume to shadow ray if needed
560
+ int volidx = Path_GetVolumeIdx (path );
561
+ if (volidx != -1 )
562
+ {
563
+ radiance *= Volume_Transmittance (& volumes [volidx ], & shadowrays [2 * globalid + 1 ], shadow_ray_length );
564
+ radiance += Volume_Emission (& volumes [volidx ], & shadowrays [2 * globalid + 1 ], shadow_ray_length ) * throughput ;
565
+ }
566
+
567
+ // And write the light sample
568
+ lightsamples [2 * globalid + 1 ] = REASONABLE_RADIANCE (radiance );
569
+ }
570
+ else
571
+ {
572
+ // Otherwise save some intersector cycles
573
+ Ray_SetInactive (shadowrays + 2 * globalid + 1 );
574
+ lightsamples [2 * globalid + 1 ] = 0 ;
575
+ }
576
+ }
577
+ else
578
+ {
579
+ // Otherwise save some intersector cycles
580
+ Ray_SetInactive (shadowrays + 2 * globalid + 1 );
581
+ lightsamples [2 * globalid + 1 ] = 0 ;
582
+ }
496
583
}
497
584
else
498
585
{
499
586
// Otherwise save some intersector cycles
500
- Ray_SetInactive (shadowrays + globalid );
501
- lightsamples [globalid ] = 0 ;
587
+ Ray_SetInactive (shadowrays + 2 * globalid + 1 );
588
+ lightsamples [2 * globalid + 1 ] = 0 ;
502
589
}
503
-
590
+
504
591
// Apply Russian roulette
505
592
float q = max (min (0.5f ,
506
593
// Luminance
@@ -514,13 +601,9 @@ __kernel void ShadeSurface(
514
601
Path_MulThroughput (path , 1.f / q );
515
602
}
516
603
517
- if (Bxdf_IsSingular (& diffgeo ))
518
- {
519
- bxdfweight = 1.f ;
520
- }
521
-
604
+ float3 bxdf = Bxdf_Sample (& diffgeo , wi , TEXTURE_ARGS , sample5 , & bxdfwo , & bxdfpdf );
522
605
bxdfwo = normalize (bxdfwo );
523
- float3 t = bxdf * fabs (dot (diffgeo .n , bxdfwo )) * bxdfweight ;
606
+ float3 t = bxdf * fabs (dot (diffgeo .n , bxdfwo ));
524
607
525
608
// Only continue if we have non-zero throughput & pdf
526
609
if (NON_BLACK (t ) && bxdfpdf > 0.f && !rr_stop )
@@ -626,10 +709,16 @@ __kernel void GatherLightSamples(
626
709
// Start collecting samples
627
710
{
628
711
// If shadow ray didn't hit anything and reached skydome
629
- if (shadowhits [globalid ] == -1 )
712
+ if (shadowhits [2 * globalid ] == -1 )
713
+ {
714
+ // Add its contribution to radiance accumulator
715
+ radiance += lightsamples [2 * globalid ];
716
+ }
717
+
718
+ if (shadowhits [2 * globalid + 1 ] == -1 )
630
719
{
631
720
// Add its contribution to radiance accumulator
632
- radiance += lightsamples [globalid ];
721
+ radiance += lightsamples [2 * globalid + 1 ];
633
722
}
634
723
}
635
724
0 commit comments