     T ret = d[0];
     for (int i = 1; i < dim; i++) {
       ret += d[i];
+    }
+    return ret;
+  }
+
+  TC_FORCE_INLINE T __device__ __host__ min() const {
+    T ret = d[0];
+    for (int i = 1; i < dim; i++) {
+      ret = ret < d[i] ? ret : d[i];
+    }
+    return ret;
+  }
+
+  TC_FORCE_INLINE T __device__ __host__ max() const {
+    T ret = d[0];
+    for (int i = 1; i < dim; i++) {
+      ret = ret > d[i] ? ret : d[i];
     }
     return ret;
   }

  return ret;
}

TC_FORCE_INLINE __host__ __device__ float fract(float x) {
  return x - floorf(x);
}

    TVector<int, dim> res,
    TVector<int, dim> idx) {
  int ret = idx[0];
  for (int i = 1; i < dim; i++) {
    ret *= res[i];
    ret += idx[i];
