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

47 files changed, 2046 insertions, 69 deletions

diff --git a/intern/cycles/blender/CMakeLists.txt b/intern/cycles/blender/CMakeLists.txt
index 1f05b1aa234..fe7d0b89bb0 100644
--- a/intern/cycles/blender/CMakeLists.txt
+++ b/intern/cycles/blender/CMakeLists.txt
@@ -40,6 +40,7 @@ set(SRC
   object_cull.cpp
   output_driver.cpp
   particles.cpp
+  pointcloud.cpp
   curves.cpp
   logging.cpp
   python.cpp
diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index fd86d75a301..e4fbc898070 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -1020,7 +1020,7 @@ class CYCLES_OBJECT_PT_motion_blur(CyclesButtonsPanel, Panel):
     def poll(cls, context):
         ob = context.object
         if CyclesButtonsPanel.poll(context) and ob:
-            if ob.type in {'MESH', 'CURVE', 'CURVE', 'SURFACE', 'FONT', 'META', 'CAMERA'}:
+            if ob.type in {'MESH', 'CURVE', 'CURVE', 'SURFACE', 'FONT', 'META', 'CAMERA', 'HAIR', 'POINTCLOUD'}:
                 return True
             if ob.instance_type == 'COLLECTION' and ob.instance_collection:
                 return True
diff --git a/intern/cycles/blender/geometry.cpp b/intern/cycles/blender/geometry.cpp
index a5f36a3a17d..78c803b7adb 100644
--- a/intern/cycles/blender/geometry.cpp
+++ b/intern/cycles/blender/geometry.cpp
@@ -19,6 +19,7 @@
 #include "scene/hair.h"
 #include "scene/mesh.h"
 #include "scene/object.h"
+#include "scene/pointcloud.h"
 #include "scene/volume.h"
 
 #include "blender/sync.h"
@@ -39,6 +40,10 @@ static Geometry::Type determine_geom_type(BObjectInfo &b_ob_info, bool use_parti
     return Geometry::HAIR;
   }
 
+  if (b_ob_info.object_data.is_a(&RNA_PointCloud)) {
+    return Geometry::POINTCLOUD;
+  }
+
   if (b_ob_info.object_data.is_a(&RNA_Volume) ||
       (b_ob_info.object_data == b_ob_info.real_object.data() &&
        object_fluid_gas_domain_find(b_ob_info.real_object))) {
@@ -111,6 +116,9 @@ Geometry *BlenderSync::sync_geometry(BL::Depsgraph &b_depsgraph,
     else if (geom_type == Geometry::VOLUME) {
       geom = scene->create_node<Volume>();
     }
+    else if (geom_type == Geometry::POINTCLOUD) {
+      geom = scene->create_node<PointCloud>();
+    }
     else {
       geom = scene->create_node<Mesh>();
     }
@@ -170,6 +178,10 @@ Geometry *BlenderSync::sync_geometry(BL::Depsgraph &b_depsgraph,
       Volume *volume = static_cast<Volume *>(geom);
       sync_volume(b_ob_info, volume);
     }
+    else if (geom_type == Geometry::POINTCLOUD) {
+      PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+      sync_pointcloud(pointcloud, b_ob_info);
+    }
     else {
       Mesh *mesh = static_cast<Mesh *>(geom);
       sync_mesh(b_depsgraph, b_ob_info, mesh);
@@ -231,6 +243,10 @@ void BlenderSync::sync_geometry_motion(BL::Depsgraph &b_depsgraph,
              object_fluid_gas_domain_find(b_ob_info.real_object)) {
       /* No volume motion blur support yet. */
     }
+    else if (b_ob_info.object_data.is_a(&RNA_PointCloud)) {
+      PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+      sync_pointcloud_motion(pointcloud, b_ob_info, motion_step);
+    }
     else {
       Mesh *mesh = static_cast<Mesh *>(geom);
       sync_mesh_motion(b_depsgraph, b_ob_info, mesh, motion_step);
diff --git a/intern/cycles/blender/object.cpp b/intern/cycles/blender/object.cpp
index d6c52f7faf9..86314d3b196 100644
--- a/intern/cycles/blender/object.cpp
+++ b/intern/cycles/blender/object.cpp
@@ -72,7 +72,8 @@ bool BlenderSync::object_is_geometry(BObjectInfo &b_ob_info)
 
   BL::Object::type_enum type = b_ob_info.iter_object.type();
 
-  if (type == BL::Object::type_VOLUME || type == BL::Object::type_HAIR) {
+  if (type == BL::Object::type_VOLUME || type == BL::Object::type_HAIR ||
+      type == BL::Object::type_POINTCLOUD) {
     /* Will be exported attached to mesh. */
     return true;
   }
@@ -206,7 +207,7 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph,
     return NULL;
   }
 
-  /* only interested in object that we can create meshes from */
+  /* only interested in object that we can create geometry from */
   if (!object_is_geometry(b_ob_info)) {
     return NULL;
   }
diff --git a/intern/cycles/blender/pointcloud.cpp b/intern/cycles/blender/pointcloud.cpp
new file mode 100644
index 00000000000..a9e616a468f
--- /dev/null
+++ b/intern/cycles/blender/pointcloud.cpp
@@ -0,0 +1,253 @@
+/*
+ * Copyright 2011-2013 Blender Foundation
+ *
+ * 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.
+ */
+
+#include "scene/pointcloud.h"
+#include "scene/attribute.h"
+#include "scene/scene.h"
+
+#include "blender/sync.h"
+#include "blender/util.h"
+
+#include "util/foreach.h"
+#include "util/hash.h"
+
+CCL_NAMESPACE_BEGIN
+
+template<typename TypeInCycles, typename GetValueAtIndex>
+static void fill_generic_attribute(BL::PointCloud &b_pointcloud,
+                                   TypeInCycles *data,
+                                   const GetValueAtIndex &get_value_at_index)
+{
+  const int num_points = b_pointcloud.points.length();
+  for (int i = 0; i < num_points; i++) {
+    data[i] = get_value_at_index(i);
+  }
+}
+
+static void copy_attributes(PointCloud *pointcloud, BL::PointCloud b_pointcloud)
+{
+  AttributeSet &attributes = pointcloud->attributes;
+  for (BL::Attribute &b_attribute : b_pointcloud.attributes) {
+    const ustring name{b_attribute.name().c_str()};
+
+    if (attributes.find(name)) {
+      continue;
+    }
+
+    const AttributeElement element = ATTR_ELEMENT_VERTEX;
+    const BL::Attribute::data_type_enum b_data_type = b_attribute.data_type();
+    switch (b_data_type) {
+      case BL::Attribute::data_type_FLOAT: {
+        BL::FloatAttribute b_float_attribute{b_attribute};
+        Attribute *attr = attributes.add(name, TypeFloat, element);
+        float *data = attr->data_float();
+        fill_generic_attribute(
+            b_pointcloud, data, [&](int i) { return b_float_attribute.data[i].value(); });
+        break;
+      }
+      case BL::Attribute::data_type_BOOLEAN: {
+        BL::BoolAttribute b_bool_attribute{b_attribute};
+        Attribute *attr = attributes.add(name, TypeFloat, element);
+        float *data = attr->data_float();
+        fill_generic_attribute(
+            b_pointcloud, data, [&](int i) { return (float)b_bool_attribute.data[i].value(); });
+        break;
+      }
+      case BL::Attribute::data_type_INT: {
+        BL::IntAttribute b_int_attribute{b_attribute};
+        Attribute *attr = attributes.add(name, TypeFloat, element);
+        float *data = attr->data_float();
+        fill_generic_attribute(
+            b_pointcloud, data, [&](int i) { return (float)b_int_attribute.data[i].value(); });
+        break;
+      }
+      case BL::Attribute::data_type_FLOAT_VECTOR: {
+        BL::FloatVectorAttribute b_vector_attribute{b_attribute};
+        Attribute *attr = attributes.add(name, TypeVector, element);
+        float3 *data = attr->data_float3();
+        fill_generic_attribute(b_pointcloud, data, [&](int i) {
+          BL::Array<float, 3> v = b_vector_attribute.data[i].vector();
+          return make_float3(v[0], v[1], v[2]);
+        });
+        break;
+      }
+      case BL::Attribute::data_type_FLOAT_COLOR: {
+        BL::FloatColorAttribute b_color_attribute{b_attribute};
+        Attribute *attr = attributes.add(name, TypeRGBA, element);
+        float4 *data = attr->data_float4();
+        fill_generic_attribute(b_pointcloud, data, [&](int i) {
+          BL::Array<float, 4> v = b_color_attribute.data[i].color();
+          return make_float4(v[0], v[1], v[2], v[3]);
+        });
+        break;
+      }
+      case BL::Attribute::data_type_FLOAT2: {
+        BL::Float2Attribute b_float2_attribute{b_attribute};
+        Attribute *attr = attributes.add(name, TypeFloat2, element);
+        float2 *data = attr->data_float2();
+        fill_generic_attribute(b_pointcloud, data, [&](int i) {
+          BL::Array<float, 2> v = b_float2_attribute.data[i].vector();
+          return make_float2(v[0], v[1]);
+        });
+        break;
+      }
+      default:
+        /* Not supported. */
+        break;
+    }
+  }
+}
+
+static void export_pointcloud(Scene *scene, PointCloud *pointcloud, BL::PointCloud b_pointcloud)
+{
+  /* TODO: optimize so we can straight memcpy arrays from Blender? */
+
+  /* Add requested attributes. */
+  Attribute *attr_random = NULL;
+  if (pointcloud->need_attribute(scene, ATTR_STD_POINT_RANDOM)) {
+    attr_random = pointcloud->attributes.add(ATTR_STD_POINT_RANDOM);
+  }
+
+  /* Reserve memory. */
+  const int num_points = b_pointcloud.points.length();
+  pointcloud->reserve(num_points);
+
+  /* Export points. */
+  BL::PointCloud::points_iterator b_point_iter;
+  for (b_pointcloud.points.begin(b_point_iter); b_point_iter != b_pointcloud.points.end();
+       ++b_point_iter) {
+    BL::Point b_point = *b_point_iter;
+    const float3 co = get_float3(b_point.co());
+    const float radius = b_point.radius();
+    pointcloud->add_point(co, radius);
+
+    /* Random number per point. */
+    if (attr_random != NULL) {
+      attr_random->add(hash_uint2_to_float(b_point.index(), 0));
+    }
+  }
+
+  /* Export attributes */
+  copy_attributes(pointcloud, b_pointcloud);
+}
+
+static void export_pointcloud_motion(PointCloud *pointcloud,
+                                     BL::PointCloud b_pointcloud,
+                                     int motion_step)
+{
+  /* Find or add attribute. */
+  Attribute *attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+  bool new_attribute = false;
+
+  if (!attr_mP) {
+    attr_mP = pointcloud->attributes.add(ATTR_STD_MOTION_VERTEX_POSITION);
+    new_attribute = true;
+  }
+
+  /* Export motion points. */
+  const int num_points = pointcloud->num_points();
+  float3 *mP = attr_mP->data_float3() + motion_step * num_points;
+  bool have_motion = false;
+  int num_motion_points = 0;
+  const array<float3> &pointcloud_points = pointcloud->get_points();
+
+  BL::PointCloud::points_iterator b_point_iter;
+  for (b_pointcloud.points.begin(b_point_iter); b_point_iter != b_pointcloud.points.end();
+       ++b_point_iter) {
+    BL::Point b_point = *b_point_iter;
+
+    if (num_motion_points < num_points) {
+      float3 P = get_float3(b_point.co());
+      P.w = b_point.radius();
+      mP[num_motion_points] = P;
+      have_motion = have_motion || (P != pointcloud_points[num_motion_points]);
+      num_motion_points++;
+    }
+  }
+
+  /* In case of new attribute, we verify if there really was any motion. */
+  if (new_attribute) {
+    if (num_motion_points != num_points || !have_motion) {
+      pointcloud->attributes.remove(ATTR_STD_MOTION_VERTEX_POSITION);
+    }
+    else if (motion_step > 0) {
+      /* Motion, fill up previous steps that we might have skipped because
+       * they had no motion, but we need them anyway now. */
+      for (int step = 0; step < motion_step; step++) {
+        pointcloud->copy_center_to_motion_step(step);
+      }
+    }
+  }
+
+  /* Export attributes */
+  copy_attributes(pointcloud, b_pointcloud);
+}
+
+void BlenderSync::sync_pointcloud(PointCloud *pointcloud, BObjectInfo &b_ob_info)
+{
+  size_t old_numpoints = pointcloud->num_points();
+
+  array<Node *> used_shaders = pointcloud->get_used_shaders();
+
+  PointCloud new_pointcloud;
+  new_pointcloud.set_used_shaders(used_shaders);
+
+  /* TODO: add option to filter out points in the view layer. */
+  BL::PointCloud b_pointcloud(b_ob_info.object_data);
+  export_pointcloud(scene, &new_pointcloud, b_pointcloud);
+
+  /* update original sockets */
+  for (const SocketType &socket : new_pointcloud.type->inputs) {
+    /* Those sockets are updated in sync_object, so do not modify them. */
+    if (socket.name == "use_motion_blur" || socket.name == "motion_steps" ||
+        socket.name == "used_shaders") {
+      continue;
+    }
+    pointcloud->set_value(socket, new_pointcloud, socket);
+  }
+
+  pointcloud->attributes.clear();
+  foreach (Attribute &attr, new_pointcloud.attributes.attributes) {
+    pointcloud->attributes.attributes.push_back(std::move(attr));
+  }
+
+  /* tag update */
+  const bool rebuild = (pointcloud && old_numpoints != pointcloud->num_points());
+  pointcloud->tag_update(scene, rebuild);
+}
+
+void BlenderSync::sync_pointcloud_motion(PointCloud *pointcloud,
+                                         BObjectInfo &b_ob_info,
+                                         int motion_step)
+{
+  /* Skip if nothing exported. */
+  if (pointcloud->num_points() == 0) {
+    return;
+  }
+
+  /* Export deformed coordinates. */
+  if (ccl::BKE_object_is_deform_modified(b_ob_info, b_scene, preview)) {
+    /* PointCloud object. */
+    BL::PointCloud b_pointcloud(b_ob_info.object_data);
+    export_pointcloud_motion(pointcloud, b_pointcloud, motion_step);
+  }
+  else {
+    /* No deformation on this frame, copy coordinates if other frames did have it. */
+    pointcloud->copy_center_to_motion_step(motion_step);
+  }
+}
+
+CCL_NAMESPACE_END
diff --git a/intern/cycles/blender/sync.h b/intern/cycles/blender/sync.h
index 2db4125ba1c..d074f90bb1b 100644
--- a/intern/cycles/blender/sync.h
+++ b/intern/cycles/blender/sync.h
@@ -169,12 +169,16 @@ class BlenderSync {
       Hair *hair, BL::Mesh &b_mesh, BObjectInfo &b_ob_info, bool motion, int motion_step = 0);
   bool object_has_particle_hair(BL::Object b_ob);
 
+  /* Point Cloud */
+  void sync_pointcloud(PointCloud *pointcloud, BObjectInfo &b_ob_info);
+  void sync_pointcloud_motion(PointCloud *pointcloud, BObjectInfo &b_ob_info, int motion_step = 0);
+
   /* Camera */
   void sync_camera_motion(
       BL::RenderSettings &b_render, BL::Object &b_ob, int width, int height, float motion_time);
 
   /* Geometry */
-  Geometry *sync_geometry(BL::Depsgraph &b_depsgrpah,
+  Geometry *sync_geometry(BL::Depsgraph &b_depsgraph,
                           BObjectInfo &b_ob_info,
                           bool object_updated,
                           bool use_particle_hair,
diff --git a/intern/cycles/bvh/build.cpp b/intern/cycles/bvh/build.cpp
index 3ce268dfb25..91198e4e2a2 100644
--- a/intern/cycles/bvh/build.cpp
+++ b/intern/cycles/bvh/build.cpp
@@ -26,6 +26,7 @@
 #include "scene/hair.h"
 #include "scene/mesh.h"
 #include "scene/object.h"
+#include "scene/pointcloud.h"
 #include "scene/scene.h"
 
 #include "util/algorithm.h"
@@ -113,9 +114,9 @@ void BVHBuild::add_reference_triangles(BoundBox &root,
     else {
       /* Motion triangles, trace optimized case:  we split triangle
        * primitives into separate nodes for each of the time steps.
-       * This way we minimize overlap of neighbor curve primitives.
+       * This way we minimize overlap of neighbor triangle primitives.
        */
-      const int num_bvh_steps = params.num_motion_curve_steps * 2 + 1;
+      const int num_bvh_steps = params.num_motion_triangle_steps * 2 + 1;
       const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1);
       const size_t num_verts = mesh->verts.size();
       const size_t num_steps = mesh->motion_steps;
@@ -269,6 +270,101 @@ void BVHBuild::add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair
   }
 }
 
+void BVHBuild::add_reference_points(BoundBox &root,
+                                    BoundBox &center,
+                                    PointCloud *pointcloud,
+                                    int i)
+{
+  const Attribute *point_attr_mP = NULL;
+  if (pointcloud->has_motion_blur()) {
+    point_attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+  }
+
+  const float3 *points_data = &pointcloud->points[0];
+  const float *radius_data = &pointcloud->radius[0];
+  const size_t num_points = pointcloud->num_points();
+  const float3 *motion_data = (point_attr_mP) ? point_attr_mP->data_float3() : NULL;
+  const size_t num_steps = pointcloud->get_motion_steps();
+
+  if (point_attr_mP == NULL) {
+    /* Really simple logic for static points. */
+    for (uint j = 0; j < num_points; j++) {
+      const PointCloud::Point point = pointcloud->get_point(j);
+      BoundBox bounds = BoundBox::empty;
+      point.bounds_grow(points_data, radius_data, bounds);
+      if (bounds.valid()) {
+        references.push_back(BVHReference(bounds, j, i, PRIMITIVE_POINT));
+        root.grow(bounds);
+        center.grow(bounds.center2());
+      }
+    }
+  }
+  else if (params.num_motion_point_steps == 0 || params.use_spatial_split) {
+    /* Simple case of motion points: single node for the whole
+     * shutter time. Lowest memory usage but less optimal
+     * rendering.
+     */
+    /* TODO(sergey): Support motion steps for spatially split BVH. */
+    for (uint j = 0; j < num_points; j++) {
+      const PointCloud::Point point = pointcloud->get_point(j);
+      BoundBox bounds = BoundBox::empty;
+      point.bounds_grow(points_data, radius_data, bounds);
+      for (size_t step = 0; step < num_steps - 1; step++) {
+        point.bounds_grow(motion_data + step * num_points, radius_data, bounds);
+      }
+      if (bounds.valid()) {
+        references.push_back(BVHReference(bounds, j, i, PRIMITIVE_MOTION_POINT));
+        root.grow(bounds);
+        center.grow(bounds.center2());
+      }
+    }
+  }
+  else {
+    /* Motion points, trace optimized case:  we split point
+     * primitives into separate nodes for each of the time steps.
+     * This way we minimize overlap of neighbor point primitives.
+     */
+    const int num_bvh_steps = params.num_motion_point_steps * 2 + 1;
+    const float num_bvh_steps_inv_1 = 1.0f / (num_bvh_steps - 1);
+
+    for (uint j = 0; j < num_points; j++) {
+      const PointCloud::Point point = pointcloud->get_point(j);
+      const size_t num_steps = pointcloud->get_motion_steps();
+      const float3 *point_steps = point_attr_mP->data_float3();
+
+      /* Calculate bounding box of the previous time step.
+       * Will be reused later to avoid duplicated work on
+       * calculating BVH time step boundbox.
+       */
+      float4 prev_key = point.motion_key(
+          points_data, radius_data, point_steps, num_points, num_steps, 0.0f, j);
+      BoundBox prev_bounds = BoundBox::empty;
+      point.bounds_grow(prev_key, prev_bounds);
+      /* Create all primitive time steps, */
+      for (int bvh_step = 1; bvh_step < num_bvh_steps; ++bvh_step) {
+        const float curr_time = (float)(bvh_step)*num_bvh_steps_inv_1;
+        float4 curr_key = point.motion_key(
+            points_data, radius_data, point_steps, num_points, num_steps, curr_time, j);
+        BoundBox curr_bounds = BoundBox::empty;
+        point.bounds_grow(curr_key, curr_bounds);
+        BoundBox bounds = prev_bounds;
+        bounds.grow(curr_bounds);
+        if (bounds.valid()) {
+          const float prev_time = (float)(bvh_step - 1) * num_bvh_steps_inv_1;
+          references.push_back(
+              BVHReference(bounds, j, i, PRIMITIVE_MOTION_POINT, prev_time, curr_time));
+          root.grow(bounds);
+          center.grow(bounds.center2());
+        }
+        /* Current time boundbox becomes previous one for the
+         * next time step.
+         */
+        prev_bounds = curr_bounds;
+      }
+    }
+  }
+}
+
 void BVHBuild::add_reference_geometry(BoundBox &root,
                                       BoundBox &center,
                                       Geometry *geom,
@@ -282,6 +378,10 @@ void BVHBuild::add_reference_geometry(BoundBox &root,
     Hair *hair = static_cast<Hair *>(geom);
     add_reference_curves(root, center, hair, object_index);
   }
+  else if (geom->geometry_type == Geometry::POINTCLOUD) {
+    PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+    add_reference_points(root, center, pointcloud, object_index);
+  }
 }
 
 void BVHBuild::add_reference_object(BoundBox &root, BoundBox &center, Object *ob, int i)
@@ -311,6 +411,10 @@ static size_t count_primitives(Geometry *geom)
     Hair *hair = static_cast<Hair *>(geom);
     return count_curve_segments(hair);
   }
+  else if (geom->geometry_type == Geometry::POINTCLOUD) {
+    PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+    return pointcloud->num_points();
+  }
 
   return 0;
 }
@@ -328,8 +432,9 @@ void BVHBuild::add_references(BVHRange &root)
       if (!ob->get_geometry()->is_instanced()) {
         num_alloc_references += count_primitives(ob->get_geometry());
       }
-      else
+      else {
         num_alloc_references++;
+      }
     }
     else {
       num_alloc_references += count_primitives(ob->get_geometry());
@@ -394,7 +499,7 @@ BVHNode *BVHBuild::run()
   spatial_min_overlap = root.bounds().safe_area() * params.spatial_split_alpha;
   spatial_free_index = 0;
 
-  need_prim_time = params.num_motion_curve_steps > 0 || params.num_motion_triangle_steps > 0;
+  need_prim_time = params.use_motion_steps();
 
   /* init progress updates */
   double build_start_time;
@@ -535,7 +640,8 @@ bool BVHBuild::range_within_max_leaf_size(const BVHRange &range,
                                           const vector<BVHReference> &references) const
 {
   size_t size = range.size();
-  size_t max_leaf_size = max(params.max_triangle_leaf_size, params.max_curve_leaf_size);
+  size_t max_leaf_size = max(max(params.max_triangle_leaf_size, params.max_curve_leaf_size),
+                             params.max_point_leaf_size);
 
   if (size > max_leaf_size)
     return false;
@@ -544,6 +650,8 @@ bool BVHBuild::range_within_max_leaf_size(const BVHRange &range,
   size_t num_motion_triangles = 0;
   size_t num_curves = 0;
   size_t num_motion_curves = 0;
+  size_t num_points = 0;
+  size_t num_motion_points = 0;
 
   for (int i = 0; i < size; i++) {
     const BVHReference &ref = references[range.start() + i];
@@ -564,12 +672,22 @@ bool BVHBuild::range_within_max_leaf_size(const BVHRange &range,
         num_triangles++;
       }
     }
+    else if (ref.prim_type() & PRIMITIVE_ALL_POINT) {
+      if (ref.prim_type() & PRIMITIVE_ALL_MOTION) {
+        num_motion_points++;
+      }
+      else {
+        num_points++;
+      }
+    }
   }
 
   return (num_triangles <= params.max_triangle_leaf_size) &&
          (num_motion_triangles <= params.max_motion_triangle_leaf_size) &&
          (num_curves <= params.max_curve_leaf_size) &&
-         (num_motion_curves <= params.max_motion_curve_leaf_size);
+         (num_motion_curves <= params.max_motion_curve_leaf_size) &&
+         (num_points <= params.max_point_leaf_size) &&
+         (num_motion_points <= params.max_motion_point_leaf_size);
 }
 
 /* multithreaded binning builder */
diff --git a/intern/cycles/bvh/build.h b/intern/cycles/bvh/build.h
index 06b318f1ee0..5b9bb59d9f8 100644
--- a/intern/cycles/bvh/build.h
+++ b/intern/cycles/bvh/build.h
@@ -39,6 +39,7 @@ class Geometry;
 class Hair;
 class Mesh;
 class Object;
+class PointCloud;
 class Progress;
 
 /* BVH Builder */
@@ -68,6 +69,7 @@ class BVHBuild {
   /* Adding references. */
   void add_reference_triangles(BoundBox &root, BoundBox &center, Mesh *mesh, int i);
   void add_reference_curves(BoundBox &root, BoundBox &center, Hair *hair, int i);
+  void add_reference_points(BoundBox &root, BoundBox &center, PointCloud *pointcloud, int i);
   void add_reference_geometry(BoundBox &root, BoundBox &center, Geometry *geom, int i);
   void add_reference_object(BoundBox &root, BoundBox &center, Object *ob, int i);
   void add_references(BVHRange &root);
diff --git a/intern/cycles/bvh/bvh2.cpp b/intern/cycles/bvh/bvh2.cpp
index 04290602145..744e7fa9898 100644
--- a/intern/cycles/bvh/bvh2.cpp
+++ b/intern/cycles/bvh/bvh2.cpp
@@ -20,6 +20,7 @@
 #include "scene/hair.h"
 #include "scene/mesh.h"
 #include "scene/object.h"
+#include "scene/pointcloud.h"
 
 #include "bvh/build.h"
 #include "bvh/node.h"
@@ -409,6 +410,30 @@ void BVH2::refit_primitives(int start, int end, BoundBox &bbox, uint &visibility
           }
         }
       }
+      else if (pack.prim_type[prim] & PRIMITIVE_ALL_POINT) {
+        /* Points. */
+        const PointCloud *pointcloud = static_cast<const PointCloud *>(ob->get_geometry());
+        int prim_offset = (params.top_level) ? pointcloud->prim_offset : 0;
+        const float3 *points = &pointcloud->points[0];
+        const float *radius = &pointcloud->radius[0];
+        PointCloud::Point point = pointcloud->get_point(pidx - prim_offset);
+
+        point.bounds_grow(points, radius, bbox);
+
+        /* Motion points. */
+        if (pointcloud->get_use_motion_blur()) {
+          Attribute *attr = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+
+          if (attr) {
+            size_t pointcloud_size = pointcloud->points.size();
+            size_t steps = pointcloud->get_motion_steps() - 1;
+            float3 *point_steps = attr->data_float3();
+
+            for (size_t i = 0; i < steps; i++)
+              point.bounds_grow(point_steps + i * pointcloud_size, radius, bbox);
+          }
+        }
+      }
       else {
         /* Triangles. */
         const Mesh *mesh = static_cast<const Mesh *>(ob->get_geometry());
@@ -505,7 +530,8 @@ void BVH2::pack_instances(size_t nodes_size, size_t leaf_nodes_size)
   pack.leaf_nodes.resize(leaf_nodes_size);
   pack.object_node.resize(objects.size());
 
-  if (params.num_motion_curve_steps > 0 || params.num_motion_triangle_steps > 0) {
+  if (params.num_motion_curve_steps > 0 || params.num_motion_triangle_steps > 0 ||
+      params.num_motion_point_steps > 0) {
     pack.prim_time.resize(prim_index_size);
   }
 
diff --git a/intern/cycles/bvh/embree.cpp b/intern/cycles/bvh/embree.cpp
index b54b38f2798..eab193f45cb 100644
--- a/intern/cycles/bvh/embree.cpp
+++ b/intern/cycles/bvh/embree.cpp
@@ -45,6 +45,7 @@
 #  include "scene/hair.h"
 #  include "scene/mesh.h"
 #  include "scene/object.h"
+#  include "scene/pointcloud.h"
 
 #  include "util/foreach.h"
 #  include "util/log.h"
@@ -245,7 +246,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
   }
 }
 
-static void rtc_filter_func_thick_curve(const RTCFilterFunctionNArguments *args)
+static void rtc_filter_func_backface_cull(const RTCFilterFunctionNArguments *args)
 {
   const RTCRay *ray = (RTCRay *)args->ray;
   RTCHit *hit = (RTCHit *)args->hit;
@@ -258,7 +259,7 @@ static void rtc_filter_func_thick_curve(const RTCFilterFunctionNArguments *args)
   }
 }
 
-static void rtc_filter_occluded_func_thick_curve(const RTCFilterFunctionNArguments *args)
+static void rtc_filter_occluded_func_backface_cull(const RTCFilterFunctionNArguments *args)
 {
   const RTCRay *ray = (RTCRay *)args->ray;
   RTCHit *hit = (RTCHit *)args->hit;
@@ -410,6 +411,12 @@ void BVHEmbree::add_object(Object *ob, int i)
       add_curves(ob, hair, i);
     }
   }
+  else if (geom->geometry_type == Geometry::POINTCLOUD) {
+    PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+    if (pointcloud->num_points() > 0) {
+      add_points(ob, pointcloud, i);
+    }
+  }
 }
 
 void BVHEmbree::add_instance(Object *ob, int i)
