Cycles: Add support for shader raytracing in OptiX

Support for the AO and bevel shader nodes requires calling "optixTrace" from within the shading
VM, which is only allowed from inlined functions to the raygen program or callables. This patch
therefore converts the shading VM to use direct callables to make it work. To prevent performance
regressions a separate kernel module is compiled and used for this purpose.

Reviewed By: brecht

Differential Revision: https://developer.blender.org/D9733

1 files changed, 22 insertions, 7 deletions

diff --git a/intern/cycles/kernel/kernel_subsurface.h b/intern/cycles/kernel/kernel_subsurface.h
index ed8572467ea..917f35d37dc 100644
--- a/intern/cycles/kernel/kernel_subsurface.h
+++ b/intern/cycles/kernel/kernel_subsurface.h
@@ -281,13 +281,28 @@ ccl_device_inline int subsurface_scatter_disk(KernelGlobals *kg,
   return num_eval_hits;
 }
 
-ccl_device_noinline void subsurface_scatter_multi_setup(KernelGlobals *kg,
-                                                        LocalIntersection *ss_isect,
-                                                        int hit,
-                                                        ShaderData *sd,
-                                                        ccl_addr_space PathState *state,
-                                                        ClosureType type,
-                                                        float roughness)
+#if defined(__KERNEL_OPTIX__) && defined(__SHADER_RAYTRACE__)
+ccl_device_inline void subsurface_scatter_multi_setup(KernelGlobals *kg,
+                                                      LocalIntersection *ss_isect,
+                                                      int hit,
+                                                      ShaderData *sd,
+                                                      ccl_addr_space PathState *state,
+                                                      ClosureType type,
+                                                      float roughness)
+{
+  optixDirectCall<void>(2, kg, ss_isect, hit, sd, state, type, roughness);
+}
+extern "C" __device__ void __direct_callable__subsurface_scatter_multi_setup(
+#else
+ccl_device_noinline void subsurface_scatter_multi_setup(
+#endif
+    KernelGlobals *kg,
+    LocalIntersection *ss_isect,
+    int hit,
+    ShaderData *sd,
+    ccl_addr_space PathState *state,
+    ClosureType type,
+    float roughness)
 {
 #ifdef __SPLIT_KERNEL__
   Ray ray_object = ss_isect->ray;