   ptrdiff_t strides_[ndim];
 
 public:
+  // Constructor supports pointers or initializer lists and strides is optional
+  // as long as _c_contiguous=true.
+  //    CArray<T, 3> ca(data, shape, strides);
+  //    CArray<T, 3, true> ca(data, shape);
+  //    CArray<T, 3> ca(data, {1, 2, 3}, {48, 24, 8});
+  //    CArray<T, 3, true> ca(data, {1, 2, 3});
+  // Initializer lists and optional strides are only supported with -std=c++11
+  // or higher.
+
+  template <typename Int1, typename Int2>
+  __device__ CArray(T* data, const Int1* shape, const Int2* strides)
+      : data_(data), size_(1)
+  {
+    if (_c_contiguous) {
+      assert(strides[_ndim-1] == sizeof(T));
+      for (int i = _ndim-1; i > 0; i--) {
+        assert(strides[i-1] == shape[i] * strides[i]);
+      }
+    }
+    for (int i = 0; i < _ndim; i++) {
+      this->size_ *= shape[i];
+      this->shape_[i] = shape[i];
+      this->strides_[i] = strides[i];
+    }
+  }
+
+#if __cplusplus >= 201103
+  template <typename Int, typename U=T>
+  __device__ CArray(typename std::enable_if<_c_contiguous, U>::type* data,
+                    const Int* shape)
+      : data_(data), size_(1)
+  {
+    for (int i = 0; i < _ndim; i++) {
+      this->size_ *= shape[i];
+      this->shape_[i] = shape[i];
+    }
+    this->strides_[_ndim-1] = sizeof(T);
+    for (int i = _ndim-1; i > 0; i--) {
+      this->strides_[i-1] = shape[i] * this->strides_[i];
+    }
+  }
+
+  template <typename Int, typename U=T>
+  __device__ CArray(typename std::enable_if<_c_contiguous, U>::type* data,
+                    const std::initializer_list<Int> shape)
+      : CArray(data, shape.begin())
+  {
+    assert(shape.size() == _ndim);
+  }
+
+  template <typename Int1, typename Int2>
+  __device__ CArray(T* data,
+                    const std::initializer_list<Int1> shape,
+                    const std::initializer_list<Int2> strides)
+      : CArray(data, shape.begin(), strides.begin())
+  {
+    assert(shape.size() == _ndim);
+    assert(strides.size() == _ndim);
+  }
+#endif
+
   __device__ ptrdiff_t size() const {
     return size_;
   }

public:
  static const int ndim = 0;

  __device__ ptrdiff_t size() const {
    return size_;
  }

  __device__ const ptrdiff_t* strides() const {
    return strides_;
  }

#if __cplusplus >= 201103
  template <typename Int>
  __device__ T& operator[](const std::initializer_list<Int> idx_) {
    assert(idx_.size() == _ndim);
    Int idx[ndim];
    memcpy(idx, idx_.begin(), ndim*sizeof(Int));
    return this->operator[](idx);
  }

  template <typename Int>
  __device__ const T& operator[](const std::initializer_list<Int> idx_) const {
    assert(idx_.size() == _ndim);
    Int idx[ndim];
    memcpy(idx, idx_.begin(), ndim*sizeof(Int));
    return this->operator[](idx);
  }
#endif

  template <typename Int>
  __device__ T& operator[](const Int (&idx)[ndim]) {