@@ -624,6 +631,89 @@ void BVHEmbree::set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, c
   }
 }
 
+void BVHEmbree::set_point_vertex_buffer(RTCGeometry geom_id,
+                                        const PointCloud *pointcloud,
+                                        const bool update)
+{
+  const Attribute *attr_mP = NULL;
+  size_t num_motion_steps = 1;
+  if (pointcloud->has_motion_blur()) {
+    attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+    if (attr_mP) {
+      num_motion_steps = pointcloud->get_motion_steps();
+    }
+  }
+
+  const size_t num_points = pointcloud->num_points();
+
+  /* Copy the point data to Embree */
+  const int t_mid = (num_motion_steps - 1) / 2;
+  const float *radius = pointcloud->get_radius().data();
+  for (int t = 0; t < num_motion_steps; ++t) {
+    const float3 *verts;
+    if (t == t_mid || attr_mP == NULL) {
+      verts = pointcloud->get_points().data();
+    }
+    else {
+      int t_ = (t > t_mid) ? (t - 1) : t;
+      verts = &attr_mP->data_float3()[t_ * num_points];
+    }
+
+    float4 *rtc_verts = (update) ? (float4 *)rtcGetGeometryBufferData(
+                                       geom_id, RTC_BUFFER_TYPE_VERTEX, t) :
+                                   (float4 *)rtcSetNewGeometryBuffer(geom_id,
+                                                                     RTC_BUFFER_TYPE_VERTEX,
+                                                                     t,
+                                                                     RTC_FORMAT_FLOAT4,
+                                                                     sizeof(float) * 4,
+                                                                     num_points);
+
+    assert(rtc_verts);
+    if (rtc_verts) {
+      for (size_t j = 0; j < num_points; ++j) {
+        rtc_verts[j] = float3_to_float4(verts[j]);
+        rtc_verts[j].w = radius[j];
+      }
+    }
+
+    if (update) {
+      rtcUpdateGeometryBuffer(geom_id, RTC_BUFFER_TYPE_VERTEX, t);
+    }
+  }
+}
+
+void BVHEmbree::add_points(const Object *ob, const PointCloud *pointcloud, int i)
+{
+  size_t prim_offset = pointcloud->prim_offset;
+
+  const Attribute *attr_mP = NULL;
+  size_t num_motion_steps = 1;
+  if (pointcloud->has_motion_blur()) {
+    attr_mP = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+    if (attr_mP) {
+      num_motion_steps = pointcloud->get_motion_steps();
+    }
+  }
+
+  enum RTCGeometryType type = RTC_GEOMETRY_TYPE_SPHERE_POINT;
+
+  RTCGeometry geom_id = rtcNewGeometry(rtc_device, type);
+
+  rtcSetGeometryBuildQuality(geom_id, build_quality);
+  rtcSetGeometryTimeStepCount(geom_id, num_motion_steps);
+
+  set_point_vertex_buffer(geom_id, pointcloud, false);
+
+  rtcSetGeometryUserData(geom_id, (void *)prim_offset);
+  rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_func_backface_cull);
+  rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func_backface_cull);
+  rtcSetGeometryMask(geom_id, ob->visibility_for_tracing());
+
+  rtcCommitGeometry(geom_id);
+  rtcAttachGeometryByID(scene, geom_id, i * 2);
+  rtcReleaseGeometry(geom_id);
+}
+
 void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
 {
   size_t prim_offset = hair->curve_segment_offset;
@@ -678,8 +768,8 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
     rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func);
   }
   else {
-    rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_func_thick_curve);
-    rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func_thick_curve);
+    rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_func_backface_cull);
+    rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func_backface_cull);
   }
   rtcSetGeometryMask(geom_id, ob->visibility_for_tracing());
 
