       // contiguous arrays can be directly addressed by the
       // numeric value, avoiding expensive 64 bit operations in cuda
       return data_[idx];
-    }
-    // 64-bit mults and divs are pretty expensive and can 
-    // lead to severe perforamance degradation in computation bound
-    // kernels
-    const char* ptr = reinterpret_cast<const char*>(data_);
-    if (use_32bit_indexing) {
-      int i = static_cast<int>(idx);
+    } else {
+      // 64-bit mults and divs are pretty expensive and can lead to severe
+      // performance degradation in computation bound kernels
+      const char* ptr = reinterpret_cast<const char*>(data_);
+      index_t i = static_cast<index_t>(idx);
       for (int dim = ndim; --dim > 0; ) {
-        int shape_dim = static_cast<int>(shape_[dim]);
-        ptr += static_cast<int>(strides_[dim]) * (i % shape_dim);
-        i /= shape_dim;
+          index_t shape_dim = static_cast<index_t>(shape_[dim]);
+          ptr += static_cast<index_t>(strides_[dim]) * (i % shape_dim);
+          i /= shape_dim;
       }
-      if (ndim > 0) {
-        ptr += static_cast<int>(strides_[0]) * i;
-      }
-    } else {
-      ptrdiff_t i = idx;
-      for (int dim = ndim; --dim > 0; ) {
-        ptr += static_cast<ptrdiff_t>(strides_[dim]) * (i % shape_[dim]);
-        i /= shape_[dim];
-      }
-      if (ndim > 0) {
-        ptr += static_cast<ptrdiff_t>(strides_[0]) * i;
-      }
-    }
-    return *reinterpret_cast<const T*>(ptr);
+      ptr += static_cast<index_t>(strides_[0]) * i;
+      return *reinterpret_cast<const T*>(ptr);
+    }
   }
 };
 

  template <typename Int>
  __device__ const T& operator[](const Int (&idx)[ndim]) const {
    const char* ptr = reinterpret_cast<const char*>(data_);
    if (use_32bit_indexing) {
      for (int dim = 0; dim < ndim; ++dim) {
        int i = static_cast<int>(idx[dim]);
        ptr += static_cast<int>(strides_[dim]) * i;
      }
    } else {
      for (int dim = 0; dim < ndim; ++dim) {
        ptr += static_cast<ptrdiff_t>(strides_[dim]) * idx[dim];
      }
    }
    return *reinterpret_cast<const T*>(ptr);
  }

// in NVRTC std:initializer_list is pre-defined (no need to include it)
#include <initializer_list>
#endif
#else
// Basic implementation of std::conditional when not using C++11
namespace std {
  template<bool B, class T, class F>
  struct conditional { typedef T type; };
  template<class T, class F>
  struct conditional<false, T, F> { typedef F type; };
}
#endif

// math
