diff --git a/include/xgboost/data.h b/include/xgboost/data.h
index 1678f4b1f4f1..a5210890da57 100644
--- a/include/xgboost/data.h
+++ b/include/xgboost/data.h
@@ -11,12 +11,14 @@
 #include <dmlc/data.h>
 #include <dmlc/serializer.h>
 #include <xgboost/base.h>
-#include <xgboost/span.h>
 #include <xgboost/host_device_vector.h>
+#include <xgboost/linalg.h>
+#include <xgboost/span.h>
+#include <xgboost/string_view.h>
 
+#include <algorithm>
 #include <memory>
 #include <numeric>
-#include <algorithm>
 #include <string>
 #include <utility>
 #include <vector>
@@ -157,7 +159,7 @@ class MetaInfo {
    *
    *        Right now only 1 column is permitted.
    */
-  void SetInfo(const char* key, std::string const& interface_str);
+  void SetInfo(StringView key, std::string const& interface_str);
 
   void GetInfo(char const* key, bst_ulong* out_len, DataType dtype,
                const void** out_dptr) const;
diff --git a/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu b/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu
index de7a1fc41495..4ecf8b0f1da1 100644
--- a/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu
+++ b/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu
@@ -35,13 +35,12 @@ template <typename T> T CheckJvmCall(T const &v, JNIEnv *jenv) {
 }
 
 template <typename VCont>
-void CopyColumnMask(xgboost::ArrayInterface const &interface,
+void CopyColumnMask(xgboost::ArrayInterface<1> const &interface,
                     std::vector<Json> const &columns, cudaMemcpyKind kind,
                     size_t c, VCont *p_mask, Json *p_out, cudaStream_t stream) {
   auto &mask = *p_mask;
   auto &out = *p_out;
-  auto size = sizeof(typename VCont::value_type) * interface.num_rows *
-              interface.num_cols;
+  auto size = sizeof(typename VCont::value_type) * interface.n;
   mask.resize(size);
   CHECK(RawPtr(mask));
   CHECK(size);
@@ -67,11 +66,11 @@ void CopyColumnMask(xgboost::ArrayInterface const &interface,
     LOG(FATAL) << "Invalid shape of mask";
   }
   out["mask"]["typestr"] = String("<t1");
-  out["mask"]["version"] = Integer(1);
+  out["mask"]["version"] = Integer(3);
 }
 
 template <typename DCont, typename VCont>
-void CopyInterface(std::vector<xgboost::ArrayInterface> &interface_arr,
+void CopyInterface(std::vector<xgboost::ArrayInterface<1>> &interface_arr,
                    std::vector<Json> const &columns, cudaMemcpyKind kind,
                    std::vector<DCont> *p_data, std::vector<VCont> *p_mask,
                    std::vector<xgboost::Json> *p_out, cudaStream_t stream) {
@@ -81,7 +80,7 @@ void CopyInterface(std::vector<xgboost::ArrayInterface> &interface_arr,
   for (size_t c = 0; c < interface_arr.size(); ++c) {
     auto &interface = interface_arr.at(c);
     size_t element_size = interface.ElementSize();
-    size_t size = element_size * interface.num_rows * interface.num_cols;
+    size_t size = element_size * interface.n;
 
     auto &data = (*p_data)[c];
     auto &mask = (*p_mask)[c];
@@ -95,14 +94,13 @@ void CopyInterface(std::vector<xgboost::ArrayInterface> &interface_arr,
         Json{Boolean{false}}};
 
     out["data"] = Array(std::move(j_data));
-    out["shape"] = Array(std::vector<Json>{Json(Integer(interface.num_rows)),
-                                           Json(Integer(interface.num_cols))});
+    out["shape"] = Array(std::vector<Json>{Json(Integer(interface.Shape(0)))});
 
     if (interface.valid.Data()) {
       CopyColumnMask(interface, columns, kind, c, &mask, &out, stream);
     }
     out["typestr"] = String("<f4");
-    out["version"] = Integer(1);
+    out["version"] = Integer(3);
   }
 }
 
@@ -110,10 +108,10 @@ void CopyMetaInfo(Json *p_interface, dh::device_vector<float> *out, cudaStream_t
   auto &j_interface = *p_interface;
   CHECK_EQ(get<Array const>(j_interface).size(), 1);
   auto object = get<Object>(get<Array>(j_interface)[0]);
-  ArrayInterface interface(object);
-  out->resize(interface.num_rows);
+  ArrayInterface<1> interface(object);
+  out->resize(interface.Shape(0));
   size_t element_size = interface.ElementSize();
-  size_t size = element_size * interface.num_rows;
+  size_t size = element_size * interface.n;
   dh::safe_cuda(cudaMemcpyAsync(RawPtr(*out), interface.data, size,
                                 cudaMemcpyDeviceToDevice, stream));
   j_interface[0]["data"][0] = reinterpret_cast<Integer::Int>(RawPtr(*out));
@@ -285,11 +283,11 @@ class DataIteratorProxy {
 
     Json features = json_interface["features_str"];
     auto json_columns = get<Array const>(features);
-    std::vector<ArrayInterface> interfaces;
+    std::vector<ArrayInterface<1>> interfaces;
 
     // Stage the data
     for (auto &json_col : json_columns) {
-      auto column = ArrayInterface(get<Object const>(json_col));
+      auto column = ArrayInterface<1>(get<Object const>(json_col));
       interfaces.emplace_back(column);
     }
     Json::Dump(features, &interface_str);
@@ -342,9 +340,9 @@ class DataIteratorProxy {
     // Data
     auto const &json_interface = host_columns_.at(it_)->interfaces;
 
-    std::vector<ArrayInterface> in;
+    std::vector<ArrayInterface<1>> in;
     for (auto interface : json_interface) {
-      auto column = ArrayInterface(get<Object const>(interface));
+      auto column = ArrayInterface<1>(get<Object const>(interface));
       in.emplace_back(column);
     }
     std::vector<Json> out;
diff --git a/src/common/device_helpers.cuh b/src/common/device_helpers.cuh
index f35af3b5e647..b316453f5d2c 100644
--- a/src/common/device_helpers.cuh
+++ b/src/common/device_helpers.cuh
@@ -136,7 +136,7 @@ inline ncclResult_t ThrowOnNcclError(ncclResult_t code, const char *file,
 }
 #endif
 
-inline int32_t CudaGetPointerDevice(void* ptr) {
+inline int32_t CudaGetPointerDevice(void const *ptr) {
   int32_t device = -1;
   cudaPointerAttributes attr;
   dh::safe_cuda(cudaPointerGetAttributes(&attr, ptr));
diff --git a/src/data/adapter.h b/src/data/adapter.h
index 27da8c6e3b36..4214171e8b7e 100644
--- a/src/data/adapter.h
+++ b/src/data/adapter.h
@@ -254,20 +254,20 @@ class ArrayAdapterBatch : public detail::NoMetaInfo {
   static constexpr bool kIsRowMajor = true;
 
  private:
-  ArrayInterface array_interface_;
+  ArrayInterface<2> array_interface_;
 
   class Line {
-    ArrayInterface array_interface_;
+    ArrayInterface<2> array_interface_;
     size_t ridx_;
 
    public:
-    Line(ArrayInterface array_interface, size_t ridx)
+    Line(ArrayInterface<2> array_interface, size_t ridx)
         : array_interface_{std::move(array_interface)}, ridx_{ridx} {}
 
-    size_t Size() const { return array_interface_.num_cols; }
+    size_t Size() const { return array_interface_.Shape(1); }
 
     COOTuple GetElement(size_t idx) const {
-      return {ridx_, idx, array_interface_.GetElement(ridx_, idx)};
+      return {ridx_, idx, array_interface_(ridx_, idx)};
     }
   };
 
@@ -277,11 +277,11 @@ class ArrayAdapterBatch : public detail::NoMetaInfo {
     return Line{array_interface_, idx};
   }
 
-  size_t NumRows() const { return array_interface_.num_rows; }
-  size_t NumCols() const { return array_interface_.num_cols; }
+  size_t NumRows() const { return array_interface_.Shape(0); }
+  size_t NumCols() const { return array_interface_.Shape(1); }
   size_t Size() const { return this->NumRows(); }
 
-  explicit ArrayAdapterBatch(ArrayInterface array_interface)
+  explicit ArrayAdapterBatch(ArrayInterface<2> array_interface)
       : array_interface_{std::move(array_interface)} {}
 };
 
@@ -294,43 +294,42 @@ class ArrayAdapter : public detail::SingleBatchDataIter<ArrayAdapterBatch> {
  public:
   explicit ArrayAdapter(StringView array_interface) {
     auto j = Json::Load(array_interface);
-    array_interface_ = ArrayInterface(get<Object const>(j));
+    array_interface_ = ArrayInterface<2>(get<Object const>(j));
     batch_ = ArrayAdapterBatch{array_interface_};
   }
   ArrayAdapterBatch const& Value() const override { return batch_; }
-  size_t NumRows() const { return array_interface_.num_rows; }
-  size_t NumColumns() const { return array_interface_.num_cols; }
+  size_t NumRows() const { return array_interface_.Shape(0); }
+  size_t NumColumns() const { return array_interface_.Shape(1); }
 
  private:
   ArrayAdapterBatch batch_;
-  ArrayInterface array_interface_;
+  ArrayInterface<2> array_interface_;
 };
 
 class CSRArrayAdapterBatch : public detail::NoMetaInfo {
-  ArrayInterface indptr_;
-  ArrayInterface indices_;
-  ArrayInterface values_;
+  ArrayInterface<1> indptr_;
+  ArrayInterface<1> indices_;
+  ArrayInterface<1> values_;
   bst_feature_t n_features_;
 
   class Line {
-    ArrayInterface indices_;
-    ArrayInterface values_;
+    ArrayInterface<1> indices_;
+    ArrayInterface<1> values_;
     size_t ridx_;
     size_t offset_;
 
    public:
-    Line(ArrayInterface indices, ArrayInterface values, size_t ridx,
+    Line(ArrayInterface<1> indices, ArrayInterface<1> values, size_t ridx,
          size_t offset)
         : indices_{std::move(indices)}, values_{std::move(values)}, ridx_{ridx},
           offset_{offset} {}
 
     COOTuple GetElement(size_t idx) const {
-      return {ridx_, indices_.GetElement<size_t>(offset_ + idx, 0),
-              values_.GetElement(offset_ + idx, 0)};
+      return {ridx_, TypedIndex<size_t, 1>{indices_}(offset_ + idx), values_(offset_ + idx)};
     }
 
     size_t Size() const {
-      return values_.num_rows * values_.num_cols;
+      return values_.Shape(0);
     }
   };
 
@@ -339,17 +338,16 @@ class CSRArrayAdapterBatch : public detail::NoMetaInfo {
 
  public:
   CSRArrayAdapterBatch() = default;
-  CSRArrayAdapterBatch(ArrayInterface indptr, ArrayInterface indices,
-                       ArrayInterface values, bst_feature_t n_features)
-      : indptr_{std::move(indptr)}, indices_{std::move(indices)},
-        values_{std::move(values)}, n_features_{n_features} {
-    indptr_.AsColumnVector();
-    values_.AsColumnVector();
-    indices_.AsColumnVector();
+  CSRArrayAdapterBatch(ArrayInterface<1> indptr, ArrayInterface<1> indices,
+                       ArrayInterface<1> values, bst_feature_t n_features)
+      : indptr_{std::move(indptr)},
+        indices_{std::move(indices)},
+        values_{std::move(values)},
+        n_features_{n_features} {
   }
 
   size_t NumRows() const {
-    size_t size = indptr_.num_rows * indptr_.num_cols;
+    size_t size = indptr_.Shape(0);
     size = size == 0 ? 0 : size - 1;
     return size;
   }
@@ -357,19 +355,19 @@ class CSRArrayAdapterBatch : public detail::NoMetaInfo {
   size_t Size() const { return this->NumRows(); }
 
   Line const GetLine(size_t idx) const {
-    auto begin_offset = indptr_.GetElement<size_t>(idx, 0);
-    auto end_offset = indptr_.GetElement<size_t>(idx + 1, 0);
+    auto begin_no_stride = TypedIndex<size_t, 1>{indptr_}(idx);
+    auto end_no_stride = TypedIndex<size_t, 1>{indptr_}(idx + 1);
 
     auto indices = indices_;
     auto values = values_;
+    // Slice indices and values, stride remains unchanged since this is slicing by
+    // specific index.
+    auto offset = indices.strides[0] * begin_no_stride;
 
-    values.num_cols = end_offset - begin_offset;
-    values.num_rows = 1;
+    indices.shape[0] = end_no_stride - begin_no_stride;
+    values.shape[0] = end_no_stride - begin_no_stride;
 
-    indices.num_cols = values.num_cols;
-    indices.num_rows = values.num_rows;
-
-    return Line{indices, values, idx, begin_offset};
+    return Line{indices, values, idx, offset};
   }
 };
 
@@ -391,7 +389,7 @@ class CSRArrayAdapter : public detail::SingleBatchDataIter<CSRArrayAdapterBatch>
     return batch_;
   }
   size_t NumRows() const {
-    size_t size = indptr_.num_cols * indptr_.num_rows;
+    size_t size = indptr_.Shape(0);
     size = size == 0 ? 0 : size - 1;
     return  size;
   }
@@ -399,9 +397,9 @@ class CSRArrayAdapter : public detail::SingleBatchDataIter<CSRArrayAdapterBatch>
 
  private:
   CSRArrayAdapterBatch batch_;
-  ArrayInterface indptr_;
-  ArrayInterface indices_;
-  ArrayInterface values_;
+  ArrayInterface<1> indptr_;
+  ArrayInterface<1> indices_;
+  ArrayInterface<1> values_;
   size_t num_cols_;
 };
 
diff --git a/src/data/array_interface.cu b/src/data/array_interface.cu
index def4de195523..b1a80251ecc4 100644
--- a/src/data/array_interface.cu
+++ b/src/data/array_interface.cu
@@ -1,21 +1,56 @@
 /*!
  * Copyright 2021 by Contributors
  */
-#include "array_interface.h"
 #include "../common/common.h"
+#include "array_interface.h"
 
 namespace xgboost {
 void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) {
   switch (stream) {
-  case 0:
-    LOG(FATAL) << "Invalid stream ID in array interface: " << stream;
-  case 1:
-    // default legacy stream
-    break;
-  case 2:
-    // default per-thread stream
-  default:
-    dh::safe_cuda(cudaStreamSynchronize(reinterpret_cast<cudaStream_t>(stream)));
+    case 0:
+      /**
+       * disallowed by the `__cuda_array_interface__`.  Quote:
+       *
+       *   This is disallowed as it would be ambiguous between None and the default
+       *   stream, and also between the legacy and per-thread default streams. Any use
+       *   case where 0 might be given should either use None, 1, or 2 instead for
+       *   clarity.
+       */
+      LOG(FATAL) << "Invalid stream ID in array interface: " << stream;
+    case 1:
+      // default legacy stream
+      break;
+    case 2:
+      // default per-thread stream
+    default:
+      dh::safe_cuda(cudaStreamSynchronize(reinterpret_cast<cudaStream_t>(stream)));
+  }
+}
+
+bool ArrayInterfaceHandler::IsCudaPtr(void const* ptr) {
+  if (!ptr) {
+    return false;
+  }
+  cudaPointerAttributes attr;
+  auto err = cudaPointerGetAttributes(&attr, ptr);
+  // reset error
+  CHECK_EQ(err, cudaGetLastError());
+  if (err == cudaErrorInvalidValue) {
+    // CUDA < 11
+    return false;
+  } else if (err == cudaSuccess) {
+    // CUDA >= 11
+    switch (attr.type) {
+      case cudaMemoryTypeUnregistered:
+      case cudaMemoryTypeHost:
+        return false;
+      default:
+        return true;
+    }
+    return true;
+  } else {
+    // other errors, `cudaErrorNoDevice`, `cudaErrorInsufficientDriver` etc.
+    return false;
   }
 }
 }  // namespace xgboost
diff --git a/src/data/array_interface.h b/src/data/array_interface.h
index 6524f4512407..1b2545a57d62 100644
--- a/src/data/array_interface.h
+++ b/src/data/array_interface.h
@@ -13,24 +13,23 @@
 #include <utility>
 #include <vector>
 
+#include "../common/bitfield.h"
+#include "../common/common.h"
 #include "xgboost/base.h"
 #include "xgboost/data.h"
 #include "xgboost/json.h"
+#include "xgboost/linalg.h"
 #include "xgboost/logging.h"
 #include "xgboost/span.h"
-#include "../common/bitfield.h"
-#include "../common/common.h"
 
 namespace xgboost {
 // Common errors in parsing columnar format.
 struct ArrayInterfaceErrors {
-  static char const* Contigious() {
-    return "Memory should be contigious.";
-  }
-  static char const* TypestrFormat() {
+  static char const *Contiguous() { return "Memory should be contiguous."; }
+  static char const *TypestrFormat() {
     return "`typestr' should be of format <endian><type><size of type in bytes>.";
   }
-  static char const* Dimension(int32_t d) {
+  static char const *Dimension(int32_t d) {
     static std::string str;
     str.clear();
     str += "Only ";
@@ -38,11 +37,11 @@ struct ArrayInterfaceErrors {
     str += " dimensional array is valid.";
     return str.c_str();
   }
-  static char const* Version() {
-    return "Only version <= 3 of "
-           "`__cuda_array_interface__/__array_interface__' are supported.";
+  static char const *Version() {
+    return "Only version <= 3 of `__cuda_array_interface__' and `__array_interface__' are "
+           "supported.";
   }
-  static char const* OfType(std::string const& type) {
+  static char const *OfType(std::string const &type) {
     static std::string str;
     str.clear();
     str += " should be of ";
@@ -92,49 +91,39 @@ struct ArrayInterfaceErrors {
   }
 };
 
-// TODO(trivialfis): Abstract this into a class that accept a json
-// object and turn it into an array (for cupy and numba).
+/**
+ * Utilities for consuming array interface.
+ */
 class ArrayInterfaceHandler {
  public:
-  template <typename T>
-  static constexpr char TypeChar() {
-    return
-        (std::is_floating_point<T>::value ? 'f' :
-         (std::is_integral<T>::value ?
-          (std::is_signed<T>::value ? 'i' : 'u') : '\0'));
-  }
+  enum Type : std::int8_t { kF4, kF8, kF16, kI1, kI2, kI4, kI8, kU1, kU2, kU4, kU8 };
 
   template <typename PtrType>
-  static PtrType GetPtrFromArrayData(std::map<std::string, Json> const& obj) {
-    if (obj.find("data") == obj.cend()) {
+  static PtrType GetPtrFromArrayData(std::map<std::string, Json> const &obj) {
+    auto data_it = obj.find("data");
+    if (data_it == obj.cend()) {
       LOG(FATAL) << "Empty data passed in.";
     }
-    auto p_data = reinterpret_cast<PtrType>(static_cast<size_t>(
-        get<Integer const>(
-            get<Array const>(
-                obj.at("data"))
-            .at(0))));
+    auto p_data = reinterpret_cast<PtrType>(
+        static_cast<size_t>(get<Integer const>(get<Array const>(data_it->second).at(0))));
     return p_data;
   }
 
-  static void Validate(std::map<std::string, Json> const& array) {
+  static void Validate(std::map<std::string, Json> const &array) {
     auto version_it = array.find("version");
     if (version_it == array.cend()) {
       LOG(FATAL) << "Missing `version' field for array interface";
     }
-    auto stream_it = array.find("stream");
-    if (stream_it != array.cend() && !IsA<Null>(stream_it->second)) {
-      // is cuda, check the version.
-      if (get<Integer const>(version_it->second) > 3) {
-        LOG(FATAL) << ArrayInterfaceErrors::Version();
-      }
+    if (get<Integer const>(version_it->second) > 3) {
+      LOG(FATAL) << ArrayInterfaceErrors::Version();
     }
 
-    if (array.find("typestr") == array.cend()) {
+    auto typestr_it = array.find("typestr");
+    if (typestr_it == array.cend()) {
       LOG(FATAL) << "Missing `typestr' field for array interface";
     }
 
-    auto typestr = get<String const>(array.at("typestr"));
+    auto typestr = get<String const>(typestr_it->second);
     CHECK(typestr.size() == 3 || typestr.size() == 4) << ArrayInterfaceErrors::TypestrFormat();
 
     if (array.find("shape") == array.cend()) {
@@ -149,12 +138,12 @@ class ArrayInterfaceHandler {
   // Mask object is also an array interface, but with different requirements.
   static size_t ExtractMask(std::map<std::string, Json> const &column,
                             common::Span<RBitField8::value_type> *p_out) {
-    auto& s_mask = *p_out;
+    auto &s_mask = *p_out;
     if (column.find("mask") != column.cend()) {
-      auto const& j_mask = get<Object const>(column.at("mask"));
+      auto const &j_mask = get<Object const>(column.at("mask"));
       Validate(j_mask);
 
-      auto p_mask = GetPtrFromArrayData<RBitField8::value_type*>(j_mask);
+      auto p_mask = GetPtrFromArrayData<RBitField8::value_type *>(j_mask);
 
       auto j_shape = get<Array const>(j_mask.at("shape"));
       CHECK_EQ(j_shape.size(), 1) << ArrayInterfaceErrors::Dimension(1);
@@ -186,8 +175,8 @@ class ArrayInterfaceHandler {
 
       if (j_mask.find("strides") != j_mask.cend()) {
         auto strides = get<Array const>(column.at("strides"));
-        CHECK_EQ(strides.size(),                        1) << ArrayInterfaceErrors::Dimension(1);
-        CHECK_EQ(get<Integer>(strides.at(0)), type_length) << ArrayInterfaceErrors::Contigious();
+        CHECK_EQ(strides.size(), 1) << ArrayInterfaceErrors::Dimension(1);
+        CHECK_EQ(get<Integer>(strides.at(0)), type_length) << ArrayInterfaceErrors::Contiguous();
       }
 
       s_mask = {p_mask, span_size};
@@ -195,77 +184,212 @@ class ArrayInterfaceHandler {
     }
     return 0;
   }
-
-  static std::pair<bst_row_t, bst_feature_t> ExtractShape(
-      std::map<std::string, Json> const& column) {
-    auto j_shape = get<Array const>(column.at("shape"));
-    auto typestr = get<String const>(column.at("typestr"));
-    if (j_shape.size() == 1) {
-      return {static_cast<bst_row_t>(get<Integer const>(j_shape.at(0))), 1};
-    } else {
-      CHECK_EQ(j_shape.size(), 2) << "Only 1-D and 2-D arrays are supported.";
-      return {static_cast<bst_row_t>(get<Integer const>(j_shape.at(0))),
-              static_cast<bst_feature_t>(get<Integer const>(j_shape.at(1)))};
+  /**
+   * \brief Handle vector inputs.  For higher dimension, we require strictly correct shape.
+   */
+  template <int32_t D>
+  static void HandleRowVector(std::vector<size_t> const &shape, std::vector<size_t> *p_out) {
+    auto &out = *p_out;
+    if (shape.size() == 2 && D == 1) {
+      auto m = shape[0];
+      auto n = shape[1];
+      CHECK(m == 1 || n == 1);
+      if (m == 1) {
+        // keep the number of columns
+        out[0] = out[1];
+        out.resize(1);
+      } else if (n == 1) {
+        // keep the number of rows.
+        out.resize(1);
+      }
+      // when both m and n are 1, above logic keeps the column.
+      // when neither m nor n is 1, caller should throw an error about Dimension.
     }
   }
 
-  static void ExtractStride(std::map<std::string, Json> const &column,
-                            size_t *stride_r, size_t *stride_c, size_t rows,
-                            size_t cols, size_t itemsize) {
-    auto strides_it = column.find("strides");
-    if (strides_it == column.cend() || IsA<Null>(strides_it->second)) {
-      // default strides
-      *stride_r = cols;
-      *stride_c = 1;
-    } else {
-      // strides specified by the array interface
-      auto const &j_strides = get<Array const>(strides_it->second);
-      CHECK_LE(j_strides.size(), 2) << ArrayInterfaceErrors::Dimension(2);
-      *stride_r = get<Integer const>(j_strides[0]) / itemsize;
-      size_t n = 1;
-      if (j_strides.size() == 2) {
-        n = get<Integer const>(j_strides[1]) / itemsize;
-      }
-      *stride_c = n;
+  template <int32_t D>
+  static void ExtractShape(std::map<std::string, Json> const &array, size_t (&out_shape)[D]) {
+    auto const &j_shape = get<Array const>(array.at("shape"));
+    std::vector<size_t> shape_arr(j_shape.size(), 0);
+    std::transform(j_shape.cbegin(), j_shape.cend(), shape_arr.begin(),
+                   [](Json in) { return get<Integer const>(in); });
+    // handle column vector vs. row vector
+    HandleRowVector<D>(shape_arr, &shape_arr);
+    // Copy shape.
+    size_t i;
+    for (i = 0; i < shape_arr.size(); ++i) {
+      CHECK_LT(i, D) << ArrayInterfaceErrors::Dimension(D);
+      out_shape[i] = shape_arr[i];
     }
+    // Fill the remaining dimensions
+    std::fill(out_shape + i, out_shape + D, 1);
+  }
 
-    auto valid = rows * (*stride_r) + cols * (*stride_c) >= (rows * cols);
-    CHECK(valid) << "Invalid strides in array."
-                 << "  strides: (" << (*stride_r) << "," << (*stride_c)
-                 << "), shape: (" << rows << ", " << cols << ")";
+  /**
+   * \brief Extracts the optiona `strides' field and returns whether the array is c-contiguous.
+   */
+  template <int32_t D>
+  static bool ExtractStride(std::map<std::string, Json> const &array, size_t itemsize,
+                            size_t (&shape)[D], size_t (&stride)[D]) {
+    auto strides_it = array.find("strides");
+    // No stride is provided
+    if (strides_it == array.cend() || IsA<Null>(strides_it->second)) {
+      // No stride is provided, we can calculate it from shape.
+      linalg::detail::CalcStride(shape, stride);
+      // Quote:
+      //
+      //   strides: Either None to indicate a C-style contiguous array or a Tuple of
+      //            strides which provides the number of bytes
+      return true;
+    }
+    // Get shape, we need to make changes to handle row vector, so some duplicated code
+    // from `ExtractShape` for copying out the shape.
+    auto const &j_shape = get<Array const>(array.at("shape"));
+    std::vector<size_t> shape_arr(j_shape.size(), 0);
+    std::transform(j_shape.cbegin(), j_shape.cend(), shape_arr.begin(),
+                   [](Json in) { return get<Integer const>(in); });
+    // Get stride
+    auto const &j_strides = get<Array const>(strides_it->second);
+    CHECK_EQ(j_strides.size(), j_shape.size()) << "stride and shape don't match.";
+    std::vector<size_t> stride_arr(j_strides.size(), 0);
+    std::transform(j_strides.cbegin(), j_strides.cend(), stride_arr.begin(),
+                   [](Json in) { return get<Integer const>(in); });
+
+    // Handle column vector vs. row vector
+    HandleRowVector<D>(shape_arr, &stride_arr);
+    size_t i;
+    for (i = 0; i < stride_arr.size(); ++i) {
+      // If one of the dim has shape 0 then total size is 0, stride is meaningless, but we
+      // set it to 0 here just to be consistent
+      CHECK_LT(i, D) << ArrayInterfaceErrors::Dimension(D);
+      // We use number of items instead of number of bytes
+      stride[i] = stride_arr[i] / itemsize;
+    }
+    std::fill(stride + i, stride + D, 1);
+    // If the stride can be calculated from shape then it's contiguous.
+    size_t stride_tmp[D];
+    linalg::detail::CalcStride(shape, stride_tmp);
+    return std::equal(stride_tmp, stride_tmp + D, stride);
   }
 
-  static void* ExtractData(std::map<std::string, Json> const &column,
-                                     std::pair<size_t, size_t> shape) {
-    Validate(column);
-    void* p_data = ArrayInterfaceHandler::GetPtrFromArrayData<void*>(column);
+  static void *ExtractData(std::map<std::string, Json> const &array, size_t size) {
+    Validate(array);
+    void *p_data = ArrayInterfaceHandler::GetPtrFromArrayData<void *>(array);
     if (!p_data) {
-      CHECK_EQ(shape.first * shape.second, 0) << "Empty data with non-zero shape.";
+      CHECK_EQ(size, 0) << "Empty data with non-zero shape.";
     }
     return p_data;
   }
-
+  /**
+   * \brief Whether the ptr is allocated by CUDA.
+   */
+  static bool IsCudaPtr(void const *ptr);
+  /**
+   * \brief Sync the CUDA stream.
+   */
   static void SyncCudaStream(int64_t stream);
 };
 
+/**
+ * Dispatch compile time type to runtime type.
+ */
+template <typename T, typename E = void>
+struct ToDType;
+// float
+template <>
+struct ToDType<float> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF4;
+};
+template <>
+struct ToDType<double> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF8;
+};
+template <typename T>
+struct ToDType<T,
+               std::enable_if_t<std::is_same<T, long double>::value && sizeof(long double) == 16>> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kF16;
+};
+// uint
+template <>
+struct ToDType<uint8_t> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU1;
+};
+template <>
+struct ToDType<uint16_t> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU2;
+};
+template <>
+struct ToDType<uint32_t> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU4;
+};
+template <>
+struct ToDType<uint64_t> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kU8;
+};
+// int
+template <>
+struct ToDType<int8_t> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI1;
+};
+template <>
+struct ToDType<int16_t> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI2;
+};
+template <>
+struct ToDType<int32_t> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI4;
+};
+template <>
+struct ToDType<int64_t> {
+  static constexpr ArrayInterfaceHandler::Type kType = ArrayInterfaceHandler::kI8;
+};
+
 #if !defined(XGBOOST_USE_CUDA)
-inline void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) {
-  common::AssertGPUSupport();
-}
+inline void ArrayInterfaceHandler::SyncCudaStream(int64_t stream) { common::AssertGPUSupport(); }
+inline bool ArrayInterfaceHandler::IsCudaPtr(void const *ptr) { return false; }
 #endif  // !defined(XGBOOST_USE_CUDA)
 
-// A view over __array_interface__
+/**
+ * \brief A type erased view over __array_interface__ protocol defined by numpy
+ *
+ *   <a href="https://numpy.org/doc/stable/reference/arrays.interface.html">numpy</a>.
+ *
+ * \tparam D The number of maximum dimension.
+
+ *   User input array must have dim <= D for all non-trivial dimensions.  During
+ *   construction, the ctor can automatically remove those trivial dimensions.
+ *
+ * \tparam allow_mask Whether masked array is accepted.
+ *
+ *   Currently this only supported for 1-dim vector, which is used by cuDF column
+ *   (apache arrow format).  For general masked array, as the time of writting, only
+ *   numpy has the proper support even though it's in the __cuda_array_interface__
+ *   protocol defined by numba.
+ */
+template <int32_t D, bool allow_mask = (D == 1)>
 class ArrayInterface {
-  void Initialize(std::map<std::string, Json> const &array,
-                  bool allow_mask = true) {
+  static_assert(D > 0, "Invalid dimension for array interface.");
+
+  /**
+   * \brief Initialize the object, by extracting shape, stride and type.
+   *
+   *   The function also perform some basic validation for input array.  Lastly it will
+   *   also remove trivial dimensions like converting a matrix with shape (n_samples, 1)
+   *   to a vector of size n_samples.  For for inputs like weights, this should be a 1
+   *   dimension column vector even though user might provide a matrix.
+   */
+  void Initialize(std::map<std::string, Json> const &array) {
     ArrayInterfaceHandler::Validate(array);
+
     auto typestr = get<String const>(array.at("typestr"));
     this->AssignType(StringView{typestr});
+    ArrayInterfaceHandler::ExtractShape(array, shape);
+    size_t itemsize = typestr[2] - '0';
+    is_contiguous = ArrayInterfaceHandler::ExtractStride(array, itemsize, shape, strides);
+    n = linalg::detail::CalcSize(shape);
 
-    std::tie(num_rows, num_cols) = ArrayInterfaceHandler::ExtractShape(array);
-    data = ArrayInterfaceHandler::ExtractData(
-        array, std::make_pair(num_rows, num_cols));
-
+    data = ArrayInterfaceHandler::ExtractData(array, n);
+    static_assert(allow_mask ? D == 1 : D >= 1, "Masked ndarray is not supported.");
     if (allow_mask) {
       common::Span<RBitField8::value_type> s_mask;
       size_t n_bits = ArrayInterfaceHandler::ExtractMask(array, &s_mask);
@@ -273,18 +397,13 @@ class ArrayInterface {
       valid = RBitField8(s_mask);
 
       if (s_mask.data()) {
-        CHECK_EQ(n_bits, num_rows)
-            << "Shape of bit mask doesn't match data shape. "
-            << "XGBoost doesn't support internal broadcasting.";
+        CHECK_EQ(n_bits, n) << "Shape of bit mask doesn't match data shape. "
+                            << "XGBoost doesn't support internal broadcasting.";
       }
     } else {
-      CHECK(array.find("mask") == array.cend())
-          << "Masked array is not yet supported.";
+      CHECK(array.find("mask") == array.cend()) << "Masked array is not yet supported.";
     }
 
-    ArrayInterfaceHandler::ExtractStride(array, &stride_row, &stride_col,
-                                         num_rows, num_cols, typestr[2] - '0');
-
     auto stream_it = array.find("stream");
     if (stream_it != array.cend() && !IsA<Null>(stream_it->second)) {
       int64_t stream = get<Integer const>(stream_it->second);
@@ -292,151 +411,147 @@ class ArrayInterface {
     }
   }
 
- public:
-  enum Type : std::int8_t { kF4, kF8, kF16, kI1, kI2, kI4, kI8, kU1, kU2, kU4, kU8 };
-
  public:
   ArrayInterface() = default;
-  explicit ArrayInterface(std::string const &str, bool allow_mask = true)
-      : ArrayInterface{StringView{str.c_str(), str.size()}, allow_mask} {}
+  explicit ArrayInterface(std::map<std::string, Json> const &array) { this->Initialize(array); }
 
-  explicit ArrayInterface(std::map<std::string, Json> const &column,
-                          bool allow_mask = true) {
-    this->Initialize(column, allow_mask);
-  }
-
-  explicit ArrayInterface(StringView str, bool allow_mask = true) {
-    auto jinterface = Json::Load(str);
-    if (IsA<Object>(jinterface)) {
-      this->Initialize(get<Object const>(jinterface), allow_mask);
+  explicit ArrayInterface(Json const &array) {
+    if (IsA<Object>(array)) {
+      this->Initialize(get<Object const>(array));
       return;
     }
-    if (IsA<Array>(jinterface)) {
-      CHECK_EQ(get<Array const>(jinterface).size(), 1)
+    if (IsA<Array>(array)) {
+      CHECK_EQ(get<Array const>(array).size(), 1)
           << "Column: " << ArrayInterfaceErrors::Dimension(1);
-      this->Initialize(get<Object const>(get<Array const>(jinterface)[0]), allow_mask);
+      this->Initialize(get<Object const>(get<Array const>(array)[0]));
       return;
     }
   }
 
-  void AsColumnVector() {
-    CHECK(num_rows == 1 || num_cols == 1) << "Array should be a vector instead of matrix.";
-    num_rows = std::max(num_rows, static_cast<size_t>(num_cols));
-    num_cols = 1;
+  explicit ArrayInterface(std::string const &str) : ArrayInterface{StringView{str}} {}
 
-    stride_row = std::max(stride_row, stride_col);
-    stride_col = 1;
-  }
+  explicit ArrayInterface(StringView str) : ArrayInterface<D>{Json::Load(str)} {}
 
   void AssignType(StringView typestr) {
-    if (typestr.size() == 4 && typestr[1] == 'f' && typestr[2] == '1' &&
-        typestr[3] == '6') {
-      type = kF16;
+    using T = ArrayInterfaceHandler::Type;
+    if (typestr.size() == 4 && typestr[1] == 'f' && typestr[2] == '1' && typestr[3] == '6') {
+      type = T::kF16;
       CHECK(sizeof(long double) == 16)
           << "128-bit floating point is not supported on current platform.";
     } else if (typestr[1] == 'f' && typestr[2] == '4') {
-      type = kF4;
+      type = T::kF4;
     } else if (typestr[1] == 'f' && typestr[2] == '8') {
-      type = kF8;
+      type = T::kF8;
     } else if (typestr[1] == 'i' && typestr[2] == '1') {
-      type = kI1;
+      type = T::kI1;
     } else if (typestr[1] == 'i' && typestr[2] == '2') {
-      type = kI2;
+      type = T::kI2;
     } else if (typestr[1] == 'i' && typestr[2] == '4') {
-      type = kI4;
+      type = T::kI4;
     } else if (typestr[1] == 'i' && typestr[2] == '8') {
-      type = kI8;
+      type = T::kI8;
     } else if (typestr[1] == 'u' && typestr[2] == '1') {
-      type = kU1;
+      type = T::kU1;
     } else if (typestr[1] == 'u' && typestr[2] == '2') {
-      type = kU2;
+      type = T::kU2;
     } else if (typestr[1] == 'u' && typestr[2] == '4') {
-      type = kU4;
+      type = T::kU4;
     } else if (typestr[1] == 'u' && typestr[2] == '8') {
-      type = kU8;
+      type = T::kU8;
     } else {
       LOG(FATAL) << ArrayInterfaceErrors::UnSupportedType(typestr);
       return;
     }
   }
 
+  XGBOOST_DEVICE size_t Shape(size_t i) const { return shape[i]; }
+  XGBOOST_DEVICE size_t Stride(size_t i) const { return strides[i]; }
+
   template <typename Fn>
-  XGBOOST_HOST_DEV_INLINE decltype(auto) DispatchCall(Fn func) const {
+  XGBOOST_HOST_DEV_INLINE constexpr decltype(auto) DispatchCall(Fn func) const {
+    using T = ArrayInterfaceHandler::Type;
     switch (type) {
-    case kF4:
-      return func(reinterpret_cast<float *>(data));
-    case kF8:
-      return func(reinterpret_cast<double *>(data));
+      case T::kF4:
+        return func(reinterpret_cast<float const *>(data));
+      case T::kF8:
+        return func(reinterpret_cast<double const *>(data));
 #ifdef __CUDA_ARCH__
-    case kF16: {
-      // CUDA device code doesn't support long double.
-      SPAN_CHECK(false);
-      return func(reinterpret_cast<double *>(data));
-    }
+      case T::kF16: {
+        // CUDA device code doesn't support long double.
+        SPAN_CHECK(false);
+        return func(reinterpret_cast<double const *>(data));
+      }
 #else
-    case kF16:
-      return func(reinterpret_cast<long double *>(data));
+      case T::kF16:
+        return func(reinterpret_cast<long double const *>(data));
 #endif
-    case kI1:
-      return func(reinterpret_cast<int8_t *>(data));
-    case kI2:
-      return func(reinterpret_cast<int16_t *>(data));
-    case kI4:
-      return func(reinterpret_cast<int32_t *>(data));
-    case kI8:
-      return func(reinterpret_cast<int64_t *>(data));
-    case kU1:
-      return func(reinterpret_cast<uint8_t *>(data));
-    case kU2:
-      return func(reinterpret_cast<uint16_t *>(data));
-    case kU4:
-      return func(reinterpret_cast<uint32_t *>(data));
-    case kU8:
-      return func(reinterpret_cast<uint64_t *>(data));
+      case T::kI1:
+        return func(reinterpret_cast<int8_t const *>(data));
+      case T::kI2:
+        return func(reinterpret_cast<int16_t const *>(data));
+      case T::kI4:
+        return func(reinterpret_cast<int32_t const *>(data));
+      case T::kI8:
+        return func(reinterpret_cast<int64_t const *>(data));
+      case T::kU1:
+        return func(reinterpret_cast<uint8_t const *>(data));
+      case T::kU2:
+        return func(reinterpret_cast<uint16_t const *>(data));
+      case T::kU4:
+        return func(reinterpret_cast<uint32_t const *>(data));
+      case T::kU8:
+        return func(reinterpret_cast<uint64_t const *>(data));
     }
     SPAN_CHECK(false);
-    return func(reinterpret_cast<uint64_t *>(data));
+    return func(reinterpret_cast<uint64_t const *>(data));
   }
 
-  XGBOOST_DEVICE size_t ElementSize() {
-    return this->DispatchCall([](auto* p_values) {
-      return sizeof(std::remove_pointer_t<decltype(p_values)>);
-    });
+  XGBOOST_DEVICE size_t constexpr ElementSize() {
+    return this->DispatchCall(
+        [](auto *p_values) { return sizeof(std::remove_pointer_t<decltype(p_values)>); });
   }
 
-  template <typename T = float>
-  XGBOOST_DEVICE T GetElement(size_t r, size_t c) const {
-    return this->DispatchCall(
-        [=](auto *p_values) -> T { return p_values[stride_row * r + stride_col * c]; });
+  template <typename T = float, typename... Index>
+  XGBOOST_DEVICE T operator()(Index &&...index) const {
+    static_assert(sizeof...(index) <= D, "Invalid index.");
+    return this->DispatchCall([=](auto const *p_values) -> T {
+      size_t offset = linalg::detail::Offset<0ul>(strides, 0ul, index...);
+      return static_cast<T>(p_values[offset]);
+    });
   }
 
+  // Used only by columnar format.
   RBitField8 valid;
-  bst_row_t num_rows;
-  bst_feature_t num_cols;
-  size_t stride_row{0};
-  size_t stride_col{0};
-  void* data;
-  Type type;
+  // Array stride
+  size_t strides[D]{0};
+  // Array shape
+  size_t shape[D]{0};
+  // Type earsed pointer referencing the data.
+  void const *data{nullptr};
+  // Total number of items
+  size_t n{0};
+  // Whether the memory is c-contiguous
+  bool is_contiguous{false};
+  // RTTI, initialized to the f16 to avoid masking potential bugs in initialization.
+  ArrayInterfaceHandler::Type type{ArrayInterfaceHandler::kF16};
 };
 
-template <typename T> std::string MakeArrayInterface(T const *data, size_t n) {
-  Json arr{Object{}};
-  arr["data"] = Array(std::vector<Json>{
-      Json{Integer{reinterpret_cast<int64_t>(data)}}, Json{Boolean{false}}});
-  arr["shape"] = Array{std::vector<Json>{Json{Integer{n}}, Json{Integer{1}}}};
-  std::string typestr;
-  if (DMLC_LITTLE_ENDIAN) {
-    typestr.push_back('<');
-  } else {
-    typestr.push_back('>');
+/**
+ * \brief Helper for type casting.
+ */
+template <typename T, int32_t D>
+struct TypedIndex {
+  ArrayInterface<D> const &array;
+  template <typename... I>
+  XGBOOST_DEVICE T operator()(I &&...ind) const {
+    static_assert(sizeof...(ind) <= D, "Invalid index.");
+    return array.template operator()<T>(ind...);
   }
-  typestr.push_back(ArrayInterfaceHandler::TypeChar<T>());
-  typestr += std::to_string(sizeof(T));
-  arr["typestr"] = typestr;
-  arr["version"] = 3;
-  std::string str;
-  Json::Dump(arr, &str);
-  return str;
+};
+
+template <int32_t D>
+inline void CheckArrayInterface(StringView key, ArrayInterface<D> const &array) {
+  CHECK(!array.valid.Data()) << "Meta info " << key << " should be dense, found validity mask";
 }
 }  // namespace xgboost
 #endif  // XGBOOST_DATA_ARRAY_INTERFACE_H_
diff --git a/src/data/data.cc b/src/data/data.cc
index acf6e47b9d21..a6b76ee2a39c 100644
--- a/src/data/data.cc
+++ b/src/data/data.cc
@@ -1,5 +1,5 @@
 /*!
- * Copyright 2015-2020 by Contributors
+ * Copyright 2015-2021 by Contributors
  * \file data.cc
  */
 #include <dmlc/registry.h>
@@ -24,6 +24,7 @@
 #include "../data/iterative_device_dmatrix.h"
 #include "file_iterator.h"
 
+#include "validation.h"
 #include "./sparse_page_source.h"
 #include "./sparse_page_dmatrix.h"
 
@@ -337,17 +338,6 @@ inline bool MetaTryLoadFloatInfo(const std::string& fname,
   return true;
 }
 
-void ValidateQueryGroup(std::vector<bst_group_t> const &group_ptr_) {
-  bool valid_query_group = true;
-  for (size_t i = 1; i < group_ptr_.size(); ++i) {
-    valid_query_group = valid_query_group && group_ptr_[i] >= group_ptr_[i - 1];
-    if (!valid_query_group) {
-      break;
-    }
-  }
-  CHECK(valid_query_group) << "Invalid group structure.";
-}
-
 // macro to dispatch according to specified pointer types
 #define DISPATCH_CONST_PTR(dtype, old_ptr, cast_ptr, proc)              \
   switch (dtype) {                                                      \
@@ -398,7 +388,7 @@ void MetaInfo::SetInfo(const char* key, const void* dptr, DataType dtype, size_t
     for (size_t i = 1; i < group_ptr_.size(); ++i) {
       group_ptr_[i] = group_ptr_[i - 1] + group_ptr_[i];
     }
-    ValidateQueryGroup(group_ptr_);
+    data::ValidateQueryGroup(group_ptr_);
   } else if (!std::strcmp(key, "qid")) {
     std::vector<uint32_t> query_ids(num, 0);
     DISPATCH_CONST_PTR(dtype, dptr, cast_dptr,
@@ -632,7 +622,7 @@ void MetaInfo::Validate(int32_t device) const {
 }
 
 #if !defined(XGBOOST_USE_CUDA)
-void MetaInfo::SetInfo(const char * c_key, std::string const& interface_str) {
+void MetaInfo::SetInfo(StringView key, std::string const& interface_str) {
   common::AssertGPUSupport();
 }
 #endif  // !defined(XGBOOST_USE_CUDA)
diff --git a/src/data/data.cu b/src/data/data.cu
index aee62d1b7fad..8e13db9ce751 100644
--- a/src/data/data.cu
+++ b/src/data/data.cu
@@ -9,84 +9,81 @@
 #include "xgboost/json.h"
 #include "array_interface.h"
 #include "../common/device_helpers.cuh"
+#include "../common/linalg_op.cuh"
 #include "device_adapter.cuh"
 #include "simple_dmatrix.h"
+#include "validation.h"
 
 namespace xgboost {
-
-void CopyInfoImpl(ArrayInterface column, HostDeviceVector<float>* out) {
-  auto SetDeviceToPtr = [](void* ptr) {
-    cudaPointerAttributes attr;
-    dh::safe_cuda(cudaPointerGetAttributes(&attr, ptr));
-    int32_t ptr_device = attr.device;
-    if (ptr_device >= 0) {
-      dh::safe_cuda(cudaSetDevice(ptr_device));
-    }
-    return ptr_device;
-  };
-  auto ptr_device = SetDeviceToPtr(column.data);
-
-  if (column.num_rows == 0) {
-    return;
-  }
-  out->SetDevice(ptr_device);
-
-  size_t size = column.num_rows * column.num_cols;
-  CHECK_NE(size, 0);
-  out->Resize(size);
-
-  auto p_dst = thrust::device_pointer_cast(out->DevicePointer());
-  dh::LaunchN(size, [=] __device__(size_t idx) {
-    size_t ridx = idx / column.num_cols;
-    size_t cidx = idx - (ridx * column.num_cols);
-    p_dst[idx] = column.GetElement(ridx, cidx);
-  });
-}
-
 namespace {
-auto SetDeviceToPtr(void *ptr) {
+auto SetDeviceToPtr(void const* ptr) {
   cudaPointerAttributes attr;
   dh::safe_cuda(cudaPointerGetAttributes(&attr, ptr));
   int32_t ptr_device = attr.device;
   dh::safe_cuda(cudaSetDevice(ptr_device));
   return ptr_device;
 }
-}  // anonymous namespace
 
-void CopyGroupInfoImpl(ArrayInterface column, std::vector<bst_group_t>* out) {
-  CHECK(column.type != ArrayInterface::kF4 && column.type != ArrayInterface::kF8)
+template <typename T, int32_t D>
+void CopyTensorInfoImpl(Json arr_interface, linalg::Tensor<T, D>* p_out) {
+  ArrayInterface<D> array(arr_interface);
+  if (array.n == 0) {
+    p_out->SetDevice(0);
+    return;
+  }
+  CHECK(array.valid.Size() == 0) << "Meta info like label or weight can not have missing value.";
+  auto ptr_device = SetDeviceToPtr(array.data);
+  p_out->SetDevice(ptr_device);
+
+  if (array.is_contiguous && array.type == ToDType<T>::kType) {
+    p_out->ModifyInplace([&](HostDeviceVector<T>* data, common::Span<size_t, D> shape) {
+      // set shape
+      std::copy(array.shape, array.shape + D, shape.data());
+      // set data
+      data->Resize(array.n);
+      dh::safe_cuda(cudaMemcpyAsync(data->DevicePointer(), array.data, array.n * sizeof(T),
+                                    cudaMemcpyDefault));
+    });
+    return;
+  }
+  p_out->Reshape(array.shape);
+  auto t = p_out->View(ptr_device);
+  linalg::ElementWiseKernelDevice(t, [=] __device__(size_t i, T) {
+    return linalg::detail::Apply(TypedIndex<T, D>{array}, linalg::UnravelIndex<D>(i, array.shape));
+  });
+}
+
+void CopyGroupInfoImpl(ArrayInterface<1> column, std::vector<bst_group_t>* out) {
+  CHECK(column.type != ArrayInterfaceHandler::kF4 && column.type != ArrayInterfaceHandler::kF8)
       << "Expected integer for group info.";
 
   auto ptr_device = SetDeviceToPtr(column.data);
   CHECK_EQ(ptr_device, dh::CurrentDevice());
-  dh::TemporaryArray<bst_group_t> temp(column.num_rows);
-  auto d_tmp = temp.data();
+  dh::TemporaryArray<bst_group_t> temp(column.Shape(0));
+  auto d_tmp = temp.data().get();
 
-  dh::LaunchN(column.num_rows, [=] __device__(size_t idx) {
-    d_tmp[idx] = column.GetElement<size_t>(idx, 0);
-  });
-  auto length = column.num_rows;
+  dh::LaunchN(column.Shape(0),
+              [=] __device__(size_t idx) { d_tmp[idx] = TypedIndex<size_t, 1>{column}(idx); });
+  auto length = column.Shape(0);
   out->resize(length + 1);
   out->at(0) = 0;
   thrust::copy(temp.data(), temp.data() + length, out->begin() + 1);
   std::partial_sum(out->begin(), out->end(), out->begin());
 }
 
-void CopyQidImpl(ArrayInterface array_interface,
-                 std::vector<bst_group_t> *p_group_ptr) {
+void CopyQidImpl(ArrayInterface<1> array_interface, std::vector<bst_group_t>* p_group_ptr) {
   auto &group_ptr_ = *p_group_ptr;
   auto it = dh::MakeTransformIterator<uint32_t>(
-      thrust::make_counting_iterator(0ul),
-      [array_interface] __device__(size_t i) {
-        return array_interface.GetElement<uint32_t>(i, 0);
+      thrust::make_counting_iterator(0ul), [array_interface] __device__(size_t i) {
+        return TypedIndex<uint32_t, 1>{array_interface}(i);
       });
   dh::caching_device_vector<bool> flag(1);
   auto d_flag = dh::ToSpan(flag);
   auto d = SetDeviceToPtr(array_interface.data);
   dh::LaunchN(1, [=] __device__(size_t) { d_flag[0] = true; });
-  dh::LaunchN(array_interface.num_rows - 1, [=] __device__(size_t i) {
-    if (array_interface.GetElement<uint32_t>(i, 0) >
-        array_interface.GetElement<uint32_t>(i + 1, 0)) {
+  dh::LaunchN(array_interface.Shape(0) - 1, [=] __device__(size_t i) {
+    auto typed = TypedIndex<uint32_t, 1>{array_interface};
+    if (typed(i) > typed(i + 1)) {
       d_flag[0] = false;
     }
   });
@@ -95,16 +92,16 @@ void CopyQidImpl(ArrayInterface array_interface,
                            cudaMemcpyDeviceToHost));
   CHECK(non_dec) << "`qid` must be sorted in increasing order along with data.";
   size_t bytes = 0;
-  dh::caching_device_vector<uint32_t> out(array_interface.num_rows);
-  dh::caching_device_vector<uint32_t> cnt(array_interface.num_rows);
+  dh::caching_device_vector<uint32_t> out(array_interface.Shape(0));
+  dh::caching_device_vector<uint32_t> cnt(array_interface.Shape(0));
   HostDeviceVector<int> d_num_runs_out(1, 0, d);
   cub::DeviceRunLengthEncode::Encode(
       nullptr, bytes, it, out.begin(), cnt.begin(),
-      d_num_runs_out.DevicePointer(), array_interface.num_rows);
+      d_num_runs_out.DevicePointer(), array_interface.Shape(0));
   dh::caching_device_vector<char> tmp(bytes);
   cub::DeviceRunLengthEncode::Encode(
       tmp.data().get(), bytes, it, out.begin(), cnt.begin(),
-      d_num_runs_out.DevicePointer(), array_interface.num_rows);
+      d_num_runs_out.DevicePointer(), array_interface.Shape(0));
 
   auto h_num_runs_out = d_num_runs_out.HostSpan()[0];
   group_ptr_.clear();
@@ -115,77 +112,56 @@ void CopyQidImpl(ArrayInterface array_interface,
   thrust::copy(cnt.begin(), cnt.begin() + h_num_runs_out,
                group_ptr_.begin() + 1);
 }
+}  // namespace
 
-namespace {
-// thrust::all_of tries to copy lambda function.
-struct LabelsCheck {
-  __device__ bool operator()(float y) { return ::isnan(y) || ::isinf(y); }
-};
-struct WeightsCheck {
-  __device__ bool operator()(float w) { return LabelsCheck{}(w) || w < 0; }  // NOLINT
-};
-}  // anonymous namespace
-
-void ValidateQueryGroup(std::vector<bst_group_t> const &group_ptr_);
-
-void MetaInfo::SetInfo(const char * c_key, std::string const& interface_str) {
-  Json j_interface = Json::Load({interface_str.c_str(), interface_str.size()});
-  ArrayInterface array_interface(interface_str);
-  std::string key{c_key};
-
-  CHECK(!array_interface.valid.Data())
-      << "Meta info " << key << " should be dense, found validity mask";
-  if (array_interface.num_rows == 0) {
-    return;
-  }
-
+void MetaInfo::SetInfo(StringView key, std::string const& interface_str) {
+  Json array = Json::Load(StringView{interface_str});
+  // multi-dim float info
   if (key == "base_margin") {
-    CopyInfoImpl(array_interface, &base_margin_);
+    // FIXME(jiamingy): This is temporary until #7405 can be fully merged
+    linalg::Tensor<float, 3> t;
+    CopyTensorInfoImpl(array, &t);
+    base_margin_ = std::move(*t.Data());
     return;
   }
-
-  CHECK(array_interface.num_cols == 1 || array_interface.num_rows == 1)
-      << "MetaInfo: " << c_key << " has invalid shape";
-  if (!((array_interface.num_cols == 1 && array_interface.num_rows == 0) ||
-        (array_interface.num_cols == 0 && array_interface.num_rows == 1))) {
-    // Not an empty column, transform it.
-    array_interface.AsColumnVector();
+  // uint info
+  if (key == "group") {
+    auto array_interface{ArrayInterface<1>(array)};
+    CopyGroupInfoImpl(array_interface, &group_ptr_);
+    data::ValidateQueryGroup(group_ptr_);
+    return;
+  } else if (key == "qid") {
+    auto array_interface{ArrayInterface<1>(array)};
+    CopyQidImpl(array_interface, &group_ptr_);
+    data::ValidateQueryGroup(group_ptr_);
+    return;
   }
+  // float info
+  linalg::Tensor<float, 1> t;
+  CopyTensorInfoImpl(array, &t);
   if (key == "label") {
-    CopyInfoImpl(array_interface, &labels_);
+    this->labels_ = std::move(*t.Data());
     auto ptr = labels_.ConstDevicePointer();
-    auto valid = thrust::none_of(thrust::device, ptr, ptr + labels_.Size(),
-                                 LabelsCheck{});
+    auto valid = thrust::none_of(thrust::device, ptr, ptr + labels_.Size(), data::LabelsCheck{});
     CHECK(valid) << "Label contains NaN, infinity or a value too large.";
   } else if (key == "weight") {
-    CopyInfoImpl(array_interface, &weights_);
+    this->weights_ = std::move(*t.Data());
     auto ptr = weights_.ConstDevicePointer();
-    auto valid = thrust::none_of(thrust::device, ptr, ptr + weights_.Size(),
-                                 WeightsCheck{});
+    auto valid = thrust::none_of(thrust::device, ptr, ptr + weights_.Size(), data::WeightsCheck{});
     CHECK(valid) << "Weights must be positive values.";
-  } else if (key == "group") {
-    CopyGroupInfoImpl(array_interface, &group_ptr_);
-    ValidateQueryGroup(group_ptr_);
-    return;
-  } else if (key == "qid") {
-    CopyQidImpl(array_interface, &group_ptr_);
-    return;
   } else if (key == "label_lower_bound") {
-    CopyInfoImpl(array_interface, &labels_lower_bound_);
-    return;
+    this->labels_lower_bound_ = std::move(*t.Data());
   } else if (key == "label_upper_bound") {
-    CopyInfoImpl(array_interface, &labels_upper_bound_);
-    return;
+    this->labels_upper_bound_ = std::move(*t.Data());
   } else if (key == "feature_weights") {
-    CopyInfoImpl(array_interface, &feature_weights);
+    this->feature_weights = std::move(*t.Data());
     auto d_feature_weights = feature_weights.ConstDeviceSpan();
-    auto valid = thrust::none_of(
-        thrust::device, d_feature_weights.data(),
-        d_feature_weights.data() + d_feature_weights.size(), WeightsCheck{});
+    auto valid =
+        thrust::none_of(thrust::device, d_feature_weights.data(),
+                        d_feature_weights.data() + d_feature_weights.size(), data::WeightsCheck{});
     CHECK(valid) << "Feature weight must be greater than 0.";
-    return;
   } else {
-    LOG(FATAL) << "Unknown metainfo: " << key;
+    LOG(FATAL) << "Unknown key for MetaInfo: " << key;
   }
 }
 
diff --git a/src/data/device_adapter.cuh b/src/data/device_adapter.cuh
index 628878f319f1..d1bda280a7d5 100644
--- a/src/data/device_adapter.cuh
+++ b/src/data/device_adapter.cuh
@@ -20,7 +20,7 @@ class CudfAdapterBatch : public detail::NoMetaInfo {
 
  public:
   CudfAdapterBatch() = default;
-  CudfAdapterBatch(common::Span<ArrayInterface> columns, size_t num_rows)
+  CudfAdapterBatch(common::Span<ArrayInterface<1>> columns, size_t num_rows)
       : columns_(columns),
         num_rows_(num_rows) {}
   size_t Size() const { return num_rows_ * columns_.size(); }
@@ -29,7 +29,7 @@ class CudfAdapterBatch : public detail::NoMetaInfo {
     size_t row_idx = idx / columns_.size();
     auto const& column = columns_[column_idx];
     float value = column.valid.Data() == nullptr || column.valid.Check(row_idx)
-                  ? column.GetElement(row_idx, 0)
+                      ? column(row_idx)
                       : std::numeric_limits<float>::quiet_NaN();
     return {row_idx, column_idx, value};
   }
@@ -38,7 +38,7 @@ class CudfAdapterBatch : public detail::NoMetaInfo {
   XGBOOST_DEVICE bst_row_t NumCols() const { return columns_.size(); }
 
  private:
-  common::Span<ArrayInterface> columns_;
+  common::Span<ArrayInterface<1>> columns_;
   size_t num_rows_;
 };
 
@@ -101,9 +101,9 @@ class CudfAdapter : public detail::SingleBatchDataIter<CudfAdapterBatch> {
 
     auto const& typestr = get<String const>(json_columns[0]["typestr"]);
     CHECK_EQ(typestr.size(), 3) << ArrayInterfaceErrors::TypestrFormat();
-    std::vector<ArrayInterface> columns;
-    auto first_column = ArrayInterface(get<Object const>(json_columns[0]));
-    num_rows_ = first_column.num_rows;
+    std::vector<ArrayInterface<1>> columns;
+    auto first_column = ArrayInterface<1>(get<Object const>(json_columns[0]));
+    num_rows_ = first_column.Shape(0);
     if (num_rows_ == 0) {
       return;
     }
@@ -112,13 +112,12 @@ class CudfAdapter : public detail::SingleBatchDataIter<CudfAdapterBatch> {
     CHECK_NE(device_idx_, -1);
     dh::safe_cuda(cudaSetDevice(device_idx_));
     for (auto& json_col : json_columns) {
-      auto column = ArrayInterface(get<Object const>(json_col));
+      auto column = ArrayInterface<1>(get<Object const>(json_col));
       columns.push_back(column);
-      CHECK_EQ(column.num_cols, 1);
-      num_rows_ = std::max(num_rows_, size_t(column.num_rows));
+      num_rows_ = std::max(num_rows_, size_t(column.Shape(0)));
       CHECK_EQ(device_idx_, dh::CudaGetPointerDevice(column.data))
           << "All columns should use the same device.";
-      CHECK_EQ(num_rows_, column.num_rows)
+      CHECK_EQ(num_rows_, column.Shape(0))
           << "All columns should have same number of rows.";
     }
     columns_ = columns;
@@ -135,7 +134,7 @@ class CudfAdapter : public detail::SingleBatchDataIter<CudfAdapterBatch> {
 
  private:
   CudfAdapterBatch batch_;
-  dh::device_vector<ArrayInterface> columns_;
+  dh::device_vector<ArrayInterface<1>> columns_;
   size_t num_rows_{0};
   int device_idx_;
 };
@@ -143,23 +142,23 @@ class CudfAdapter : public detail::SingleBatchDataIter<CudfAdapterBatch> {
 class CupyAdapterBatch : public detail::NoMetaInfo {
  public:
   CupyAdapterBatch() = default;
-  explicit CupyAdapterBatch(ArrayInterface array_interface)
+  explicit CupyAdapterBatch(ArrayInterface<2> array_interface)
     : array_interface_(std::move(array_interface)) {}
   size_t Size() const {
-    return array_interface_.num_rows * array_interface_.num_cols;
+    return array_interface_.Shape(0) * array_interface_.Shape(1);
   }
   __device__ COOTuple GetElement(size_t idx) const {
-    size_t column_idx = idx % array_interface_.num_cols;
-    size_t row_idx = idx / array_interface_.num_cols;
-    float value = array_interface_.GetElement(row_idx, column_idx);
+    size_t column_idx = idx % array_interface_.Shape(1);
+    size_t row_idx = idx / array_interface_.Shape(1);
+    float value = array_interface_(row_idx, column_idx);
     return {row_idx, column_idx, value};
   }
 
-  XGBOOST_DEVICE bst_row_t NumRows() const { return array_interface_.num_rows; }
-  XGBOOST_DEVICE bst_row_t NumCols() const { return array_interface_.num_cols; }
+  XGBOOST_DEVICE bst_row_t NumRows() const { return array_interface_.Shape(0); }
+  XGBOOST_DEVICE bst_row_t NumCols() const { return array_interface_.Shape(1); }
 
  private:
-  ArrayInterface array_interface_;
+  ArrayInterface<2> array_interface_;
 };
 
 class CupyAdapter : public detail::SingleBatchDataIter<CupyAdapterBatch> {
@@ -167,9 +166,9 @@ class CupyAdapter : public detail::SingleBatchDataIter<CupyAdapterBatch> {
   explicit CupyAdapter(std::string cuda_interface_str) {
     Json json_array_interface =
         Json::Load({cuda_interface_str.c_str(), cuda_interface_str.size()});
-    array_interface_ = ArrayInterface(get<Object const>(json_array_interface), false);
+    array_interface_ = ArrayInterface<2>(get<Object const>(json_array_interface));
     batch_ = CupyAdapterBatch(array_interface_);
-    if (array_interface_.num_rows == 0) {
+    if (array_interface_.Shape(0) == 0) {
       return;
     }
     device_idx_ = dh::CudaGetPointerDevice(array_interface_.data);
@@ -177,12 +176,12 @@ class CupyAdapter : public detail::SingleBatchDataIter<CupyAdapterBatch> {
   }
   const CupyAdapterBatch& Value() const override { return batch_; }
 
-  size_t NumRows() const { return array_interface_.num_rows; }
-  size_t NumColumns() const { return array_interface_.num_cols; }
+  size_t NumRows() const { return array_interface_.Shape(0); }
+  size_t NumColumns() const { return array_interface_.Shape(1); }
   int32_t DeviceIdx() const { return device_idx_; }
 
  private:
-  ArrayInterface array_interface_;
+  ArrayInterface<2> array_interface_;
   CupyAdapterBatch batch_;
   int32_t device_idx_ {-1};
 };
diff --git a/src/data/file_iterator.h b/src/data/file_iterator.h
index 6d6adb62b008..70a5d51c30b9 100644
--- a/src/data/file_iterator.h
+++ b/src/data/file_iterator.h
@@ -12,6 +12,7 @@
 #include "dmlc/data.h"
 #include "xgboost/c_api.h"
 #include "xgboost/json.h"
+#include "xgboost/linalg.h"
 #include "array_interface.h"
 
 namespace xgboost {
@@ -58,16 +59,14 @@ class FileIterator {
     CHECK(parser_);
     if (parser_->Next()) {
       row_block_ = parser_->Value();
+      using linalg::MakeVec;
 
-      indptr_ = MakeArrayInterface(row_block_.offset, row_block_.size + 1);
-      values_ = MakeArrayInterface(row_block_.value,
-                                   row_block_.offset[row_block_.size]);
-      indices_ = MakeArrayInterface(row_block_.index,
-                                    row_block_.offset[row_block_.size]);
+      indptr_ = MakeVec(row_block_.offset, row_block_.size + 1).ArrayInterfaceStr();
+      values_ = MakeVec(row_block_.value, row_block_.offset[row_block_.size]).ArrayInterfaceStr();
+      indices_ = MakeVec(row_block_.index, row_block_.offset[row_block_.size]).ArrayInterfaceStr();
 
-      size_t n_columns = *std::max_element(
-          row_block_.index,
-          row_block_.index + row_block_.offset[row_block_.size]);
+      size_t n_columns = *std::max_element(row_block_.index,
+                                           row_block_.index + row_block_.offset[row_block_.size]);
       // dmlc parser converts 1-based indexing back to 0-based indexing so we can ignore
       // this condition and just add 1 to n_columns
       n_columns += 1;
diff --git a/src/data/validation.h b/src/data/validation.h
new file mode 100644
index 000000000000..6d3701114886
--- /dev/null
+++ b/src/data/validation.h
@@ -0,0 +1,40 @@
+/*!
+ * Copyright 2021 by XGBoost Contributors
+ */
+#ifndef XGBOOST_DATA_VALIDATION_H_
+#define XGBOOST_DATA_VALIDATION_H_
+#include <cmath>
+#include <vector>
+
+#include "xgboost/base.h"
+#include "xgboost/logging.h"
+
+namespace xgboost {
+namespace data {
+struct LabelsCheck {
+  XGBOOST_DEVICE bool operator()(float y) {
+#if defined(__CUDA_ARCH__)
+    return ::isnan(y) || ::isinf(y);
+#else
+    return std::isnan(y) || std::isinf(y);
+#endif
+  }
+};
+
+struct WeightsCheck {
+  XGBOOST_DEVICE bool operator()(float w) { return LabelsCheck{}(w) || w < 0; }  // NOLINT
+};
+
+inline void ValidateQueryGroup(std::vector<bst_group_t> const &group_ptr_) {
+  bool valid_query_group = true;
+  for (size_t i = 1; i < group_ptr_.size(); ++i) {
+    valid_query_group = valid_query_group && group_ptr_[i] >= group_ptr_[i - 1];
+    if (XGBOOST_EXPECT(!valid_query_group, false)) {
+      break;
+    }
+  }
+  CHECK(valid_query_group) << "Invalid group structure.";
+}
+}  // namespace data
+}  // namespace xgboost
+#endif  // XGBOOST_DATA_VALIDATION_H_
diff --git a/tests/cpp/data/test_adapter.cc b/tests/cpp/data/test_adapter.cc
index ccb19de71a74..c6de226194f6 100644
--- a/tests/cpp/data/test_adapter.cc
+++ b/tests/cpp/data/test_adapter.cc
@@ -41,9 +41,10 @@ TEST(Adapter, CSRArrayAdapter) {
   HostDeviceVector<bst_feature_t> indices;
   size_t n_features = 100, n_samples = 10;
   RandomDataGenerator{n_samples, n_features, 0.5}.GenerateCSR(&values, &indptr, &indices);
-  auto indptr_arr = MakeArrayInterface(indptr.HostPointer(), indptr.Size());
-  auto values_arr = MakeArrayInterface(values.HostPointer(), values.Size());
-  auto indices_arr = MakeArrayInterface(indices.HostPointer(), indices.Size());
+  using linalg::MakeVec;
+  auto indptr_arr = MakeVec(indptr.HostPointer(), indptr.Size()).ArrayInterfaceStr();
+  auto values_arr = MakeVec(values.HostPointer(), values.Size()).ArrayInterfaceStr();
+  auto indices_arr = MakeVec(indices.HostPointer(), indices.Size()).ArrayInterfaceStr();
   auto adapter = data::CSRArrayAdapter(
       StringView{indptr_arr.c_str(), indptr_arr.size()},
       StringView{values_arr.c_str(), values_arr.size()},
diff --git a/tests/cpp/data/test_array_interface.cc b/tests/cpp/data/test_array_interface.cc
index 875858855ed5..8efd30eb75b2 100644
--- a/tests/cpp/data/test_array_interface.cc
+++ b/tests/cpp/data/test_array_interface.cc
@@ -11,21 +11,22 @@ TEST(ArrayInterface, Initialize) {
   size_t constexpr kRows = 10, kCols = 10;
   HostDeviceVector<float> storage;
   auto array = RandomDataGenerator{kRows, kCols, 0}.GenerateArrayInterface(&storage);
-  auto arr_interface = ArrayInterface(array);
-  ASSERT_EQ(arr_interface.num_rows, kRows);
-  ASSERT_EQ(arr_interface.num_cols, kCols);
+  auto arr_interface = ArrayInterface<2>(StringView{array});
+  ASSERT_EQ(arr_interface.Shape(0), kRows);
+  ASSERT_EQ(arr_interface.Shape(1), kCols);
   ASSERT_EQ(arr_interface.data, storage.ConstHostPointer());
   ASSERT_EQ(arr_interface.ElementSize(), 4);
-  ASSERT_EQ(arr_interface.type, ArrayInterface::kF4);
+  ASSERT_EQ(arr_interface.type, ArrayInterfaceHandler::kF4);
 
   HostDeviceVector<size_t> u64_storage(storage.Size());
-  std::string u64_arr_str;
-  Json::Dump(GetArrayInterface(&u64_storage, kRows, kCols), &u64_arr_str);
+  std::string u64_arr_str{linalg::TensorView<size_t const, 2>{
+      u64_storage.ConstHostSpan(), {kRows, kCols}, GenericParameter::kCpuId}
+                              .ArrayInterfaceStr()};
   std::copy(storage.ConstHostVector().cbegin(), storage.ConstHostVector().cend(),
             u64_storage.HostSpan().begin());
-  auto u64_arr = ArrayInterface{u64_arr_str};
+  auto u64_arr = ArrayInterface<2>{u64_arr_str};
   ASSERT_EQ(u64_arr.ElementSize(), 8);
-  ASSERT_EQ(u64_arr.type, ArrayInterface::kU8);
+  ASSERT_EQ(u64_arr.type, ArrayInterfaceHandler::kU8);
 }
 
 TEST(ArrayInterface, Error) {
@@ -38,23 +39,22 @@ TEST(ArrayInterface, Error) {
         Json(Boolean(false))};
 
   auto const& column_obj = get<Object>(column);
-  std::pair<size_t, size_t> shape{kRows, kCols};
   std::string typestr{"<f4"};
+  size_t n = kRows * kCols;
 
   // missing version
-  EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, shape),
-               dmlc::Error);
-  column["version"] = Integer(static_cast<Integer::Int>(1));
+  EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, n), dmlc::Error);
+  column["version"] = 3;
   // missing data
-  EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, shape),
+  EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, n),
                dmlc::Error);
   column["data"] = j_data;
   // missing typestr
-  EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, shape),
+  EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, n),
                dmlc::Error);
   column["typestr"] = String("<f4");
   // nullptr is not valid
-  EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, shape),
+  EXPECT_THROW(ArrayInterfaceHandler::ExtractData(column_obj, n),
                dmlc::Error);
 
   HostDeviceVector<float> storage;
@@ -63,22 +63,52 @@ TEST(ArrayInterface, Error) {
       Json(Integer(reinterpret_cast<Integer::Int>(storage.ConstHostPointer()))),
       Json(Boolean(false))};
   column["data"] = j_data;
-  EXPECT_NO_THROW(ArrayInterfaceHandler::ExtractData(column_obj, shape));
+  EXPECT_NO_THROW(ArrayInterfaceHandler::ExtractData(column_obj, n));
 }
 
 TEST(ArrayInterface, GetElement) {
   size_t kRows = 4, kCols = 2;
   HostDeviceVector<float> storage;
   auto intefrace_str = RandomDataGenerator{kRows, kCols, 0}.GenerateArrayInterface(&storage);
-  ArrayInterface array_interface{intefrace_str};
+  ArrayInterface<2> array_interface{intefrace_str};
 
   auto const& h_storage = storage.ConstHostVector();
   for (size_t i = 0; i < kRows; ++i) {
     for (size_t j = 0; j < kCols; ++j) {
-      float v0 = array_interface.GetElement(i, j);
+      float v0 = array_interface(i, j);
       float v1 = h_storage.at(i * kCols + j);
       ASSERT_EQ(v0, v1);
     }
   }
 }
+
+TEST(ArrayInterface, TrivialDim) {
+  size_t kRows{1000}, kCols = 1;
+  HostDeviceVector<float> storage;
+  auto interface_str = RandomDataGenerator{kRows, kCols, 0}.GenerateArrayInterface(&storage);
+  {
+    ArrayInterface<1> arr_i{interface_str};
+    ASSERT_EQ(arr_i.n, kRows);
+    ASSERT_EQ(arr_i.Shape(0), kRows);
+  }
+
+  std::swap(kRows, kCols);
+  interface_str = RandomDataGenerator{kRows, kCols, 0}.GenerateArrayInterface(&storage);
+  {
+    ArrayInterface<1> arr_i{interface_str};
+    ASSERT_EQ(arr_i.n, kCols);
+    ASSERT_EQ(arr_i.Shape(0), kCols);
+  }
+}
+
+TEST(ArrayInterface, ToDType) {
+  static_assert(ToDType<float>::kType == ArrayInterfaceHandler::kF4, "");
+  static_assert(ToDType<double>::kType == ArrayInterfaceHandler::kF8, "");
+
+  static_assert(ToDType<uint32_t>::kType == ArrayInterfaceHandler::kU4, "");
+  static_assert(ToDType<uint64_t>::kType == ArrayInterfaceHandler::kU8, "");
+
+  static_assert(ToDType<int32_t>::kType == ArrayInterfaceHandler::kI4, "");
+  static_assert(ToDType<int64_t>::kType == ArrayInterfaceHandler::kI8, "");
+}
 }  // namespace xgboost
diff --git a/tests/cpp/data/test_array_interface.cu b/tests/cpp/data/test_array_interface.cu
index 75923e74ba1a..c8e07852534b 100644
--- a/tests/cpp/data/test_array_interface.cu
+++ b/tests/cpp/data/test_array_interface.cu
@@ -32,11 +32,24 @@ TEST(ArrayInterface, Stream) {
   dh::caching_device_vector<uint64_t> out(1, 0);
   uint64_t dur = 1e9;
   dh::LaunchKernel{1, 1, 0, stream}(SleepForTest, out.data().get(), dur);
-  ArrayInterface arr(arr_str);
+  ArrayInterface<2> arr(arr_str);
 
   auto t = out[0];
   CHECK_GE(t, dur);
 
   cudaStreamDestroy(stream);
 }
+
+TEST(ArrayInterface, Ptr) {
+  std::vector<float> h_data(10);
+  ASSERT_FALSE(ArrayInterfaceHandler::IsCudaPtr(h_data.data()));
+  dh::safe_cuda(cudaGetLastError());
+
+  dh::device_vector<float> d_data(10);
+  ASSERT_TRUE(ArrayInterfaceHandler::IsCudaPtr(d_data.data().get()));
+  dh::safe_cuda(cudaGetLastError());
+
+  ASSERT_FALSE(ArrayInterfaceHandler::IsCudaPtr(nullptr));
+  dh::safe_cuda(cudaGetLastError());
+}
 }  // namespace xgboost
diff --git a/tests/cpp/data/test_array_interface.h b/tests/cpp/data/test_array_interface.h
index 7872a9507aa5..78bce76f53e7 100644
--- a/tests/cpp/data/test_array_interface.h
+++ b/tests/cpp/data/test_array_interface.h
@@ -19,6 +19,7 @@ Json GenerateDenseColumn(std::string const& typestr, size_t kRows,
   std::vector<Json> j_shape {Json(Integer(static_cast<Integer::Int>(kRows)))};
   column["shape"] = Array(j_shape);
   column["strides"] = Array(std::vector<Json>{Json(Integer(static_cast<Integer::Int>(sizeof(T))))});
+  column["stream"] = nullptr;
 
   d_data.resize(kRows);
   thrust::sequence(thrust::device, d_data.begin(), d_data.end(), 0.0f, 2.0f);
@@ -30,7 +31,7 @@ Json GenerateDenseColumn(std::string const& typestr, size_t kRows,
         Json(Boolean(false))};
   column["data"] = j_data;
 
-  column["version"] = Integer(static_cast<Integer::Int>(1));
+  column["version"] = 3;
   column["typestr"] = String(typestr);
   return column;
 }
@@ -43,6 +44,7 @@ Json GenerateSparseColumn(std::string const& typestr, size_t kRows,
   std::vector<Json> j_shape {Json(Integer(static_cast<Integer::Int>(kRows)))};
   column["shape"] = Array(j_shape);
   column["strides"] = Array(std::vector<Json>{Json(Integer(static_cast<Integer::Int>(sizeof(T))))});
+  column["stream"] = nullptr;
 
   d_data.resize(kRows);
   for (size_t i = 0; i < d_data.size(); ++i) {
@@ -56,7 +58,7 @@ Json GenerateSparseColumn(std::string const& typestr, size_t kRows,
         Json(Boolean(false))};
   column["data"] = j_data;
 
-  column["version"] = Integer(static_cast<Integer::Int>(1));
+  column["version"] = 3;
   column["typestr"] = String(typestr);
   return column;
 }
@@ -75,9 +77,9 @@ Json Generate2dArrayInterface(int rows, int cols, std::string typestr,
       Json(Integer(reinterpret_cast<Integer::Int>(data.data().get()))),
       Json(Boolean(false))};
   array_interface["data"] = j_data;
-  array_interface["version"] = Integer(static_cast<Integer::Int>(1));
+  array_interface["version"] = 3;
   array_interface["typestr"] = String(typestr);
+  array_interface["stream"] = nullptr;
   return array_interface;
 }
-
 }  // namespace xgboost
diff --git a/tests/cpp/data/test_iterative_device_dmatrix.cu b/tests/cpp/data/test_iterative_device_dmatrix.cu
index cb64a3b5cdb2..0fc992f24187 100644
--- a/tests/cpp/data/test_iterative_device_dmatrix.cu
+++ b/tests/cpp/data/test_iterative_device_dmatrix.cu
@@ -103,9 +103,9 @@ TEST(IterativeDeviceDMatrix, RowMajor) {
 
     auto j_interface =
         Json::Load({interface_str.c_str(), interface_str.size()});
-    ArrayInterface loaded {get<Object const>(j_interface)};
+    ArrayInterface<2> loaded {get<Object const>(j_interface)};
     std::vector<float> h_data(cols * rows);
-    common::Span<float> s_data{static_cast<float*>(loaded.data), cols * rows};
+    common::Span<float const> s_data{static_cast<float const*>(loaded.data), cols * rows};
     dh::CopyDeviceSpanToVector(&h_data, s_data);
 
     for(auto i = 0ull; i < rows * cols; i++) {
@@ -128,9 +128,9 @@ TEST(IterativeDeviceDMatrix, RowMajorMissing) {
   std::string interface_str = iter.AsArray();
   auto j_interface =
       Json::Load({interface_str.c_str(), interface_str.size()});
-  ArrayInterface loaded {get<Object const>(j_interface)};
+  ArrayInterface<2> loaded {get<Object const>(j_interface)};
   std::vector<float> h_data(cols * rows);
-  common::Span<float> s_data{static_cast<float*>(loaded.data), cols * rows};
+  common::Span<float const> s_data{static_cast<float const*>(loaded.data), cols * rows};
   dh::CopyDeviceSpanToVector(&h_data, s_data);
   h_data[1] = kMissing;
   h_data[5] = kMissing;
diff --git a/tests/cpp/data/test_metainfo.cu b/tests/cpp/data/test_metainfo.cu
index 090374b913d6..205844a5e961 100644
--- a/tests/cpp/data/test_metainfo.cu
+++ b/tests/cpp/data/test_metainfo.cu
@@ -1,4 +1,4 @@
-/*! Copyright 2019 by Contributors */
+/*! Copyright 2019-2021 by XGBoost Contributors */
 
 #include <gtest/gtest.h>
 #include <xgboost/data.h>
diff --git a/tests/cpp/helpers.h b/tests/cpp/helpers.h
index c424d65ced05..2ccbd8a77e83 100644
--- a/tests/cpp/helpers.h
+++ b/tests/cpp/helpers.h
@@ -198,7 +198,7 @@ Json GetArrayInterface(HostDeviceVector<T> *storage, size_t rows, size_t cols) {
   array_interface["shape"][0] = rows;
   array_interface["shape"][1] = cols;
 
-  char t = ArrayInterfaceHandler::TypeChar<T>();
+  char t = linalg::detail::ArrayInterfaceHandler::TypeChar<T>();
   array_interface["typestr"] = String(std::string{"<"} + t + std::to_string(sizeof(T)));
   array_interface["version"] = 1;
   return array_interface;