Cycles: pointcloud rendering

This add support for rendering of the point cloud object in Blender, as a native
geometry type in Cycles that is more memory and time efficient than instancing
sphere meshes. This can be useful for rendering sand, water splashes, particles,
motion graphics, etc.

Points are currently always rendered as spheres, with backface culling. More
shapes are likely to be added later, but this is the most important one and can
be customized with shaders.

For CPU rendering the Embree primitive is used, for GPU there is our own
intersection code. Motion blur is suppored. Volumes inside points are not
currently supported.

Implemented with help from:
* Kévin Dietrich: Alembic procedural integration
* Patrick Mourse: OptiX integration
* Josh Whelchel: update for cycles-x changes

Ref T92573

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

18 files changed, 559 insertions, 38 deletions

diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index 674eb702814..51158b86c5a 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -179,11 +179,14 @@ set(SRC_KERNEL_GEOM_HEADERS
   geom/curve.h
   geom/curve_intersect.h
   geom/motion_curve.h
+  geom/motion_point.h
   geom/motion_triangle.h
   geom/motion_triangle_intersect.h
   geom/motion_triangle_shader.h
   geom/object.h
   geom/patch.h
+  geom/point.h
+  geom/point_intersect.h
   geom/primitive.h
   geom/shader_data.h
   geom/subd_triangle.h
diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h
index 37cfac25060..67804fb1d0d 100644
--- a/intern/cycles/kernel/bvh/bvh.h
+++ b/intern/cycles/kernel/bvh/bvh.h
@@ -49,24 +49,24 @@ CCL_NAMESPACE_BEGIN
 #  include "kernel/bvh/nodes.h"
 
 #  define BVH_FUNCTION_NAME bvh_intersect
-#  define BVH_FUNCTION_FEATURES 0
+#  define BVH_FUNCTION_FEATURES BVH_POINTCLOUD
 #  include "kernel/bvh/traversal.h"
 
 #  if defined(__HAIR__)
 #    define BVH_FUNCTION_NAME bvh_intersect_hair
-#    define BVH_FUNCTION_FEATURES BVH_HAIR
+#    define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_POINTCLOUD
 #    include "kernel/bvh/traversal.h"
 #  endif
 
 #  if defined(__OBJECT_MOTION__)
 #    define BVH_FUNCTION_NAME bvh_intersect_motion
-#    define BVH_FUNCTION_FEATURES BVH_MOTION
+#    define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_POINTCLOUD
 #    include "kernel/bvh/traversal.h"
 #  endif
 
 #  if defined(__HAIR__) && defined(__OBJECT_MOTION__)
 #    define BVH_FUNCTION_NAME bvh_intersect_hair_motion
-#    define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION
+#    define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION | BVH_POINTCLOUD
 #    include "kernel/bvh/traversal.h"
 #  endif
 
@@ -102,26 +102,27 @@ CCL_NAMESPACE_BEGIN
 
 #  if defined(__SHADOW_RECORD_ALL__)
 #    define BVH_FUNCTION_NAME bvh_intersect_shadow_all
-#    define BVH_FUNCTION_FEATURES 0
+#    define BVH_FUNCTION_FEATURES BVH_POINTCLOUD
 #    include "kernel/bvh/shadow_all.h"
 
 #    if defined(__HAIR__)
 #      define BVH_FUNCTION_NAME bvh_intersect_shadow_all_hair
-#      define BVH_FUNCTION_FEATURES BVH_HAIR
+#      define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_POINTCLOUD
 #      include "kernel/bvh/shadow_all.h"
 #    endif
 
 #    if defined(__OBJECT_MOTION__)
 #      define BVH_FUNCTION_NAME bvh_intersect_shadow_all_motion
-#      define BVH_FUNCTION_FEATURES BVH_MOTION
+#      define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_POINTCLOUD
 #      include "kernel/bvh/shadow_all.h"
 #    endif
 
 #    if defined(__HAIR__) && defined(__OBJECT_MOTION__)
 #      define BVH_FUNCTION_NAME bvh_intersect_shadow_all_hair_motion
-#      define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION
+#      define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION | BVH_POINTCLOUD
 #      include "kernel/bvh/shadow_all.h"
 #    endif
+
 #  endif /* __SHADOW_RECORD_ALL__ */
 
 /* Record all intersections - Volume BVH traversal. */
diff --git a/intern/cycles/kernel/bvh/shadow_all.h b/intern/cycles/kernel/bvh/shadow_all.h
index 049c6a03fe0..caca85aac1a 100644
--- a/intern/cycles/kernel/bvh/shadow_all.h
+++ b/intern/cycles/kernel/bvh/shadow_all.h
@@ -28,6 +28,7 @@
  * without new features slowing things down.
  *
  * BVH_HAIR: hair curve rendering
+ * BVH_POINTCLOUD: point cloud rendering
  * BVH_MOTION: motion blur rendering
  */
 
@@ -199,6 +200,34 @@ ccl_device_inline
                 break;
               }
 #endif
