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,
@@ -1374,6 +1402,86 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
       build_input.triangleArray.numSbtRecords = 1;
       build_input.triangleArray.primitiveIndexOffset = mesh->prim_offset;
 
+      if (!build_optix_bvh(bvh_optix, operation, build_input, num_motion_steps)) {
+        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");
       }
@@ -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
@@ -47,6 +47,11 @@ ccl_device_inline float primitive_surface_attribute_float(KernelGlobals kg,
   else if (sd->type & PRIMITIVE_ALL_CURVE) {
     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)
@@ -73,6 +78,11 @@ ccl_device_inline float2 primitive_surface_attribute_float2(KernelGlobals kg,
   else if (sd->type & PRIMITIVE_ALL_CURVE) {
     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)
@@ -99,6 +109,11 @@ ccl_device_inline float3 primitive_surface_attribute_float3(KernelGlobals kg,
   else if (sd->type & PRIMITIVE_ALL_CURVE) {
     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)
@@ -125,6 +140,11 @@ ccl_device_forceinline float4 primitive_surface_attribute_float4(KernelGlobals k
   else if (sd->type & PRIMITIVE_ALL_CURVE) {
     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)
@@ -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
@@ -74,6 +74,13 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
     curve_shader_setup(kg, sd, ray->P, ray->D, isect->t, isect->object, isect->prim);
   }
   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 */
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 << 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_FLOAT_NEEDS_REALLOC = (1 << 10),
+  ATTR_FLOAT2_NEEDS_REALLOC = (1 << 11),
+  ATTR_FLOAT3_NEEDS_REALLOC = (1 << 12),
+  ATTR_FLOAT4_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);
diff --git a/tests/python/CMakeLists.txt b/tests/python/CMakeLists.txt
index a2c2d44fe20..7f221e1c4eb 100644
--- a/tests/python/CMakeLists.txt
+++ b/tests/python/CMakeLists.txt
@@ -649,6 +649,7 @@ if(WITH_CYCLES OR WITH_OPENGL_RENDER_TESTS)
       integrator
       light
       mesh
+      pointcloud
       shader
       shadow_catcher
       sss