New time-fused LSTM API. Performance improvements.

This change adds a new LSTM API that fuses operations across RNN
time steps. It performs significantly faster than the existing
iterative API.

The existing iterative LSTM API also received performance
improvements.

BREAKING CHANGES:
  Previously, callers were expected to transpose `h` before passing
  it to `BackwardPass`. Callers must not transpose `h` anymore.

  The `dv` parameter in `BackwardPass` has been removed and `v` must
  now be mutable.

Author: sharvil

examples/device_ptr.h

@@ -23,6 +23,10 @@ template<typename T>
 struct device_ptr {
   static constexpr size_t ElemSize = sizeof(typename T::Scalar);
 
+  static device_ptr<T> NewByteSized(size_t bytes) {
+    return device_ptr<T>((bytes + ElemSize - 1) / ElemSize);
+  }
+
   explicit device_ptr(size_t size_)
       : data(nullptr), size(size_) {
     void* tmp;
@@ -38,6 +42,24 @@ struct device_ptr {
     ToDevice(elem);
   }
 
+  device_ptr(device_ptr<T>&& other) : data(other.data), size(other.size) {
+    other.data = nullptr;
+    other.size = 0;
+  }
+
+  device_ptr& operator=(const device_ptr<T>&& other) {
+    if (&other != this) {
+      data = other.data;
+      size = other.size;
+      other.data = nullptr;
+      other.size = 0;
+    }
+    return *this;
+  }
+
+  device_ptr(const device_ptr<T>& other) = delete;
+  device_ptr& operator=(const device_ptr<T>& other) = delete;
+
   void ToDevice(const T& src) {
     assert(size == src.size());
     cudaMemcpy(data, src.data(), src.size() * ElemSize, cudaMemcpyHostToDevice);
@@ -48,6 +70,10 @@ struct device_ptr {
     cudaMemcpy(target.data(), data, target.size() * ElemSize, cudaMemcpyDeviceToHost);
   }
 
+  size_t Size() const {
+    return size;
+  }
+
   void zero() {
     cudaMemset(data, 0, size * ElemSize);
   }
@@ -57,5 +83,5 @@ struct device_ptr {
   }
 
   typename T::Scalar* data;
-  const size_t size;
+  size_t size;
 };