Blender  V3.3
ao.h
Go to the documentation of this file.
1 /* SPDX-License-Identifier: Apache-2.0
2  * Copyright 2011-2022 Blender Foundation */
3 
4 #pragma once
5 
6 #include "kernel/bvh/bvh.h"
7 
9 
10 #ifdef __SHADER_RAYTRACE__
11 
12 # ifdef __KERNEL_OPTIX__
13 extern "C" __device__ float __direct_callable__svm_node_ao(
14 # else
15 ccl_device float svm_ao(
16 # endif
17  KernelGlobals kg,
20  float3 N,
21  float max_dist,
22  int num_samples,
23  int flags)
24 {
25  if (flags & NODE_AO_GLOBAL_RADIUS) {
26  max_dist = kernel_data.integrator.ao_bounces_distance;
27  }
28 
29  /* Early out if no sampling needed. */
30  if (max_dist <= 0.0f || num_samples < 1 || sd->object == OBJECT_NONE) {
31  return 1.0f;
32  }
33 
34  /* Can't ray-trace from shaders like displacement, before BVH exists. */
35  if (kernel_data.bvh.bvh_layout == BVH_LAYOUT_NONE) {
36  return 1.0f;
37  }
38 
39  if (flags & NODE_AO_INSIDE) {
40  N = -N;
41  }
42 
43  float3 T, B;
44  make_orthonormals(N, &T, &B);
45 
46  /* TODO: support ray-tracing in shadow shader evaluation? */
47  RNGState rng_state;
48  path_state_rng_load(state, &rng_state);
49 
50  int unoccluded = 0;
51  for (int sample = 0; sample < num_samples; sample++) {
52  float disk_u, disk_v;
53  path_branched_rng_2D(kg, &rng_state, sample, num_samples, PRNG_BEVEL_U, &disk_u, &disk_v);
54 
55  float2 d = concentric_sample_disk(disk_u, disk_v);
56  float3 D = make_float3(d.x, d.y, safe_sqrtf(1.0f - dot(d, d)));
57 
58  /* Create ray. */
59  Ray ray;
60  ray.P = sd->P;
61  ray.D = D.x * T + D.y * B + D.z * N;
62  ray.tmin = 0.0f;
63  ray.tmax = max_dist;
64  ray.time = sd->time;
65  ray.self.object = sd->object;
66  ray.self.prim = sd->prim;
71 
72  if (flags & NODE_AO_ONLY_LOCAL) {
73  if (!scene_intersect_local(kg, &ray, NULL, sd->object, NULL, 0)) {
74  unoccluded++;
75  }
76  }
77  else {
78  Intersection isect;
79  if (!scene_intersect(kg, &ray, PATH_RAY_SHADOW_OPAQUE, &isect)) {
80  unoccluded++;
81  }
82  }
83  }
84 
85  return ((float)unoccluded) / num_samples;
86 }
87 
88 template<uint node_feature_mask, typename ConstIntegratorGenericState>
89 # if defined(__KERNEL_OPTIX__)
91 # else
93 # endif
94  void
95  svm_node_ao(KernelGlobals kg,
96  ConstIntegratorGenericState state,
98  ccl_private float *stack,
99  uint4 node)
100 {
101  uint flags, dist_offset, normal_offset, out_ao_offset;
102  svm_unpack_node_uchar4(node.y, &flags, &dist_offset, &normal_offset, &out_ao_offset);
103 
104  uint color_offset, out_color_offset, samples;
105  svm_unpack_node_uchar3(node.z, &color_offset, &out_color_offset, &samples);
106 
107  float ao = 1.0f;
108 
109  IF_KERNEL_NODES_FEATURE(RAYTRACE)
110  {
111  float dist = stack_load_float_default(stack, dist_offset, node.w);
112  float3 normal = stack_valid(normal_offset) ? stack_load_float3(stack, normal_offset) : sd->N;
113 
114 # ifdef __KERNEL_OPTIX__
115  ao = optixDirectCall<float>(0, kg, state, sd, normal, dist, samples, flags);
116 # else
117  ao = svm_ao(kg, state, sd, normal, dist, samples, flags);
118 # endif
119  }
120 
121  if (stack_valid(out_ao_offset)) {
122  stack_store_float(stack, out_ao_offset, ao);
123  }
124 
125  if (stack_valid(out_color_offset)) {
126  float3 color = stack_load_float3(stack, color_offset);
127  stack_store_float3(stack, out_color_offset, ao * color);
128  }
129 }
130 
131 #endif /* __SHADER_RAYTRACE__ */
132 
MINLINE float safe_sqrtf(float a)
unsigned int uint
Definition: BLI_sys_types.h:67
Group Output data from inside of a node group A color picker Mix two input colors RGB to Convert a color s luminance to a grayscale value Generate a normal vector and a dot product Bright Control the brightness and contrast of the input color Vector Map an input vectors to used to fine tune the interpolation of the input Camera Retrieve information about the camera and how it relates to the current shading point s position Clamp a value between a minimum and a maximum Vector Perform vector math operation Invert a color
#define ccl_device
Definition: cuda/compat.h:32
#define ccl_private
Definition: cuda/compat.h:48
#define ccl_device_inline
Definition: cuda/compat.h:34
#define ccl_device_noinline
Definition: cuda/compat.h:40
#define CCL_NAMESPACE_END
Definition: cuda/compat.h:9
OperationNode * node
#define kernel_data
const KernelGlobalsCPU *ccl_restrict KernelGlobals
ccl_device_forceinline float differential_zero_compact()
Definition: differential.h:112
IconTextureDrawCall normal
ccl_device_intersect bool scene_intersect(KernelGlobals kg, ccl_private const Ray *ray, const uint visibility, ccl_private Intersection *isect)
const int state
ccl_gpu_kernel_postfix ccl_global float int int int int ccl_global const float int int int int int int int int int int int int num_samples
ccl_device_inline void stack_store_float3(ccl_private float *stack, uint a, float3 f)
CCL_NAMESPACE_BEGIN ccl_device_inline float3 stack_load_float3(ccl_private float *stack, uint a)
ccl_device_forceinline void svm_unpack_node_uchar3(uint i, ccl_private uint *x, ccl_private uint *y, ccl_private uint *z)
ccl_device_inline float stack_load_float_default(ccl_private float *stack, uint a, uint value)
ccl_device_inline void stack_store_float(ccl_private float *stack, uint a, float f)
ccl_device_forceinline void svm_unpack_node_uchar4(uint i, ccl_private uint *x, ccl_private uint *y, ccl_private uint *z, ccl_private uint *w)
ccl_device_inline bool stack_valid(uint a)
@ NODE_AO_INSIDE
@ NODE_AO_GLOBAL_RADIUS
@ NODE_AO_ONLY_LOCAL
#define IF_KERNEL_NODES_FEATURE(feature)
@ PRNG_BEVEL_U
Definition: kernel/types.h:174
#define PRIM_NONE
Definition: kernel/types.h:41
@ PATH_RAY_SHADOW_OPAQUE
Definition: kernel/types.h:204
#define OBJECT_NONE
Definition: kernel/types.h:40
ShaderData
Definition: kernel/types.h:925
@ BVH_LAYOUT_NONE
#define N
#define T
#define B
#define __device__
Definition: metal/compat.h:248
#define make_float3(x, y, z)
Definition: metal/compat.h:204
T dot(const vec_base< T, Size > &a, const vec_base< T, Size > &b)
ccl_device_inline void path_branched_rng_2D(KernelGlobals kg, ccl_private const RNGState *rng_state, int branch, int num_branches, int dimension, ccl_private float *fx, ccl_private float *fy)
Definition: path_state.h:340
ccl_device_inline void path_state_rng_load(ConstIntegratorState state, ccl_private RNGState *rng_state)
Definition: path_state.h:283
ccl_device float2 concentric_sample_disk(float u1, float u2)
const IntegratorStateCPU *ccl_restrict ConstIntegratorState
Definition: state.h:148
float tmax
Definition: kernel/types.h:527
float tmin
Definition: kernel/types.h:526
float dD
Definition: kernel/types.h:534
float3 P
Definition: kernel/types.h:524
float time
Definition: kernel/types.h:528
float dP
Definition: kernel/types.h:533
RaySelfPrimitives self
Definition: kernel/types.h:530
float3 D
Definition: kernel/types.h:525
ccl_device_inline void make_orthonormals(const float3 N, ccl_private float3 *a, ccl_private float3 *b)
Definition: util/math.h:566
BLI_INLINE float D(const float *data, const int res[3], int x, int y, int z)
Definition: voxel.c:13