     return ret;
   }
 
-  __device__ T dot(TVector &other) const {
+  TC_FORCE_INLINE __host__ __device__ T dot(TVector &other) const {
     T ret = 0;
     for (int i = 0; i < dim; i++) {
       ret += d[i] * other[i];
+    }
+    return ret;
+  }
+
+  TC_FORCE_INLINE T __host__ __device__ __host__ prod() {
+    T ret = d[0];
+    for (int i = 1; i < dim; i++) {
+      ret *= d[i];
+    }
+    return ret;
+  }
+
+  TC_FORCE_INLINE T __device__ __host__ sum() {
+    T ret = d[0];
+    for (int i = 1; i < dim; i++) {
+      ret += d[i];
     }
     return ret;
   }

  }

  template <int dim__ = dim_>
  TC_FORCE_INLINE __device__ TVector(T x, T y) {
    static_assert(dim__ == 2, "");
    d[0] = x;
    d[1] = y;
  }

  template <int dim__ = dim_>
  TC_FORCE_INLINE __device__ TVector(T x, T y, T z) {
    static_assert(dim__ == 3, "");
    d[0] = x;
    d[1] = y;
    d[2] = z;
  }

  TC_FORCE_INLINE __device__ TVector(T x = 0) {
    for (int i = 0; i < dim; i++) {
      d[i] = x;
    }

    return d[i];
  }

  TC_FORCE_INLINE __host__ __device__ TVector &operator+=(const TVector &o) {
    for (int i = 0; i < dim; i++) {
      d[i] += o[i];
    }
    return *this;
  }

  TC_FORCE_INLINE __host__ __device__ TVector operator+(const TVector &o) {
    TVector ret;
    for (int i = 0; i < dim; i++) {
      ret[i] = d[i] + o[i];
