diff options
24 files changed, 499 insertions, 109 deletions
diff --git a/intern/cycles/bvh/bvh_embree.cpp b/intern/cycles/bvh/bvh_embree.cpp index 343d62dedf4..cd19e009bf3 100644 --- a/intern/cycles/bvh/bvh_embree.cpp +++ b/intern/cycles/bvh/bvh_embree.cpp @@ -80,31 +80,49 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args) Intersection current_isect; kernel_embree_convert_hit(kg, ray, hit, ¤t_isect); - /* If no transparent shadows, all light is blocked. */ + /* If no transparent shadows or max number of hits exceeded, all light is blocked. */ const int flags = intersection_get_shader_flags(kg, current_isect.prim, current_isect.type); - if (!(flags & (SD_HAS_TRANSPARENT_SHADOW)) || ctx->max_hits == 0) { + if (!(flags & (SD_HAS_TRANSPARENT_SHADOW)) || ctx->num_hits >= ctx->max_hits) { ctx->opaque_hit = true; return; } + ++ctx->num_hits; + + /* Always use baked shadow transparency for curves. */ + if (current_isect.type & PRIMITIVE_ALL_CURVE) { + ctx->throughput *= intersection_curve_shadow_transparency( + kg, current_isect.object, current_isect.prim, current_isect.u); + + if (ctx->throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { + ctx->opaque_hit = true; + return; + } + else { + *args->valid = 0; + return; + } + } + /* Test if we need to record this transparent intersection. */ - if (ctx->num_hits < ctx->max_hits || ray->tfar < ctx->max_t) { + const uint max_record_hits = min(ctx->max_hits, INTEGRATOR_SHADOW_ISECT_SIZE); + if (ctx->num_recorded_hits < max_record_hits || ray->tfar < ctx->max_t) { /* If maximum number of hits was reached, replace the intersection with the * highest distance. We want to find the N closest intersections. */ - const int num_recorded_hits = min(ctx->num_hits, ctx->max_hits); - int isect_index = num_recorded_hits; - if (num_recorded_hits + 1 >= ctx->max_hits) { + const uint num_recorded_hits = min(ctx->num_recorded_hits, max_record_hits); + uint isect_index = num_recorded_hits; + if (num_recorded_hits + 1 >= max_record_hits) { float max_t = ctx->isect_s[0].t; - int max_recorded_hit = 0; + uint max_recorded_hit = 0; - for (int i = 1; i < num_recorded_hits; ++i) { + for (uint i = 1; i < num_recorded_hits; ++i) { if (ctx->isect_s[i].t > max_t) { max_recorded_hit = i; max_t = ctx->isect_s[i].t; } } - if (num_recorded_hits >= ctx->max_hits) { + if (num_recorded_hits >= max_record_hits) { isect_index = max_recorded_hit; } @@ -118,10 +136,9 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args) ctx->isect_s[isect_index] = current_isect; } - /* Always increase the number of hits, even beyond ray.max_hits so that - * the caller can detect this as and consider it opaque, or trace another - * ray. */ - ++ctx->num_hits; + /* Always increase the number of recorded hits, even beyond the maximum, + * so that we can detect this and trace another ray if needed. */ + ++ctx->num_recorded_hits; /* This tells Embree to continue tracing. */ *args->valid = 0; @@ -160,7 +177,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args) if (ctx->lcg_state) { /* See triangle_intersect_subsurface() for the native equivalent. */ - for (int i = min(ctx->max_hits, local_isect->num_hits) - 1; i >= 0; --i) { + for (int i = min((int)ctx->max_hits, local_isect->num_hits) - 1; i >= 0; --i) { if (local_isect->hits[i].t == ray->tfar) { /* This tells Embree to continue tracing. */ *args->valid = 0; diff --git a/intern/cycles/device/cpu/kernel.cpp b/intern/cycles/device/cpu/kernel.cpp index 91282390e27..bbad2f3147d 100644 --- a/intern/cycles/device/cpu/kernel.cpp +++ b/intern/cycles/device/cpu/kernel.cpp @@ -44,6 +44,7 @@ CPUKernels::CPUKernels() /* Shader evaluation. */ REGISTER_KERNEL(shader_eval_displace), REGISTER_KERNEL(shader_eval_background), + REGISTER_KERNEL(shader_eval_curve_shadow_transparency), /* Adaptive sampling. */ REGISTER_KERNEL(adaptive_sampling_convergence_check), REGISTER_KERNEL(adaptive_sampling_filter_x), diff --git a/intern/cycles/device/cpu/kernel.h b/intern/cycles/device/cpu/kernel.h index 2db09057e44..3787fe37a33 100644 --- a/intern/cycles/device/cpu/kernel.h +++ b/intern/cycles/device/cpu/kernel.h @@ -58,6 +58,7 @@ class CPUKernels { ShaderEvalFunction shader_eval_displace; ShaderEvalFunction shader_eval_background; + ShaderEvalFunction shader_eval_curve_shadow_transparency; /* Adaptive stopping. */ diff --git a/intern/cycles/device/device_kernel.cpp b/intern/cycles/device/device_kernel.cpp index ceaddee4756..e0833331b77 100644 --- a/intern/cycles/device/device_kernel.cpp +++ b/intern/cycles/device/device_kernel.cpp @@ -74,6 +74,8 @@ const char *device_kernel_as_string(DeviceKernel kernel) return "shader_eval_displace"; case DEVICE_KERNEL_SHADER_EVAL_BACKGROUND: return "shader_eval_background"; + case DEVICE_KERNEL_SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY: + return "shader_eval_curve_shadow_transparency"; /* Film. */ diff --git a/intern/cycles/integrator/shader_eval.cpp b/intern/cycles/integrator/shader_eval.cpp index cfc30056f7d..3de7bb6fd16 100644 --- a/intern/cycles/integrator/shader_eval.cpp +++ b/intern/cycles/integrator/shader_eval.cpp @@ -122,6 +122,9 @@ bool ShaderEval::eval_cpu(Device *device, case SHADER_EVAL_BACKGROUND: kernels.shader_eval_background(kg, input_data, output_data, work_index); break; + case SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY: + kernels.shader_eval_curve_shadow_transparency(kg, input_data, output_data, work_index); + break; } }); }); @@ -144,6 +147,9 @@ bool ShaderEval::eval_gpu(Device *device, case SHADER_EVAL_BACKGROUND: kernel = DEVICE_KERNEL_SHADER_EVAL_BACKGROUND; break; + case SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY: + kernel = DEVICE_KERNEL_SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY; + break; }; /* Create device queue. */ diff --git a/intern/cycles/integrator/shader_eval.h b/intern/cycles/integrator/shader_eval.h index 013fad17d4f..43b6b1bdd47 100644 --- a/intern/cycles/integrator/shader_eval.h +++ b/intern/cycles/integrator/shader_eval.h @@ -30,6 +30,7 @@ class Progress; enum ShaderEvalType { SHADER_EVAL_DISPLACE, SHADER_EVAL_BACKGROUND, + SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY, }; /* ShaderEval class performs shader evaluation for background light and displacement. */ diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h index 0d9ba7e6369..813ac15711e 100644 --- a/intern/cycles/kernel/bvh/bvh.h +++ b/intern/cycles/kernel/bvh/bvh.h @@ -367,12 +367,13 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, ccl_private const Ray *ray, uint visibility, uint max_hits, - ccl_private uint *num_hits) + ccl_private uint *num_recorded_hits, + ccl_private float *throughput) { # ifdef __KERNEL_OPTIX__ uint p0 = state; - uint p1 = 0; /* Unused */ - uint p2 = 0; /* Number of hits. */ + uint p1 = __float_as_uint(1.0f); /* Throughput. */ + uint p2 = 0; /* Number of hits. */ uint p3 = max_hits; uint p4 = visibility; uint p5 = false; @@ -382,7 +383,6 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, ray_mask = 0xFF; } - *num_hits = 0; /* Initialize hit count to zero. */ optixTrace(scene_intersect_valid(ray) ? kernel_data.bvh.scene : 0, ray->P, ray->D, @@ -402,12 +402,14 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, p4, p5); - *num_hits = p2; + *num_recorded_hits = uint16_unpack_from_uint_0(p2); + *throughput = __uint_as_float(p1); return p5; # else /* __KERNEL_OPTIX__ */ if (!scene_intersect_valid(ray)) { - *num_hits = 0; + *num_recorded_hits = 0; + *throughput = 1.0f; return false; } @@ -422,7 +424,8 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, kernel_embree_setup_ray(*ray, rtc_ray, visibility); rtcOccluded1(kernel_data.bvh.scene, &rtc_ctx.context, &rtc_ray); - *num_hits = ctx.num_hits; + *num_recorded_hits = ctx.num_recorded_hits; + *throughput = ctx.throughput; return ctx.opaque_hit; } # endif /* __EMBREE__ */ @@ -431,21 +434,25 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, if (kernel_data.bvh.have_motion) { # ifdef __HAIR__ if (kernel_data.bvh.have_curves) { - return bvh_intersect_shadow_all_hair_motion(kg, ray, state, visibility, max_hits, num_hits); + return bvh_intersect_shadow_all_hair_motion( + kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); } # endif /* __HAIR__ */ - return bvh_intersect_shadow_all_motion(kg, ray, state, visibility, max_hits, num_hits); + return bvh_intersect_shadow_all_motion( + kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); } # endif /* __OBJECT_MOTION__ */ # ifdef __HAIR__ if (kernel_data.bvh.have_curves) { - return bvh_intersect_shadow_all_hair(kg, ray, state, visibility, max_hits, num_hits); + return bvh_intersect_shadow_all_hair( + kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); } # endif /* __HAIR__ */ - return bvh_intersect_shadow_all(kg, ray, state, visibility, max_hits, num_hits); + return bvh_intersect_shadow_all( + kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); # endif /* __KERNEL_OPTIX__ */ } #endif /* __SHADOW_RECORD_ALL__ */ diff --git a/intern/cycles/kernel/bvh/bvh_embree.h b/intern/cycles/kernel/bvh/bvh_embree.h index 4f85e8bee4b..321e0f28dae 100644 --- a/intern/cycles/kernel/bvh/bvh_embree.h +++ b/intern/cycles/kernel/bvh/bvh_embree.h @@ -40,8 +40,10 @@ struct CCLIntersectContext { /* for shadow rays */ Intersection *isect_s; - int max_hits; - int num_hits; + uint max_hits; + uint num_hits; + uint num_recorded_hits; + float throughput; float max_t; bool opaque_hit; @@ -56,6 +58,8 @@ struct CCLIntersectContext { type = type_; max_hits = 1; num_hits = 0; + num_recorded_hits = 0; + throughput = 1.0f; max_t = FLT_MAX; opaque_hit = false; isect_s = NULL; diff --git a/intern/cycles/kernel/bvh/bvh_shadow_all.h b/intern/cycles/kernel/bvh/bvh_shadow_all.h index b997235b6e4..049c6a03fe0 100644 --- a/intern/cycles/kernel/bvh/bvh_shadow_all.h +++ b/intern/cycles/kernel/bvh/bvh_shadow_all.h @@ -41,7 +41,8 @@ ccl_device_inline IntegratorShadowState state, const uint visibility, const uint max_hits, - ccl_private uint *num_hits) + ccl_private uint *num_recorded_hits, + ccl_private float *throughput) { /* todo: * - likely and unlikely for if() statements @@ -61,6 +62,7 @@ ccl_device_inline float3 dir = bvh_clamp_direction(ray->D); float3 idir = bvh_inverse_direction(dir); int object = OBJECT_NONE; + uint num_hits = 0; #if BVH_FEATURE(BVH_MOTION) Transform ob_itfm; @@ -77,7 +79,8 @@ ccl_device_inline * otherwise. */ float t_world_to_instance = 1.0f; - *num_hits = 0; + *num_recorded_hits = 0; + *throughput = 1.0f; /* traversal loop */ do { @@ -212,42 +215,62 @@ ccl_device_inline * the primitive has a transparent shadow shader? */ const int flags = intersection_get_shader_flags(kg, isect.prim, isect.type); - if (!(flags & SD_HAS_TRANSPARENT_SHADOW) || max_hits == 0) { + if (!(flags & SD_HAS_TRANSPARENT_SHADOW) || num_hits >= max_hits) { /* If no transparent shadows, all light is blocked and we can * stop immediately. */ return true; } - /* Increase the number of hits, possibly beyond max_hits, we will - * simply not record those and only keep the max_hits closest. */ - uint record_index = (*num_hits)++; - - if (record_index >= max_hits - 1) { - /* If maximum number of hits reached, find the intersection with - * the largest distance to potentially replace when another hit - * is found. */ - const int num_recorded_hits = min(max_hits, record_index); - float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t); - int max_recorded_hit = 0; - - for (int i = 1; i < num_recorded_hits; i++) { - const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t); - if (isect_t > max_recorded_t) { - max_recorded_t = isect_t; - max_recorded_hit = i; - } + num_hits++; + + bool record_intersection = true; + + /* Always use baked shadow transparency for curves. */ + if (isect.type & PRIMITIVE_ALL_CURVE) { + *throughput *= intersection_curve_shadow_transparency( + kg, isect.object, isect.prim, isect.u); + + if (*throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { + return true; + } + else { + record_intersection = false; } + } - if (record_index >= max_hits) { - record_index = max_recorded_hit; + if (record_intersection) { + /* Increase the number of hits, possibly beyond max_hits, we will + * simply not record those and only keep the max_hits closest. */ + uint record_index = (*num_recorded_hits)++; + + const uint max_record_hits = min(max_hits, INTEGRATOR_SHADOW_ISECT_SIZE); + if (record_index >= max_record_hits - 1) { + /* If maximum number of hits reached, find the intersection with + * the largest distance to potentially replace when another hit + * is found. */ + const int num_recorded_hits = min(max_record_hits, record_index); + float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t); + int max_recorded_hit = 0; + + for (int i = 1; i < num_recorded_hits; i++) { + const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t); + if (isect_t > max_recorded_t) { + max_recorded_t = isect_t; + max_recorded_hit = i; + } + } + + if (record_index >= max_record_hits) { + record_index = max_recorded_hit; + } + + /* Limit the ray distance and stop counting hits beyond this. */ + t_max_world = max(max_recorded_t, isect.t); + t_max_current = t_max_world * t_world_to_instance; } - /* Limit the ray distance and stop counting hits beyond this. */ - t_max_world = max(max_recorded_t, isect.t); - t_max_current = t_max_world * t_world_to_instance; + integrator_state_write_shadow_isect(state, &isect, record_index); } - - integrator_state_write_shadow_isect(state, &isect, record_index); } prim_addr++; @@ -304,9 +327,11 @@ ccl_device_inline bool BVH_FUNCTION_NAME(KernelGlobals kg, IntegratorShadowState state, const uint visibility, const uint max_hits, - ccl_private uint *num_hits) + ccl_private uint *num_recorded_hits, + ccl_private float *throughput) { - return BVH_FUNCTION_FULL_NAME(BVH)(kg, ray, state, visibility, max_hits, num_hits); + return BVH_FUNCTION_FULL_NAME(BVH)( + kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); } #undef BVH_FUNCTION_NAME diff --git a/intern/cycles/kernel/bvh/bvh_util.h b/intern/cycles/kernel/bvh/bvh_util.h index 869311b38e2..309f0eeb1e2 100644 --- a/intern/cycles/kernel/bvh/bvh_util.h +++ b/intern/cycles/kernel/bvh/bvh_util.h @@ -195,4 +195,32 @@ ccl_device_inline int intersection_find_attribute(KernelGlobals kg, return (attr_map.y == ATTR_ELEMENT_NONE) ? (int)ATTR_STD_NOT_FOUND : (int)attr_map.z; } +/* Transparent Shadows */ + +/* Cut-off value to stop transparent shadow tracing when practically opaque. */ +#define CURVE_SHADOW_TRANSPARENCY_CUTOFF 0.001f + +ccl_device_inline float intersection_curve_shadow_transparency(KernelGlobals kg, + const int object, + const int prim, + const float u) +{ + /* Find attribute. */ + const int offset = intersection_find_attribute(kg, object, ATTR_STD_SHADOW_TRANSPARENCY); + if (offset == ATTR_STD_NOT_FOUND) { + /* If no shadow transparency attribute, assume opaque. */ + return 0.0f; + } + + /* Interpolate transparency between curve keys. */ + const KernelCurve kcurve = kernel_tex_fetch(__curves, prim); + const int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(kcurve.type); + const int k1 = k0 + 1; + + const float f0 = kernel_tex_fetch(__attributes_float, offset + k0); + const float f1 = kernel_tex_fetch(__attributes_float, offset + k1); + + return (1.0f - u) * f0 + u * f1; +} + CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/device/cpu/kernel_arch.h b/intern/cycles/kernel/device/cpu/kernel_arch.h index ae7fab65100..7a438b58e73 100644 --- a/intern/cycles/kernel/device/cpu/kernel_arch.h +++ b/intern/cycles/kernel/device/cpu/kernel_arch.h @@ -64,6 +64,11 @@ void KERNEL_FUNCTION_FULL_NAME(shader_eval_displace)(const KernelGlobalsCPU *kg, const KernelShaderEvalInput *input, float *output, const int offset); +void KERNEL_FUNCTION_FULL_NAME(shader_eval_curve_shadow_transparency)( + const KernelGlobalsCPU *kg, + const KernelShaderEvalInput *input, + float *output, + const int offset); /* -------------------------------------------------------------------- * Adaptive sampling. diff --git a/intern/cycles/kernel/device/cpu/kernel_arch_impl.h b/intern/cycles/kernel/device/cpu/kernel_arch_impl.h index 2b0eea4fb61..ac606c768db 100644 --- a/intern/cycles/kernel/device/cpu/kernel_arch_impl.h +++ b/intern/cycles/kernel/device/cpu/kernel_arch_impl.h @@ -150,6 +150,19 @@ void KERNEL_FUNCTION_FULL_NAME(shader_eval_background)(const KernelGlobalsCPU *k #endif } +void KERNEL_FUNCTION_FULL_NAME(shader_eval_curve_shadow_transparency)( + const KernelGlobalsCPU *kg, + const KernelShaderEvalInput *input, + float *output, + const int offset) +{ +#ifdef KERNEL_STUB + STUB_ASSERT(KERNEL_ARCH, shader_eval_curve_shadow_transparency); +#else + kernel_curve_shadow_transparency_evaluate(kg, input, output, offset); +#endif +} + /* -------------------------------------------------------------------- * Adaptive sampling. */ diff --git a/intern/cycles/kernel/device/gpu/kernel.h b/intern/cycles/kernel/device/gpu/kernel.h index 6b4d79ed5b7..b6df74e835a 100644 --- a/intern/cycles/kernel/device/gpu/kernel.h +++ b/intern/cycles/kernel/device/gpu/kernel.h @@ -621,7 +621,7 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) } } -/* Background Shader Evaluation */ +/* Background */ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) kernel_gpu_shader_eval_background(KernelShaderEvalInput *input, @@ -635,6 +635,20 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) } } +/* Curve Shadow Transparency */ + +ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS) + kernel_gpu_shader_eval_curve_shadow_transparency(KernelShaderEvalInput *input, + float *output, + const int offset, + const int work_size) +{ + int i = ccl_gpu_global_id_x(); + if (i < work_size) { + kernel_curve_shadow_transparency_evaluate(NULL, input, output, offset + i); + } +} + /* -------------------------------------------------------------------- * Denoising. */ diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu index 574f66ab708..a3bafb9846c 100644 --- a/intern/cycles/kernel/device/optix/kernel.cu +++ b/intern/cycles/kernel/device/optix/kernel.cu @@ -210,29 +210,50 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() optixSetPayload_5(true); return optixTerminateRay(); # else - const int max_hits = optixGetPayload_3(); + const uint max_hits = optixGetPayload_3(); + const uint num_hits_packed = optixGetPayload_2(); + const uint num_recorded_hits = uint16_unpack_from_uint_0(num_hits_packed); + const uint num_hits = uint16_unpack_from_uint_1(num_hits_packed); /* If no transparent shadows, all light is blocked and we can stop immediately. */ - if (max_hits == 0 || + if (num_hits >= max_hits || !(intersection_get_shader_flags(NULL, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) { optixSetPayload_5(true); return optixTerminateRay(); } + /* Always use baked shadow transparency for curves. */ + if (type & PRIMITIVE_ALL_CURVE) { + float throughput = __uint_as_float(optixGetPayload_1()); + throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, u); + optixSetPayload_1(__float_as_uint(throughput)); + optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits, num_hits + 1)); + + if (throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { + optixSetPayload_4(true); + return optixTerminateRay(); + } + else { + /* Continue tracing. */ + optixIgnoreIntersection(); + return; + } + } + /* Record transparent intersection. */ - const int num_hits = optixGetPayload_2(); - int record_index = num_hits; + optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits + 1, num_hits + 1)); - optixSetPayload_2(num_hits + 1); + uint record_index = num_recorded_hits; const IntegratorShadowState state = optixGetPayload_0(); - if (record_index >= max_hits) { + const uint max_record_hits = min(max_hits, INTEGRATOR_SHADOW_ISECT_SIZE); + if (record_index >= max_record_hits) { /* If maximum number of hits reached, find a hit to replace. */ float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t); - int max_recorded_hit = 0; + uint max_recorded_hit = 0; - for (int i = 1; i < max_hits; i++) { + for (int i = 1; i < max_record_hits; i++) { const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t); if (isect_t > max_recorded_t) { max_recorded_t = isect_t; @@ -256,6 +277,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object; INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type; + /* Continue tracing. */ optixIgnoreIntersection(); # endif /* __TRANSPARENT_SHADOWS__ */ #endif /* __SHADOW_RECORD_ALL__ */ diff --git a/intern/cycles/kernel/geom/geom_shader_data.h b/intern/cycles/kernel/geom/geom_shader_data.h index e6a5b8f7923..46bda2b656c 100644 --- a/intern/cycles/kernel/geom/geom_shader_data.h +++ b/intern/cycles/kernel/geom/geom_shader_data.h @@ -279,6 +279,81 @@ ccl_device void shader_setup_from_displace(KernelGlobals kg, LAMP_NONE); } +/* ShaderData setup for point on curve. */ + +ccl_device void shader_setup_from_curve(KernelGlobals kg, + ccl_private ShaderData *ccl_restrict sd, + int object, + int prim, + int segment, + float u) +{ + /* Primitive */ + sd->type = PRIMITIVE_PACK_SEGMENT(PRIMITIVE_CURVE_THICK, segment); + sd->lamp = LAMP_NONE; + sd->prim = prim; + sd->u = u; + sd->v = 0.0f; + sd->time = 0.5f; + sd->ray_length = 0.0f; + + /* Shader */ + sd->shader = kernel_tex_fetch(__curves, prim).shader_id; + sd->flag = kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags; + + /* Object */ + sd->object = object; + sd->object_flag = kernel_tex_fetch(__object_flag, sd->object); +#ifdef __OBJECT_MOTION__ + shader_setup_object_transforms(kg, sd, sd->time); +#endif + + /* Get control points. */ + KernelCurve kcurve = kernel_tex_fetch(__curves, prim); + + int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type); + int k1 = k0 + 1; + int ka = max(k0 - 1, kcurve.first_key); + int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1); + + float4 P_curve[4]; + + P_curve[0] = kernel_tex_fetch(__curve_keys, ka); + P_curve[1] = kernel_tex_fetch(__curve_keys, k0); + P_curve[2] = kernel_tex_fetch(__curve_keys, k1); + P_curve[3] = kernel_tex_fetch(__curve_keys, kb); + + /* Interpolate position and tangent. */ + sd->P = float4_to_float3(catmull_rom_basis_derivative(P_curve, sd->u)); +#ifdef __DPDU__ + sd->dPdu = float4_to_float3(catmull_rom_basis_derivative(P_curve, sd->u)); +#endif + + /* Transform into world space */ + if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { + object_position_transform_auto(kg, sd, &sd->P); +#ifdef __DPDU__ + object_dir_transform_auto(kg, sd, &sd->dPdu); +#endif + } + + /* No view direction, normals or bitangent. */ + sd->I = zero_float3(); + sd->N = zero_float3(); + sd->Ng = zero_float3(); +#ifdef __DPDU__ + sd->dPdv = zero_float3(); +#endif + + /* No ray differentials currently. */ +#ifdef __RAY_DIFFERENTIALS__ + sd->dP = differential3_zero(); + sd->dI = differential3_zero(); + sd->du = differential_zero(); + sd->dv = differential_zero(); +#endif +} + /* ShaderData setup from ray into background */ ccl_device_inline void shader_setup_from_background(KernelGlobals kg, diff --git a/intern/cycles/kernel/integrator/integrator_intersect_shadow.h b/intern/cycles/kernel/integrator/integrator_intersect_shadow.h index d5c6ec145f0..90422445fad 100644 --- a/intern/cycles/kernel/integrator/integrator_intersect_shadow.h +++ b/intern/cycles/kernel/integrator/integrator_intersect_shadow.h @@ -115,18 +115,25 @@ ccl_device bool integrate_intersect_shadow_transparent(KernelGlobals kg, { /* Limit the number hits to the max transparent bounces allowed and the size that we * have available in the integrator state. */ - const uint max_transparent_hits = integrate_shadow_max_transparent_hits(kg, state); - const uint max_hits = min(max_transparent_hits, (uint)INTEGRATOR_SHADOW_ISECT_SIZE); + const uint max_hits = integrate_shadow_max_transparent_hits(kg, state); uint num_hits = 0; - bool opaque_hit = scene_intersect_shadow_all(kg, state, ray, visibility, max_hits, &num_hits); + float throughput = 1.0f; + bool opaque_hit = scene_intersect_shadow_all( + kg, state, ray, visibility, max_hits, &num_hits, &throughput); + + /* Computed throughput from baked shadow transparency, where we can bypass recording + * intersections and shader evaluation. */ + if (throughput != 1.0f) { + INTEGRATOR_STATE_WRITE(state, shadow_path, throughput) *= throughput; + } /* If number of hits exceed the transparent bounces limit, make opaque. */ - if (num_hits > max_transparent_hits) { + if (num_hits > max_hits) { opaque_hit = true; } if (!opaque_hit) { - uint num_recorded_hits = min(num_hits, max_hits); + const uint num_recorded_hits = min(num_hits, min(max_hits, INTEGRATOR_SHADOW_ISECT_SIZE)); if (num_recorded_hits > 0) { sort_shadow_intersections(state, num_recorded_hits); diff --git a/intern/cycles/kernel/integrator/integrator_shade_shadow.h b/intern/cycles/kernel/integrator/integrator_shade_shadow.h index 94900754b76..2d056a0b76f 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_shadow.h +++ b/intern/cycles/kernel/integrator/integrator_shade_shadow.h @@ -23,7 +23,7 @@ CCL_NAMESPACE_BEGIN -ccl_device_inline bool shadow_intersections_has_remaining(const int num_hits) +ccl_device_inline bool shadow_intersections_has_remaining(const uint num_hits) { return num_hits >= INTEGRATOR_SHADOW_ISECT_SIZE; } @@ -105,12 +105,12 @@ ccl_device_inline void integrate_transparent_volume_shadow(KernelGlobals kg, ccl_device_inline bool integrate_transparent_shadow(KernelGlobals kg, IntegratorShadowState state, - const int num_hits) + const uint num_hits) { /* Accumulate shadow for transparent surfaces. */ - const int num_recorded_hits = min(num_hits, INTEGRATOR_SHADOW_ISECT_SIZE); + const uint num_recorded_hits = min(num_hits, INTEGRATOR_SHADOW_ISECT_SIZE); - for (int hit = 0; hit < num_recorded_hits + 1; hit++) { + for (uint hit = 0; hit < num_recorded_hits + 1; hit++) { /* Volume shaders. */ if (hit < num_recorded_hits || !shadow_intersections_has_remaining(num_hits)) { # ifdef __VOLUME__ @@ -162,7 +162,7 @@ ccl_device void integrator_shade_shadow(KernelGlobals kg, ccl_global float *ccl_restrict render_buffer) { PROFILING_INIT(kg, PROFILING_SHADE_SHADOW_SETUP); - const int num_hits = INTEGRATOR_STATE(state, shadow_path, num_hits); + const uint num_hits = INTEGRATOR_STATE(state, shadow_path, num_hits); #ifdef __TRANSPARENT_SHADOWS__ /* Evaluate transparent shadows. */ diff --git a/intern/cycles/kernel/integrator/integrator_state.h b/intern/cycles/kernel/integrator/integrator_state.h index 84f34c6b986..4f21ab35d1f 100644 --- a/intern/cycles/kernel/integrator/integrator_state.h +++ b/intern/cycles/kernel/integrator/integrator_state.h @@ -48,19 +48,6 @@ CCL_NAMESPACE_BEGIN -/* Constants - * - * TODO: these could be made dynamic depending on the features used in the scene. */ - -#define INTEGRATOR_SHADOW_ISECT_SIZE_CPU 1024 -#define INTEGRATOR_SHADOW_ISECT_SIZE_GPU 4 - -#ifdef __KERNEL_CPU__ -# define INTEGRATOR_SHADOW_ISECT_SIZE INTEGRATOR_SHADOW_ISECT_SIZE_CPU -#else -# define INTEGRATOR_SHADOW_ISECT_SIZE INTEGRATOR_SHADOW_ISECT_SIZE_GPU -#endif - /* Data structures */ /* Integrator State diff --git a/intern/cycles/kernel/kernel_bake.h b/intern/cycles/kernel/kernel_bake.h index 1473e9ba8bf..30a41b9d3e3 100644 --- a/intern/cycles/kernel/kernel_bake.h +++ b/intern/cycles/kernel/kernel_bake.h @@ -96,4 +96,25 @@ ccl_device void kernel_background_evaluate(KernelGlobals kg, output[offset * 3 + 2] += color.z; } +ccl_device void kernel_curve_shadow_transparency_evaluate( + KernelGlobals kg, + ccl_global const KernelShaderEvalInput *input, + ccl_global float *output, + const int offset) +{ + /* Setup shader data. */ + const KernelShaderEvalInput in = input[offset]; + + ShaderData sd; + shader_setup_from_curve(kg, &sd, in.object, in.prim, __float_as_int(in.v), in.u); + + /* Evaluate transparency. */ + shader_eval_surface<KERNEL_FEATURE_NODE_MASK_SURFACE_SHADOW & + ~(KERNEL_FEATURE_NODE_RAYTRACE | KERNEL_FEATURE_NODE_LIGHT_PATH)>( + kg, INTEGRATOR_STATE_NULL, &sd, NULL, PATH_RAY_SHADOW); + + /* Write output. */ + output[offset] = clamp(average(shader_bsdf_transparency(kg, &sd)), 0.0f, 1.0f); +} + CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h index edae158f403..fa8453b99cb 100644 --- a/intern/cycles/kernel/kernel_types.h +++ b/intern/cycles/kernel/kernel_types.h @@ -61,6 +61,15 @@ CCL_NAMESPACE_BEGIN #define ID_NONE (0.0f) #define PASS_UNUSED (~0) +#define INTEGRATOR_SHADOW_ISECT_SIZE_CPU 1024U +#define INTEGRATOR_SHADOW_ISECT_SIZE_GPU 4U + +#ifdef __KERNEL_CPU__ +# define INTEGRATOR_SHADOW_ISECT_SIZE INTEGRATOR_SHADOW_ISECT_SIZE_CPU +#else +# define INTEGRATOR_SHADOW_ISECT_SIZE INTEGRATOR_SHADOW_ISECT_SIZE_GPU +#endif + /* Kernel features */ #define __SOBOL__ #define __DPDU__ @@ -582,6 +591,7 @@ typedef enum AttributeStandard { ATTR_STD_VOLUME_VELOCITY, ATTR_STD_POINTINESS, ATTR_STD_RANDOM_PER_ISLAND, + ATTR_STD_SHADOW_TRANSPARENCY, ATTR_STD_NUM, ATTR_STD_NOT_FOUND = ~0 @@ -1452,6 +1462,7 @@ typedef enum DeviceKernel { DEVICE_KERNEL_SHADER_EVAL_DISPLACE, DEVICE_KERNEL_SHADER_EVAL_BACKGROUND, + DEVICE_KERNEL_SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY, #define DECLARE_FILM_CONVERT_KERNEL(variant) \ DEVICE_KERNEL_FILM_CONVERT_##variant, DEVICE_KERNEL_FILM_CONVERT_##variant##_HALF_RGBA diff --git a/intern/cycles/render/attribute.cpp b/intern/cycles/render/attribute.cpp index aaf21ad9fd2..d7e6939cd80 100644 --- a/intern/cycles/render/attribute.cpp +++ b/intern/cycles/render/attribute.cpp @@ -366,6 +366,8 @@ const char *Attribute::standard_name(AttributeStandard std) return "pointiness"; case ATTR_STD_RANDOM_PER_ISLAND: return "random_per_island"; + case ATTR_STD_SHADOW_TRANSPARENCY: + return "shadow_transparency"; case ATTR_STD_NOT_FOUND: case ATTR_STD_NONE: case ATTR_STD_NUM: @@ -603,6 +605,9 @@ Attribute *AttributeSet::add(AttributeStandard std, ustring name) case ATTR_STD_RANDOM_PER_ISLAND: attr = add(name, TypeDesc::TypeFloat, ATTR_ELEMENT_FACE); break; + case ATTR_STD_SHADOW_TRANSPARENCY: + attr = add(name, TypeDesc::TypeFloat, ATTR_ELEMENT_CURVE_KEY); + break; default: assert(0); break; diff --git a/intern/cycles/render/geometry.cpp b/intern/cycles/render/geometry.cpp index 5d89060c1a1..5cedab24ceb 100644 --- a/intern/cycles/render/geometry.cpp +++ b/intern/cycles/render/geometry.cpp @@ -734,6 +734,10 @@ void GeometryManager::device_update_attributes(Device *device, Shader *shader = static_cast<Shader *>(node); geom_attributes[i].add(shader->attributes); } + + if (geom->is_hair() && static_cast<Hair *>(geom)->need_shadow_transparency()) { + geom_attributes[i].add(ATTR_STD_SHADOW_TRANSPARENCY); + } } /* convert object attributes to use the same data structures as geometry ones */ @@ -1659,6 +1663,7 @@ void GeometryManager::device_update(Device *device, VLOG(1) << "Total " << scene->geometry.size() << " meshes."; bool true_displacement_used = false; + bool curve_shadow_transparency_used = false; size_t total_tess_needed = 0; { @@ -1669,26 +1674,33 @@ void GeometryManager::device_update(Device *device, }); foreach (Geometry *geom, scene->geometry) { - if (geom->is_modified() && - (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME)) { - Mesh *mesh = static_cast<Mesh *>(geom); + if (geom->is_modified()) { + if ((geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME)) { + Mesh *mesh = static_cast<Mesh *>(geom); - /* Update normals. */ - mesh->add_face_normals(); - mesh->add_vertex_normals(); + /* Update normals. */ + mesh->add_face_normals(); + mesh->add_vertex_normals(); - if (mesh->need_attribute(scene, ATTR_STD_POSITION_UNDISPLACED)) { - mesh->add_undisplaced(); - } + if (mesh->need_attribute(scene, ATTR_STD_POSITION_UNDISPLACED)) { + mesh->add_undisplaced(); + } - /* Test if we need tessellation. */ - if (mesh->need_tesselation()) { - total_tess_needed++; - } + /* Test if we need tessellation. */ + if (mesh->need_tesselation()) { + total_tess_needed++; + } - /* Test if we need displacement. */ - if (mesh->has_true_displacement()) { - true_displacement_used = true; + /* Test if we need displacement. */ + if (mesh->has_true_displacement()) { + true_displacement_used = true; + } + } + else if (geom->geometry_type == Geometry::HAIR) { + Hair *hair = static_cast<Hair *>(geom); + if (hair->need_shadow_transparency()) { + curve_shadow_transparency_used = true; + } } if (progress.get_cancel()) { @@ -1752,7 +1764,7 @@ void GeometryManager::device_update(Device *device, /* Update images needed for true displacement. */ bool old_need_object_flags_update = false; - if (true_displacement_used) { + if (true_displacement_used || curve_shadow_transparency_used) { scoped_callback_timer timer([scene](double time) { if (scene->update_stats) { scene->update_stats->geometry.times.add_entry( @@ -1770,7 +1782,7 @@ void GeometryManager::device_update(Device *device, const BVHLayout bvh_layout = BVHParams::best_bvh_layout(scene->params.bvh_layout, device->get_bvh_layout_mask()); mesh_calc_offset(scene, bvh_layout); - if (true_displacement_used) { + if (true_displacement_used || curve_shadow_transparency_used) { scoped_callback_timer timer([scene](double time) { if (scene->update_stats) { scene->update_stats->geometry.times.add_entry( @@ -1795,8 +1807,9 @@ void GeometryManager::device_update(Device *device, } } - /* Update displacement. */ + /* Update displacement and hair shadow transparency. */ bool displacement_done = false; + bool curve_shadow_transparency_done = false; size_t num_bvh = 0; { @@ -1817,6 +1830,12 @@ void GeometryManager::device_update(Device *device, displacement_done = true; } } + else if (geom->geometry_type == Geometry::HAIR) { + Hair *hair = static_cast<Hair *>(geom); + if (hair->update_shadow_transparency(device, scene, progress)) { + curve_shadow_transparency_done = true; + } + } } if (geom->is_modified() || geom->need_update_bvh_for_offset) { @@ -1836,7 +1855,7 @@ void GeometryManager::device_update(Device *device, } /* Device re-update after displacement. */ - if (displacement_done) { + if (displacement_done || curve_shadow_transparency_done) { scoped_callback_timer timer([scene](double time) { if (scene->update_stats) { scene->update_stats->geometry.times.add_entry( diff --git a/intern/cycles/render/hair.cpp b/intern/cycles/render/hair.cpp index e757e3fd3e0..4656148119a 100644 --- a/intern/cycles/render/hair.cpp +++ b/intern/cycles/render/hair.cpp @@ -18,8 +18,13 @@ #include "render/curves.h" #include "render/hair.h" +#include "render/object.h" #include "render/scene.h" +#include "integrator/shader_eval.h" + +#include "util/util_progress.h" + CCL_NAMESPACE_BEGIN /* Hair Curve */ @@ -514,4 +519,114 @@ PrimitiveType Hair::primitive_type() const ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON : PRIMITIVE_CURVE_THICK); } +/* Fill in coordinates for curve transparency shader evaluation on device. */ +static int fill_shader_input(const Hair *hair, + const int object_index, + device_vector<KernelShaderEvalInput> &d_input) +{ + int d_input_size = 0; + KernelShaderEvalInput *d_input_data = d_input.data(); + + const int num_curves = hair->num_curves(); + for (int i = 0; i < num_curves; i++) { + const Hair::Curve curve = hair->get_curve(i); + const int num_segments = curve.num_segments(); + + for (int j = 0; j < num_segments + 1; j++) { + KernelShaderEvalInput in; + in.object = object_index; + in.prim = hair->prim_offset + i; + in.u = (j < num_segments) ? 0.0f : 1.0f; + in.v = (j < num_segments) ? __int_as_float(j) : __int_as_float(j - 1); + d_input_data[d_input_size++] = in; + } + } + + return d_input_size; +} + +/* Read back curve transparency shader output. */ +static void read_shader_output(float *shadow_transparency, + bool &is_fully_opaque, + const device_vector<float> &d_output) +{ + const int num_keys = d_output.size(); + const float *output_data = d_output.data(); + bool is_opaque = true; + + for (int i = 0; i < num_keys; i++) { + shadow_transparency[i] = output_data[i]; + if (shadow_transparency[i] > 0.0f) { + is_opaque = false; + } + } + + is_fully_opaque = is_opaque; +} + +bool Hair::need_shadow_transparency() +{ + for (const Node *node : used_shaders) { + const Shader *shader = static_cast<const Shader *>(node); + if (shader->has_surface_transparent && shader->get_use_transparent_shadow()) { + return true; + } + } + + return false; +} + +bool Hair::update_shadow_transparency(Device *device, Scene *scene, Progress &progress) +{ + if (!need_shadow_transparency()) { + /* If no shaders with shadow transparency, remove attribute. */ + Attribute *attr = attributes.find(ATTR_STD_SHADOW_TRANSPARENCY); + if (attr) { + attributes.remove(attr); + return true; + } + else { + return false; + } + } + + string msg = string_printf("Computing Shadow Transparency %s", name.c_str()); + progress.set_status("Updating Hair", msg); + + /* Create shadow transparency attribute. */ + Attribute *attr = attributes.find(ATTR_STD_SHADOW_TRANSPARENCY); + const bool attribute_exists = (attr != nullptr); + if (!attribute_exists) { + attr = attributes.add(ATTR_STD_SHADOW_TRANSPARENCY); + } + + float *attr_data = attr->data_float(); + + /* Find object index. */ + size_t object_index = OBJECT_NONE; + + for (size_t i = 0; i < scene->objects.size(); i++) { + if (scene->objects[i]->get_geometry() == this) { + object_index = i; + break; + } + } + + /* Evaluate shader on device. */ + ShaderEval shader_eval(device, progress); + bool is_fully_opaque = false; + shader_eval.eval(SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY, + num_keys(), + 1, + function_bind(&fill_shader_input, this, object_index, _1), + function_bind(&read_shader_output, attr_data, is_fully_opaque, _1)); + + if (is_fully_opaque) { + attributes.remove(attr); + return attribute_exists; + } + + return true; +} + CCL_NAMESPACE_END diff --git a/intern/cycles/render/hair.h b/intern/cycles/render/hair.h index 920e9601b35..3e91fc3dcbb 100644 --- a/intern/cycles/render/hair.h +++ b/intern/cycles/render/hair.h @@ -153,6 +153,10 @@ class Hair : public Geometry { KernelCurveSegment *curve_segments); PrimitiveType primitive_type() const override; + + /* Attributes */ + bool need_shadow_transparency(); + bool update_shadow_transparency(Device *device, Scene *scene, Progress &progress); }; CCL_NAMESPACE_END |