     }
   }
 
-  __device__ const ptrdiff_t* get() const {
+  __device__ const index_t& get() const {
     return index_;
   }
 };

    return *reinterpret_cast<T*>(ptr);
  }

  __device__ T operator[](const int* idx) const {
    return (*const_cast<CArray<T, ndim>*>(this))[idx];
  }

  __device__ T& operator[](const ptrdiff_t* idx) {
    char* ptr = reinterpret_cast<char*>(data_);
    for (int dim = 0; dim < ndim; ++dim) {
      ptr += strides_[dim] * idx[dim];
    }
    return *reinterpret_cast<T*>(ptr);
  }

  __device__ T operator[](const ptrdiff_t* idx) const {
    return (*const_cast<CArray<T, ndim>*>(this))[idx];
  }


  ptrdiff_t size_;
  ptrdiff_t shape_[ndim];
  ptrdiff_t index_[ndim];

  typedef ptrdiff_t index_t[ndim];

public:
  __device__ ptrdiff_t size() const {