+#if BVH_FEATURE(BVH_POINTCLOUD)
+              case PRIMITIVE_POINT:
+              case PRIMITIVE_MOTION_POINT: {
+                if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) {
+                  const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
+                  if (ray->time < prim_time.x || ray->time > prim_time.y) {
+                    hit = false;
+                    break;
+                  }
+                }
+
+                const int point_object = (object == OBJECT_NONE) ?
+                                             kernel_tex_fetch(__prim_object, prim_addr) :
+                                             object;
+                const int point_prim = kernel_tex_fetch(__prim_index, prim_addr);
+                const int point_type = kernel_tex_fetch(__prim_type, prim_addr);
+                hit = point_intersect(kg,
+                                      &isect,
+                                      P,
+                                      dir,
+                                      t_max_current,
+                                      point_object,
+                                      point_prim,
+                                      ray->time,
+                                      point_type);
+                break;
+              }
+#endif /* BVH_FEATURE(BVH_POINTCLOUD) */
               default: {
                 hit = false;
                 break;
diff --git a/intern/cycles/kernel/bvh/traversal.h b/intern/cycles/kernel/bvh/traversal.h
index 1c17ebf767f..180f19d11c5 100644
--- a/intern/cycles/kernel/bvh/traversal.h
+++ b/intern/cycles/kernel/bvh/traversal.h
@@ -28,6 +28,7 @@
  * without new features slowing things down.
  *
  * BVH_HAIR: hair curve rendering
+ * BVH_POINTCLOUD: point cloud rendering
  * BVH_MOTION: motion blur rendering
  */
 
@@ -188,6 +189,33 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
               break;
             }
 #endif /* BVH_FEATURE(BVH_HAIR) */
