Blender V5.0
bevel.h
Go to the documentation of this file.
1/* SPDX-FileCopyrightText: 2011-2022 Blender Foundation
2 *
3 * SPDX-License-Identifier: Apache-2.0 */
4
5#pragma once
6
7#include "kernel/bvh/bvh.h"
8
12
14
15#include "kernel/svm/util.h"
16
18
19#ifdef __SHADER_RAYTRACE__
20
21/* Planar Cubic BSSRDF falloff, reused for bevel.
22 *
23 * This is basically (Rm - x)^3, with some factors to normalize it. For sampling
24 * we integrate 2*pi*x * (Rm - x)^3, which gives us a quintic equation that as
25 * far as I can tell has no closed form solution. So we get an iterative solution
26 * instead with newton-raphson. */
27
28ccl_device float svm_bevel_cubic_eval(const float radius, const float r)
29{
30 const float Rm = radius;
31
32 if (r >= Rm) {
33 return 0.0f;
34 }
35
36 /* integrate (2*pi*r * 10*(R - r)^3)/(pi * R^5) from 0 to R = 1 */
37 const float Rm5 = (Rm * Rm) * (Rm * Rm) * Rm;
38 const float f = Rm - r;
39 const float num = f * f * f;
40
41 return (10.0f * num) / (Rm5 * M_PI_F);
42}
43
44ccl_device float svm_bevel_cubic_pdf(const float radius, const float r)
45{
46 return svm_bevel_cubic_eval(radius, r);
47}
48
49/* solve 10x^2 - 20x^3 + 15x^4 - 4x^5 - xi == 0 */
50ccl_device_forceinline float svm_bevel_cubic_quintic_root_find(const float xi)
51{
52 /* newton-raphson iteration, usually succeeds in 2-4 iterations, except
53 * outside 0.02 ... 0.98 where it can go up to 10, so overall performance
54 * should not be too bad */
55 const float tolerance = 1e-6f;
56 const int max_iteration_count = 10;
57 float x = 0.25f;
58 int i;
59
60 for (i = 0; i < max_iteration_count; i++) {
61 const float x2 = x * x;
62 const float x3 = x2 * x;
63 const float nx = (1.0f - x);
64
65 const float f = 10.0f * x2 - 20.0f * x3 + 15.0f * x2 * x2 - 4.0f * x2 * x3 - xi;
66 const float f_ = 20.0f * (x * nx) * (nx * nx);
67
68 if (fabsf(f) < tolerance || f_ == 0.0f) {
69 break;
70 }
71
72 x = saturatef(x - f / f_);
73 }
74
75 return x;
76}
77
78ccl_device void svm_bevel_cubic_sample(const float radius,
79 const float xi,
80 ccl_private float *r,
81 ccl_private float *h)
82{
83 const float Rm = radius;
84 float r_ = svm_bevel_cubic_quintic_root_find(xi);
85
86 r_ *= Rm;
87 *r = r_;
88
89 /* h^2 + r^2 = Rm^2 */
90 *h = safe_sqrtf(Rm * Rm - r_ * r_);
91}
92
93/* Bevel shader averaging normals from nearby surfaces.
94 *
95 * Sampling strategy from: BSSRDF Importance Sampling, SIGGRAPH 2013
96 * http://library.imageworks.com/pdfs/imageworks-library-BSSRDF-sampling.pdf
97 */
98
99# ifdef __KERNEL_OPTIX__
100extern "C" __device__ float3 __direct_callable__svm_node_bevel(
101# else
102ccl_device float3 svm_bevel(
103# endif
104 KernelGlobals kg,
106 ccl_private ShaderData *sd,
107 const float radius,
108 const int num_samples)
109{
110 /* Early out if no sampling needed. */
111 if (radius <= 0.0f || num_samples < 1 || sd->object == OBJECT_NONE) {
112 return sd->N;
113 }
114
115 /* Can't ray-trace from shaders like displacement, before BVH exists. */
116 if (kernel_data.bvh.bvh_layout == BVH_LAYOUT_NONE) {
117 return sd->N;
118 }
119
120 /* Don't bevel for blurry indirect rays. */
121 if (INTEGRATOR_STATE(state, path, min_ray_pdf) < 8.0f) {
122 return sd->N;
123 }
124
125 /* Setup for multi intersection. */
126 LocalIntersection isect;
127 uint lcg_state = lcg_state_init(INTEGRATOR_STATE(state, path, rng_pixel),
128 INTEGRATOR_STATE(state, path, rng_offset),
130 0x64c6a40e);
131
132 /* Sample normals from surrounding points on surface. */
133 float3 sum_N = make_float3(0.0f, 0.0f, 0.0f);
134
135 /* TODO: support ray-tracing in shadow shader evaluation? */
136 RNGState rng_state;
137 path_state_rng_load(state, &rng_state);
138
139 for (int sample = 0; sample < num_samples; sample++) {
140 float2 rand_disk = path_branched_rng_2D(
141 kg, &rng_state, sample, num_samples, PRNG_SURFACE_BEVEL);
142
143 /* Pick random axis in local frame and point on disk. */
144 float3 disk_N;
145 float3 disk_T;
146 float3 disk_B;
147 float pick_pdf_N;
148 float pick_pdf_T;
149 float pick_pdf_B;
150
151 disk_N = sd->Ng;
152 make_orthonormals(disk_N, &disk_T, &disk_B);
153
154 const float axisu = rand_disk.x;
155
156 if (axisu < 0.5f) {
157 pick_pdf_N = 0.5f;
158 pick_pdf_T = 0.25f;
159 pick_pdf_B = 0.25f;
160 rand_disk.x *= 2.0f;
161 }
162 else if (axisu < 0.75f) {
163 const float3 tmp = disk_N;
164 disk_N = disk_T;
165 disk_T = tmp;
166 pick_pdf_N = 0.25f;
167 pick_pdf_T = 0.5f;
168 pick_pdf_B = 0.25f;
169 rand_disk.x = (rand_disk.x - 0.5f) * 4.0f;
170 }
171 else {
172 const float3 tmp = disk_N;
173 disk_N = disk_B;
174 disk_B = tmp;
175 pick_pdf_N = 0.25f;
176 pick_pdf_T = 0.25f;
177 pick_pdf_B = 0.5f;
178 rand_disk.x = (rand_disk.x - 0.75f) * 4.0f;
179 }
180
181 /* Sample point on disk. */
182 const float phi = M_2PI_F * rand_disk.x;
183 float disk_r = rand_disk.y;
184 float disk_height;
185
186 /* Perhaps find something better than Cubic BSSRDF, but happens to work well. */
187 svm_bevel_cubic_sample(radius, disk_r, &disk_r, &disk_height);
188
189 const float3 disk_P = to_global(polar_to_cartesian(disk_r, phi), disk_T, disk_B);
190
191 /* Create ray. */
193 ray.P = sd->P + disk_N * disk_height + disk_P;
194 ray.D = -disk_N;
195 ray.tmin = 0.0f;
196 ray.tmax = 2.0f * disk_height;
197 ray.dP = differential_zero_compact();
198 ray.dD = differential_zero_compact();
199 ray.time = sd->time;
200 ray.self.object = OBJECT_NONE;
201 ray.self.prim = PRIM_NONE;
202 ray.self.light_object = OBJECT_NONE;
203 ray.self.light_prim = PRIM_NONE;
204
205 /* Intersect with the same object. if multiple intersections are found it
206 * will use at most LOCAL_MAX_HITS hits, a random subset of all hits. */
207 scene_intersect_local(kg, &ray, &isect, sd->object, &lcg_state, LOCAL_MAX_HITS);
208
209 const int num_eval_hits = min(isect.num_hits, LOCAL_MAX_HITS);
210
211 for (int hit = 0; hit < num_eval_hits; hit++) {
212 /* Quickly retrieve P and Ng without setting up ShaderData. */
213 float3 hit_P;
214 if (sd->type == PRIMITIVE_TRIANGLE) {
216 kg, sd, isect.hits[hit].prim, isect.hits[hit].u, isect.hits[hit].v);
217 }
218# ifdef __OBJECT_MOTION__
219 else if (sd->type == PRIMITIVE_MOTION_TRIANGLE) {
220 float3 verts[3];
221 motion_triangle_vertices(kg, sd->object, isect.hits[hit].prim, sd->time, verts);
222 hit_P = motion_triangle_point_from_uv(kg, sd, isect.hits[hit].u, isect.hits[hit].v, verts);
223 }
224# endif /* __OBJECT_MOTION__ */
225
226 /* Get geometric normal. */
227 float3 hit_Ng = isect.Ng[hit];
228 const int object = isect.hits[hit].object;
229 const int object_flag = kernel_data_fetch(object_flag, object);
230 if (object_negative_scale_applied(object_flag)) {
231 hit_Ng = -hit_Ng;
232 }
233
234 /* Compute smooth normal. */
235 float3 N = hit_Ng;
236 const int prim = isect.hits[hit].prim;
237 const int shader = kernel_data_fetch(tri_shader, prim);
238
239 if (shader & SHADER_SMOOTH_NORMAL) {
240 const float u = isect.hits[hit].u;
241 const float v = isect.hits[hit].v;
242
243 if (sd->type == PRIMITIVE_TRIANGLE) {
244 N = triangle_smooth_normal(kg, N, prim, u, v);
245 }
246# ifdef __OBJECT_MOTION__
247 else if (sd->type == PRIMITIVE_MOTION_TRIANGLE) {
248 N = motion_triangle_smooth_normal(kg, N, sd->object, prim, u, v, sd->time);
249 }
250# endif /* __OBJECT_MOTION__ */
251 }
252
253 /* Transform normals to world space. */
254 if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
255 object_normal_transform(kg, sd, &N);
256 object_normal_transform(kg, sd, &hit_Ng);
257 }
258
259 /* Probability densities for local frame axes. */
260 const float pdf_N = pick_pdf_N * fabsf(dot(disk_N, hit_Ng));
261 const float pdf_T = pick_pdf_T * fabsf(dot(disk_T, hit_Ng));
262 const float pdf_B = pick_pdf_B * fabsf(dot(disk_B, hit_Ng));
263
264 /* Multiple importance sample between 3 axes, power heuristic
265 * found to be slightly better than balance heuristic. pdf_N
266 * in the MIS weight and denominator canceled out. */
267 float w = pdf_N / (sqr(pdf_N) + sqr(pdf_T) + sqr(pdf_B));
268 if (isect.num_hits > LOCAL_MAX_HITS) {
269 w *= isect.num_hits / (float)LOCAL_MAX_HITS;
270 }
271
272 /* Real distance to sampled point. */
273 const float r = len(hit_P - sd->P);
274
275 /* Compute weight. */
276 const float pdf = svm_bevel_cubic_pdf(radius, r);
277 const float disk_pdf = svm_bevel_cubic_pdf(radius, disk_r);
278
279 w *= pdf / disk_pdf;
280
281 /* Sum normal and weight. */
282 sum_N += w * N;
283 }
284 }
285
286 /* Normalize. */
287 const float3 N = safe_normalize(sum_N);
288 return is_zero(N) ? sd->N : (sd->flag & SD_BACKFACING) ? -N : N;
289}
290
291template<uint node_feature_mask, typename ConstIntegratorGenericState>
292# if defined(__KERNEL_OPTIX__)
294# else
296# endif
297 void
298 svm_node_bevel(KernelGlobals kg,
299 ConstIntegratorGenericState state,
300 ccl_private ShaderData *sd,
301 ccl_private float *stack,
302 const uint4 node)
303{
304 uint num_samples;
305 uint radius_offset;
306 uint normal_offset;
307 uint out_offset;
308 svm_unpack_node_uchar4(node.y, &num_samples, &radius_offset, &normal_offset, &out_offset);
309
310 float3 bevel_N = sd->N;
311
313 {
314 float radius = stack_load_float(stack, radius_offset);
315
316# ifdef __KERNEL_OPTIX__
317 bevel_N = optixDirectCall<float3>(1, kg, state, sd, radius, num_samples);
318# else
319 bevel_N = svm_bevel(kg, state, sd, radius, num_samples);
320# endif
321
322 if (stack_valid(normal_offset)) {
323 /* Preserve input normal. */
324 const float3 ref_N = stack_load_float3(stack, normal_offset);
325 bevel_N = normalize(ref_N + (bevel_N - sd->N));
326 }
327 }
328
329 stack_store_float3(stack, out_offset, bevel_N);
330}
331
332#endif /* __SHADER_RAYTRACE__ */
333
MINLINE float safe_sqrtf(float a)
ATTR_WARN_UNUSED_RESULT const size_t num
unsigned int uint
ATTR_WARN_UNUSED_RESULT const BMVert * v
SIMD_FORCE_INLINE const btScalar & w() const
Return the w value.
Definition btQuadWord.h:119
nullptr float
dot(value.rgb, luminance_coefficients)") DEFINE_VALUE("REDUCE(lhs
ccl_device_inline float stack_load_float(const ccl_private float *stack, const uint a)
ccl_device_inline void stack_store_float3(ccl_private float *stack, const uint a, const float3 f)
ccl_device_forceinline void svm_unpack_node_uchar4(const uint i, ccl_private uint *x, ccl_private uint *y, ccl_private uint *z, ccl_private uint *w)
ccl_device_inline bool stack_valid(const uint a)
CCL_NAMESPACE_BEGIN ccl_device_inline float3 stack_load_float3(const ccl_private float *stack, const uint a)
ccl_device_inline T to_global(const float2 p, const T X, const T Y)
ccl_device_inline float2 polar_to_cartesian(const float r, const float phi)
#define kernel_data
#define ccl_device_forceinline
#define IF_KERNEL_NODES_FEATURE(feature)
#define ccl_optional_struct_init
#define kernel_data_fetch(name, index)
#define PRIM_NONE
#define OBJECT_NONE
#define LOCAL_MAX_HITS
#define ccl_private
const ThreadKernelGlobalsCPU * KernelGlobals
#define ccl_device_inline
#define ccl_device_noinline
#define CCL_NAMESPACE_END
#define saturatef(x)
ccl_device_forceinline float3 make_float3(const float x, const float y, const float z)
#define __device__
ccl_device_forceinline float differential_zero_compact()
ccl_device_inline float3 triangle_smooth_normal(KernelGlobals kg, const float3 Ng, const int prim, const float u, float v)
static float verts[][3]
VecBase< float, D > normalize(VecOp< float, D >) RET
ccl_device_inline bool object_negative_scale_applied(const int object_flag)
ccl_device_inline void object_normal_transform(KernelGlobals kg, const ccl_private ShaderData *sd, ccl_private float3 *N)
@ SD_BACKFACING
@ PRIMITIVE_MOTION_TRIANGLE
@ PRIMITIVE_TRIANGLE
@ PRNG_SURFACE_BEVEL
@ SHADER_SMOOTH_NORMAL
@ SD_OBJECT_TRANSFORM_APPLIED
@ BVH_LAYOUT_NONE
ccl_device_inline uint lcg_state_init(const uint rng_hash, const uint rng_offset, const uint sample, const uint scramble)
Definition lcg.h:45
ccl_device_inline bool is_zero(const float2 a)
ccl_device_inline float2 safe_normalize(const float2 a)
ccl_device_inline void make_orthonormals(const float3 N, ccl_private float3 *a, ccl_private float3 *b)
static ulong state[N]
#define N
ccl_device_inline void motion_triangle_vertices(KernelGlobals kg, const int object, const uint3 tri_vindex, const int numsteps, const int numverts, const int step, const float t, float3 verts[3])
ccl_device_inline float3 motion_triangle_smooth_normal(KernelGlobals kg, const float3 Ng, const int object, const uint3 tri_vindex, const int numsteps, const int step, const float t, const float u, const float v)
CCL_NAMESPACE_BEGIN ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg, ccl_private ShaderData *sd, const float u, const float v, const float3 verts[3])
ccl_device_inline void path_state_rng_load(ConstIntegratorState state, ccl_private RNGState *rng_state)
Definition path_state.h:329
ccl_device_inline float2 path_branched_rng_2D(KernelGlobals kg, const ccl_private RNGState *rng_state, const int branch, const int num_branches, const int dimension)
Definition path_state.h:389
#define sqr
#define fabsf
#define ccl_device
#define M_2PI_F
#define M_PI_F
#define min(a, b)
Definition sort.cc:36
#define INTEGRATOR_STATE(state, nested_struct, member)
Definition state.h:235
const IntegratorStateCPU * ConstIntegratorState
Definition state.h:229
struct Intersection hits[LOCAL_MAX_HITS]
float3 Ng[LOCAL_MAX_HITS]
float x
float y
uint y
Definition types_uint4.h:13
i
Definition text_draw.cc:230
ccl_device_inline float3 triangle_point_from_uv(KernelGlobals kg, ccl_private ShaderData *sd, const int isect_prim, const float u, const float v)
uint len