@@ -716,6 +806,14 @@ void BVHEmbree::refit(Progress &progress)
           rtcCommitGeometry(geom);
         }
       }
+      else if (geom->geometry_type == Geometry::POINTCLOUD) {
+        PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+        if (pointcloud->num_points() > 0) {
+          RTCGeometry geom = rtcGetGeometry(scene, geom_id);
+          set_point_vertex_buffer(geom, pointcloud, true);
+          rtcCommitGeometry(geom);
+        }
+      }
     }
     geom_id += 2;
   }
diff --git a/intern/cycles/bvh/embree.h b/intern/cycles/bvh/embree.h
index 746ca97b504..3b30b2bbcf7 100644
--- a/intern/cycles/bvh/embree.h
+++ b/intern/cycles/bvh/embree.h
@@ -33,6 +33,7 @@ CCL_NAMESPACE_BEGIN
 
 class Hair;
 class Mesh;
+class PointCloud;
 
 class BVHEmbree : public BVH {
  public:
@@ -51,11 +52,15 @@ class BVHEmbree : public BVH {
   void add_object(Object *ob, int i);
   void add_instance(Object *ob, int i);
   void add_curves(const Object *ob, const Hair *hair, int i);
+  void add_points(const Object *ob, const PointCloud *pointcloud, int i);
   void add_triangles(const Object *ob, const Mesh *mesh, int i);
 
  private:
   void set_tri_vertex_buffer(RTCGeometry geom_id, const Mesh *mesh, const bool update);
   void set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, const bool update);
+  void set_point_vertex_buffer(RTCGeometry geom_id,
+                               const PointCloud *pointcloud,
+                               const bool update);
 
   RTCDevice rtc_device;
   enum RTCBuildQuality build_quality;
diff --git a/intern/cycles/bvh/params.h b/intern/cycles/bvh/params.h
index 8f185a2640f..16edf2e88e4 100644
--- a/intern/cycles/bvh/params.h
+++ b/intern/cycles/bvh/params.h
@@ -83,6 +83,8 @@ class BVHParams {
   int max_motion_triangle_leaf_size;
   int max_curve_leaf_size;
   int max_motion_curve_leaf_size;
+  int max_point_leaf_size;
+  int max_motion_point_leaf_size;
 
   /* object or mesh level bvh */
   bool top_level;
@@ -98,13 +100,13 @@ class BVHParams {
   /* Split time range to this number of steps and create leaf node for each
    * of this time steps.
    *
-   * Speeds up rendering of motion curve primitives in the cost of higher
-   * memory usage.
+   * Speeds up rendering of motion primitives in the cost of higher memory usage.
    */
-  int num_motion_curve_steps;
 
   /* Same as above, but for triangle primitives. */
   int num_motion_triangle_steps;
+  int num_motion_curve_steps;
+  int num_motion_point_steps;
 
   /* Same as in SceneParams. */
   int bvh_type;
@@ -132,6 +134,8 @@ class BVHParams {
     max_motion_triangle_leaf_size = 8;
     max_curve_leaf_size = 1;
     max_motion_curve_leaf_size = 4;
+    max_point_leaf_size = 8;
+    max_motion_point_leaf_size = 8;
 
     top_level = false;
     bvh_layout = BVH_LAYOUT_BVH2;
@@ -139,6 +143,7 @@ class BVHParams {
 
     num_motion_curve_steps = 0;
     num_motion_triangle_steps = 0;
+    num_motion_point_steps = 0;
 
     bvh_type = 0;
 
@@ -166,6 +171,12 @@ class BVHParams {
     return (size <= min_leaf_size || level >= MAX_DEPTH);
   }
 
+  bool use_motion_steps()
+  {
+    return num_motion_curve_steps > 0 || num_motion_triangle_steps > 0 ||
+           num_motion_point_steps > 0;
+  }
+
   /* Gets best matching BVH.
    *
    * If the requested layout is supported by the device, it will be used.
diff --git a/intern/cycles/bvh/split.cpp b/intern/cycles/bvh/split.cpp
index 102c50e2979..34d12de97c0 100644
--- a/intern/cycles/bvh/split.cpp
+++ b/intern/cycles/bvh/split.cpp
@@ -23,6 +23,7 @@
 #include "scene/hair.h"
 #include "scene/mesh.h"
 #include "scene/object.h"
+#include "scene/pointcloud.h"
 
 #include "util/algorithm.h"
 
@@ -426,6 +427,32 @@ void BVHSpatialSplit::split_curve_primitive(const Hair *hair,
   }
 }
 
+void BVHSpatialSplit::split_point_primitive(const PointCloud *pointcloud,
+                                            const Transform *tfm,
+                                            int prim_index,
+                                            int dim,
+                                            float pos,
+                                            BoundBox &left_bounds,
+                                            BoundBox &right_bounds)
+{
+  /* No real splitting support for points, assume they are small enough for it
+   * not to matter. */
+  float3 point = pointcloud->get_points()[prim_index];
+
+  if (tfm != NULL) {
+    point = transform_point(tfm, point);
+  }
+  point = get_unaligned_point(point);
+
+  if (point[dim] <= pos) {
+    left_bounds.grow(point);
+  }
+
+  if (point[dim] >= pos) {
+    right_bounds.grow(point);
+  }
+}
+
 void BVHSpatialSplit::split_triangle_reference(const BVHReference &ref,
                                                const Mesh *mesh,
                                                int dim,
@@ -453,6 +480,16 @@ void BVHSpatialSplit::split_curve_reference(const BVHReference &ref,
                         right_bounds);
 }
 
+void BVHSpatialSplit::split_point_reference(const BVHReference &ref,
+                                            const PointCloud *pointcloud,
+                                            int dim,
+                                            float pos,
+                                            BoundBox &left_bounds,
+                                            BoundBox &right_bounds)
+{
+  split_point_primitive(pointcloud, NULL, ref.prim_index(), dim, pos, left_bounds, right_bounds);
+}
+
 void BVHSpatialSplit::split_object_reference(
     const Object *object, int dim, float pos, BoundBox &left_bounds, BoundBox &right_bounds)
 {
@@ -475,6 +512,13 @@ void BVHSpatialSplit::split_object_reference(
       }
     }
   }
+  else if (geom->geometry_type == Geometry::POINTCLOUD) {
+    PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+    for (int point_idx = 0; point_idx < pointcloud->num_points(); ++point_idx) {
+      split_point_primitive(
+          pointcloud, &object->get_tfm(), point_idx, dim, pos, left_bounds, right_bounds);
+    }
+  }
 }
 
 void BVHSpatialSplit::split_reference(const BVHBuild &builder,
@@ -499,6 +543,10 @@ void BVHSpatialSplit::split_reference(const BVHBuild &builder,
     Hair *hair = static_cast<Hair *>(ob->get_geometry());
     split_curve_reference(ref, hair, dim, pos, left_bounds, right_bounds);
   }
+  else if (ref.prim_type() & PRIMITIVE_ALL_POINT) {
+    PointCloud *pointcloud = static_cast<PointCloud *>(ob->get_geometry());
+    split_point_reference(ref, pointcloud, dim, pos, left_bounds, right_bounds);
+  }
   else {
     split_object_reference(ob, dim, pos, left_bounds, right_bounds);
   }
diff --git a/intern/cycles/bvh/split.h b/intern/cycles/bvh/split.h
index 2650a500ea9..92953c40040 100644
--- a/intern/cycles/bvh/split.h
+++ b/intern/cycles/bvh/split.h
@@ -26,6 +26,7 @@ CCL_NAMESPACE_BEGIN
 class BVHBuild;
 class Hair;
 class Mesh;
+class PointCloud;
 struct Transform;
 
 /* Object Split */
@@ -123,6 +124,13 @@ class BVHSpatialSplit {
                              float pos,
                              BoundBox &left_bounds,
                              BoundBox &right_bounds);
+  void split_point_primitive(const PointCloud *pointcloud,
+                             const Transform *tfm,
+                             int prim_index,
+                             int dim,
+                             float pos,
+                             BoundBox &left_bounds,
+                             BoundBox &right_bounds);
 
   /* Lower-level functions which calculates boundaries of left and right nodes
    * needed for spatial split.
@@ -141,6 +149,12 @@ class BVHSpatialSplit {
                              float pos,
                              BoundBox &left_bounds,
                              BoundBox &right_bounds);
+  void split_point_reference(const BVHReference &ref,
+                             const PointCloud *pointcloud,
+                             int dim,
+                             float pos,
+                             BoundBox &left_bounds,
+                             BoundBox &right_bounds);
   void split_object_reference(
       const Object *object, int dim, float pos, BoundBox &left_bounds, BoundBox &right_bounds);
 
diff --git a/intern/cycles/device/optix/device_impl.cpp b/intern/cycles/device/optix/device_impl.cpp
index 11786f4a6e2..38cc3330ebd 100644
--- a/intern/cycles/device/optix/device_impl.cpp
+++ b/intern/cycles/device/optix/device_impl.cpp
@@ -28,6 +28,7 @@
 #  include "scene/mesh.h"
 #  include "scene/object.h"
 #  include "scene/pass.h"
+#  include "scene/pointcloud.h"
 #  include "scene/scene.h"
 
 #  include "util/debug.h"
@@ -242,6 +243,9 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
     else
       pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_CUSTOM;
   }
+  if (kernel_features & KERNEL_FEATURE_POINTCLOUD) {
+    pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_CUSTOM;
+  }
 
   /* Keep track of whether motion blur is enabled, so to enable/disable motion in BVH builds
    * This is necessary since objects may be reported to have motion if the Vector pass is
@@ -372,6 +376,18 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
     }
   }
 
+  /* Pointclouds */
+  if (kernel_features & KERNEL_FEATURE_POINTCLOUD) {
+    group_descs[PG_HITD_POINTCLOUD] = group_descs[PG_HITD];
+    group_descs[PG_HITD_POINTCLOUD].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
+    group_descs[PG_HITD_POINTCLOUD].hitgroup.moduleIS = optix_module;
+    group_descs[PG_HITD_POINTCLOUD].hitgroup.entryFunctionNameIS = "__intersection__point";
+    group_descs[PG_HITS_POINTCLOUD] = group_descs[PG_HITS];
+    group_descs[PG_HITS_POINTCLOUD].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
+    group_descs[PG_HITS_POINTCLOUD].hitgroup.moduleIS = optix_module;
+    group_descs[PG_HITS_POINTCLOUD].hitgroup.entryFunctionNameIS = "__intersection__point";
+  }
+
   if (kernel_features & (KERNEL_FEATURE_SUBSURFACE | KERNEL_FEATURE_NODE_RAYTRACE)) {
     /* Add hit group for local intersections. */
     group_descs[PG_HITL].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
@@ -419,6 +435,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
                        stack_size[PG_HITD_MOTION].cssIS + stack_size[PG_HITD_MOTION].cssAH);
   trace_css = std::max(trace_css,
                        stack_size[PG_HITS_MOTION].cssIS + stack_size[PG_HITS_MOTION].cssAH);
+  trace_css = std::max(
+      trace_css, stack_size[PG_HITD_POINTCLOUD].cssIS + stack_size[PG_HITD_POINTCLOUD].cssAH);
+  trace_css = std::max(
+      trace_css, stack_size[PG_HITS_POINTCLOUD].cssIS + stack_size[PG_HITS_POINTCLOUD].cssAH);
 
   OptixPipelineLinkOptions link_options = {};
   link_options.maxTraceDepth = 1;
@@ -444,6 +464,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
       pipeline_groups.push_back(groups[PG_HITD_MOTION]);
       pipeline_groups.push_back(groups[PG_HITS_MOTION]);
     }
+    if (kernel_features & KERNEL_FEATURE_POINTCLOUD) {
+      pipeline_groups.push_back(groups[PG_HITD_POINTCLOUD]);
+      pipeline_groups.push_back(groups[PG_HITS_POINTCLOUD]);
+    }
     pipeline_groups.push_back(groups[PG_CALL_SVM_AO]);
     pipeline_groups.push_back(groups[PG_CALL_SVM_BEVEL]);
 
@@ -483,6 +507,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
       pipeline_groups.push_back(groups[PG_HITD_MOTION]);
       pipeline_groups.push_back(groups[PG_HITS_MOTION]);
     }
+    if (kernel_features & KERNEL_FEATURE_POINTCLOUD) {
+      pipeline_groups.push_back(groups[PG_HITD_POINTCLOUD]);
+      pipeline_groups.push_back(groups[PG_HITS_POINTCLOUD]);
+    }
 
     optix_assert(optixPipelineCreate(context,
                                      &pipeline_options,
@@ -1378,6 +1406,86 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         progress.set_error("Failed to build OptiX acceleration structure");
       }
     }
+    else if (geom->geometry_type == Geometry::POINTCLOUD) {
+      /* Build BLAS for points primitives. */
+      PointCloud *const pointcloud = static_cast<PointCloud *const>(geom);
+      const size_t num_points = pointcloud->num_points();
+      if (num_points == 0) {
+        return;
+      }
+
+      size_t num_motion_steps = 1;
+      Attribute *motion_points = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+      if (motion_blur && pointcloud->get_use_motion_blur() && motion_points) {
+        num_motion_steps = pointcloud->get_motion_steps();
+      }
+
+      device_vector<OptixAabb> aabb_data(this, "optix temp aabb data", MEM_READ_ONLY);
+      aabb_data.alloc(num_points * num_motion_steps);
+
+      /* Get AABBs for each motion step. */
+      for (size_t step = 0; step < num_motion_steps; ++step) {
+        /* The center step for motion vertices is not stored in the attribute. */
+        const float3 *points = pointcloud->get_points().data();
+        const float *radius = pointcloud->get_radius().data();
+        size_t center_step = (num_motion_steps - 1) / 2;
+        if (step != center_step) {
+          size_t attr_offset = (step > center_step) ? step - 1 : step;
+          /* Technically this is a float4 array, but sizeof(float3) == sizeof(float4). */
+          points = motion_points->data_float3() + attr_offset * num_points;
+        }
+
+        for (size_t i = 0; i < num_points; ++i) {
+          const PointCloud::Point point = pointcloud->get_point(i);
+          BoundBox bounds = BoundBox::empty;
+          point.bounds_grow(points, radius, bounds);
+
+          const size_t index = step * num_points + i;
+          aabb_data[index].minX = bounds.min.x;
+          aabb_data[index].minY = bounds.min.y;
+          aabb_data[index].minZ = bounds.min.z;
+          aabb_data[index].maxX = bounds.max.x;
+          aabb_data[index].maxY = bounds.max.y;
+          aabb_data[index].maxZ = bounds.max.z;
+        }
+      }
+
+      /* Upload AABB data to GPU. */
+      aabb_data.copy_to_device();
+
+      vector<device_ptr> aabb_ptrs;
+      aabb_ptrs.reserve(num_motion_steps);
+      for (size_t step = 0; step < num_motion_steps; ++step) {
+        aabb_ptrs.push_back(aabb_data.device_pointer + step * num_points * sizeof(OptixAabb));
+      }
+
+      /* Disable visibility test any-hit program, since it is already checked during
+       * intersection. Those trace calls that require anyhit can force it with a ray flag.
+       * For those, force a single any-hit call, so shadow record-all behavior works correctly. */
+      unsigned int build_flags = OPTIX_GEOMETRY_FLAG_DISABLE_ANYHIT |
+                                 OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;
+      OptixBuildInput build_input = {};
+      build_input.type = OPTIX_BUILD_INPUT_TYPE_CUSTOM_PRIMITIVES;
+#  if OPTIX_ABI_VERSION < 23
+      build_input.aabbArray.aabbBuffers = (CUdeviceptr *)aabb_ptrs.data();
+      build_input.aabbArray.numPrimitives = num_points;
+      build_input.aabbArray.strideInBytes = sizeof(OptixAabb);
+      build_input.aabbArray.flags = &build_flags;
+      build_input.aabbArray.numSbtRecords = 1;
+      build_input.aabbArray.primitiveIndexOffset = pointcloud->prim_offset;
+#  else
+      build_input.customPrimitiveArray.aabbBuffers = (CUdeviceptr *)aabb_ptrs.data();
+      build_input.customPrimitiveArray.numPrimitives = num_points;
+      build_input.customPrimitiveArray.strideInBytes = sizeof(OptixAabb);
+      build_input.customPrimitiveArray.flags = &build_flags;
+      build_input.customPrimitiveArray.numSbtRecords = 1;
+      build_input.customPrimitiveArray.primitiveIndexOffset = pointcloud->prim_offset;
+#  endif
+
+      if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
+        progress.set_error("Failed to build OptiX acceleration structure");
+      }
+    }
   }
   else {
     unsigned int num_instances = 0;
@@ -1461,12 +1569,22 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
           instance.sbtOffset = PG_HITD_MOTION - PG_HITD;
         }
       }
