 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];
   }
 };

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

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

  __device__ T& operator[](ssize_t i) {
    char* ptr = reinterpret_cast<char*>(data_);
    for (int dim = ndim; --dim > 0; ) {
      ptr += strides_[dim] * (i % shape_[dim]);


// CArray
#define CUPY_FOR(i, n) \
    for (ptrdiff_t i = blockIdx.x * blockDim.x + threadIdx.x; \
         i < (n); \
         i += blockDim.x * gridDim.x)

