diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index d566c4df421..abad4d7bbca 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -157,6 +157,7 @@ add_library(cudf
     src/column/column.cu
     src/column/column_device_view.cu
     src/column/column_factories.cpp
+    src/column/column_factories.cu
     src/column/column_view.cpp
     src/comms/ipc/ipc.cpp
     src/copying/concatenate.cu
@@ -375,7 +376,7 @@ add_library(cudf
     src/structs/copying/concatenate.cu
     src/structs/structs_column_factories.cu
     src/structs/structs_column_view.cpp
-    src/structs/utilities.cu
+    src/structs/utilities.cpp
     src/table/table.cpp
     src/table/table_device_view.cu
     src/table/table_view.cpp
diff --git a/cpp/include/cudf/scalar/scalar.hpp b/cpp/include/cudf/scalar/scalar.hpp
index 3de8762c763..3025c01d747 100644
--- a/cpp/include/cudf/scalar/scalar.hpp
+++ b/cpp/include/cudf/scalar/scalar.hpp
@@ -16,6 +16,7 @@
 #pragma once
 
 #include <cudf/column/column.hpp>
+#include <cudf/table/table.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/traits.hpp>
 
@@ -572,5 +573,59 @@ class list_scalar : public scalar {
   cudf::column _data;
 };
 
+/**
+ * @brief An owning class to represent a struct value in device memory
+ */
+class struct_scalar : public scalar {
+ public:
+  struct_scalar();
+  ~struct_scalar()                          = default;
+  struct_scalar(struct_scalar&& other)      = default;
+  struct_scalar(struct_scalar const& other) = default;
+  struct_scalar& operator=(struct_scalar const& other) = delete;
+  struct_scalar& operator=(struct_scalar&& other) = delete;
+
+  /**
+   * @brief Construct a new struct scalar object from table_view
+   *
+   * The input table_view is deep-copied.
+   *
+   * @param data The table data to copy.
+   * @param is_valid Whether the value held by the scalar is valid
+   * @param stream CUDA stream used for device memory operations.
+   * @param mr Device memory resource to use for device memory allocation
+   */
+  struct_scalar(table_view const& data,
+                bool is_valid                       = true,
+                rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+                rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+  /**
+   * @brief Construct a new struct scalar object from a host_span of column_views
+   *
+   * The input column_views are deep-copied.
+   *
+   * @param data The column_views to copy.
+   * @param is_valid Whether the value held by the scalar is valid
+   * @param stream CUDA stream used for device memory operations.
+   * @param mr Device memory resource to use for device memory allocation
+   */
+  struct_scalar(host_span<column_view const> data,
+                bool is_valid                       = true,
+                rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+                rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+  /**
+   * @brief Returns a non-owning, immutable view to underlying device data
+   */
+  table_view view() const;
+
+ private:
+  table _data;
+
+  void init(bool is_valid, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr);
+  void superimpose_nulls(rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr);
+};
+
 /** @} */  // end of group
 }  // namespace cudf
diff --git a/cpp/include/cudf/scalar/scalar_factories.hpp b/cpp/include/cudf/scalar/scalar_factories.hpp
index 8e10e122571..b96a8c65a04 100644
--- a/cpp/include/cudf/scalar/scalar_factories.hpp
+++ b/cpp/include/cudf/scalar/scalar_factories.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -168,5 +168,33 @@ std::unique_ptr<scalar> make_list_scalar(
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+/**
+ * @brief Construct a struct scalar using the given table_view.
+ *
+ * The columns must have 1 row.
+ *
+ * @param data The columnar data to store in the scalar object
+ * @param stream CUDA stream used for device memory operations.
+ * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool.
+ */
+std::unique_ptr<scalar> make_struct_scalar(
+  table_view const& data,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+/**
+ * @brief Construct a struct scalar using the given span of column views.
+ *
+ * The columns must have 1 row.
+ *
+ * @param value The columnar data to store in the scalar object
+ * @param stream CUDA stream used for device memory operations.
+ * @param mr Device memory resource used to allocate the scalar's `data` and `is_valid` bool.
+ */
+std::unique_ptr<scalar> make_struct_scalar(
+  host_span<column_view const> data,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 /** @} */  // end of group
 }  // namespace cudf
diff --git a/cpp/include/cudf/types.hpp b/cpp/include/cudf/types.hpp
index b08fccc0d66..ddba575cb07 100644
--- a/cpp/include/cudf/types.hpp
+++ b/cpp/include/cudf/types.hpp
@@ -61,6 +61,7 @@ class scalar;
 
 // clang-format off
 class list_scalar;
+class struct_scalar;
 class string_scalar;
 template <typename T> class numeric_scalar;
 template <typename T> class fixed_point_scalar;
@@ -74,8 +75,6 @@ template <typename T> class timestamp_scalar_device_view;
 template <typename T> class duration_scalar_device_view;
 // clang-format on
 
-class struct_scalar;
-
 class table;
 class table_view;
 class mutable_table_view;
diff --git a/cpp/src/column/column_factories.cpp b/cpp/src/column/column_factories.cpp
index 03339c2e0a8..86059a72e8f 100644
--- a/cpp/src/column/column_factories.cpp
+++ b/cpp/src/column/column_factories.cpp
@@ -20,7 +20,6 @@
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/dictionary/dictionary_factories.hpp>
 #include <cudf/fixed_point/fixed_point.hpp>
-#include <cudf/scalar/scalar.hpp>
 #include <cudf/scalar/scalar_factories.hpp>
 #include <cudf/strings/detail/fill.hpp>
 #include <cudf/utilities/error.hpp>
@@ -159,90 +158,6 @@ std::unique_ptr<column> make_fixed_width_column(data_type type,
   /// clang-format on
 }
 
-struct column_from_scalar_dispatch {
-  template <typename T>
-  std::unique_ptr<cudf::column> operator()(scalar const& value,
-                                           size_type size,
-                                           rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr) const
-  {
-    if (!value.is_valid())
-      return make_fixed_width_column(value.type(), size, mask_state::ALL_NULL, stream, mr);
-    auto output_column =
-      make_fixed_width_column(value.type(), size, mask_state::UNALLOCATED, stream, mr);
-    auto view = output_column->mutable_view();
-    detail::fill_in_place(view, 0, size, value, stream);
-    return output_column;
-  }
-};
-
-template <>
-std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::string_view>(
-  scalar const& value,
-  size_type size,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr) const
-{
-  auto null_mask = detail::create_null_mask(size, mask_state::ALL_NULL, stream, mr);
-
-  if (!value.is_valid())
-    return std::make_unique<column>(value.type(),
-                                    size,
-                                    rmm::device_buffer{0, stream, mr},
-                                    null_mask,
-                                    size);
-
-  // Create a strings column_view with all nulls and no children.
-  // Since we are setting every row to the scalar, the fill() never needs to access
-  // any of the children in the strings column which would otherwise cause an exception.
-  column_view sc{
-    data_type{type_id::STRING}, size, nullptr, static_cast<bitmask_type*>(null_mask.data()), size};
-  auto sv = static_cast<scalar_type_t<cudf::string_view> const&>(value);
-  // fill the column with the scalar
-  auto output = strings::detail::fill(strings_column_view(sc), 0, size, sv, stream, mr);
-  output->set_null_mask(rmm::device_buffer{0, stream, mr}, 0);  // should be no nulls
-  return output;
-}
-
-template <>
-std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::dictionary32>(
-  scalar const& value,
-  size_type size,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr) const
-{
-  CUDF_FAIL("dictionary not supported when creating from scalar");
-}
-
-template <>
-std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::list_view>(
-  scalar const& value,
-  size_type size,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr) const
-{
-  CUDF_FAIL("TODO");
-}
-
-template <>
-std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::struct_view>(
-  scalar const& value,
-  size_type size,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr) const
-{
-  CUDF_FAIL("TODO. struct_view currently not supported.");
-}
-
-std::unique_ptr<column> make_column_from_scalar(scalar const& s,
-                                                size_type size,
-                                                rmm::cuda_stream_view stream,
-                                                rmm::mr::device_memory_resource* mr)
-{
-  if (size == 0) return make_empty_column(s.type());
-  return type_dispatcher(s.type(), column_from_scalar_dispatch{}, s, size, stream, mr);
-}
-
 std::unique_ptr<column> make_dictionary_from_scalar(scalar const& s,
                                                     size_type size,
                                                     rmm::cuda_stream_view stream,
diff --git a/cpp/src/column/column_factories.cu b/cpp/src/column/column_factories.cu
new file mode 100644
index 00000000000..60e642ea3d5
--- /dev/null
+++ b/cpp/src/column/column_factories.cu
@@ -0,0 +1,124 @@
+/*
+ * Copyright (c) 2021, NVIDIA CORPORATION.
+ *
+ * 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 <cudf/column/column_factories.hpp>
+#include <cudf/detail/fill.hpp>
+#include <cudf/detail/gather.cuh>
+#include <cudf/dictionary/dictionary_factories.hpp>
+#include <cudf/scalar/scalar.hpp>
+#include <cudf/strings/detail/fill.hpp>
+
+namespace cudf {
+
+namespace {
+
+struct column_from_scalar_dispatch {
+  template <typename T>
+  std::unique_ptr<cudf::column> operator()(scalar const& value,
+                                           size_type size,
+                                           rmm::cuda_stream_view stream,
+                                           rmm::mr::device_memory_resource* mr) const
+  {
+    if (!value.is_valid())
+      return make_fixed_width_column(value.type(), size, mask_state::ALL_NULL, stream, mr);
+    auto output_column =
+      make_fixed_width_column(value.type(), size, mask_state::UNALLOCATED, stream, mr);
+    auto view = output_column->mutable_view();
+    detail::fill_in_place(view, 0, size, value, stream);
+    return output_column;
+  }
+};
+
+template <>
+std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::string_view>(
+  scalar const& value,
+  size_type size,
+  rmm::cuda_stream_view stream,
+  rmm::mr::device_memory_resource* mr) const
+{
+  auto null_mask = detail::create_null_mask(size, mask_state::ALL_NULL, stream, mr);
+
+  if (!value.is_valid())
+    return std::make_unique<column>(
+      value.type(), size, rmm::device_buffer{0, stream, mr}, null_mask, size);
+
+  // Create a strings column_view with all nulls and no children.
+  // Since we are setting every row to the scalar, the fill() never needs to access
+  // any of the children in the strings column which would otherwise cause an exception.
+  column_view sc{
+    data_type{type_id::STRING}, size, nullptr, static_cast<bitmask_type*>(null_mask.data()), size};
+  auto sv = static_cast<scalar_type_t<cudf::string_view> const&>(value);
+  // fill the column with the scalar
+  auto output = strings::detail::fill(strings_column_view(sc), 0, size, sv, stream, mr);
+  output->set_null_mask(rmm::device_buffer{0, stream, mr}, 0);  // should be no nulls
+  return output;
+}
+
+template <>
+std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::dictionary32>(
+  scalar const& value,
+  size_type size,
+  rmm::cuda_stream_view stream,
+  rmm::mr::device_memory_resource* mr) const
+{
+  CUDF_FAIL("dictionary not supported when creating from scalar");
+}
+
+template <>
+std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::list_view>(
+  scalar const& value,
+  size_type size,
+  rmm::cuda_stream_view stream,
+  rmm::mr::device_memory_resource* mr) const
+{
+  CUDF_FAIL("TODO");
+}
+
+template <>
+std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::struct_view>(
+  scalar const& value,
+  size_type size,
+  rmm::cuda_stream_view stream,
+  rmm::mr::device_memory_resource* mr) const
+{
+  auto ss   = static_cast<scalar_type_t<cudf::struct_view> const&>(value);
+  auto iter = thrust::make_constant_iterator(0);
+
+  auto children =
+    detail::gather(ss.view(), iter, iter + size, out_of_bounds_policy::NULLIFY, stream, mr);
+  auto const is_valid = ss.is_valid();
+  return make_structs_column(size,
+                             std::move(children->release()),
+                             is_valid ? 0 : size,
+                             is_valid
+                               ? rmm::device_buffer{}
+                               : detail::create_null_mask(size, mask_state::ALL_NULL, stream, mr),
+                             stream,
+                             mr);
+}
+
+}  // anonymous namespace
+
+std::unique_ptr<column> make_column_from_scalar(scalar const& s,
+                                                size_type size,
+                                                rmm::cuda_stream_view stream,
+                                                rmm::mr::device_memory_resource* mr)
+{
+  if (size == 0) return make_empty_column(s.type());
+  return type_dispatcher(s.type(), column_from_scalar_dispatch{}, s, size, stream, mr);
+}
+
+}  // namespace cudf
diff --git a/cpp/src/scalar/scalar.cpp b/cpp/src/scalar/scalar.cpp
index 858d2c063b3..f21b1c7ca20 100644
--- a/cpp/src/scalar/scalar.cpp
+++ b/cpp/src/scalar/scalar.cpp
@@ -14,7 +14,10 @@
  * limitations under the License.
  */
 
+#include <structs/utilities.hpp>
+
 #include <cudf/column/column.hpp>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/fixed_point/fixed_point.hpp>
 #include <cudf/scalar/scalar.hpp>
 #include <cudf/strings/string_view.hpp>
@@ -446,4 +449,53 @@ list_scalar::list_scalar(cudf::column&& data,
 
 column_view list_scalar::view() const { return _data.view(); }
 
+struct_scalar::struct_scalar() : scalar(data_type(type_id::STRUCT)) {}
+
+struct_scalar::struct_scalar(table_view const& data,
+                             bool is_valid,
+                             rmm::cuda_stream_view stream,
+                             rmm::mr::device_memory_resource* mr)
+  : scalar(data_type(type_id::STRUCT), is_valid, stream, mr), _data(data, stream, mr)
+{
+  init(is_valid, stream, mr);
+}
+
+struct_scalar::struct_scalar(host_span<column_view const> data,
+                             bool is_valid,
+                             rmm::cuda_stream_view stream,
+                             rmm::mr::device_memory_resource* mr)
+  : scalar(data_type(type_id::STRUCT), is_valid, stream, mr),
+    _data(table_view{std::vector<column_view>{data.begin(), data.end()}}, stream, mr)
+{
+  init(is_valid, stream, mr);
+}
+
+table_view struct_scalar::view() const { return _data.view(); }
+
+void struct_scalar::init(bool is_valid,
+                         rmm::cuda_stream_view stream,
+                         rmm::mr::device_memory_resource* mr)
+{
+  table_view tv = static_cast<table_view>(_data);
+  CUDF_EXPECTS(
+    std::all_of(tv.begin(), tv.end(), [](column_view const& col) { return col.size() == 1; }),
+    "Struct scalar inputs must have exactly 1 row");
+
+  // validity pushdown
+  if (!is_valid) { superimpose_nulls(stream, mr); }
+}
+
+void struct_scalar::superimpose_nulls(rmm::cuda_stream_view stream,
+                                      rmm::mr::device_memory_resource* mr)
+{
+  // push validity mask down
+  std::vector<bitmask_type> host_validity({0});
+  auto validity = cudf::detail::make_device_uvector_sync(host_validity, stream, mr);
+  auto iter     = thrust::make_counting_iterator(0);
+  std::for_each(iter, iter + _data.num_columns(), [&](size_type i) {
+    cudf::structs::detail::superimpose_parent_nulls(
+      validity.data(), 1, _data.get_column(i), stream, mr);
+  });
+}
+
 }  // namespace cudf
diff --git a/cpp/src/scalar/scalar_factories.cpp b/cpp/src/scalar/scalar_factories.cpp
index 7054e2b5c8f..e1d71b279d6 100644
--- a/cpp/src/scalar/scalar_factories.cpp
+++ b/cpp/src/scalar/scalar_factories.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -104,6 +104,20 @@ std::unique_ptr<scalar> make_list_scalar(column_view elements,
   return std::make_unique<list_scalar>(elements, true, stream, mr);
 }
 
+std::unique_ptr<scalar> make_struct_scalar(table_view const& data,
+                                           rmm::cuda_stream_view stream,
+                                           rmm::mr::device_memory_resource* mr)
+{
+  return std::make_unique<struct_scalar>(data, true, stream, mr);
+}
+
+std::unique_ptr<scalar> make_struct_scalar(host_span<column_view const> data,
+                                           rmm::cuda_stream_view stream,
+                                           rmm::mr::device_memory_resource* mr)
+{
+  return std::make_unique<struct_scalar>(data, true, stream, mr);
+}
+
 namespace {
 struct default_scalar_functor {
   template <typename T>
diff --git a/cpp/src/structs/structs_column_factories.cu b/cpp/src/structs/structs_column_factories.cu
index 330cecd1815..d8b94d1c448 100644
--- a/cpp/src/structs/structs_column_factories.cu
+++ b/cpp/src/structs/structs_column_factories.cu
@@ -14,10 +14,10 @@
  * limitations under the License.
  */
 
+#include <structs/utilities.hpp>
+
 #include <cudf/column/column_factories.hpp>
-#include <cudf/detail/null_mask.hpp>
 #include <cudf/types.hpp>
-
 #include <rmm/cuda_stream_view.hpp>
 
 #include <thrust/iterator/counting_iterator.h>
@@ -25,56 +25,6 @@
 #include <algorithm>
 #include <memory>
 namespace cudf {
-namespace {
-// Helper function to superimpose validity of parent struct
-// over the specified member (child) column.
-void superimpose_parent_nullmask(bitmask_type const* parent_null_mask,
-                                 std::size_t parent_null_mask_size,
-                                 size_type parent_null_count,
-                                 column& child,
-                                 rmm::cuda_stream_view stream,
-                                 rmm::mr::device_memory_resource* mr)
-{
-  if (!child.nullable()) {
-    // Child currently has no null mask. Copy parent's null mask.
-    child.set_null_mask(rmm::device_buffer{parent_null_mask, parent_null_mask_size, stream, mr});
-    child.set_null_count(parent_null_count);
-  } else {
-    // Child should have a null mask.
-    // `AND` the child's null mask with the parent's.
-
-    auto current_child_mask = child.mutable_view().null_mask();
-
-    std::vector<bitmask_type const*> masks{
-      reinterpret_cast<bitmask_type const*>(parent_null_mask),
-      reinterpret_cast<bitmask_type const*>(current_child_mask)};
-    std::vector<size_type> begin_bits{0, 0};
-    cudf::detail::inplace_bitmask_and(
-      device_span<bitmask_type>(current_child_mask, num_bitmask_words(child.size())),
-      masks,
-      begin_bits,
-      child.size(),
-      stream,
-      mr);
-    child.set_null_count(UNKNOWN_NULL_COUNT);
-  }
-
-  // If the child is also a struct, repeat for all grandchildren.
-  if (child.type().id() == cudf::type_id::STRUCT) {
-    const auto current_child_mask = child.mutable_view().null_mask();
-    std::for_each(thrust::make_counting_iterator(0),
-                  thrust::make_counting_iterator(child.num_children()),
-                  [&current_child_mask, &child, parent_null_mask_size, stream, mr](auto i) {
-                    superimpose_parent_nullmask(current_child_mask,
-                                                parent_null_mask_size,
-                                                UNKNOWN_NULL_COUNT,
-                                                child.child(i),
-                                                stream,
-                                                mr);
-                  });
-  }
-}
-}  // namespace
 
 /// Column factory that adopts child columns.
 std::unique_ptr<cudf::column> make_structs_column(
@@ -95,12 +45,8 @@ std::unique_ptr<cudf::column> make_structs_column(
 
   if (!null_mask.is_empty()) {
     for (auto& child : child_columns) {
-      superimpose_parent_nullmask(static_cast<bitmask_type const*>(null_mask.data()),
-                                  null_mask.size(),
-                                  null_count,
-                                  *child,
-                                  stream,
-                                  mr);
+      cudf::structs::detail::superimpose_parent_nulls(
+        static_cast<bitmask_type const*>(null_mask.data()), null_count, *child, stream, mr);
     }
   }
 
diff --git a/cpp/src/structs/utilities.cu b/cpp/src/structs/utilities.cpp
similarity index 77%
rename from cpp/src/structs/utilities.cu
rename to cpp/src/structs/utilities.cpp
index 0e944cd975c..6cc537d2042 100644
--- a/cpp/src/structs/utilities.cu
+++ b/cpp/src/structs/utilities.cpp
@@ -16,6 +16,7 @@
 
 #include <thrust/iterator/counting_iterator.h>
 
+#include <cudf/detail/null_mask.hpp>
 #include <cudf/structs/structs_column_view.hpp>
 #include <cudf/table/table_view.hpp>
 #include <cudf/unary.hpp>
@@ -167,6 +168,51 @@ flatten_nested_columns(table_view const& input,
   return flattened_table{input, column_order, null_precedence, nullability}();
 }
 
+// Helper function to superimpose validity of parent struct
+// over the specified member (child) column.
+void superimpose_parent_nulls(bitmask_type const* parent_null_mask,
+                              size_type parent_null_count,
+                              column& child,
+                              rmm::cuda_stream_view stream,
+                              rmm::mr::device_memory_resource* mr)
+{
+  if (!child.nullable()) {
+    // Child currently has no null mask. Copy parent's null mask.
+    child.set_null_mask(rmm::device_buffer{
+      parent_null_mask, cudf::bitmask_allocation_size_bytes(child.size()), stream, mr});
+    child.set_null_count(parent_null_count);
+  } else {
+    // Child should have a null mask.
+    // `AND` the child's null mask with the parent's.
+
+    auto current_child_mask = child.mutable_view().null_mask();
+
+    std::vector<bitmask_type const*> masks{
+      reinterpret_cast<bitmask_type const*>(parent_null_mask),
+      reinterpret_cast<bitmask_type const*>(current_child_mask)};
+    std::vector<size_type> begin_bits{0, 0};
+    cudf::detail::inplace_bitmask_and(
+      device_span<bitmask_type>(current_child_mask, num_bitmask_words(child.size())),
+      masks,
+      begin_bits,
+      child.size(),
+      stream,
+      mr);
+    child.set_null_count(UNKNOWN_NULL_COUNT);
+  }
+
+  // If the child is also a struct, repeat for all grandchildren.
+  if (child.type().id() == cudf::type_id::STRUCT) {
+    const auto current_child_mask = child.mutable_view().null_mask();
+    std::for_each(thrust::make_counting_iterator(0),
+                  thrust::make_counting_iterator(child.num_children()),
+                  [&current_child_mask, &child, stream, mr](auto i) {
+                    superimpose_parent_nulls(
+                      current_child_mask, UNKNOWN_NULL_COUNT, child.child(i), stream, mr);
+                  });
+  }
+}
+
 }  // namespace detail
 }  // namespace structs
 }  // namespace cudf
diff --git a/cpp/src/structs/utilities.hpp b/cpp/src/structs/utilities.hpp
index ddd28fe70be..eee9ca63146 100644
--- a/cpp/src/structs/utilities.hpp
+++ b/cpp/src/structs/utilities.hpp
@@ -76,6 +76,24 @@ flatten_nested_columns(table_view const& input,
                        std::vector<null_order> const& null_precedence,
                        column_nullability nullability = column_nullability::MATCH_INCOMING);
 
+/**
+ * @brief Pushdown nulls from a parent mask into a child column, using AND.
+ *
+ * This function will recurse through all struct descendants. It is expected that
+ * the size of `parent_null_mask` in bits is the same as `child.size()`
+ *
+ * @param parent_null_mask The mask to be applied to descendants
+ * @param parent_null_count Null count in the null mask
+ * @param column Column to apply the null mask to.
+ * @param stream CUDA stream used for device memory operations and kernel launches.
+ * @param mr     Device memory resource used to allocate new device memory.
+ */
+void superimpose_parent_nulls(bitmask_type const* parent_null_mask,
+                              size_type parent_null_count,
+                              column& child,
+                              rmm::cuda_stream_view stream,
+                              rmm::mr::device_memory_resource* mr);
+
 }  // namespace detail
 }  // namespace structs
 }  // namespace cudf
diff --git a/cpp/tests/column/factories_test.cpp b/cpp/tests/column/factories_test.cpp
index d30929b90c6..71f65eedd91 100644
--- a/cpp/tests/column/factories_test.cpp
+++ b/cpp/tests/column/factories_test.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -15,6 +15,7 @@
  */
 
 #include <cudf_test/base_fixture.hpp>
+#include <cudf_test/column_wrapper.hpp>
 #include <cudf_test/type_lists.hpp>
 
 #include <cudf/column/column.hpp>
@@ -460,3 +461,34 @@ TEST_F(ColumnFactoryTest, DictionaryFromStringScalarError)
   cudf::string_scalar value("hello", false);
   EXPECT_THROW(cudf::make_dictionary_from_scalar(value, 1), cudf::logic_error);
 }
+
+void struct_from_scalar(bool is_valid)
+{
+  using LCW = cudf::test::lists_column_wrapper<int>;
+
+  cudf::test::fixed_width_column_wrapper<int> col0{1};
+  cudf::test::strings_column_wrapper col1{"abc"};
+  cudf::test::lists_column_wrapper<int> col2{{1, 2, 3}};
+  cudf::test::lists_column_wrapper<int> col3{LCW{}};
+
+  std::vector<cudf::column_view> src_children({col0, col1, col2, col3});
+  auto value = cudf::struct_scalar(src_children, is_valid);
+  cudf::test::structs_column_wrapper struct_col({col0, col1, col2, col3}, {is_valid});
+
+  auto const num_rows = 32;
+  auto result         = cudf::make_column_from_scalar(value, num_rows);
+
+  // generate a column of size num_rows
+  std::vector<cudf::column_view> cols;
+  auto iter = thrust::make_counting_iterator(0);
+  std::transform(iter, iter + num_rows, std::back_inserter(cols), [&](int i) {
+    return static_cast<cudf::column_view>(struct_col);
+  });
+  auto expected = cudf::concatenate(cols);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*result, *expected);
+}
+
+TEST_F(ColumnFactoryTest, FromStructScalar) { struct_from_scalar(true); }
+
+TEST_F(ColumnFactoryTest, FromStructScalarNull) { struct_from_scalar(false); }
diff --git a/cpp/tests/scalar/factories_test.cpp b/cpp/tests/scalar/factories_test.cpp
index fd0a92a0168..e2f2c26a16e 100644
--- a/cpp/tests/scalar/factories_test.cpp
+++ b/cpp/tests/scalar/factories_test.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -15,6 +15,8 @@
  */
 
 #include <cudf_test/base_fixture.hpp>
+#include <cudf_test/column_wrapper.hpp>
+#include <cudf_test/table_utilities.hpp>
 #include <cudf_test/type_lists.hpp>
 
 #include <cudf/scalar/scalar_factories.hpp>
@@ -161,4 +163,33 @@ TYPED_TEST(FixedPointScalarFactory, ValueProvided)
   EXPECT_TRUE(s->is_valid());
 }
 
+struct StructScalarFactory : public ScalarFactoryTest {
+};
+
+TEST_F(StructScalarFactory, Basic)
+{
+  cudf::test::fixed_width_column_wrapper<int> col0{1};
+  cudf::test::strings_column_wrapper col1{"abc"};
+  cudf::test::lists_column_wrapper<int> col2{{1, 2, 3}};
+  cudf::test::structs_column_wrapper struct_col({col0, col1, col2});
+  cudf::column_view cv = static_cast<cudf::column_view>(struct_col);
+  std::vector<cudf::column_view> children(cv.child_begin(), cv.child_end());
+
+  // table_view constructor
+  {
+    auto sc = cudf::make_struct_scalar(cudf::table_view{children});
+    auto s  = static_cast<cudf::scalar_type_t<cudf::struct_view>*>(sc.get());
+    EXPECT_TRUE(s->is_valid());
+    CUDF_TEST_EXPECT_TABLES_EQUIVALENT(cudf::table_view{children}, s->view());
+  }
+
+  // host_span constructor
+  {
+    auto sc = cudf::make_struct_scalar(cudf::host_span<cudf::column_view const>{children});
+    auto s  = static_cast<cudf::scalar_type_t<cudf::struct_view>*>(sc.get());
+    EXPECT_TRUE(s->is_valid());
+    CUDF_TEST_EXPECT_TABLES_EQUIVALENT(cudf::table_view{children}, s->view());
+  }
+}
+
 CUDF_TEST_PROGRAM_MAIN()
diff --git a/cpp/tests/scalar/scalar_test.cpp b/cpp/tests/scalar/scalar_test.cpp
index b869569bea3..7a12c2fd27d 100644
--- a/cpp/tests/scalar/scalar_test.cpp
+++ b/cpp/tests/scalar/scalar_test.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@
 #include <cudf_test/base_fixture.hpp>
 #include <cudf_test/column_wrapper.hpp>
 #include <cudf_test/cudf_gtest.hpp>
+#include <cudf_test/table_utilities.hpp>
 #include <cudf_test/type_list_utilities.hpp>
 #include <cudf_test/type_lists.hpp>
 
@@ -231,4 +232,89 @@ TEST_F(ListScalarTest, MoveConstructorNested)
   EXPECT_EQ(s.view().num_children(), 0);
 }
 
+struct StructScalarTest : public cudf::test::BaseFixture {
+};
+
+TEST_F(StructScalarTest, Basic)
+{
+  cudf::test::fixed_width_column_wrapper<int> col0{1};
+  cudf::test::strings_column_wrapper col1{"abc"};
+  cudf::test::lists_column_wrapper<int> col2{{1, 2, 3}};
+  cudf::test::structs_column_wrapper struct_col({col0, col1, col2});
+  cudf::column_view cv = static_cast<cudf::column_view>(struct_col);
+  std::vector<cudf::column_view> children(cv.child_begin(), cv.child_end());
+
+  // table_view constructor
+  {
+    auto s = cudf::struct_scalar(children, true);
+    EXPECT_TRUE(s.is_valid());
+    CUDF_TEST_EXPECT_TABLES_EQUIVALENT(cudf::table_view{children}, s.view());
+  }
+
+  // host_span constructor
+  {
+    auto s = cudf::struct_scalar(cudf::host_span<cudf::column_view const>{children}, true);
+    EXPECT_TRUE(s.is_valid());
+    CUDF_TEST_EXPECT_TABLES_EQUIVALENT(cudf::table_view{children}, s.view());
+  }
+}
+
+TEST_F(StructScalarTest, BasicNulls)
+{
+  cudf::test::fixed_width_column_wrapper<int> col0{1};
+  cudf::test::strings_column_wrapper col1{"abc"};
+  cudf::test::lists_column_wrapper<int> col2{{1, 2, 3}};
+  std::vector<cudf::column_view> src_children({col0, col1, col2});
+
+  std::vector<std::unique_ptr<cudf::column>> src_columns;
+
+  // structs_column_wrapper takes ownership of the incoming columns, so make a copy
+  src_columns.push_back(std::make_unique<cudf::column>(src_children[0]));
+  src_columns.push_back(std::make_unique<cudf::column>(src_children[1]));
+  src_columns.push_back(std::make_unique<cudf::column>(src_children[2]));
+  cudf::test::structs_column_wrapper valid_struct_col(std::move(src_columns), {1});
+  cudf::column_view vcv = static_cast<cudf::column_view>(valid_struct_col);
+  std::vector<cudf::column_view> valid_children(vcv.child_begin(), vcv.child_end());
+
+  // structs_column_wrapper takes ownership of the incoming columns, so make a copy
+  src_columns.push_back(std::make_unique<cudf::column>(src_children[0]));
+  src_columns.push_back(std::make_unique<cudf::column>(src_children[1]));
+  src_columns.push_back(std::make_unique<cudf::column>(src_children[2]));
+  cudf::test::structs_column_wrapper invalid_struct_col(std::move(src_columns), {0});
+  cudf::column_view icv = static_cast<cudf::column_view>(invalid_struct_col);
+  std::vector<cudf::column_view> invalid_children(icv.child_begin(), icv.child_end());
+
+  // table_view constructor
+  {
+    auto s = cudf::struct_scalar(cudf::table_view{src_children}, true);
+    EXPECT_TRUE(s.is_valid());
+    CUDF_TEST_EXPECT_TABLES_EQUIVALENT(cudf::table_view{valid_children}, s.view());
+  }
+  // host_span constructor
+  {
+    auto s = cudf::struct_scalar(cudf::host_span<cudf::column_view const>{src_children}, true);
+    EXPECT_TRUE(s.is_valid());
+    CUDF_TEST_EXPECT_TABLES_EQUIVALENT(cudf::table_view{valid_children}, s.view());
+  }
+
+  // with nulls, we expect the incoming children to get nullified by passing false to
+  // the scalar constructor itself. so we use the unmodified `children` as the input, but
+  // we compare against the modified `invalid_children` produced by the source column as
+  // proof that the scalar did the validity pushdown.
+
+  // table_view constructor
+  {
+    auto s = cudf::struct_scalar(cudf::table_view{src_children}, false);
+    EXPECT_TRUE(!s.is_valid());
+    CUDF_TEST_EXPECT_TABLES_EQUIVALENT(cudf::table_view{invalid_children}, s.view());
+  }
+
+  // host_span constructor
+  {
+    auto s = cudf::struct_scalar(cudf::host_span<cudf::column_view const>{src_children}, false);
+    EXPECT_TRUE(!s.is_valid());
+    CUDF_TEST_EXPECT_TABLES_EQUIVALENT(cudf::table_view{invalid_children}, s.view());
+  }
+}
+
 CUDF_TEST_PROGRAM_MAIN()