+      else if (ob->get_geometry()->geometry_type == Geometry::POINTCLOUD) {
+        /* Use the hit group that has an intersection program for point clouds. */
+        instance.sbtOffset = PG_HITD_POINTCLOUD - PG_HITD;
+
+        /* Also skip point clouds in local trace calls. */
+        instance.visibilityMask |= 4;
+      }
+
 #  if OPTIX_ABI_VERSION < 55
       /* Cannot disable any-hit program for thick curves, since it needs to filter out end-caps. */
       else
 #  endif
       {
-        /* Can disable __anyhit__kernel_optix_visibility_test by default.
+        /* Can disable __anyhit__kernel_optix_visibility_test by default (except for thick curves,
+         * since it needs to filter out end-caps there).
+
          * It is enabled where necessary (visibility mask exceeds 8 bits or the other any-hit
          * programs like __anyhit__kernel_optix_shadow_all_hit) via OPTIX_RAY_FLAG_ENFORCE_ANYHIT.
          */
diff --git a/intern/cycles/device/optix/device_impl.h b/intern/cycles/device/optix/device_impl.h
index cf4afb46280..25073c60e69 100644
--- a/intern/cycles/device/optix/device_impl.h
+++ b/intern/cycles/device/optix/device_impl.h
@@ -44,6 +44,8 @@ enum {
   PG_HITV, /* __VOLUME__ hit group. */
   PG_HITD_MOTION,
   PG_HITS_MOTION,
+  PG_HITD_POINTCLOUD,
+  PG_HITS_POINTCLOUD,
   PG_CALL_SVM_AO,
   PG_CALL_SVM_BEVEL,
   NUM_PROGRAM_GROUPS
@@ -52,9 +54,9 @@ enum {
 static const int MISS_PROGRAM_GROUP_OFFSET = PG_MISS;
 static const int NUM_MIS_PROGRAM_GROUPS = 1;
 static const int HIT_PROGAM_GROUP_OFFSET = PG_HITD;
-static const int NUM_HIT_PROGRAM_GROUPS = 6;
+static const int NUM_HIT_PROGRAM_GROUPS = 8;
 static const int CALLABLE_PROGRAM_GROUPS_BASE = PG_CALL_SVM_AO;
-static const int NUM_CALLABLE_PROGRAM_GROUPS = 3;
+static const int NUM_CALLABLE_PROGRAM_GROUPS = 2;
 
 /* List of OptiX pipelines. */
 enum { PIP_SHADE_RAYTRACE, PIP_INTERSECT, NUM_PIPELINES };
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. */
diff --git a/intern/cycles/scene/CMakeLists.txt b/intern/cycles/scene/CMakeLists.txt
index a3fde99306b..6f8b97a4998 100644
--- a/intern/cycles/scene/CMakeLists.txt
+++ b/intern/cycles/scene/CMakeLists.txt
@@ -40,6 +40,7 @@ set(SRC
   mesh_displace.cpp
   mesh_subdivision.cpp
   procedural.cpp
+  pointcloud.cpp
   object.cpp
   osl.cpp
   particles.cpp
@@ -81,6 +82,7 @@ set(SRC_HEADERS
   particles.h
   pass.h
   procedural.h
+  pointcloud.h
   curves.h
   scene.h
   shader.h
diff --git a/intern/cycles/scene/alembic.cpp b/intern/cycles/scene/alembic.cpp
index 39b5f467736..71fa1863b8b 100644
--- a/intern/cycles/scene/alembic.cpp
+++ b/intern/cycles/scene/alembic.cpp
@@ -21,6 +21,7 @@
 #include "scene/curves.h"
 #include "scene/mesh.h"
 #include "scene/object.h"
+#include "scene/pointcloud.h"
 #include "scene/scene.h"
 #include "scene/shader.h"
 
@@ -99,6 +100,9 @@ void CachedData::clear()
   triangles.clear();
   uv_loops.clear();
   vertices.clear();
+  points.clear();
+  radiuses.clear();
+  points_shader.clear();
 
   for (CachedAttribute &attr : attributes) {
     attr.data.clear();
@@ -146,6 +150,9 @@ bool CachedData::is_constant() const
   CHECK_IF_CONSTANT(triangles)
   CHECK_IF_CONSTANT(uv_loops)
   CHECK_IF_CONSTANT(vertices)
+  CHECK_IF_CONSTANT(points)
+  CHECK_IF_CONSTANT(radiuses)
+  CHECK_IF_CONSTANT(points_shader)
 
   for (const CachedAttribute &attr : attributes) {
     if (!attr.data.is_constant()) {
@@ -185,6 +192,9 @@ void CachedData::invalidate_last_loaded_time(bool attributes_only)
   triangles.invalidate_last_loaded_time();
   uv_loops.invalidate_last_loaded_time();
   vertices.invalidate_last_loaded_time();
+  points.invalidate_last_loaded_time();
+  radiuses.invalidate_last_loaded_time();
+  points_shader.invalidate_last_loaded_time();
 }
 
 void CachedData::set_time_sampling(TimeSampling time_sampling)
@@ -206,6 +216,9 @@ void CachedData::set_time_sampling(TimeSampling time_sampling)
   triangles.set_time_sampling(time_sampling);
   uv_loops.set_time_sampling(time_sampling);
   vertices.set_time_sampling(time_sampling);
+  points.set_time_sampling(time_sampling);
+  radiuses.set_time_sampling(time_sampling);
+  points_shader.set_time_sampling(time_sampling);
 
   for (CachedAttribute &attr : attributes) {
     attr.data.set_time_sampling(time_sampling);
@@ -233,6 +246,9 @@ size_t CachedData::memory_used() const
   mem_used += triangles.memory_used();
   mem_used += uv_loops.memory_used();
   mem_used += vertices.memory_used();
+  mem_used += points.memory_used();
+  mem_used += radiuses.memory_used();
+  mem_used += points_shader.memory_used();
 
   for (const CachedAttribute &attr : attributes) {
     mem_used += attr.data.memory_used();
@@ -901,6 +917,9 @@ void AlembicProcedural::generate(Scene *scene, Progress &progress)
     else if (object->schema_type == AlembicObject::CURVES) {
       read_curves(object, frame_time);
     }
+    else if (object->schema_type == AlembicObject::POINTS) {
+      read_points(object, frame_time);
+    }
     else if (object->schema_type == AlembicObject::SUBD) {
       read_subd(object, frame_time);
     }
@@ -970,6 +989,9 @@ void AlembicProcedural::load_objects(Progress &progress)
       if (abc_object->schema_type == AlembicObject::CURVES) {
         geometry = scene_->create_node<Hair>();
       }
+      else if (abc_object->schema_type == AlembicObject::POINTS) {
+        geometry = scene_->create_node<PointCloud>();
+      }
       else if (abc_object->schema_type == AlembicObject::POLY_MESH ||
                abc_object->schema_type == AlembicObject::SUBD) {
         geometry = scene_->create_node<Mesh>();
@@ -1202,6 +1224,45 @@ void AlembicProcedural::read_curves(AlembicObject *abc_object, Abc::chrono_t fra
   hair->tag_update(scene_, rebuild);
 }
 
+void AlembicProcedural::read_points(AlembicObject *abc_object, Abc::chrono_t frame_time)
+{
+  CachedData &cached_data = abc_object->get_cached_data();
+
+  /* update sockets */
+
+  Object *object = abc_object->get_object();
+  cached_data.transforms.copy_to_socket(frame_time, object, object->get_tfm_socket());
+
+  if (object->is_modified()) {
+    object->tag_update(scene_);
+  }
+
+  /* Only update sockets for the original Geometry. */
+  if (abc_object->instance_of) {
+    return;
+  }
+
+  PointCloud *point_cloud = static_cast<PointCloud *>(object->get_geometry());
+
+  /* Make sure shader ids are also updated. */
+  if (point_cloud->used_shaders_is_modified()) {
+    point_cloud->tag_shader_modified();
+  }
+
+  cached_data.points.copy_to_socket(frame_time, point_cloud, point_cloud->get_points_socket());
+  cached_data.radiuses.copy_to_socket(frame_time, point_cloud, point_cloud->get_radius_socket());
+  cached_data.points_shader.copy_to_socket(
+      frame_time, point_cloud, point_cloud->get_shader_socket());
+
+  /* update attributes */
+
+  update_attributes(point_cloud->attributes, cached_data, frame_time);
+
+  const bool rebuild = (point_cloud->points_is_modified() || point_cloud->radius_is_modified() ||
+                        point_cloud->shader_is_modified());
+  point_cloud->tag_update(scene_, rebuild);
+}
+
 void AlembicProcedural::walk_hierarchy(
     IObject parent,
     const ObjectHeader &header,
@@ -1314,7 +1375,23 @@ void AlembicProcedural::walk_hierarchy(
     // ignore the face set, it will be read along with the data
   }
   else if (IPoints::matches(header)) {
-    // unsupported for now
+    IPoints points(parent, header.getName());
+
+    unordered_map<std::string, AlembicObject *>::const_iterator iter;
+    iter = object_map.find(points.getFullName());
+
+    if (iter != object_map.end()) {
+      AlembicObject *abc_object = iter->second;
+      abc_object->iobject = points;
+      abc_object->schema_type = AlembicObject::POINTS;
+
+      if (matrix_samples_data.samples) {
+        abc_object->xform_samples = *matrix_samples_data.samples;
+        abc_object->xform_time_sampling = matrix_samples_data.time_sampling;
+      }
+    }
+
+    next_object = points;
   }
   else if (INuPatch::matches(header)) {
     // unsupported for now
@@ -1391,6 +1468,14 @@ void AlembicProcedural::build_caches(Progress &progress)
         object->load_data_in_cache(object->get_cached_data(), this, schema, progress);
       }
     }
+    else if (object->schema_type == AlembicObject::POINTS) {
+      if (!object->has_data_loaded() || default_radius_is_modified() ||
+          object->radius_scale_is_modified()) {
+        IPoints points(object->iobject, Alembic::Abc::kWrapExisting);
+        IPointsSchema schema = points.getSchema();
+        object->load_data_in_cache(object->get_cached_data(), this, schema, progress);
+      }
+    }
     else if (object->schema_type == AlembicObject::SUBD) {
       if (!object->has_data_loaded()) {
         ISubD subd_mesh(object->iobject, Alembic::Abc::kWrapExisting);
diff --git a/intern/cycles/scene/alembic.h b/intern/cycles/scene/alembic.h
index 77aafd0ab32..3a4d37da3ff 100644
--- a/intern/cycles/scene/alembic.h
+++ b/intern/cycles/scene/alembic.h
@@ -327,6 +327,11 @@ struct CachedData {
   DataStore<array<int>> curve_first_key;
   DataStore<array<int>> curve_shader;
 
+  /* point data */
+  DataStore<array<float3>> points;
+  DataStore<array<float>> radiuses;
+  DataStore<array<int>> points_shader;
+
   struct CachedAttribute {
     AttributeStandard std;
     AttributeElement element;
@@ -414,6 +419,7 @@ class AlembicObject : public Node {
     POLY_MESH,
     SUBD,
     CURVES,
+    POINTS,
   };
 
   bool need_shader_update = true;
@@ -550,6 +556,10 @@ class AlembicProcedural : public Procedural {
    * Object Nodes in the Cycles scene if none exist yet. */
   void read_curves(AlembicObject *abc_object, Alembic::AbcGeom::Abc::chrono_t frame_time);
 
+  /* Read the data for an IPoints at the specified frame_time. Creates corresponding Geometry and
+   * Object Nodes in the Cycles scene if none exist yet. */
+  void read_points(AlembicObject *abc_object, Alembic::AbcGeom::Abc::chrono_t frame_time);
+
   /* Read the data for an ISubD at the specified frame_time. Creates corresponding Geometry and
    * Object Nodes in the Cycles scene if none exist yet. */
   void read_subd(AlembicObject *abc_object, Alembic::AbcGeom::Abc::chrono_t frame_time);
diff --git a/intern/cycles/scene/alembic_read.cpp b/intern/cycles/scene/alembic_read.cpp
index 35f4854127a..07874b7528e 100644
--- a/intern/cycles/scene/alembic_read.cpp
+++ b/intern/cycles/scene/alembic_read.cpp
@@ -609,6 +609,55 @@ void read_geometry_data(AlembicProcedural *proc,
   read_data_loop(proc, cached_data, data, read_curves_data, progress);
 }
 
+/* Points Geometries. */
+
+static void read_points_data(CachedData &cached_data, const PointsSchemaData &data, chrono_t time)
+{
+  const ISampleSelector iss = ISampleSelector(time);
+
+  const P3fArraySamplePtr position = data.positions.getValue(iss);
+  FloatArraySamplePtr radiuses;
+
+  array<float3> a_positions;
+  array<float> a_radius;
+  array<int> a_shader;
+  a_positions.reserve(position->size());
+  a_radius.reserve(position->size());
+  a_shader.reserve(position->size());
+
+  if (data.radiuses.valid()) {
+    IFloatGeomParam::Sample wsample = data.radiuses.getExpandedValue(iss);
+    radiuses = wsample.getVals();
+  }
+
+  const bool do_radius = (radiuses != nullptr) && (radiuses->size() > 1);
+  float radius = (radiuses && radiuses->size() == 1) ? (*radiuses)[0] : data.default_radius;
+
+  int offset = 0;
+  for (size_t i = 0; i < position->size(); i++) {
+    const V3f &f = position->get()[offset + i];
+    a_positions.push_back_slow(make_float3_from_yup(f));
+
+    if (do_radius) {
+      radius = (*radiuses)[offset + i];
+      a_radius.push_back_slow(radius);
+    }
+
+    a_shader.push_back_slow((int)0);
+  }
+
+  cached_data.points.add_data(a_positions, time);
+  cached_data.radiuses.add_data(a_radius, time);
+  cached_data.points_shader.add_data(a_shader, time);
+}
+
+void read_geometry_data(AlembicProcedural *proc,
+                        CachedData &cached_data,
+                        const PointsSchemaData &data,
+                        Progress &progress)
+{
+  read_data_loop(proc, cached_data, data, read_points_data, progress);
+}
 /* Attributes conversions. */
 
 /* Type traits for converting between Alembic and Cycles types.
diff --git a/intern/cycles/scene/alembic_read.h b/intern/cycles/scene/alembic_read.h
index 6b656b59481..e95750ce691 100644
--- a/intern/cycles/scene/alembic_read.h
+++ b/intern/cycles/scene/alembic_read.h
@@ -122,6 +122,26 @@ void read_geometry_data(AlembicProcedural *proc,
                         const CurvesSchemaData &data,
                         Progress &progress);
 
+/* Data of a ICurvesSchema that we need to read. */
+struct PointsSchemaData {
+  Alembic::AbcGeom::TimeSamplingPtr time_sampling;
+  size_t num_samples;
+
+  float default_radius;
+  float radius_scale;
+
+  Alembic::AbcGeom::IP3fArrayProperty positions;
+  Alembic::AbcGeom::IInt32ArrayProperty num_points;
+  Alembic::AbcGeom::IFloatGeomParam radiuses;
+  // Those are unsupported for now.
+  Alembic::AbcGeom::IV3fArrayProperty velocities;
+};
+
+void read_geometry_data(AlembicProcedural *proc,
+                        CachedData &cached_data,
+                        const PointsSchemaData &data,
+                        Progress &progress);
+
 void read_attributes(AlembicProcedural *proc,
                      CachedData &cache,
                      const Alembic::AbcGeom::ICompoundProperty &arb_geom_params,
diff --git a/intern/cycles/scene/attribute.cpp b/intern/cycles/scene/attribute.cpp
index 6d15f3325f7..f1021a26e8a 100644
--- a/intern/cycles/scene/attribute.cpp
+++ b/intern/cycles/scene/attribute.cpp
@@ -18,6 +18,7 @@
 #include "scene/hair.h"
 #include "scene/image.h"
 #include "scene/mesh.h"
+#include "scene/pointcloud.h"
 
 #include "util/foreach.h"
 #include "util/log.h"
@@ -205,6 +206,10 @@ size_t Attribute::element_size(Geometry *geom, AttributePrimitive prim) const
           size -= mesh->get_num_subd_verts();
         }
       }
+      else if (geom->geometry_type == Geometry::POINTCLOUD) {
+        PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+        size = pointcloud->num_points();
+      }
       break;
     case ATTR_ELEMENT_VERTEX_MOTION:
       if (geom->geometry_type == Geometry::MESH) {
@@ -215,6 +220,10 @@ size_t Attribute::element_size(Geometry *geom, AttributePrimitive prim) const
           size -= mesh->get_num_subd_verts() * (mesh->get_motion_steps() - 1);
         }
       }
+      else if (geom->geometry_type == Geometry::POINTCLOUD) {
+        PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+        size = pointcloud->num_points() * (pointcloud->get_motion_steps() - 1);
+      }
       break;
     case ATTR_ELEMENT_FACE:
       if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
@@ -346,6 +355,8 @@ const char *Attribute::standard_name(AttributeStandard std)
       return "curve_length";
     case ATTR_STD_CURVE_RANDOM:
       return "curve_random";
+    case ATTR_STD_POINT_RANDOM:
+      return "point_random";
     case ATTR_STD_PTEX_FACE_ID:
       return "ptex_face_id";
     case ATTR_STD_PTEX_UV:
@@ -555,6 +566,28 @@ Attribute *AttributeSet::add(AttributeStandard std, ustring name)
         break;
     }
   }
+  else if (geometry->geometry_type == Geometry::POINTCLOUD) {
+    switch (std) {
+      case ATTR_STD_UV:
+        attr = add(name, TypeFloat2, ATTR_ELEMENT_VERTEX);
+        break;
+      case ATTR_STD_GENERATED:
+        attr = add(name, TypeDesc::TypePoint, ATTR_ELEMENT_VERTEX);
+        break;
+      case ATTR_STD_MOTION_VERTEX_POSITION:
+        attr = add(name, TypeDesc::TypeFloat4, ATTR_ELEMENT_VERTEX_MOTION);
+        break;
+      case ATTR_STD_POINT_RANDOM:
+        attr = add(name, TypeDesc::TypeFloat, ATTR_ELEMENT_VERTEX);
+        break;
+      case ATTR_STD_GENERATED_TRANSFORM:
+        attr = add(name, TypeDesc::TypeMatrix, ATTR_ELEMENT_MESH);
+        break;
+      default:
+        assert(0);
+        break;
+    }
+  }
   else if (geometry->geometry_type == Geometry::VOLUME) {
     switch (std) {
       case ATTR_STD_VERTEX_NORMAL:
diff --git a/intern/cycles/scene/geometry.cpp b/intern/cycles/scene/geometry.cpp
index ca330c5544d..558b0e13b0f 100644
--- a/intern/cycles/scene/geometry.cpp
+++ b/intern/cycles/scene/geometry.cpp
@@ -26,6 +26,7 @@
 #include "scene/light.h"
 #include "scene/mesh.h"
 #include "scene/object.h"
+#include "scene/pointcloud.h"
 #include "scene/scene.h"
 #include "scene/shader.h"
 #include "scene/shader_nodes.h"
@@ -240,6 +241,7 @@ void Geometry::compute_bvh(
                                     params->use_bvh_unaligned_nodes;
       bparams.num_motion_triangle_steps = params->num_bvh_time_steps;
       bparams.num_motion_curve_steps = params->num_bvh_time_steps;
+      bparams.num_motion_point_steps = params->num_bvh_time_steps;
       bparams.bvh_type = params->bvh_type;
       bparams.curve_subdivisions = params->curve_subdivisions();
 
@@ -723,6 +725,12 @@ void GeometryManager::update_attribute_element_offset(Geometry *geom,
       else if (element == ATTR_ELEMENT_CURVE_KEY_MOTION)
         offset -= hair->curve_key_offset;
     }
+    else if (geom->is_pointcloud()) {
+      if (element == ATTR_ELEMENT_VERTEX)
+        offset -= geom->prim_offset;
+      else if (element == ATTR_ELEMENT_VERTEX_MOTION)
+        offset -= geom->prim_offset;
+    }
   }
   else {
     /* attribute not found */
@@ -1006,6 +1014,8 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout)
   size_t curve_key_size = 0;
   size_t curve_segment_size = 0;
 
+  size_t point_size = 0;
+
   size_t patch_size = 0;
   size_t face_size = 0;
   size_t corner_size = 0;
@@ -1054,6 +1064,14 @@ void GeometryManager::mesh_calc_offset(Scene *scene, BVHLayout bvh_layout)
       curve_key_size += hair->get_curve_keys().size();
       curve_segment_size += hair->num_segments();
     }
+    else if (geom->is_pointcloud()) {
+      PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+
+      prim_offset_changed = (pointcloud->prim_offset != point_size);
+
+      pointcloud->prim_offset = point_size;
+      point_size += pointcloud->num_points();
+    }
 
     if (prim_offset_changed) {
       /* Need to rebuild BVH in OptiX, since refit only allows modified mesh data there */
@@ -1079,6 +1097,8 @@ void GeometryManager::device_update_mesh(Device *,
   size_t curve_size = 0;
   size_t curve_segment_size = 0;
 
+  size_t point_size = 0;
+
   size_t patch_size = 0;
 
   foreach (Geometry *geom, scene->geometry) {
@@ -1106,6 +1126,10 @@ void GeometryManager::device_update_mesh(Device *,
       curve_size += hair->num_curves();
       curve_segment_size += hair->num_segments();
     }
+    else if (geom->is_pointcloud()) {
+      PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+      point_size += pointcloud->num_points();
+    }
   }
 
   /* Fill in all the arrays. */
@@ -1201,6 +1225,26 @@ void GeometryManager::device_update_mesh(Device *,
     dscene->curve_segments.copy_to_device_if_modified();
   }
 
+  if (point_size != 0) {
+    progress.set_status("Updating Mesh", "Copying Point clouds to device");
+
+    float4 *points = dscene->points.alloc(point_size);
+    uint *points_shader = dscene->points_shader.alloc(point_size);
+
+    foreach (Geometry *geom, scene->geometry) {
+      if (geom->is_pointcloud()) {
+        PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+        pointcloud->pack(
+            scene, &points[pointcloud->prim_offset], &points_shader[pointcloud->prim_offset]);
+        if (progress.get_cancel())
+          return;
+      }
+    }
+
+    dscene->points.copy_to_device();
+    dscene->points_shader.copy_to_device();
+  }
+
   if (patch_size != 0 && dscene->patches.need_realloc()) {
     progress.set_status("Updating Mesh", "Copying Patches to device");
 
@@ -1242,6 +1286,7 @@ void GeometryManager::device_update_bvh(Device *device,
                                 scene->params.use_bvh_unaligned_nodes;
   bparams.num_motion_triangle_steps = scene->params.num_bvh_time_steps;
   bparams.num_motion_curve_steps = scene->params.num_bvh_time_steps;
+  bparams.num_motion_point_steps = scene->params.num_bvh_time_steps;
   bparams.bvh_type = scene->params.bvh_type;
   bparams.curve_subdivisions = scene->params.curve_subdivisions();
 
@@ -1323,6 +1368,7 @@ void GeometryManager::device_update_bvh(Device *device,
 enum {
   DEVICE_CURVE_DATA_MODIFIED = (1 << 0),
   DEVICE_MESH_DATA_MODIFIED = (1 << 1),
+  DEVICE_POINT_DATA_MODIFIED = (1 << 2),
 
   ATTR_FLOAT_MODIFIED = (1 << 2),
   ATTR_FLOAT2_MODIFIED = (1 << 3),
@@ -1332,17 +1378,20 @@ enum {
 
   CURVE_DATA_NEED_REALLOC = (1 << 7),
   MESH_DATA_NEED_REALLOC = (1 << 8),
+  POINT_DATA_NEED_REALLOC = (1 << 9),
+
+  ATTR_FLOAT_NEEDS_REALLOC = (1 << 10),
+  ATTR_FLOAT2_NEEDS_REALLOC = (1 << 11),
+  ATTR_FLOAT3_NEEDS_REALLOC = (1 << 12),
+  ATTR_FLOAT4_NEEDS_REALLOC = (1 << 13),
 
-  ATTR_FLOAT_NEEDS_REALLOC = (1 << 9),
-  ATTR_FLOAT2_NEEDS_REALLOC = (1 << 10),
-  ATTR_FLOAT3_NEEDS_REALLOC = (1 << 11),
-  ATTR_FLOAT4_NEEDS_REALLOC = (1 << 12),
-  ATTR_UCHAR4_NEEDS_REALLOC = (1 << 13),
+  ATTR_UCHAR4_NEEDS_REALLOC = (1 << 14),
 
   ATTRS_NEED_REALLOC = (ATTR_FLOAT_NEEDS_REALLOC | ATTR_FLOAT2_NEEDS_REALLOC |
                         ATTR_FLOAT3_NEEDS_REALLOC | ATTR_FLOAT4_NEEDS_REALLOC |
                         ATTR_UCHAR4_NEEDS_REALLOC),
   DEVICE_MESH_DATA_NEEDS_REALLOC = (MESH_DATA_NEED_REALLOC | ATTRS_NEED_REALLOC),
+  DEVICE_POINT_DATA_NEEDS_REALLOC = (POINT_DATA_NEED_REALLOC | ATTRS_NEED_REALLOC),
   DEVICE_CURVE_DATA_NEEDS_REALLOC = (CURVE_DATA_NEED_REALLOC | ATTRS_NEED_REALLOC),
 };
 
@@ -1529,6 +1578,17 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
         device_update_flags |= DEVICE_MESH_DATA_MODIFIED;
       }
     }
+
+    if (geom->is_pointcloud()) {
+      PointCloud *pointcloud = static_cast<PointCloud *>(geom);
+
+      if (pointcloud->need_update_rebuild) {
+        device_update_flags |= DEVICE_POINT_DATA_NEEDS_REALLOC;
+      }
+      else if (pointcloud->is_modified()) {
+        device_update_flags |= DEVICE_POINT_DATA_MODIFIED;
+      }
+    }
   }
 
   if (update_flags & (MESH_ADDED | MESH_REMOVED)) {
@@ -1539,10 +1599,15 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
     device_update_flags |= DEVICE_CURVE_DATA_NEEDS_REALLOC;
   }
 
+  if (update_flags & (POINT_ADDED | POINT_REMOVED)) {
+    device_update_flags |= DEVICE_POINT_DATA_NEEDS_REALLOC;
+  }
+
   /* tag the device arrays for reallocation or modification */
   DeviceScene *dscene = &scene->dscene;
 
-  if (device_update_flags & (DEVICE_MESH_DATA_NEEDS_REALLOC | DEVICE_CURVE_DATA_NEEDS_REALLOC)) {
+  if (device_update_flags & (DEVICE_MESH_DATA_NEEDS_REALLOC | DEVICE_CURVE_DATA_NEEDS_REALLOC |
+                             DEVICE_POINT_DATA_NEEDS_REALLOC)) {
     delete scene->bvh;
     scene->bvh = nullptr;
 
@@ -1570,6 +1635,11 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
       dscene->curve_keys.tag_realloc();
       dscene->curve_segments.tag_realloc();
     }
+
+    if (device_update_flags & DEVICE_POINT_DATA_NEEDS_REALLOC) {
+      dscene->points.tag_realloc();
+      dscene->points_shader.tag_realloc();
+    }
   }
 
   if ((update_flags & VISIBILITY_MODIFIED) != 0) {
@@ -1630,6 +1700,11 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
     dscene->curve_segments.tag_modified();
   }
 
+  if (device_update_flags & DEVICE_POINT_DATA_MODIFIED) {
+    dscene->points.tag_modified();
+    dscene->points_shader.tag_modified();
+  }
+
   need_flags_update = false;
 }
 
@@ -2064,6 +2139,8 @@ void GeometryManager::device_update(Device *device,
   dscene->curves.clear_modified();
   dscene->curve_keys.clear_modified();
   dscene->curve_segments.clear_modified();
+  dscene->points.clear_modified();
+  dscene->points_shader.clear_modified();
   dscene->patches.clear_modified();
   dscene->attributes_map.clear_modified();
   dscene->attributes_float.clear_modified();
@@ -2092,6 +2169,8 @@ void GeometryManager::device_free(Device *device, DeviceScene *dscene, bool forc
   dscene->curves.free_if_need_realloc(force_free);
   dscene->curve_keys.free_if_need_realloc(force_free);
   dscene->curve_segments.free_if_need_realloc(force_free);
+  dscene->points.free_if_need_realloc(force_free);
+  dscene->points_shader.free_if_need_realloc(force_free);
   dscene->patches.free_if_need_realloc(force_free);
   dscene->attributes_map.free_if_need_realloc(force_free);
   dscene->attributes_float.free_if_need_realloc(force_free);
diff --git a/intern/cycles/scene/geometry.h b/intern/cycles/scene/geometry.h
index 91799d7fde8..b02387c3020 100644
--- a/intern/cycles/scene/geometry.h
+++ b/intern/cycles/scene/geometry.h
@@ -55,6 +55,7 @@ class Geometry : public Node {
     MESH,
     HAIR,
     VOLUME,
+    POINTCLOUD,
   };
 
   Type geometry_type;
@@ -155,6 +156,11 @@ class Geometry : public Node {
     return geometry_type == HAIR;
   }
 
+  bool is_pointcloud() const
+  {
+    return geometry_type == POINTCLOUD;
+  }
+
   bool is_volume() const
   {
     return geometry_type == VOLUME;
@@ -181,12 +187,14 @@ class GeometryManager {
     MESH_REMOVED = (1 << 5),
     HAIR_ADDED = (1 << 6),
     HAIR_REMOVED = (1 << 7),
+    POINT_ADDED = (1 << 12),
+    POINT_REMOVED = (1 << 13),
 
     SHADER_ATTRIBUTE_MODIFIED = (1 << 8),
     SHADER_DISPLACEMENT_MODIFIED = (1 << 9),
 
-    GEOMETRY_ADDED = MESH_ADDED | HAIR_ADDED,
-    GEOMETRY_REMOVED = MESH_REMOVED | HAIR_REMOVED,
+    GEOMETRY_ADDED = MESH_ADDED | HAIR_ADDED | POINT_ADDED,
+    GEOMETRY_REMOVED = MESH_REMOVED | HAIR_REMOVED | POINT_REMOVED,
 
     TRANSFORM_MODIFIED = (1 << 10),
 
diff --git a/intern/cycles/scene/object.cpp b/intern/cycles/scene/object.cpp
index 05f6f6ce197..f8110b20d6e 100644
--- a/intern/cycles/scene/object.cpp
+++ b/intern/cycles/scene/object.cpp
@@ -23,6 +23,7 @@
 #include "scene/light.h"
 #include "scene/mesh.h"
 #include "scene/particles.h"
+#include "scene/pointcloud.h"
 #include "scene/scene.h"
 #include "scene/stats.h"
 #include "scene/volume.h"
@@ -69,6 +70,7 @@ struct UpdateObjectTransformState {
   /* Flags which will be synchronized to Integrator. */
   bool have_motion;
   bool have_curves;
+  // bool have_points;
 
   /* ** Scheduling queue. ** */
   Scene *scene;
@@ -435,7 +437,7 @@ void ObjectManager::device_update_object_transform(UpdateObjectTransformState *s
     state->have_motion = true;
   }
 
-  if (geom->geometry_type == Geometry::MESH) {
+  if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::POINTCLOUD) {
     /* TODO: why only mesh? */
     Mesh *mesh = static_cast<Mesh *>(geom);
     if (mesh->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION)) {
@@ -491,6 +493,8 @@ void ObjectManager::device_update_object_transform(UpdateObjectTransformState *s
   kobject.dupli_generated[2] = ob->dupli_generated[2];
   kobject.numkeys = (geom->geometry_type == Geometry::HAIR) ?
                         static_cast<Hair *>(geom)->get_curve_keys().size() :
+                    (geom->geometry_type == Geometry::POINTCLOUD) ?
+                        static_cast<PointCloud *>(geom)->num_points() :
                         0;
   kobject.dupli_uv[0] = ob->dupli_uv[0];
   kobject.dupli_uv[1] = ob->dupli_uv[1];
diff --git a/intern/cycles/scene/pointcloud.cpp b/intern/cycles/scene/pointcloud.cpp
new file mode 100644
index 00000000000..4f88fe9db3d
--- /dev/null
+++ b/intern/cycles/scene/pointcloud.cpp
@@ -0,0 +1,304 @@
+/*
+ * Copyright 2011-2020 Blender Foundation
+ *
+ * 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.
+ */
+
+#include "bvh/bvh.h"
+
+#include "scene/pointcloud.h"
+#include "scene/scene.h"
+
+CCL_NAMESPACE_BEGIN
+
+/* PointCloud Point */
+
+void PointCloud::Point::bounds_grow(const float3 *points,
+                                    const float *radius,
+                                    BoundBox &bounds) const
+{
+  bounds.grow(points[index], radius[index]);
+}
+
+void PointCloud::Point::bounds_grow(const float3 *points,
+                                    const float *radius,
+                                    const Transform &aligned_space,
+                                    BoundBox &bounds) const
+{
+  float3 P = transform_point(&aligned_space, points[index]);
+  bounds.grow(P, radius[index]);
+}
+
+void PointCloud::Point::bounds_grow(const float4 &point, BoundBox &bounds) const
+{
+  bounds.grow(float4_to_float3(point), point.w);
+}
+
+float4 PointCloud::Point::motion_key(const float3 *points,
+                                     const float *radius,
+                                     const float3 *point_steps,
+                                     size_t num_points,
+                                     size_t num_steps,
+                                     float time,
+                                     size_t p) const
+{
+  /* Figure out which steps we need to fetch and their
+   * interpolation factor. */
+  const size_t max_step = num_steps - 1;
+  const size_t step = min((int)(time * max_step), max_step - 1);
+  const float t = time * max_step - step;
+  /* Fetch vertex coordinates. */
+  const float4 curr_key = point_for_step(
+      points, radius, point_steps, num_points, num_steps, step, p);
+  const float4 next_key = point_for_step(
+      points, radius, point_steps, num_points, num_steps, step + 1, p);
+  /* Interpolate between steps. */
+  return (1.0f - t) * curr_key + t * next_key;
+}
+
+float4 PointCloud::Point::point_for_step(const float3 *points,
+                                         const float *radius,
+                                         const float3 *point_steps,
+                                         size_t num_points,
+                                         size_t num_steps,
+                                         size_t step,
+                                         size_t p) const
+{
+  const size_t center_step = ((num_steps - 1) / 2);
+  if (step == center_step) {
+    /* Center step: regular key location. */
+    return make_float4(points[p].x, points[p].y, points[p].z, radius[p]);
+  }
+  else {
+    /* Center step is not stored in this array. */
+    if (step > center_step) {
+      step--;
+    }
+    const size_t offset = step * num_points;
+    return make_float4(point_steps[offset + p].x,
+                       point_steps[offset + p].y,
+                       point_steps[offset + p].z,
+                       radius[offset + p]);
+  }
+}
+
+/* PointCloud */
+
+NODE_DEFINE(PointCloud)
+{
+  NodeType *type = NodeType::add(
+      "pointcloud", create, NodeType::NONE, Geometry::get_node_base_type());
+
+  SOCKET_POINT_ARRAY(points, "Points", array<float3>());
+  SOCKET_FLOAT_ARRAY(radius, "Radius", array<float>());
+  SOCKET_INT_ARRAY(shader, "Shader", array<int>());
+
+  return type;
+}
+
+PointCloud::PointCloud() : Geometry(node_type, Geometry::POINTCLOUD)
+{
+}
+
+PointCloud::~PointCloud()
+{
+}
+
+void PointCloud::resize(int numpoints)
+{
+  points.resize(numpoints);
+  radius.resize(numpoints);
+  shader.resize(numpoints);
+  attributes.resize();
+
+  tag_points_modified();
+  tag_radius_modified();
+  tag_shader_modified();
+}
+
+void PointCloud::reserve(int numpoints)
+{
+  points.reserve(numpoints);
+  radius.reserve(numpoints);
+  shader.reserve(numpoints);
+  attributes.resize(true);
+}
+
+void PointCloud::clear(const bool preserve_shaders)
+{
+  Geometry::clear(preserve_shaders);
+
+  points.clear();
+  radius.clear();
+  shader.clear();
+  attributes.clear();
+
+  tag_points_modified();
+  tag_radius_modified();
+  tag_shader_modified();
+}
+
+void PointCloud::add_point(float3 co, float r, int shader_index)
+{
+  points.push_back_reserved(co);
+  radius.push_back_reserved(r);
+  shader.push_back_reserved(shader_index);
+
+  tag_points_modified();
+  tag_radius_modified();
+  tag_shader_modified();
+}
+
+void PointCloud::copy_center_to_motion_step(const int motion_step)
+{
+  Attribute *attr_mP = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+  if (attr_mP) {
+    float3 *points_data = points.data();
+    size_t numpoints = points.size();
+    memcpy(
+        attr_mP->data_float3() + motion_step * numpoints, points_data, sizeof(float3) * numpoints);
+  }
+}
+
+void PointCloud::get_uv_tiles(ustring map, unordered_set<int> &tiles)
+{
+  Attribute *attr;
+
+  if (map.empty()) {
+    attr = attributes.find(ATTR_STD_UV);
+  }
+  else {
+    attr = attributes.find(map);
+  }
+
+  if (attr) {
+    attr->get_uv_tiles(this, ATTR_PRIM_GEOMETRY, tiles);
+  }
+}
+
+void PointCloud::compute_bounds()
+{
+  BoundBox bnds = BoundBox::empty;
+  size_t numpoints = points.size();
+
+  if (numpoints > 0) {
+    for (size_t i = 0; i < numpoints; i++) {
+      bnds.grow(points[i], radius[i]);
+    }
+
+    Attribute *attr = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+    if (use_motion_blur && attr) {
+      size_t steps_size = points.size() * (motion_steps - 1);
+      float3 *point_steps = attr->data_float3();
+
+      for (size_t i = 0; i < steps_size; i++)
+        bnds.grow(point_steps[i]);
+    }
+
+    if (!bnds.valid()) {
+      bnds = BoundBox::empty;
+
+      /* skip nan or inf coordinates */
+      for (size_t i = 0; i < numpoints; i++)
+        bnds.grow_safe(points[i], radius[i]);
+
+      if (use_motion_blur && attr) {
+        size_t steps_size = points.size() * (motion_steps - 1);
+        float3 *point_steps = attr->data_float3();
+
+        for (size_t i = 0; i < steps_size; i++)
+          bnds.grow_safe(point_steps[i]);
+      }
+    }
+  }
+
+  if (!bnds.valid()) {
+    /* empty mesh */
+    bnds.grow(make_float3(0.0f, 0.0f, 0.0f));
+  }
+
+  bounds = bnds;
+}
+
+void PointCloud::apply_transform(const Transform &tfm, const bool apply_to_motion)
+{
+  /* compute uniform scale */
+  float3 c0 = transform_get_column(&tfm, 0);
+  float3 c1 = transform_get_column(&tfm, 1);
+  float3 c2 = transform_get_column(&tfm, 2);
+  float scalar = powf(fabsf(dot(cross(c0, c1), c2)), 1.0f / 3.0f);
+
+  /* apply transform to curve keys */
+  for (size_t i = 0; i < points.size(); i++) {
+    float3 co = transform_point(&tfm, points[i]);
+    float r = radius[i] * scalar;
+
+    /* scale for curve radius is only correct for uniform scale
+     */
+    points[i] = co;
+    radius[i] = r;
+  }
+
+  if (apply_to_motion) {
+    Attribute *attr = attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
+
+    if (attr) {
+      /* apply transform to motion curve keys */
+      size_t steps_size = points.size() * (motion_steps - 1);
+      float4 *point_steps = attr->data_float4();
+
+      for (size_t i = 0; i < steps_size; i++) {
+        float3 co = transform_point(&tfm, float4_to_float3(point_steps[i]));
+        float radius = point_steps[i].w * scalar;
+
+        /* scale for curve radius is only correct for uniform
+         * scale */
+        point_steps[i] = float3_to_float4(co);
+        point_steps[i].w = radius;
+      }
+    }
+  }
+}
+
+void PointCloud::pack(Scene *scene, float4 *packed_points, uint *packed_shader)
+{
+  size_t numpoints = points.size();
+  float3 *points_data = points.data();
+  float *radius_data = radius.data();
+  int *shader_data = shader.data();
+
+  for (size_t i = 0; i < numpoints; i++) {
+    packed_points[i] = make_float4(
+        points_data[i].x, points_data[i].y, points_data[i].z, radius_data[i]);
+  }
+
+  uint shader_id = 0;
+  uint last_shader = -1;
+  for (size_t i = 0; i < numpoints; i++) {
+    if (last_shader != shader_data[i]) {
+      last_shader = shader_data[i];
+      Shader *shader = (last_shader < used_shaders.size()) ?
+                           static_cast<Shader *>(used_shaders[last_shader]) :
+                           scene->default_surface;
+      shader_id = scene->shader_manager->get_shader_id(shader);
+    }
+    packed_shader[i] = shader_id;
+  }
+}
+
+PrimitiveType PointCloud::primitive_type() const
+{
+  return has_motion_blur() ? PRIMITIVE_MOTION_POINT : PRIMITIVE_POINT;
+}
+
+CCL_NAMESPACE_END
diff --git a/intern/cycles/scene/pointcloud.h b/intern/cycles/scene/pointcloud.h
new file mode 100644
index 00000000000..08987ee6c44
--- /dev/null
+++ b/intern/cycles/scene/pointcloud.h
@@ -0,0 +1,114 @@
+/*
+ * Copyright 2011-2020 Blender Foundation
+ *
+ * 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
+
+#ifndef __POINTCLOUD_H__
+#  define __POINTCLOUD_H__
+
+#  include "scene/geometry.h"
+
+CCL_NAMESPACE_BEGIN
+
+class PointCloud : public Geometry {
+ public:
+  NODE_DECLARE
+
+  /* PointCloud Point */
+  struct Point {
+    int index;
+
+    void bounds_grow(const float3 *points, const float *radius, BoundBox &bounds) const;
+    void bounds_grow(const float3 *points,
+                     const float *radius,
+                     const Transform &aligned_space,
+                     BoundBox &bounds) const;
+    void bounds_grow(const float4 &point, BoundBox &bounds) const;
+
+    float4 motion_key(const float3 *points,
+                      const float *radius,
+                      const float3 *point_steps,
+                      size_t num_points,
+                      size_t num_steps,
+                      float time,
+                      size_t p) const;
+    float4 point_for_step(const float3 *points,
+                          const float *radius,
+                          const float3 *point_steps,
+                          size_t num_points,
+                          size_t num_steps,
+                          size_t step,
+                          size_t p) const;
+  };
+
+  NODE_SOCKET_API_ARRAY(array<float3>, points)
+  NODE_SOCKET_API_ARRAY(array<float>, radius)
+  NODE_SOCKET_API_ARRAY(array<int>, shader)
+
+  /* Constructor/Destructor */
+  PointCloud();
+  ~PointCloud();
+
+  /* Geometry */
+  void clear(const bool preserver_shaders = false) override;
+
+  void resize(int numpoints);
+  void reserve(int numpoints);
+  void add_point(float3 loc, float radius, int shader = 0);
+
+  void copy_center_to_motion_step(const int motion_step);
+
+  void compute_bounds() override;
+  void apply_transform(const Transform &tfm, const bool apply_to_motion) override;
+
+  /* Points */
+  Point get_point(int i) const
+  {
+    Point point = {i};
+    return point;
+  }
+
+  size_t num_points() const
+  {
+    return points.size();
+  }
+
+  size_t num_attributes() const
+  {
+    return 1;
+  }
+
+  /* UDIM */
+  void get_uv_tiles(ustring map, unordered_set<int> &tiles) override;
+
+  PrimitiveType primitive_type() const override;
+
+  /* BVH */
+  void pack(Scene *scene, float4 *packed_points, uint *packed_shader);
+
+ private:
+  friend class BVH2;
+  friend class BVHBuild;
+  friend class BVHSpatialSplit;
+  friend class DiagSplit;
+  friend class EdgeDice;
+  friend class GeometryManager;
+  friend class ObjectManager;
+};
+
+CCL_NAMESPACE_END
+
+#endif /* __POINTCLOUD_H__ */
diff --git a/intern/cycles/scene/scene.cpp b/intern/cycles/scene/scene.cpp
index 452b5215836..1963ebbbb19 100644
--- a/intern/cycles/scene/scene.cpp
+++ b/intern/cycles/scene/scene.cpp
@@ -30,6 +30,7 @@
 #include "scene/object.h"
 #include "scene/osl.h"
 #include "scene/particles.h"
+#include "scene/pointcloud.h"
 #include "scene/procedural.h"
 #include "scene/scene.h"
 #include "scene/shader.h"
@@ -64,6 +65,8 @@ DeviceScene::DeviceScene(Device *device)
       curve_keys(device, "__curve_keys", MEM_GLOBAL),
       curve_segments(device, "__curve_segments", MEM_GLOBAL),
       patches(device, "__patches", MEM_GLOBAL),
+      points(device, "__points", MEM_GLOBAL),
+      points_shader(device, "__points_shader", MEM_GLOBAL),
       objects(device, "__objects", MEM_GLOBAL),
       object_motion_pass(device, "__object_motion_pass", MEM_GLOBAL),
       object_motion(device, "__object_motion", MEM_GLOBAL),
@@ -523,6 +526,9 @@ void Scene::update_kernel_features()
     else if (geom->is_hair()) {
       kernel_features |= KERNEL_FEATURE_HAIR;
     }
+    else if (geom->is_pointcloud()) {
+      kernel_features |= KERNEL_FEATURE_POINTCLOUD;
+    }
   }
 
   if (bake_manager->get_baking()) {
@@ -575,6 +581,7 @@ static void log_kernel_features(const uint features)
   VLOG(2) << "Use Path Tracing " << string_from_bool(features & KERNEL_FEATURE_PATH_TRACING)
           << "\n";
   VLOG(2) << "Use Hair " << string_from_bool(features & KERNEL_FEATURE_HAIR) << "\n";
+  VLOG(2) << "Use Pointclouds " << string_from_bool(features & KERNEL_FEATURE_POINTCLOUD) << "\n";
   VLOG(2) << "Use Object Motion " << string_from_bool(features & KERNEL_FEATURE_OBJECT_MOTION)
           << "\n";
   VLOG(2) << "Use Camera Motion " << string_from_bool(features & KERNEL_FEATURE_CAMERA_MOTION)
@@ -757,6 +764,15 @@ template<> Volume *Scene::create_node<Volume>()
   return node;
 }
 
+template<> PointCloud *Scene::create_node<PointCloud>()
+{
+  PointCloud *node = new PointCloud();
+  node->set_owner(this);
+  geometry.push_back(node);
+  geometry_manager->tag_update(this, GeometryManager::POINT_ADDED);
+  return node;
+}
+
 template<> Object *Scene::create_node<Object>()
 {
   Object *node = new Object();
@@ -844,6 +860,12 @@ template<> void Scene::delete_node_impl(Volume *node)
   geometry_manager->tag_update(this, GeometryManager::MESH_REMOVED);
 }
 
+template<> void Scene::delete_node_impl(PointCloud *node)
+{
+  delete_node_from_array(geometry, static_cast<Geometry *>(node));
+  geometry_manager->tag_update(this, GeometryManager::POINT_REMOVED);
+}
+
 template<> void Scene::delete_node_impl(Geometry *node)
 {
   uint flag;
diff --git a/intern/cycles/scene/scene.h b/intern/cycles/scene/scene.h
index f8f672a079a..ec935b41be6 100644
--- a/intern/cycles/scene/scene.h
+++ b/intern/cycles/scene/scene.h
@@ -54,6 +54,7 @@ class Object;
 class ObjectManager;
 class ParticleSystemManager;
 class ParticleSystem;
+class PointCloud;
 class Procedural;
 class ProceduralManager;
 class CurveSystemManager;
@@ -94,6 +95,10 @@ class DeviceScene {
 
   device_vector<uint> patches;
 
+  /* pointcloud */
+  device_vector<float4> points;
+  device_vector<uint> points_shader;
+
   /* objects */
   device_vector<KernelObject> objects;
   device_vector<Transform> object_motion_pass;
@@ -365,6 +370,8 @@ template<> Hair *Scene::create_node<Hair>();
 
 template<> Volume *Scene::create_node<Volume>();
 
+template<> PointCloud *Scene::create_node<PointCloud>();
+
 template<> ParticleSystem *Scene::create_node<ParticleSystem>();
 
 template<> Shader *Scene::create_node<Shader>();
@@ -379,6 +386,8 @@ template<> void Scene::delete_node_impl(Mesh *node);
 
 template<> void Scene::delete_node_impl(Volume *node);
 
+template<> void Scene::delete_node_impl(PointCloud *node);
+
 template<> void Scene::delete_node_impl(Hair *node);
 
 template<> void Scene::delete_node_impl(Geometry *node);