+#if BVH_FEATURE(BVH_POINTCLOUD)
+            case PRIMITIVE_POINT:
+            case PRIMITIVE_MOTION_POINT: {
+              for (; prim_addr < prim_addr2; prim_addr++) {
+                if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) {
+                  const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
+                  if (ray->time < prim_time.x || ray->time > prim_time.y) {
+                    continue;
+                  }
+                }
+
+                const int point_object = (object == OBJECT_NONE) ?
+                                             kernel_tex_fetch(__prim_object, prim_addr) :
+                                             object;
+                const int point_prim = kernel_tex_fetch(__prim_index, prim_addr);
+                const int point_type = kernel_tex_fetch(__prim_type, prim_addr);
+                const bool hit = point_intersect(
+                    kg, isect, P, dir, isect->t, point_object, point_prim, ray->time, point_type);
+                if (hit) {
+                  /* shadow ray early termination */
+                  if (visibility & PATH_RAY_SHADOW_OPAQUE)
+                    return true;
+                }
+              }
+              break;
+            }
+#endif /* BVH_FEATURE(BVH_POINTCLOUD) */
           }
         }
         else {
diff --git a/intern/cycles/kernel/bvh/types.h b/intern/cycles/kernel/bvh/types.h
index 6039e707fc3..f16f43333f8 100644
--- a/intern/cycles/kernel/bvh/types.h
+++ b/intern/cycles/kernel/bvh/types.h
@@ -34,6 +34,7 @@ CCL_NAMESPACE_BEGIN
 
 #define BVH_MOTION 1
 #define BVH_HAIR 2
+#define BVH_POINTCLOUD 4
 
 #define BVH_NAME_JOIN(x, y) x##_##y
 #define BVH_NAME_EVAL(x, y) BVH_NAME_JOIN(x, y)
diff --git a/intern/cycles/kernel/bvh/util.h b/intern/cycles/kernel/bvh/util.h
index 26ba136dd79..57593e42a88 100644
--- a/intern/cycles/kernel/bvh/util.h
+++ b/intern/cycles/kernel/bvh/util.h
@@ -118,14 +118,16 @@ ccl_device_forceinline int intersection_get_shader_flags(KernelGlobals kg,
 {
   int shader = 0;
 
-#ifdef __HAIR__
-  if (type & PRIMITIVE_ALL_TRIANGLE)
-#endif
-  {
+  if (type & PRIMITIVE_ALL_TRIANGLE) {
     shader = kernel_tex_fetch(__tri_shader, prim);
   }
+#ifdef __POINTCLOUD__
+  else if (type & PRIMITIVE_ALL_POINT) {
+    shader = kernel_tex_fetch(__points_shader, prim);
+  }
+#endif
 #ifdef __HAIR__
-  else {
+  else if (type & PRIMITIVE_ALL_CURVE) {
     shader = kernel_tex_fetch(__curves, prim).shader_id;
   }
 #endif
@@ -139,14 +141,16 @@ ccl_device_forceinline int intersection_get_shader_from_isect_prim(KernelGlobals
 {
   int shader = 0;
 
-#ifdef __HAIR__
-  if (isect_type & PRIMITIVE_ALL_TRIANGLE)
-#endif
-  {
+  if (isect_type & PRIMITIVE_ALL_TRIANGLE) {
     shader = kernel_tex_fetch(__tri_shader, prim);
   }
+#ifdef __POINTCLOUD__
+  else if (isect_type & PRIMITIVE_ALL_POINT) {
+    shader = kernel_tex_fetch(__points_shader, prim);
+  }
+#endif
 #ifdef __HAIR__
-  else {
+  else if (isect_type & PRIMITIVE_ALL_CURVE) {
     shader = kernel_tex_fetch(__curves, prim).shader_id;
   }
 #endif
diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu
index c9a48315544..c639dc87f35 100644
--- a/intern/cycles/kernel/device/optix/kernel.cu
+++ b/intern/cycles/kernel/device/optix/kernel.cu
@@ -97,9 +97,9 @@ extern "C" __global__ void __miss__kernel_optix_miss()
 
 extern "C" __global__ void __anyhit__kernel_optix_local_hit()
 {
-#ifdef __HAIR__
+#if defined(__HAIR__) || defined(__POINTCLOUD__)
   if (!optixIsTriangleHit()) {
-    /* Ignore curves. */
+    /* Ignore curves and points. */
     return optixIgnoreIntersection();
   }
 #endif
@@ -194,7 +194,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
     type = kernel_tex_fetch(__objects, object).primitive_type;
   }
 #  ifdef __HAIR__
-  else {
+  else if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) {
     u = __uint_as_float(optixGetAttribute_0());
     v = __uint_as_float(optixGetAttribute_1());
 
@@ -210,6 +210,11 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
 #    endif
   }
 #  endif
+  else {
+    type = kernel_tex_fetch(__objects, object).primitive_type;
+    u = 0.0f;
+    v = 0.0f;
+  }
 
 #  ifndef __TRANSPARENT_SHADOWS__
   /* No transparent shadows support compiled in, make opaque. */
@@ -291,7 +296,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
 
 extern "C" __global__ void __anyhit__kernel_optix_volume_test()
 {
-#ifdef __HAIR__
+#if defined(__HAIR__) || defined(__POINTCLOUD__)
   if (!optixIsTriangleHit()) {
     /* Ignore curves. */
     return optixIgnoreIntersection();
@@ -315,7 +320,7 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
 {
 #ifdef __HAIR__
 #  if OPTIX_ABI_VERSION < 55
-  if (!optixIsTriangleHit()) {
+  if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) {
     /* Filter out curve endcaps. */
     const float u = __uint_as_float(optixGetAttribute_0());
     if (u == 0.0f || u == 1.0f) {
@@ -354,13 +359,19 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
     optixSetPayload_3(prim);
     optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type);
   }
-  else {
+  else if (optixGetHitKind() & PRIMITIVE_ALL_CURVE) {
     const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
     optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */
     optixSetPayload_2(optixGetAttribute_1());
     optixSetPayload_3(segment.prim);
     optixSetPayload_5(segment.type);
   }
+  else {
+    optixSetPayload_1(0);
+    optixSetPayload_2(0);
+    optixSetPayload_3(prim);
+    optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type);
+  }
 }
 
 #ifdef __HAIR__
@@ -411,4 +422,45 @@ extern "C" __global__ void __intersection__curve_ribbon()
     optix_intersection_curve(prim, type);
   }
 }
+
+#endif
+
+#ifdef __POINTCLOUD__
+extern "C" __global__ void __intersection__point()
+{
+  const int prim = optixGetPrimitiveIndex();
+  const int object = get_object_id();
+  const int type = kernel_tex_fetch(__objects, object).primitive_type;
+
+#  ifdef __VISIBILITY_FLAG__
+  const uint visibility = optixGetPayload_4();
+  if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
+    return;
+  }
+#  endif
+
+  float3 P = optixGetObjectRayOrigin();
+  float3 dir = optixGetObjectRayDirection();
+
+  /* The direction is not normalized by default, the point intersection routine expects that. */
+  float len;
+  dir = normalize_len(dir, &len);
+
+#  ifdef __OBJECT_MOTION__
+  const float time = optixGetRayTime();
+#  else
+  const float time = 0.0f;
+#  endif
+
+  Intersection isect;
+  isect.t = optixGetRayTmax();
+  /* Transform maximum distance into object space. */
+  if (isect.t != FLT_MAX) {
+    isect.t *= len;
+  }
+
+  if (point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
+    optixReportIntersection(isect.t / len, type & PRIMITIVE_ALL);
+  }
+}
 #endif
diff --git a/intern/cycles/kernel/geom/geom.h b/intern/cycles/kernel/geom/geom.h
index 9d023375a35..efc296c8d39 100644
--- a/intern/cycles/kernel/geom/geom.h
+++ b/intern/cycles/kernel/geom/geom.h
@@ -29,6 +29,9 @@
 #include "kernel/geom/motion_triangle_intersect.h"
 #include "kernel/geom/motion_triangle_shader.h"
 #include "kernel/geom/motion_curve.h"
+#include "kernel/geom/motion_point.h"
+#include "kernel/geom/point.h"
+#include "kernel/geom/point_intersect.h"
 #include "kernel/geom/curve.h"
 #include "kernel/geom/curve_intersect.h"
 #include "kernel/geom/volume.h"
diff --git a/intern/cycles/kernel/geom/motion_point.h b/intern/cycles/kernel/geom/motion_point.h
new file mode 100644
index 00000000000..d7ffc80c045
--- /dev/null
+++ b/intern/cycles/kernel/geom/motion_point.h
@@ -0,0 +1,74 @@
+/*
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+CCL_NAMESPACE_BEGIN
+
+/* Motion Point Primitive
+ *
+ * These are stored as regular points, plus extra positions and radii at times
+ * other than the frame center. Computing the point at a given ray time is
+ * a matter of interpolation of the two steps between which the ray time lies.
+ *
+ * The extra points are stored as ATTR_STD_MOTION_VERTEX_POSITION.
+ */
+
+#ifdef __POINTCLOUD__
+
+ccl_device_inline float4
+motion_point_for_step(KernelGlobals kg, int offset, int numkeys, int numsteps, int step, int prim)
+{
+  if (step == numsteps) {
+    /* center step: regular key location */
+    return kernel_tex_fetch(__points, prim);
+  }
+  else {
+    /* center step is not stored in this array */
+    if (step > numsteps)
+      step--;
+
+    offset += step * numkeys;
+
+    return kernel_tex_fetch(__attributes_float4, offset + prim);
+  }
+}
+
+/* return 2 point key locations */
+ccl_device_inline float4 motion_point(KernelGlobals kg, int object, int prim, float time)
+{
+  /* get motion info */
+  int numsteps, numkeys;
+  object_motion_info(kg, object, &numsteps, NULL, &numkeys);
+
+  /* figure out which steps we need to fetch and their interpolation factor */
+  int maxstep = numsteps * 2;
+  int step = min((int)(time * maxstep), maxstep - 1);
+  float t = time * maxstep - step;
+
+  /* find attribute */
+  int offset = intersection_find_attribute(kg, object, ATTR_STD_MOTION_VERTEX_POSITION);
+  kernel_assert(offset != ATTR_STD_NOT_FOUND);
+
+  /* fetch key coordinates */
+  float4 point = motion_point_for_step(kg, offset, numkeys, numsteps, step, prim);
+  float4 next_point = motion_point_for_step(kg, offset, numkeys, numsteps, step + 1, prim);
+
+  /* interpolate between steps */
+  return (1.0f - t) * point + t * next_point;
+}
+
+#endif
+
+CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/geom/point.h b/intern/cycles/kernel/geom/point.h
new file mode 100644
index 00000000000..021135b76fb
--- /dev/null
+++ b/intern/cycles/kernel/geom/point.h
@@ -0,0 +1,124 @@
+/*
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+CCL_NAMESPACE_BEGIN
+
+/* Point Primitive
+ *
+ * Point primitive for rendering point clouds.
+ */
+
+#ifdef __POINTCLOUD__
+
+/* Reading attributes on various point elements */
+
+ccl_device float point_attribute_float(KernelGlobals kg,
+                                       ccl_private const ShaderData *sd,
+                                       const AttributeDescriptor desc,
+                                       ccl_private float *dx,
+                                       ccl_private float *dy)
+{
+#  ifdef __RAY_DIFFERENTIALS__
+  if (dx)
+    *dx = 0.0f;
+  if (dy)
+    *dy = 0.0f;
+#  endif
+
+  if (desc.element == ATTR_ELEMENT_VERTEX) {
+    return kernel_tex_fetch(__attributes_float, desc.offset + sd->prim);
+  }
+  else {
+    return 0.0f;
+  }
+}
+
+ccl_device float2 point_attribute_float2(
+    KernelGlobals kg, const ShaderData *sd, const AttributeDescriptor desc, float2 *dx, float2 *dy)
+{
+#  ifdef __RAY_DIFFERENTIALS__
+  if (dx)
+    *dx = make_float2(0.0f, 0.0f);
+  if (dy)
+    *dy = make_float2(0.0f, 0.0f);
+#  endif
+
+  if (desc.element == ATTR_ELEMENT_VERTEX) {
+    return kernel_tex_fetch(__attributes_float2, desc.offset + sd->prim);
+  }
+  else {
+    return make_float2(0.0f, 0.0f);
+  }
+}
+
+ccl_device float3 point_attribute_float3(
+    KernelGlobals kg, const ShaderData *sd, const AttributeDescriptor desc, float3 *dx, float3 *dy)
+{
+#  ifdef __RAY_DIFFERENTIALS__
+  if (dx)
+    *dx = make_float3(0.0f, 0.0f, 0.0f);
+  if (dy)
+    *dy = make_float3(0.0f, 0.0f, 0.0f);
+#  endif
+
+  if (desc.element == ATTR_ELEMENT_VERTEX) {
+    return float4_to_float3(kernel_tex_fetch(__attributes_float4, desc.offset + sd->prim));
+  }
+  else {
+    return make_float3(0.0f, 0.0f, 0.0f);
+  }
+}
+
+ccl_device float4 point_attribute_float4(
+    KernelGlobals kg, const ShaderData *sd, const AttributeDescriptor desc, float4 *dx, float4 *dy)
+{
+#  ifdef __RAY_DIFFERENTIALS__
+  if (dx)
+    *dx = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
+  if (dy)
+    *dy = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
+#  endif
+
+  if (desc.element == ATTR_ELEMENT_VERTEX) {
+    return kernel_tex_fetch(__attributes_float4, desc.offset + sd->prim);
+  }
+  else {
+    return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
+  }
+}
+
+/* Point radius */
+
+ccl_device float point_radius(KernelGlobals kg, ccl_private const ShaderData *sd)
+{
+  if (sd->type & PRIMITIVE_ALL_POINT) {
+    return kernel_tex_fetch(__points, sd->prim).w;
+  }
+
+  return 0.0f;
+}
+
+/* Point location for motion pass, linear interpolation between keys and
+ * ignoring radius because we do the same for the motion keys */
+
+ccl_device float3 point_motion_center_location(KernelGlobals kg, ccl_private const ShaderData *sd)
+{
+  return float4_to_float3(kernel_tex_fetch(__points, sd->prim));
+}
+
+#endif /* __POINTCLOUD__ */
+
+CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/geom/point_intersect.h b/intern/cycles/kernel/geom/point_intersect.h
new file mode 100644
index 00000000000..da0b31f9a8c
--- /dev/null
+++ b/intern/cycles/kernel/geom/point_intersect.h
@@ -0,0 +1,133 @@
+/*
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *
+ * Based on Embree code, copyright 2009-2020 Intel Corporation.
+ */
+
+#pragma once
+
+CCL_NAMESPACE_BEGIN
+
+/* Point primitive intersection functions. */
+
+#ifdef __POINTCLOUD__
+
+ccl_device_forceinline bool point_intersect_test(
+    const float4 point, const float3 P, const float3 dir, const float tmax, float *t)
+{
+  const float3 center = float4_to_float3(point);
+  const float radius = point.w;
+
+  const float rd2 = 1.0f / dot(dir, dir);
+
+  const float3 c0 = center - P;
+  const float projC0 = dot(c0, dir) * rd2;
+  const float3 perp = c0 - projC0 * dir;
+  const float l2 = dot(perp, perp);
+  const float r2 = radius * radius;
+  if (!(l2 <= r2)) {
+    return false;
+  }
+
+  const float td = sqrt((r2 - l2) * rd2);
+  const float t_front = projC0 - td;
+  const bool valid_front = (0.0f <= t_front) & (t_front <= tmax);
+
+  /* Always backface culling for now. */
+#  if 0
+  const float t_back = projC0 + td;
+  const bool valid_back = (0.0f <= t_back) & (t_back <= tmax);
+
+  /* check if there is a first hit */
+  const bool valid_first = valid_front | valid_back;
+  if (!valid_first) {
+    return false;
+  }
+
+  *t = (valid_front) ? t_front : t_back;
+  return true;
+#  else
+  if (!valid_front) {
+    return false;
+  }
+  *t = t_front;
+  return true;
+#  endif
+}
+
+ccl_device_forceinline bool point_intersect(KernelGlobals kg,
+                                            ccl_private Intersection *isect,
+                                            const float3 P,
+                                            const float3 dir,
+                                            const float tmax,
+                                            const int object,
+                                            const int prim,
+                                            const float time,
+                                            const int type)
+{
+  const float4 point = (type & PRIMITIVE_ALL_MOTION) ? motion_point(kg, object, prim, time) :
+                                                       kernel_tex_fetch(__points, prim);
+
+  if (!point_intersect_test(point, P, dir, tmax, &isect->t)) {
+    return false;
+  }
+
+  isect->prim = prim;
+  isect->object = object;
+  isect->type = type;
+  isect->u = 0.0f;
+  isect->v = 0.0f;
+  return true;
+}
+
+ccl_device_inline void point_shader_setup(KernelGlobals kg,
+                                          ccl_private ShaderData *sd,
+                                          ccl_private const Intersection *isect,
+                                          const Ray *ray)
+{
+  sd->shader = kernel_tex_fetch(__points_shader, isect->prim);
+  sd->P = ray->P + ray->D * isect->t;
+
+  /* Texture coordinates, zero for now. */
+#  ifdef __UV__
+  sd->u = isect->u;
+  sd->v = isect->v;
+#  endif
+
+  /* Computer point center for normal. */
+  float3 center = float4_to_float3((isect->type & PRIMITIVE_ALL_MOTION) ?
+                                       motion_point(kg, sd->object, sd->prim, sd->time) :
+                                       kernel_tex_fetch(__points, sd->prim));
+
+  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
+    const Transform tfm = object_get_transform(kg, sd);
+
+#  ifndef __KERNEL_OPTIX__
+    center = transform_point(&tfm, center);
+#  endif
+  }
+
+  /* Normal */
+  sd->Ng = normalize(sd->P - center);
+  sd->N = sd->Ng;
+
+#  ifdef __DPDU__
+  /* dPdu/dPdv */
+  sd->dPdu = make_float3(0.0f, 0.0f, 0.0f);
+  sd->dPdv = make_float3(0.0f, 0.0f, 0.0f);
+#  endif
+}
+
+#endif
+
+CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/geom/primitive.h b/intern/cycles/kernel/geom/primitive.h
index 6d7b550d82f..9416385638c 100644
--- a/intern/cycles/kernel/geom/primitive.h
+++ b/intern/cycles/kernel/geom/primitive.h
@@ -48,6 +48,11 @@ ccl_device_inline float primitive_surface_attribute_float(KernelGlobals kg,
     return curve_attribute_float(kg, sd, desc, dx, dy);
   }
 #endif
+#ifdef __POINTCLOUD__
+  else if (sd->type & PRIMITIVE_ALL_POINT) {
+    return point_attribute_float(kg, sd, desc, dx, dy);
+  }
+#endif
   else {
     if (dx)
       *dx = 0.0f;
@@ -74,6 +79,11 @@ ccl_device_inline float2 primitive_surface_attribute_float2(KernelGlobals kg,
     return curve_attribute_float2(kg, sd, desc, dx, dy);
   }
 #endif
+#ifdef __POINTCLOUD__
+  else if (sd->type & PRIMITIVE_ALL_POINT) {
+    return point_attribute_float2(kg, sd, desc, dx, dy);
+  }
+#endif
   else {
     if (dx)
       *dx = make_float2(0.0f, 0.0f);
@@ -100,6 +110,11 @@ ccl_device_inline float3 primitive_surface_attribute_float3(KernelGlobals kg,
     return curve_attribute_float3(kg, sd, desc, dx, dy);
   }
 #endif
+#ifdef __POINTCLOUD__
+  else if (sd->type & PRIMITIVE_ALL_POINT) {
+    return point_attribute_float3(kg, sd, desc, dx, dy);
+  }
+#endif
   else {
     if (dx)
       *dx = make_float3(0.0f, 0.0f, 0.0f);
@@ -126,6 +141,11 @@ ccl_device_forceinline float4 primitive_surface_attribute_float4(KernelGlobals k
     return curve_attribute_float4(kg, sd, desc, dx, dy);
   }
 #endif
+#ifdef __POINTCLOUD__
+  else if (sd->type & PRIMITIVE_ALL_POINT) {
+    return point_attribute_float4(kg, sd, desc, dx, dy);
+  }
+#endif
   else {
     if (dx)
       *dx = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
@@ -225,8 +245,8 @@ ccl_device bool primitive_ptex(KernelGlobals kg,
 
 ccl_device float3 primitive_tangent(KernelGlobals kg, ccl_private ShaderData *sd)
 {
-#ifdef __HAIR__
-  if (sd->type & PRIMITIVE_ALL_CURVE)
+#if defined(__HAIR__) || defined(__POINTCLOUD__)
+  if (sd->type & (PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_POINT))
 #  ifdef __DPDU__
     return normalize(sd->dPdu);
 #  else
@@ -261,10 +281,21 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg,
   /* center position */
   float3 center;
 
-#ifdef __HAIR__
-  bool is_curve_primitive = sd->type & PRIMITIVE_ALL_CURVE;
-  if (is_curve_primitive) {
-    center = curve_motion_center_location(kg, sd);
+#if defined(__HAIR__) || defined(__POINTCLOUD__)
+  bool is_curve_or_point = sd->type & (PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_POINT);
+  if (is_curve_or_point) {
+    center = make_float3(0.0f, 0.0f, 0.0f);
+
+    if (sd->type & PRIMITIVE_ALL_CURVE) {
+#  if defined(__HAIR__)
+      center = curve_motion_center_location(kg, sd);
+#  endif
+    }
+    else if (sd->type & PRIMITIVE_ALL_POINT) {
+#  if defined(__POINTCLOUD__)
+      center = point_motion_center_location(kg, sd);
+#  endif
+    }
 
     if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
       object_position_transform(kg, sd, &center);
@@ -272,7 +303,9 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg,
   }
   else
 #endif
+  {
     center = sd->P;
+  }
 
   float3 motion_pre = center, motion_post = center;
 
@@ -284,8 +317,8 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg,
     int numverts, numkeys;
     object_motion_info(kg, sd->object, NULL, &numverts, &numkeys);
 
-#ifdef __HAIR__
-    if (is_curve_primitive) {
+#if defined(__HAIR__) || defined(__POINTCLOUD__)
+    if (is_curve_or_point) {
       motion_pre = float4_to_float3(curve_attribute_float4(kg, sd, desc, NULL, NULL));
       desc.offset += numkeys;
       motion_post = float4_to_float3(curve_attribute_float4(kg, sd, desc, NULL, NULL));
diff --git a/intern/cycles/kernel/geom/shader_data.h b/intern/cycles/kernel/geom/shader_data.h
index 46bda2b656c..d3932545ef6 100644
--- a/intern/cycles/kernel/geom/shader_data.h
+++ b/intern/cycles/kernel/geom/shader_data.h
@@ -75,6 +75,13 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
   }
   else
 #endif
+#ifdef __POINTCLOUD__
+      if (sd->type & PRIMITIVE_ALL_POINT) {
+    /* point */
+    point_shader_setup(kg, sd, isect, ray);
+  }
+  else
+#endif
       if (sd->type & PRIMITIVE_TRIANGLE) {
     /* static triangle */
     float3 Ng = triangle_normal(kg, sd);
diff --git a/intern/cycles/kernel/osl/services.cpp b/intern/cycles/kernel/osl/services.cpp
index 389c1e2b746..4007005dee7 100644
--- a/intern/cycles/kernel/osl/services.cpp
+++ b/intern/cycles/kernel/osl/services.cpp
@@ -28,6 +28,7 @@
 #include "scene/colorspace.h"
 #include "scene/mesh.h"
 #include "scene/object.h"
+#include "scene/pointcloud.h"
 #include "scene/scene.h"
 
 #include "kernel/osl/closures.h"
@@ -113,6 +114,8 @@ ustring OSLRenderServices::u_curve_thickness("geom:curve_thickness");
 ustring OSLRenderServices::u_curve_length("geom:curve_length");
 ustring OSLRenderServices::u_curve_tangent_normal("geom:curve_tangent_normal");
 ustring OSLRenderServices::u_curve_random("geom:curve_random");
+ustring OSLRenderServices::u_is_point("geom:is_point");
+ustring OSLRenderServices::u_point_radius("geom:point_radius");
 ustring OSLRenderServices::u_normal_map_normal("geom:normal_map_normal");
 ustring OSLRenderServices::u_path_ray_length("path:ray_length");
 ustring OSLRenderServices::u_path_ray_depth("path:ray_depth");
@@ -994,6 +997,15 @@ bool OSLRenderServices::get_object_standard_attribute(const KernelGlobalsCPU *kg
     float3 f = curve_tangent_normal(kg, sd);
     return set_attribute_float3(f, type, derivatives, val);
   }
+  /* point attributes */
+  else if (name == u_is_point) {
+    float f = (sd->type & PRIMITIVE_ALL_POINT) != 0;
+    return set_attribute_float(f, type, derivatives, val);
+  }
+  else if (name == u_point_radius) {
+    float f = point_radius(kg, sd);
+    return set_attribute_float(f, type, derivatives, val);
+  }
   else if (name == u_normal_map_normal) {
     if (sd->type & PRIMITIVE_ALL_TRIANGLE) {
       float3 f = triangle_smooth_normal_unnormalized(kg, sd, sd->Ng, sd->prim, sd->u, sd->v);
diff --git a/intern/cycles/kernel/osl/services.h b/intern/cycles/kernel/osl/services.h
index d9f57c642ad..9526c92b8fb 100644
--- a/intern/cycles/kernel/osl/services.h
+++ b/intern/cycles/kernel/osl/services.h
@@ -297,6 +297,8 @@ class OSLRenderServices : public OSL::RendererServices {
   static ustring u_curve_length;
   static ustring u_curve_tangent_normal;
   static ustring u_curve_random;
+  static ustring u_is_point;
+  static ustring u_point_radius;
   static ustring u_normal_map_normal;
   static ustring u_path_ray_length;
   static ustring u_path_ray_depth;
diff --git a/intern/cycles/kernel/svm/wireframe.h b/intern/cycles/kernel/svm/wireframe.h
index 530a9601bce..645dc59f22e 100644
--- a/intern/cycles/kernel/svm/wireframe.h
+++ b/intern/cycles/kernel/svm/wireframe.h
@@ -42,7 +42,7 @@ ccl_device_inline float wireframe(KernelGlobals kg,
                                   int pixel_size,
                                   ccl_private float3 *P)
 {
-#ifdef __HAIR__
+#if defined(__HAIR__) || defined(__POINTCLOUD__)
   if (sd->prim != PRIM_NONE && sd->type & PRIMITIVE_ALL_TRIANGLE)
 #else
   if (sd->prim != PRIM_NONE)
diff --git a/intern/cycles/kernel/textures.h b/intern/cycles/kernel/textures.h
index 58edb239007..06074a1eca1 100644
--- a/intern/cycles/kernel/textures.h
+++ b/intern/cycles/kernel/textures.h
@@ -55,6 +55,10 @@ KERNEL_TEX(KernelCurveSegment, __curve_segments)
 /* patches */
 KERNEL_TEX(uint, __patches)
 
+/* pointclouds */
+KERNEL_TEX(float4, __points)
+KERNEL_TEX(uint, __points_shader)
+
 /* attributes */
 KERNEL_TEX(uint4, __attributes_map)
 KERNEL_TEX(float, __attributes_float)
diff --git a/intern/cycles/kernel/types.h b/intern/cycles/kernel/types.h
index 34f909a06d9..855cc97edbf 100644
--- a/intern/cycles/kernel/types.h
+++ b/intern/cycles/kernel/types.h
@@ -86,6 +86,7 @@ CCL_NAMESPACE_BEGIN
 #define __AO__
 #define __PASSES__
 #define __HAIR__
+#define __POINTCLOUD__
 #define __SVM__
 #define __EMISSION__
 #define __HOLDOUT__
@@ -125,6 +126,9 @@ CCL_NAMESPACE_BEGIN
 #  if !(__KERNEL_FEATURES & KERNEL_FEATURE_HAIR)
 #    undef __HAIR__
 #  endif
+#  if !(__KERNEL_FEATURES & KERNEL_FEATURE_POINTCLOUD)
+#    undef __POINTCLOUD__
+#  endif
 #  if !(__KERNEL_FEATURES & KERNEL_FEATURE_VOLUME)
 #    undef __VOLUME__
 #  endif
@@ -538,19 +542,22 @@ typedef enum PrimitiveType {
   PRIMITIVE_MOTION_CURVE_THICK = (1 << 3),
   PRIMITIVE_CURVE_RIBBON = (1 << 4),
   PRIMITIVE_MOTION_CURVE_RIBBON = (1 << 5),
-  PRIMITIVE_VOLUME = (1 << 6),
-  PRIMITIVE_LAMP = (1 << 7),
+  PRIMITIVE_POINT = (1 << 6),
+  PRIMITIVE_MOTION_POINT = (1 << 7),
+  PRIMITIVE_VOLUME = (1 << 8),
+  PRIMITIVE_LAMP = (1 << 9),
 
   PRIMITIVE_ALL_TRIANGLE = (PRIMITIVE_TRIANGLE | PRIMITIVE_MOTION_TRIANGLE),
   PRIMITIVE_ALL_CURVE = (PRIMITIVE_CURVE_THICK | PRIMITIVE_MOTION_CURVE_THICK |
                          PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON),
+  PRIMITIVE_ALL_POINT = (PRIMITIVE_POINT | PRIMITIVE_MOTION_POINT),
   PRIMITIVE_ALL_VOLUME = (PRIMITIVE_VOLUME),
   PRIMITIVE_ALL_MOTION = (PRIMITIVE_MOTION_TRIANGLE | PRIMITIVE_MOTION_CURVE_THICK |
-                          PRIMITIVE_MOTION_CURVE_RIBBON),
+                          PRIMITIVE_MOTION_CURVE_RIBBON | PRIMITIVE_MOTION_POINT),
   PRIMITIVE_ALL = (PRIMITIVE_ALL_TRIANGLE | PRIMITIVE_ALL_CURVE | PRIMITIVE_ALL_VOLUME |
-                   PRIMITIVE_LAMP),
+                   PRIMITIVE_LAMP | PRIMITIVE_ALL_POINT),
 
-  PRIMITIVE_NUM = 8,
+  PRIMITIVE_NUM = 10,
 } PrimitiveType;
 
 #define PRIMITIVE_PACK_SEGMENT(type, segment) ((segment << PRIMITIVE_NUM) | (type))
@@ -605,6 +612,7 @@ typedef enum AttributeStandard {
   ATTR_STD_CURVE_INTERCEPT,
   ATTR_STD_CURVE_LENGTH,
   ATTR_STD_CURVE_RANDOM,
+  ATTR_STD_POINT_RANDOM,
   ATTR_STD_PTEX_FACE_ID,
   ATTR_STD_PTEX_UV,
   ATTR_STD_VOLUME_DENSITY,
@@ -1604,6 +1612,9 @@ enum KernelFeatureFlag : unsigned int {
   KERNEL_FEATURE_AO_PASS = (1U << 25U),
   KERNEL_FEATURE_AO_ADDITIVE = (1U << 26U),
   KERNEL_FEATURE_AO = (KERNEL_FEATURE_AO_PASS | KERNEL_FEATURE_AO_ADDITIVE),
+
+  /* Point clouds. */
+  KERNEL_FEATURE_POINTCLOUD = (1U << 27U),
 };
 
 /* Shader node feature mask, to specialize shader evaluation for kernels. */