 class CArray<T, 0> {
 private:
   T* data_;
-  ssize_t size_;
-
-public:
-  __device__ ssize_t size() const {
-    return size_;
-  }
-
-  __device__ T& operator[](const ssize_t* idx) {
+  ptrdiff_t size_;
+
+public:
+  __device__ ptrdiff_t size() const {
+    return size_;
+  }
+
+  __device__ T& operator[](const ptrdiff_t* idx) {
     return *reinterpret_cast<T*>(data_);
   }
 
-  __device__ T operator[](const ssize_t* idx) const {
+  __device__ T operator[](const ptrdiff_t* idx) const {
     return (*const_cast<CArray<T, 0>*>(this))[idx];
   }
 
-  __device__ T& operator[](ssize_t i) {
+  __device__ T& operator[](ptrdiff_t i) {
     return *reinterpret_cast<T*>(data_);
   }
 
-  __device__ T operator[](ssize_t i) const {
+  __device__ T operator[](ptrdiff_t i) const {
     return (*const_cast<CArray<T, 0>*>(this))[i];
   }
 };

template <>
class CIndexer<0> {
private:
  ssize_t size_;

public:
  __device__ ssize_t size() const {
    return size_;
  }

  __device__ void set(ssize_t i) {
  }

  __device__ const ssize_t* get() const {
    return NULL;
  }
};

template <int ndim>
class CIndexer {
private:
  ptrdiff_t size_;
  ptrdiff_t shape_[ndim];
  ptrdiff_t index_[ndim];

public:
  __device__ ptrdiff_t size() const {
    return size_;
  }

  __device__ void set(ptrdiff_t i) {
    ptrdiff_t a = i;
    for (ptrdiff_t dim = ndim; --dim > 0; ) {
      ptrdiff_t s = shape_[dim];
      index_[dim] = (a % s);
      a /= s;
    }
