Repository: IBM/aihwkit (language: python)

Issue:
State-dependent short-term PCM like read noise 
## Description and motivation

PCM shows a state-dependency in the read noise, this is currently not modelled. 

## Proposed solution

Add a new weight noise type, that reversibly adds output noise that is proportional to the flowing current in a way reasonable for PCM devices.  

## Alternatives and other information

<!-- A clear and concise description of any alternative solutions or features you've considered. -->



Discussion / hints:

Fix patch:
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 4a3a6b6d..f7e055f5 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -21,20 +21,27 @@ The format is based on [Keep a Changelog], and this project adheres to
 * State indipendent inference noise model. (\# 284)
 * Transfer LR parameter for ``MixedPrecisionCompound``. (\#283)
 * The bias term can now be handled either by the analog or digital domain by controlling
-  the `digital_bias` layer parameter. (\#306 to be updated)
+  the `digital_bias` layer parameter. (\#307)
+* PCM short-term weight noise. (\#312)
+* IR-drop simulation across columns during analog mat-vec. (\#312)
+* Transposed-read for ``TransferCompound``. (\#312)
 
 ### Fixed
 
 * Removed GPU warning during destruction when using multiple GPUs. (\#277)
 * Fixed issue in transfer counter for mixed precision in case of GPU. (\#283)
 * Map location keyword for load / save observed. (\#293)
-* Fixed issue with CUDA buffer allocation when batch size changed. (\294)
+* Fixed issue with CUDA buffer allocation when batch size changed. (\#294)
 * Fixed missing load statedict for AnalogSequential. (\#295) 
 
 ### Changed
 
 * The inference noise models are now located in `aihwkit.inference`. (\#281)
 * Analog state dict structure `has changed (shared weight are not saved). (\#293)
+* Some of the parameter names of the``TransferCompound`` have
+  changed. (\#312)
+* New fast learning rate parameter for TransferCompound, SGD learning
+  rate then is applied on the slow matrix (\#312). 
 
 ### Deprecated
 
diff --git a/README.md b/README.md
index a9fada97..bec5c044 100644
--- a/README.md
+++ b/README.md
@@ -32,8 +32,9 @@ A series of primitives and features that allow using the toolkit within
 * Analog inference using torch inference workflow:
   * State-of-the-art statistical model of a phase-change memory (PCM) array
     calibrated on hardware measurements from a 1 million PCM devices chip.
-  * Hardware-aware training with hardware non-idealities and noise included
-    in the forward pass to make the trained models more robust during inference on Analog hardware.
+  * Hardware-aware training with hardware non-idealities and noise
+    included in the forward pass to make the trained models more
+    robust during inference on Analog hardware.
 
 ### Analog devices simulator
 
@@ -55,7 +56,8 @@ adjustable parameters. Features include:
 
 ### Other features
 
-Along with the two main components, the toolkit includes other functionalities such as:
+Along with the two main components, the toolkit includes other
+functionalities such as:
 
 * A library of device presets that are calibrated to real hardware data and
   based on models in the literature, along with configuration that specifies a particular device and optimizer choice.
diff --git a/examples/08_simple_layer_with_tiki_taka.py b/examples/08_simple_layer_with_tiki_taka.py
index 3eea978e..bd5f8f1b 100644
--- a/examples/08_simple_layer_with_tiki_taka.py
+++ b/examples/08_simple_layer_with_tiki_taka.py
@@ -47,10 +47,12 @@
         # Make some adjustments of the way Tiki-Taka is performed.
         units_in_mbatch=True,    # batch_size=1 anyway
         transfer_every=2,        # every 2 batches do a transfer-read
-        n_cols_per_transfer=1,   # one forward read for each transfer
+        n_reads_per_transfer=1,  # one forward read for each transfer
         gamma=0.0,               # all SGD weight in second device
         scale_transfer_lr=True,  # in relative terms to SGD LR
         transfer_lr=1.0,         # same transfer LR as for SGD
+        fast_lr=0.1,             # SGD update onto first matrix constant
+        transfer_columns=True    # transfer use columns (not rows)
     )
 )
 
diff --git a/src/aihwkit/nn/modules/base.py b/src/aihwkit/nn/modules/base.py
index c9c82ccc..3207db8b 100644
--- a/src/aihwkit/nn/modules/base.py
+++ b/src/aihwkit/nn/modules/base.py
@@ -138,7 +138,7 @@ def _setup_tile(
             weight_scaling_omega: the weight value where the max
                 weight will be scaled to. If zero, no weight scaling will
                 be performed
-            digital
+            digital_bias: whether to use bias in digital
 
         Returns:
             An analog tile with the requested parameters.
diff --git a/src/aihwkit/simulator/configs/devices.py b/src/aihwkit/simulator/configs/devices.py
index 0bd4de76..32ecdc5b 100644
--- a/src/aihwkit/simulator/configs/devices.py
+++ b/src/aihwkit/simulator/configs/devices.py
@@ -1076,17 +1076,17 @@ class TransferCompound(UnitCell):
     in the list of ``unit_cell_devices``) need to be of the same spec.
 
     The rate of transfer (e.g. learning rate and how often and how
-    many columns per transfer) and the type (ie. with ADC or without,
+    many columns/rows per transfer) and the type (ie. with ADC or without,
     with noise etc.) can be adjusted.
 
     Each transfer event that is triggered by counting the update
     cycles (in units of either mini-batch or single mat-vecs),
-    ``n_cols_per_transfer`` columns are read from the left device
+    ``n_reads_per_transfer`` columns/rows are read from the left device
     using the forward pass with transfer vectors as input and
     transferred to the right (taking the order of the
     ``unit_cell_devices`` list) using the outer-product update with
     the read-out vectors and the transfer vectors. Currently, transfer
-    vectors are fixed to be one-hot vectors. The columns to take are
+    vectors are fixed to be one-hot vectors. The columns/rows to take are
     in sequential order and warped around at the edge of the
     crossbar. The learning rate and forward and update specs of the
     transfer can be user-defined.
@@ -1136,7 +1136,7 @@ class TransferCompound(UnitCell):
     If not given explicitely with ``transfer_every_vec``, then the higher
     transfer cycles are geometrically scaled, the first is set to
     transfer_every. Each next transfer cycle is multiplied by ``x_size
-    / n_cols_per_transfer``.
+    / n_reads_per_transfer``.
     """
 
     no_self_transfer: bool = True
@@ -1159,24 +1159,50 @@ class TransferCompound(UnitCell):
     weight re-use during a while mini-batch.
     """
 
-    n_cols_per_transfer: int = 1
-    """Number of consecutive columns to use during transfer events.
+    n_reads_per_transfer: int = 1
+    """Number of consecutive reads to use during transfer events.
 
-    How many consecutive columns to read (from one tile) and write (to the next
+    How many consecutive columns or rows to read (from one tile) and write (to the next
     tile) every transfer event. For read, the input is a 1-hot vector. Once the
-    final column is reached, reading starts again from the first.
+    final columns or row is reached, reading starts again from the first.
+    """
+
+    transfer_columns: bool = True
+    """Whether to read and transfer columns or rows.
+
+    If set, read is done with an additional forward pass
+    determined by the ``transfer_forward`` settings. If not set, rows
+    are transferred instead, that is, the read is done internally
+    with a backward pass instead. However, the parameters defining the
+    backward are still given by setting the ``transfer_forward`` field for
+    convenience.
     """
 
     with_reset_prob: float = 0.0
     """Whether to apply reset of the columns that were transferred with a given
-    probability."""
+    probability.
+
+    Note:
+        Reset is only available in case of column reads
+        (``transfer_columns==True``).
+    """
+
+    random_selection: bool = False
+    """Whether to select a random starting column or row.
+
+    Whether to select a random starting column or row for each
+    transfer event and not take the next column or row that was
+    previously not transferred as a starting column or row (the
+    default).
+    """
 
-    random_column: bool = False
-    """Whether to select a random starting column.
+    fast_lr: float = 0.0
+    """Whether to set the `fast` tile's learning rate.
 
-    Whether to select a random starting column for each transfer event and not
-    take the next column that was previously not transferred as a starting
-    column (the default).
+    If set, then the SGD gradient update onto the first (fast) tile is
+    set to this learning rate and is kept constant even when the SGD
+    learning rate is scheduled. The SGD learning rate is then only
+    used to scale the transfer LR (see ``scale_transfer_lr``).
     """
 
     transfer_lr: float = 1.0
@@ -1206,7 +1232,7 @@ class TransferCompound(UnitCell):
     """Input-output parameters that define the read of a transfer event.
 
     :class:`~aihwkit.simulator.config.utils.AnalogTileInputOutputParameters` that define the read
-    (forward) of an transfer event. For instance the amount of noise
+    (forward or backward) of an transfer event. For instance the amount of noise
     or whether transfer is done using a ADC/DAC etc.
     """
 
diff --git a/src/aihwkit/simulator/configs/helpers.py b/src/aihwkit/simulator/configs/helpers.py
index 40af8370..85b004a8 100644
--- a/src/aihwkit/simulator/configs/helpers.py
+++ b/src/aihwkit/simulator/configs/helpers.py
@@ -49,7 +49,7 @@ def tile_parameters_to_bindings(params: Any) -> Any:
     field_map = {'forward': 'forward_io',
                  'backward': 'backward_io'}
     excluded_fields = ('device', 'noise_model', 'drift_compensation',
-                       'clip', 'modifier', 'mapping')
+                       'clip', 'modifier')
 
     result = params.bindings_class()
     for field, value in params.__dict__.items():
diff --git a/src/aihwkit/simulator/configs/utils.py b/src/aihwkit/simulator/configs/utils.py
index 80a45f59..2f52201f 100644
--- a/src/aihwkit/simulator/configs/utils.py
+++ b/src/aihwkit/simulator/configs/utils.py
@@ -353,9 +353,27 @@ class IOParameters(_PrintableMixin):
         :meth:`diffuse_weights`.
     """
 
+    ir_drop: float = 0.0
+    """Scale of IR drop along the inputs (rows of the weight matrix).
+
+    The IR-drop is calculated assuming that the first input is
+    farthest away from the output channel. The expected drop is
+    approximating the steady-state voltage distributions and depends
+    on the input current.
+    """
+
+    ir_drop_g_ratio: float = 1/(0.35*5e-6)
+    """Physical ratio of wire conductance from one cell to the next to
+    physical max conductance of a device.
+
+    Default is compute with 5mS maximal conductance set state and 0.35
+    Ohm wire resistance.
+    """
+
 
 @dataclass
 class UpdateParameters(_PrintableMixin):
+
     """Parameter that modify the update behaviour of a pulsed device."""
 
     bindings_class: ClassVar[Type] = devices.AnalogTileUpdateParameter
diff --git a/src/aihwkit/simulator/rpu_base_src/rpu_base_devices.cpp b/src/aihwkit/simulator/rpu_base_src/rpu_base_devices.cpp
index 22c1d392..06174e28 100644
--- a/src/aihwkit/simulator/rpu_base_src/rpu_base_devices.cpp
+++ b/src/aihwkit/simulator/rpu_base_src/rpu_base_devices.cpp
@@ -349,7 +349,9 @@ void declare_rpu_devices(py::module &m) {
       .def_readwrite("out_scale", &RPU::IOMetaParameter<T>::out_scale)
       .def_readwrite("out_sto_round", &RPU::IOMetaParameter<T>::out_sto_round)
       .def_readwrite("w_noise", &RPU::IOMetaParameter<T>::w_noise)
-      .def_readwrite("w_noise_type", &RPU::IOMetaParameter<T>::w_noise_type);
+      .def_readwrite("w_noise_type", &RPU::IOMetaParameter<T>::w_noise_type)
+      .def_readwrite("ir_drop", &RPU::IOMetaParameter<T>::ir_drop)
+      .def_readwrite("ir_drop_g_ratio", &RPU::IOMetaParameter<T>::ir_drop_Gw_div_gmax);
 
   py::class_<RPU::DriftParameter<T>>(m, "DriftParameter")
       .def(py::init<>())
@@ -567,10 +569,12 @@ void declare_rpu_devices(py::module &m) {
       .def_readwrite(
           "transfer_every_vec", &TransferParam::transfer_every_vec) // can this be filled?
       .def_readwrite("units_in_mbatch", &TransferParam::units_in_mbatch)
-      .def_readwrite("n_cols_per_transfer", &TransferParam::n_cols_per_transfer)
+      .def_readwrite("n_reads_per_transfer", &TransferParam::n_reads_per_transfer)
       .def_readwrite("with_reset_prob", &TransferParam::with_reset_prob)
-      .def_readwrite("random_column", &TransferParam::random_column)
+      .def_readwrite("random_selection", &TransferParam::random_selection)
+      .def_readwrite("transfer_columns", &TransferParam::transfer_columns)
       .def_readwrite("transfer_lr", &TransferParam::transfer_lr)
+      .def_readwrite("fast_lr", &TransferParam::fast_lr)
       .def_readwrite("transfer_lr_vec", &TransferParam::transfer_lr_vec)
       .def_readwrite("scale_transfer_lr", &TransferParam::scale_transfer_lr)
       .def_readwrite("transfer_forward", &TransferParam::transfer_io)
diff --git a/src/rpucuda/cuda/cuda_util.cu b/src/rpucuda/cuda/cuda_util.cu
index 2a5c4adb..638fff15 100644
--- a/src/rpucuda/cuda/cuda_util.cu
+++ b/src/rpucuda/cuda/cuda_util.cu
@@ -644,8 +644,6 @@ void CudaContext::subtractMemCounter(int64_t n_bytes) {
   global_mem_counter -= n_bytes;
 }
 
-int64_t CudaContext::getTotalMem() { return global_mem_counter; }
-
 //**********************************************************************//
 
 template <typename T>
diff --git a/src/rpucuda/cuda/cuda_util.h b/src/rpucuda/cuda/cuda_util.h
index 95713e60..19b55ab4 100644
--- a/src/rpucuda/cuda/cuda_util.h
+++ b/src/rpucuda/cuda/cuda_util.h
@@ -94,7 +94,9 @@
       switch (x) {                                                                                 \
       case CURAND_STATUS_VERSION_MISMATCH:                                                         \
         ss << "CURAND_STATUS_VERSION_MISMATCH";                                                    \
-        break case CURAND_STATUS_NOT_INITIALIZED : ss << "CURAND_STATUS_NOT_INITIALIZED";          \
+        break;                                                                                     \
+      case CURAND_STATUS_NOT_INITIALIZED:                                                          \
+        ss << "CURAND_STATUS_NOT_INITIALIZED";                                                     \
         break;                                                                                     \
       case CURAND_STATUS_ALLOCATION_FAILED:                                                        \
         ss << "CURAND_STATUS_ALLOCATION_FAILED";                                                   \
@@ -160,9 +162,6 @@
 #define MAX(a, b) (((a) > (b)) ? (a) : (b))
 #endif
 
-#define RPU_CUDA_1D_KERNEL_LOOP(i, n)                                                              \
-  for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < (n); i += blockDim.x * gridDim.x)
-
 #define RPU_GET_CUDA_BUFFER(CONTEXT, TYPE, BUFFER, SIZE)                                           \
   if (!BUFFER || BUFFER->getSize() < (SIZE)) {                                                     \
     BUFFER = RPU::make_unique<CudaArray<TYPE>>(CONTEXT, (SIZE));                                   \
@@ -329,7 +328,7 @@ class CudaContext {
 
   inline int64_t getAllocatedMem() { return allocated_mem_; }
 
-  int64_t getTotalMem();
+  int64_t getTotalMem() { return global_mem_counter; }
 
 private:
   int gpu_id_ = 0;
diff --git a/src/rpucuda/cuda/io_manager.cu b/src/rpucuda/cuda/io_manager.cu
index b474ab4f..a8956de3 100644
--- a/src/rpucuda/cuda/io_manager.cu
+++ b/src/rpucuda/cuda/io_manager.cu
@@ -495,6 +495,49 @@ __global__ void kernelElemSqrtAddNoiseBatch(
   STOCH_FINALIZE(true);
 }
 
+// compute x_j*(1-(1-j/n)^2)
+template <typename T>
+__global__ void kernelInputPositionCoding(
+    T *output,      // x' out
+    const T *input, // x in
+    const int size_in,
+    const int m_batch_in,
+    const bool trans_in) {
+  const bool trans = trans_in;
+  const int size = size_in;
+  const int m_batch = m_batch_in;
+  const int total_size = size * m_batch;
+
+  int tid = blockDim.x * blockIdx.x + threadIdx.x;
+  int total_threads = blockDim.x * gridDim.x;
+
+  STRIDE_LOOP(total_size, value,
+
+              int xidx = trans ? (idx / m_batch) : (idx % size);
+              T z = (T)1.0 - (T)xidx / (T)size; value *= (T)1.0 - z * z;);
+}
+
+// computes c_i and y_i -= c_i*GEMM_i
+// c_i = a_i*(a_i*(0.05*a_i - 0.2) + 0.5);
+template <typename T>
+__global__ void kernelElemMulCAdd(
+    T *output,
+    const T *input,
+    const T *a_values,
+    const T *pc_gemm_values,
+    const int total_size_in) {
+  const int total_size = total_size_in;
+
+  int total_threads = blockDim.x * gridDim.x;
+  int tid = blockDim.x * blockIdx.x + threadIdx.x;
+  const T K2 = 0.05;
+  const T K1 = 0.2;
+  const T K0 = 0.5;
+
+  STRIDE_LOOP(total_size, value, T a_i = a_values[idx]; T gemm_i = pc_gemm_values[idx];
+              T c_i = a_i * (a_i * (K2 * a_i - K1) + K0); value -= c_i * gemm_i;);
+}
+
 /*********************************************************************************/
 /*Output management: single batch */
 template <typename T, typename OutputIteratorT, bool noise_management>
@@ -1203,6 +1246,7 @@ bool InputOutputManager<T>::applyToOutputWithBoundManagement(
   int m_batch = temp_m_batch_;
 
   // actual bound management
+
   if (m_batch == RPU_IO_USE_SINGLE_BATCH_VERSION) {
     kernelOutputBoundManagement<<<nblocks_om_, nthreads_, 0, s>>>(
         dev_output, dev_output_applied_->getData(), out_size_, io_->out_noise,
@@ -1350,9 +1394,66 @@ void InputOutputManager<T>::applyOutputPCMReadNoise(const T *dev_weights, const
   }
 }
 
+template <typename T>
+void InputOutputManager<T>::applyIrDrop(const T *dev_weights, const bool out_trans) {
+  if (io_->ir_drop > 0.0) {
+    // computes for each output (row):
+    //
+    // a_i = sum_j(|w_ij|*|x_j|)*n/Gw*gmax
+    // c_i = a_i*(a_i*(0.05*a_i - 0.2) + 0.5);
+    // x'_j = x_j * (1 - (1-j/n)^2)
+    // y_i = y_i_ideal - c_i*sum_j(w_ij * x'_j)
+
+    int m_batch = temp_m_batch_;
+    int in_trans = temp_trans_;
+
+    T *in_temp = getInBuffer(); // this is in_size_ x m_batch or m_batch x in_size_ (if in_trans)
+    T *out_temp =
+        getOutBuffer(); // this is out_size_ x m_batch or m_batch x out_size_ (if out_trans)
+
+    // unfortunately needs a lot of extra buffer memory (CUTLASS will help)
+    RPU_GET_CUDA_BUFFER(context_, T, dev_extra_batch_buffer_, m_batch * MAX(in_size_, out_size_));
+    RPU_GET_CUDA_BUFFER(context_, T, dev_extra_weight_buffer_, MAX(in_size_, m_batch) * out_size_);
+    RPU_GET_CUDA_BUFFER(context_, T, dev_a_buffer_, m_batch * out_size_);
+
+    // compute thr abs weights current
+    RPU::math::elemabs(
+        context_, dev_extra_weight_buffer_->getData(), dev_weights, out_size_ * in_size_);
+
+    RPU::math::elemabs(context_, dev_extra_batch_buffer_->getData(), in_temp, in_size_ * m_batch);
+
+    // compute the a_i
+    detail::forwardMatrix(
+        context_, dev_extra_weight_buffer_->getData(), dev_extra_batch_buffer_->getData(), in_size_,
+        in_trans, dev_a_buffer_->getData(), out_size_, out_trans, m_batch,
+        in_size_ / io_->ir_drop_Gw_div_gmax);
+
+    // compute x_j*(1-(1-j/n)^2)
+    kernelInputPositionCoding<<<nblocks_om_batch_, nthreads_, 0, context_->getStream()>>>(
+        dev_extra_batch_buffer_->getData(), in_temp, in_size_, m_batch, in_trans);
+
+    // compute the position filtered GEMM
+    detail::forwardMatrix(
+        context_, dev_weights, dev_extra_batch_buffer_->getData(), in_size_, in_trans,
+        dev_extra_weight_buffer_->getData(), // re-use  buffer
+        out_size_, out_trans, m_batch, io_->ir_drop);
+
+    // computes c_i and y_i -= c_i*GEMM_i
+    kernelElemMulCAdd<<<nblocks_om_batch_, nthreads_, 0, context_->getStream()>>>(
+        out_temp,
+        out_temp,                            // in-place
+        dev_a_buffer_->getData(),            // this is the a_i
+        dev_extra_weight_buffer_->getData(), // this is GEMM_i
+        out_size_ * m_batch                  // all have same shape
+    );
+  }
+}
+
 template <typename T>
 void InputOutputManager<T>::applyOutputNonIdealities(const T *dev_weights, const bool out_trans) {
 
+  applyIrDrop(dev_weights, out_trans);
+
   switch (io_->w_noise_type) {
   case OutputWeightNoiseType::AdditiveConstant: {
     // overwrites inBuffer
diff --git a/src/rpucuda/cuda/io_manager.h b/src/rpucuda/cuda/io_manager.h
index 0626878e..e615118f 100644
--- a/src/rpucuda/cuda/io_manager.h
+++ b/src/rpucuda/cuda/io_manager.h
@@ -63,6 +63,7 @@ template <typename T> class InputOutputManager {
   void applyOutputWeightNoise(const bool out_trans);
   void applyOutputNonIdealities(const T *dev_weights, const bool out_trans);
   void applyOutputPCMReadNoise(const T *dev_weights, const bool out_trans);
+  void applyIrDrop(const T *dev_weights, const bool out_trans);
 
   void initializeBatchBuffer(int m_batch);
 
@@ -106,7 +107,6 @@ template <typename T> class InputOutputManager {
   std::unique_ptr<CudaArray<float>> dev_scale_values_ = nullptr;
   std::unique_ptr<CudaArray<int>> dev_bound_exceeded_ = nullptr;
   std::unique_ptr<CudaArray<int>> dev_any_exceeded_ = nullptr;
-  std::unique_ptr<CudaArray<int>> dev_channel_exceeded_ = nullptr;
   std::unique_ptr<CudaArray<char>> dev_flagged_temp_storage_ = nullptr;
   std::unique_ptr<CudaArray<int>> dev_selected_bidx_ = nullptr;
   std::unique_ptr<CudaArray<int>> dev_selected_m_batch_ = nullptr;
diff --git a/src/rpucuda/cuda/pulsed_weight_updater.cu b/src/rpucuda/cuda/pulsed_weight_updater.cu
index 04c9457d..4b04d2a5 100644
--- a/src/rpucuda/cuda/pulsed_weight_updater.cu
+++ b/src/rpucuda/cuda/pulsed_weight_updater.cu
@@ -94,17 +94,19 @@ void PulsedWeightUpdater<T>::executeUpdate(
     const bool x_trans_in,
     const bool d_trans_in) {
 
+  T pc_lr = rpucuda_device->getPulseCountLearningRate(lr);
   blm_->makeCounts(
-      x_in, d_in, up, rpucuda_device->getWeightGranularity(), lr, m_batch, x_trans_in, d_trans_in,
-      kpars->getOutTrans(), kpars->getUseBo64(), kpars->getImplicitPulses());
+      x_in, d_in, up, rpucuda_device->getWeightGranularity(), pc_lr, m_batch, x_trans_in,
+      d_trans_in, kpars->getOutTrans(), kpars->getUseBo64(), kpars->getImplicitPulses());
 
   CudaContext *c = context_;
   if (is_async_update_) {
     up_context_->recordWaitEvent(context_->getStream(), context_->getEvent());
     c = &*up_context_;
   }
+  // the original learninig rate needs to be passed
   rpucuda_device->runUpdateKernel(
-      kpars, c, dev_weights, m_batch, &*blm_, up, c->getRandomStates(kpars->getNStates()));
+      kpars, c, dev_weights, m_batch, &*blm_, up, lr, c->getRandomStates(kpars->getNStates()));
 }
 
 template <typename T>
@@ -290,12 +292,23 @@ void PulsedWeightUpdater<T>::doDirectUpdate(
 template <typename T>
 bool PulsedWeightUpdater<T>::checkForFPUpdate(
     AbstractRPUDeviceCuda<T> *rpucuda_device_in, const PulsedUpdateMetaParameter<T> &up) {
+
+  if (rpucuda_device_in == nullptr) {
+    return true;
+  }
+  if (rpucuda_device_in->implements() == DeviceUpdateType::FloatingPoint) {
+    return true;
+  }
+  if (rpucuda_device_in->isPulsedDevice() && up.pulse_type == PulseType::None) {
+    return true;
+  }
   if (rpucuda_device_in->hasDirectUpdate()) {
+    // also FP has direct, but that is handled above
     return false;
   }
-  return (up.pulse_type == PulseType::None) || (rpucuda_device_in == nullptr) ||
-         !rpucuda_device_in->isPulsedDevice() ||
-         (rpucuda_device_in->implements() == DeviceUpdateType::FloatingPoint);
+  // omitting !isPulsedDevice
+
+  return false;
 }
 
 #define FORCE_TUNING_THRES 0
diff --git a/src/rpucuda/cuda/rpucuda_onesided_device.cu b/src/rpucuda/cuda/rpucuda_onesided_device.cu
index 165edab3..a0600029 100644
--- a/src/rpucuda/cuda/rpucuda_onesided_device.cu
+++ b/src/rpucuda/cuda/rpucuda_onesided_device.cu
@@ -302,6 +302,7 @@ void OneSidedRPUDeviceCuda<T>::runUpdateKernel(
     int m_batch,
     const BitLineMaker<T> *blm,
     const PulsedUpdateMetaParameter<T> &up,
+    const T lr,
     curandState_t *dev_states,
     int one_sided,
     uint32_t *x_counts_chunk,
diff --git a/src/rpucuda/cuda/rpucuda_onesided_device.h b/src/rpucuda/cuda/rpucuda_onesided_device.h
index acdce64c..c73e8836 100644
--- a/src/rpucuda/cuda/rpucuda_onesided_device.h
+++ b/src/rpucuda/cuda/rpucuda_onesided_device.h
@@ -60,6 +60,7 @@ template <typename T> class OneSidedRPUDeviceCuda : public VectorRPUDeviceCuda<T
       int m_batch,
       const BitLineMaker<T> *blm,
       const PulsedUpdateMetaParameter<T> &up,
+      const T lr,
       curandState_t *dev_states,
       int one_sided = 0,
       uint32_t *x_counts_chunk = nullptr,
diff --git a/src/rpucuda/cuda/rpucuda_pulsed_device.cu b/src/rpucuda/cuda/rpucuda_pulsed_device.cu
index 477c8309..4ccf4909 100644
--- a/src/rpucuda/cuda/rpucuda_pulsed_device.cu
+++ b/src/rpucuda/cuda/rpucuda_pulsed_device.cu
@@ -366,6 +366,7 @@ void PulsedRPUDeviceCuda<T>::runUpdateKernel(
     int m_batch,
     const BitLineMaker<T> *blm,
     const PulsedUpdateMetaParameter<T> &up,
+    const T lr,
     curandState_t *dev_states,
     int one_sided,
     uint32_t *x_counts_chunk,
diff --git a/src/rpucuda/cuda/rpucuda_pulsed_device.h b/src/rpucuda/cuda/rpucuda_pulsed_device.h
index 431a1729..6107d8ae 100644
--- a/src/rpucuda/cuda/rpucuda_pulsed_device.h
+++ b/src/rpucuda/cuda/rpucuda_pulsed_device.h
@@ -80,6 +80,7 @@ template <typename T> class PulsedRPUDeviceCudaBase : public SimpleRPUDeviceCuda
       int m_batch,
       const BitLineMaker<T> *blm,
       const PulsedUpdateMetaParameter<T> &up,
+      const T lr,
       curandState_t *dev_states,
       int one_sided = 0,
       uint32_t *x_counts_chunk = nullptr,
@@ -94,6 +95,7 @@ template <typename T> class PulsedRPUDeviceCudaBase : public SimpleRPUDeviceCuda
   PulsedRPUDeviceCudaBase<T> *clone() const override { RPU_FATAL("Needs implementation"); };
 
   inline T getWeightGranularity() const { return weight_granularity_; };
+  virtual T getPulseCountLearningRate(T learning_rate) { return learning_rate; };
 
 protected:
   inline void setWeightGranularity(T weight_granularity) {
@@ -155,6 +157,7 @@ template <typename T> class PulsedRPUDeviceCuda : public PulsedRPUDeviceCudaBase
       int m_batch,
       const BitLineMaker<T> *blm,
       const PulsedUpdateMetaParameter<T> &up,
+      const T lr,
       curandState_t *dev_states,
       int one_sided = 0,
       uint32_t *x_counts_chunk = nullptr,
diff --git a/src/rpucuda/cuda/rpucuda_simple_device.h b/src/rpucuda/cuda/rpucuda_simple_device.h
index c5d72d96..c506ca13 100644
--- a/src/rpucuda/cuda/rpucuda_simple_device.h
+++ b/src/rpucuda/cuda/rpucuda_simple_device.h
@@ -48,7 +48,6 @@ template <typename T> class AbstractRPUDeviceCuda {
       T *x_buffer,
       T *d_buffer) = 0;
   virtual bool hasDirectUpdate() const = 0;
-
   virtual int getHiddenUpdateIdx() const { return 0; };
   virtual void setHiddenUpdateIdx(int idx){};
 
diff --git a/src/rpucuda/cuda/rpucuda_transfer_device.cu b/src/rpucuda/cuda/rpucuda_transfer_device.cu
index 5feada00..3a48cd68 100644
--- a/src/rpucuda/cuda/rpucuda_transfer_device.cu
+++ b/src/rpucuda/cuda/rpucuda_transfer_device.cu
@@ -13,6 +13,7 @@
 #include "forward_backward_pass.h"
 #include "rpu_pulsed_meta_parameter.h"
 #include "rpucuda_transfer_device.h"
+#include <limits>
 #include <memory>
 
 namespace RPU {
@@ -24,11 +25,17 @@ namespace RPU {
 
 */
 
-template <typename T> void TransferRPUDeviceCuda<T>::initialize() {
+template <typename T> void TransferRPUDeviceCuda<T>::initialize(bool transfer_columns) {
   transfer_pwu_ =
       RPU::make_unique<PulsedWeightUpdater<T>>(this->context_, this->x_size_, this->d_size_);
-  transfer_iom_ =
-      RPU::make_unique<InputOutputManager<T>>(this->context_, this->x_size_, this->d_size_);
+
+  if (transfer_columns) {
+    transfer_iom_ =
+        RPU::make_unique<InputOutputManager<T>>(this->context_, this->x_size_, this->d_size_);
+  } else {
+    transfer_iom_ =
+        RPU::make_unique<InputOutputManager<T>>(this->context_, this->d_size_, this->x_size_);
+  }
   this->context_->synchronize();
 }
 
@@ -47,10 +54,8 @@ TransferRPUDeviceCuda<T>::TransferRPUDeviceCuda(const TransferRPUDeviceCuda<T> &
   if (other.transfer_vecs_ != nullptr) {
     transfer_vecs_ = RPU::make_unique<CudaArray<T>>(*other.transfer_vecs_);
   }
-
-  initialize();
-
-  current_col_indices_ = other.current_col_indices_;
+  initialize(other.getPar().transfer_columns);
+  current_slice_indices_ = other.current_slice_indices_;
   fully_hidden_ = other.fully_hidden_;
   this->context_->synchronizeDevice();
 };
@@ -76,13 +81,14 @@ TransferRPUDeviceCuda<T> &TransferRPUDeviceCuda<T>::operator=(TransferRPUDeviceC
   VectorRPUDeviceCuda<T>::operator=(std::move(other));
 
   transfer_vecs_ = std::move(other.transfer_vecs_);
-  current_col_indices_ = other.current_col_indices_;
-  other.current_col_indices_.clear();
-  fully_hidden_ = other.fully_hidden_;
 
+  current_slice_indices_ = other.current_slice_indices_;
+  other.current_slice_indices_.clear();
+
+  fully_hidden_ = other.fully_hidden_;
   transfer_pwu_ = std::move(other.transfer_pwu_);
   transfer_iom_ = std::move(other.transfer_iom_);
-  // ignore transfer_tmp_ or RNG
+  transfer_tmp_ = std::move(other.transfer_tmp_);
 
   return *this;
 };
@@ -116,107 +122,162 @@ void TransferRPUDeviceCuda<T>::populateFrom(const AbstractRPUDevice<T> &rpu_devi
       RPU_FATAL("Later transfer periods need to be larger than first for CUDA.");
     }
   }
-
+  int in_size = par.getInSize();
   transfer_vecs_ = RPU::make_unique<CudaArray<T>>(
-      this->context_, this->x_size_ * this->x_size_, rpu_device.getTransferVecs());
+      this->context_, in_size * in_size, rpu_device.getTransferVecs());
 
-  initialize(); // pwu/iom
+  initialize(par.transfer_columns); // pwu/iom
 
-  current_col_indices_.resize(this->n_devices_ - 1);
-  std::fill(current_col_indices_.begin(), current_col_indices_.end(), (int)0);
+  current_slice_indices_.resize(this->n_devices_ - 1);
+  std::fill(current_slice_indices_.begin(), current_slice_indices_.end(), (int)0);
 
   this->current_update_idx_ = 0;
 
   fully_hidden_ = par.fullyHidden();
 }
 
+/*********************************************************************************/
+/* getPulseCountLearningRate */
+/* Here we compute the LR for the A matrix (the SGD update). Because
+   of the device properties it is beneficial to use a constant LR
+   here, but scale the buffer with the scheduled SGD learning rate
+   later*/
+template <typename T> T TransferRPUDeviceCuda<T>::getPulseCountLearningRate(T learning_rate) {
+  const auto &par = getPar();
+
+  if (par.fast_lr > 0) {
+    return par.fast_lr;
+  } else {
+    return learning_rate;
+  }
+}
+
+template <typename T>
+void TransferRPUDeviceCuda<T>::readMatrix(
+    int device_idx, const T *in_vec, T *out_vec, int m_batch, T alpha) {
+  const auto &par = getPar();
+
+  if (par.transfer_columns) {
+    // forward with transfer vectors
+    RPU::detail::forwardMatrixIteratorIOManaged(
+        this->context_, this->dev_weights_ptrs_[device_idx], in_vec, this->x_size_, false, out_vec,
+        this->d_size_, false, m_batch, alpha, *transfer_iom_, par.transfer_io, false);
+
+  } else {
+    // backward with transfer vectors
+    RPU::detail::backwardMatrixIteratorIOManaged(
+        this->context_, this->dev_weights_ptrs_[device_idx], in_vec, this->d_size_, false, out_vec,
+        this->x_size_, false, m_batch, alpha, *transfer_iom_, par.transfer_io);
+  }
+}
+
+template <typename T>
+void TransferRPUDeviceCuda<T>::writeMatrix(
+    int device_idx,
+    const T *in_vec,
+    const T *out_vec,
+    int m_batch,
+    const T lr,
+    const PulsedUpdateMetaParameter<T> &up) {
+  const auto &par = getPar();
+  // note that the ptrs might point to the current weight
+  T *W = this->dev_weights_ptrs_[device_idx];
+
+  if (par.transfer_columns) {
+    transfer_pwu_->update(
+        in_vec, out_vec, W, &*this->rpucuda_device_vec_[device_idx], up, fabs(lr), m_batch, false,
+        false);
+  } else {
+    transfer_pwu_->update(
+        out_vec, in_vec, W, &*this->rpucuda_device_vec_[device_idx], up, fabs(lr), m_batch, false,
+        false);
+  }
+}
+
+/*********************************************************************************/
 /* partially transfer using the given "readout" transfer vectors
    (with io-managed forward) and the usualy device update */
 template <typename T>
-void TransferRPUDeviceCuda<T>::forwardUpdate(
+void TransferRPUDeviceCuda<T>::readAndUpdate(
     int to_device_idx,
     int from_device_idx,
     int i_col_start,
     const T lr,
-    const T *x_input,
+    const T *vec,
     const int n_vec,
-    const bool trans,
     const PulsedUpdateMetaParameter<T> &up) {
-  if (!lr) {
+  if (lr == (T)0.0) {
     return;
   }
 
-  if ((transfer_tmp_ == nullptr) || transfer_tmp_->getSize() < n_vec * this->d_size_) {
-    transfer_tmp_ = RPU::make_unique<CudaArray<T>>(this->context_, this->d_size_ * n_vec);
+  const auto &par = getPar();
+  int out_size = par.getOutSize();
+  int t_size = n_vec * out_size; // transfer size
+  if ((transfer_tmp_ == nullptr) || transfer_tmp_->getSize() < t_size) {
+    transfer_tmp_ = RPU::make_unique<CudaArray<T>>(this->context_, t_size);
     this->context_->synchronize();
   }
-  // forward with transfer vectors
-  RPU::detail::forwardMatrixIteratorIOManaged(
-      this->context_, this->dev_weights_ptrs_[from_device_idx], x_input, this->x_size_, trans,
-      transfer_tmp_->getData(), this->d_size_, trans, n_vec,
-      (T)1.0, // additional output scaling. Obey?
-      *transfer_iom_, getPar().transfer_io, false);
+  T *out_vec = transfer_tmp_->getData();
 
+  // forward / backward with transfer vectors. Since we need *positive*
+  // update, LR needs to be negative. However, this is not supported
+  // in the PWU really. Thus we scale the output by -1 and set alpha
+  // accordingly
+  readMatrix(from_device_idx, vec, out_vec, n_vec, -1.0);
   // update according to device
-  T *W = this->dev_weights_ptrs_[to_device_idx]; /// note that the ptrs might point to the current
-                                                 /// weight
-
-  // since we need *positive* update, LR needs to be
-  // negative. However, this is not supported in the PWU
-  // really. Thus we scale the temp-vector by -1
-  RPU::math::scal(this->context_, this->d_size_ * n_vec, (T)-1.0, transfer_tmp_->getData(), 1);
-
-  transfer_pwu_->update(
-      x_input,                       // this is the transfer vector (x_size)
-      transfer_tmp_->getDataConst(), // this should be d_size
-      W, &*this->rpucuda_device_vec_[to_device_idx], up, fabs(lr), n_vec, trans, trans);
+  writeMatrix(to_device_idx, vec, out_vec, n_vec, fabs(lr), up);
 }
 
+/*********************************************************************************/
 template <typename T>
 void TransferRPUDeviceCuda<T>::transfer(
     int to_device_idx,
     int from_device_idx,
     const PulsedUpdateMetaParameter<T> &current_up,
     const T current_lr) {
-  int i_col = current_col_indices_[from_device_idx];
+  int i_slice = current_slice_indices_[from_device_idx];
   const auto &par = getPar();
-  if (par.random_column) {
-    i_col = MAX(MIN(floor(this->rw_rng_.sampleUniform() * this->x_size_), this->x_size_ - 1), 0);
+
+  int in_size = par.getInSize();
+  int out_size = par.getOutSize();
+
+  if (par.random_selection) {
+    i_slice = MAX(MIN(floor(this->rw_rng_.sampleUniform() * in_size), in_size - 1), 0);
   }
 
-  // transfer_vecs_ is always x_size-major (that is trans==false)
-  T *tvec = transfer_vecs_->getData() + i_col * this->x_size_;
-  int n_rest = this->x_size_ - i_col;
+  // transfer_vecs_ is always in_size-major (that is trans==false)
+  T *tvec = transfer_vecs_->getData() + i_slice * in_size;
+  int n_rest = in_size - i_slice;
 
   T lr = par.getTransferLR(to_device_idx, from_device_idx, current_lr);
 
   const PulsedUpdateMetaParameter<T> *up;
   up = &par.transfer_up;
 
-  int n_transfers = MIN(par.n_cols_per_transfer, this->x_size_);
+  int n_transfers = MIN(par.n_reads_per_transfer, in_size);
 
   if (n_rest < n_transfers) {
-
     // rest
-    forwardUpdate(to_device_idx, from_device_idx, i_col, lr, tvec, n_rest, false, *up);
+    readAndUpdate(to_device_idx, from_device_idx, i_slice, lr, tvec, n_rest, *up);
     // from beginning
-    forwardUpdate(
+    readAndUpdate(
         to_device_idx, from_device_idx, 0, lr, transfer_vecs_->getData(), n_transfers - n_rest,
-        false, *up);
+        *up);
 
   } else {
-    forwardUpdate(to_device_idx, from_device_idx, i_col, lr, tvec, n_transfers, false, *up);
+    readAndUpdate(to_device_idx, from_device_idx, i_slice, lr, tvec, n_transfers, *up);
   }
 
-  if (this->rw_rng_.sampleUniform() <
-      par.with_reset_prob) { // COL-wise prob!! device-wise reset_prob=1
+  if (par.transfer_columns && this->rw_rng_.sampleUniform() < par.with_reset_prob) {
+    // COL-wise prob!! device-wise reset_prob=1
     this->rpucuda_device_vec_[from_device_idx]->resetCols(
-        this->dev_weights_ptrs_[from_device_idx], i_col, n_transfers, 1);
+        this->dev_weights_ptrs_[from_device_idx], i_slice, n_transfers, 1);
   }
 
-  current_col_indices_[from_device_idx] = (i_col + n_transfers) % this->x_size_;
+  current_slice_indices_[from_device_idx] = (i_slice + n_transfers) % in_size;
 }
 
+/*********************************************************************************/
 template <typename T>
 inline int TransferRPUDeviceCuda<T>::getTransferEvery(int didx, int m_batch) const {
 
@@ -227,6 +288,7 @@ inline int TransferRPUDeviceCuda<T>::getTransferEvery(int didx, int m_batch) con
   }
 }
 
+/*********************************************************************************/
 template <typename T> inline int getNChunks(int m_batch, T every) {
   if (every <= 0) {
     return 1;
@@ -239,12 +301,22 @@ inline int getChunkSize(int m_batch, int nchunks) {
   return (m_batch + nchunks - 1) / nchunks; // to ensure not to have residual
 }
 
+inline uint64_t getNextTransfer(uint64_t current_update_idx, int transfer_every) {
+
+  if (transfer_every <= 0) {
+    return std::numeric_limits<uint64_t>::max();
+  }
+  return current_update_idx + transfer_every - (current_update_idx % transfer_every);
+}
+
+/*********************************************************************************/
 template <typename T>
 pwukpvec_t<T> TransferRPUDeviceCuda<T>::getUpdateKernels(
     int m_batch, int nK32, int use_bo64, bool out_trans, const PulsedUpdateMetaParameter<T> &up) {
 
   pwukpvec_t<T> v;
 
+  // just get approx chunk size for tuning
   int nchunks = getNChunks(m_batch, getTransferEvery(0, m_batch));
   int chunk_size = getChunkSize(m_batch, nchunks);
   // use the first device as the "FAST" device that gets updates with the true gradients.
@@ -259,6 +331,7 @@ pwukpvec_t<T> TransferRPUDeviceCuda<T>::getUpdateKernels(
   return v;
 }
 
+/*********************************************************************************/
 template <typename T>
 void TransferRPUDeviceCuda<T>::runUpdateKernel(
     pwukp_t<T> kpars,
@@ -267,6 +340,7 @@ void TransferRPUDeviceCuda<T>::runUpdateKernel(
     int m_batch,
     const BitLineMaker<T> *blm,
     const PulsedUpdateMetaParameter<T> &up,
+    const T lr,
     curandState_t *dev_states,
     int one_sided,
     uint32_t *x_counts_chunk,
@@ -287,10 +361,12 @@ void TransferRPUDeviceCuda<T>::runUpdateKernel(
     RPU_FATAL("Chunking not allowed here.");
   }
 
-  int nchunks = getNChunks(m_batch, getTransferEvery(0, m_batch));
-  T lr = blm->getCurrentLR();
+  // only look at first device here as it makes no sense to transfer
+  // the higer order devices more often
+  int transfer_every = getTransferEvery(0, m_batch);
+  auto next_transfer = getNextTransfer(this->current_update_idx_, transfer_every);
 
-  if (nchunks <= 1) {
+  if (next_transfer >= m_batch + this->current_update_idx_) {
 
     // just update the whole batch we do not call kpars directly to
     // also make possible to have non-pulsed devices. Note that only
@@ -300,7 +376,7 @@ void TransferRPUDeviceCuda<T>::runUpdateKernel(
     // during transfer, since we use a seperate PWU object
 
     this->rpucuda_device_vec_[0]->runUpdateKernel(
-        kpars, c, this->dev_weights_ptrs_[0], m_batch, blm, up, dev_states, one_sided);
+        kpars, c, this->dev_weights_ptrs_[0], m_batch, blm, up, lr, dev_states, one_sided);
 
     if (up._currently_tuning) {
       return;
@@ -308,31 +384,39 @@ void TransferRPUDeviceCuda<T>::runUpdateKernel(
 
     this->current_update_idx_ += m_batch; // first update idx
 
-    for (int j = 0; j < this->n_devices_ - 1; j++) {
-      // all transfer periods will be rounded up to batches.
-      int period = (getTransferEvery(j, m_batch) + m_batch - 1) / m_batch; // in m_batch
-      if (period > 0 && this->current_update_idx_ / m_batch % period == 0) {
+    // transfer
+    if (next_transfer == this->current_update_idx_) {
+      transfer(1, 0, up, lr);
+    }
+
+    // other higher order devices
+    for (int j = 1; j < this->n_devices_ - 1; j++) {
+      // all transfer periods will be rounded up to chunk_sizes
+      auto higher_order_next_transfer =
+          getNextTransfer(this->current_update_idx_ - m_batch, getTransferEvery(j, m_batch));
+      if (higher_order_next_transfer <= this->current_update_idx_) {
         transfer(j + 1, j, up, lr);
       }
     }
 
   } else {
+    // transfer is inbetween the mbatch, we need to chunk in some way
+
     // need to do it chunkwise
-    int chunk_size = getChunkSize(m_batch, nchunks);
     int batch_start = 0;
     int nK32 = blm->getNK32Current();
     auto x_counts = blm->getXCountsData();
     auto d_counts = blm->getDCountsData();
+    uint64_t final_update_idx = this->current_update_idx_ + m_batch;
 
-    for (int i_chunk = 0; i_chunk < nchunks; i_chunk++) {
+    while (next_transfer <= final_update_idx) {
 
-      // note that last chunk might be smaller.
-      int current_m_batch = chunk_size - MAX(batch_start + chunk_size - m_batch, 0);
+      int current_m_batch = (int)(next_transfer - this->current_update_idx_);
 
       this->rpucuda_device_vec_[0]->runUpdateKernel(
           kpars,
           c, // same context since sequence important
-          this->dev_weights_ptrs_[0], current_m_batch, blm, up, dev_states, one_sided,
+          this->dev_weights_ptrs_[0], current_m_batch, blm, up, lr, dev_states, one_sided,
           x_counts + batch_start * this->x_size_ * nK32, // always non-trans
           d_counts + batch_start * this->d_size_ * nK32);
 
@@ -344,19 +428,27 @@ void TransferRPUDeviceCuda<T>::runUpdateKernel(
       batch_start += current_m_batch;
 
       // transfer
-      for (int j = 0; j < this->n_devices_ - 1; j++) {
+      if (next_transfer == this->current_update_idx_) {
+        transfer(1, 0, up, lr);
+      }
+
+      // other higher order devices
+      for (int j = 1; j < this->n_devices_ - 1; j++) {
         // all transfer periods will be rounded up to chunk_sizes
-        int period = (getTransferEvery(j, m_batch) + chunk_size - 1) / chunk_size;
-        if (period > 0 && this->current_update_idx_ / chunk_size % period == 0) {
+        auto higher_order_next_transfer = getNextTransfer(
+            this->current_update_idx_ - current_m_batch, getTransferEvery(j, m_batch));
+        if (higher_order_next_transfer <= this->current_update_idx_) {
           transfer(j + 1, j, up, lr);
         }
       }
+      next_transfer = getNextTransfer(this->current_update_idx_, transfer_every);
     }
   }
   // only reduce at end
   this->reduceToWeights(up_context, dev_weights);
 }
 
+/*********************************************************************************/
 template <typename T>
 void TransferRPUDeviceCuda<T>::reduceToWeights(CudaContext *context, T *dev_weights) {
 
diff --git a/src/rpucuda/cuda/rpucuda_transfer_device.h b/src/rpucuda/cuda/rpucuda_transfer_device.h
index 47234874..a8538224 100644
--- a/src/rpucuda/cuda/rpucuda_transfer_device.h
+++ b/src/rpucuda/cuda/rpucuda_transfer_device.h
@@ -41,8 +41,7 @@ template <typename T> class TransferRPUDeviceCuda : public VectorRPUDeviceCuda<T
     swap(a.transfer_vecs_, b.transfer_vecs_);
     swap(a.transfer_iom_, b.transfer_iom_);
     swap(a.transfer_pwu_, b.transfer_pwu_);
-    swap(a.rw_rng_, b.rw_rng_);
-    swap(a.current_col_indices_, b.current_col_indices_);
+    swap(a.current_slice_indices_, b.current_slice_indices_);
     swap(a.fully_hidden_, b.fully_hidden_);
   };
 
@@ -61,6 +60,7 @@ template <typename T> class TransferRPUDeviceCuda : public VectorRPUDeviceCuda<T
       int m_batch,
       const BitLineMaker<T> *blm,
       const PulsedUpdateMetaParameter<T> &up,
+      const T lr,
       curandState_t *dev_states,
       int one_sided = 0,
       uint32_t *x_counts_chunk = nullptr,
@@ -81,16 +81,24 @@ template <typename T> class TransferRPUDeviceCuda : public VectorRPUDeviceCuda<T
 
   // uses the getPar().transfer_up and getPar().transfer_io to make a forward with transfer_vec and
   // update
-  virtual void forwardUpdate(
+  virtual void readAndUpdate(
       int to_device_idx,
       int from_device_idx,
-      int i_col_start,
+      int i_slice_start,
       const T lr,
       const T *x_input,
       const int n_vec,
-      const bool trans,
       const PulsedUpdateMetaParameter<T> &up);
 
+  virtual void writeMatrix(
+      int device_idx,
+      const T *in_vec,
+      const T *out_vec,
+      int m_batch,
+      const T lr,
+      const PulsedUpdateMetaParameter<T> &up);
+  virtual void readMatrix(int device_idx, const T *in_vec, T *out_vec, int m_batch, T alpha);
+
   pwukpvec_t<T> getUpdateKernels(
       int m_batch,
       int nK32,
@@ -98,19 +106,22 @@ template <typename T> class TransferRPUDeviceCuda : public VectorRPUDeviceCuda<T
       bool out_trans,
       const PulsedUpdateMetaParameter<T> &up) override;
 
+  T getPulseCountLearningRate(T learning_rate) override;
+
 protected:
   void reduceToWeights(CudaContext *c, T *dev_weights) override;
+
   std::unique_ptr<CudaArray<T>> transfer_tmp_ = nullptr; // no need to copy
   std::unique_ptr<CudaArray<T>> transfer_vecs_ = nullptr;
   std::unique_ptr<InputOutputManager<T>> transfer_iom_ = nullptr;
   std::unique_ptr<PulsedWeightUpdater<T>> transfer_pwu_ = nullptr;
-  std::vector<int> current_col_indices_;
+  std::vector<int> current_slice_indices_;
   bool fully_hidden_ = false;
 
 private:
   int getTransferEvery(int device_idx, int m_batch) const;
 
-  void initialize();
+  void initialize(bool transfer_columns);
 };
 
 } // namespace RPU
diff --git a/src/rpucuda/cuda/rpucuda_vector_device.cu b/src/rpucuda/cuda/rpucuda_vector_device.cu
index ff4becea..d1d44665 100644
--- a/src/rpucuda/cuda/rpucuda_vector_device.cu
+++ b/src/rpucuda/cuda/rpucuda_vector_device.cu
@@ -124,6 +124,7 @@ VectorRPUDeviceCuda<T> &VectorRPUDeviceCuda<T>::operator=(VectorRPUDeviceCuda<T>
   context_vec_ = std::move(other.context_vec_);
   other.context_vec_.clear();
 
+  rw_rng_ = std::move(other.rw_rng_);
   return *this;
 };
 
@@ -275,6 +276,7 @@ void VectorRPUDeviceCuda<T>::runUpdateKernel(
     int m_batch,
     const BitLineMaker<T> *blm,
     const PulsedUpdateMetaParameter<T> &up,
+    const T lr,
     curandState_t *dev_states,
     int one_sided,
     uint32_t *x_counts_chunk,
diff --git a/src/rpucuda/cuda/rpucuda_vector_device.h b/src/rpucuda/cuda/rpucuda_vector_device.h
index 3df9953d..84785d42 100644
--- a/src/rpucuda/cuda/rpucuda_vector_device.h
+++ b/src/rpucuda/cuda/rpucuda_vector_device.h
@@ -42,6 +42,7 @@ template <typename T> class VectorRPUDeviceCuda : public PulsedRPUDeviceCudaBase
     swap(a.current_device_idx_, b.current_device_idx_);
     swap(a.current_update_idx_, b.current_update_idx_);
     swap(a.dev_reduce_weightening_, b.dev_reduce_weightening_);
+    swap(a.rw_rng_, b.rw_rng_);
   };
 
   // implement abstract functions
@@ -74,6 +75,7 @@ template <typename T> class VectorRPUDeviceCuda : public PulsedRPUDeviceCudaBase
       int m_batch,
       const BitLineMaker<T> *blm,
       const PulsedUpdateMetaParameter<T> &up,
+      const T lr,
       curandState_t *dev_states,
       int one_sided = 0,
       uint32_t *x_counts_chunk = nullptr,
@@ -89,9 +91,9 @@ template <typename T> class VectorRPUDeviceCuda : public PulsedRPUDeviceCudaBase
   std::vector<T> getReduceWeightening() const;
 
 protected:
-  int n_devices_ = 0;
-
   virtual void reduceToWeights(CudaContext *c, T *dev_weights);
+
+  int n_devices_ = 0;
   RealWorldRNG<T> rw_rng_{0};
   std::vector<T *> dev_weights_ptrs_;
   std::unique_ptr<CudaArray<T>> dev_weights_vec_ = nullptr;
diff --git a/src/rpucuda/dense_bit_line_maker.cpp b/src/rpucuda/dense_bit_line_maker.cpp
index bce44390..13009964 100644
--- a/src/rpucuda/dense_bit_line_maker.cpp
+++ b/src/rpucuda/dense_bit_line_maker.cpp
@@ -282,14 +282,16 @@ int *DenseBitLineMaker<T>::makeCoincidences(
   T A = 0;
   T B = 0;
   int BL = 0;
+  // negative lr allowed in the below, thus fabs(lr)
   if (up.update_bl_management || up.update_management) {
 
     T x_abs_max = Find_Absolute_Max<T>(x_in, x_size_, x_inc);
     T d_abs_max = Find_Absolute_Max<T>(d_in, d_size_, d_inc);
 
-    up.performUpdateManagement(BL, A, B, up.desired_BL, x_abs_max, d_abs_max, lr, dw_min);
+    up.performUpdateManagement(BL, A, B, up.desired_BL, x_abs_max, d_abs_max, fabs(lr), dw_min);
   } else {
-    up.calculateBlAB(BL, A, B, lr, dw_min);
+
+    up.calculateBlAB(BL, A, B, fabs(lr), dw_min);
   }
 
   if (!containers_allocated_) {
diff --git a/src/rpucuda/rpu_forward_backward_pass.cpp b/src/rpucuda/rpu_forward_backward_pass.cpp
index f391a186..beb489cb 100644
--- a/src/rpucuda/rpu_forward_backward_pass.cpp
+++ b/src/rpucuda/rpu_forward_backward_pass.cpp
@@ -118,108 +118,15 @@ inline T computeNoiseManagement(
 /*********************************************************************/
 /* Noisy forward / backward pass with ADC/DAC and IO management*/
 
-template <typename T> void ForwardBackwardPassIOManaged<T>::allocateContainers() {
-
-  if (!containers_allocated_) {
-    freeContainers();
-
-    tmp_x_values_ = new T[this->x_size_]();
-    tmp_d_values_ = new T[this->d_size_]();
-    containers_allocated_ = true;
-  }
-}
-
-template <typename T> void ForwardBackwardPassIOManaged<T>::freeContainers() {
-
-  if (containers_allocated_) {
-
-    delete[] tmp_d_values_;
-    delete[] tmp_x_values_;
-
-    tmp_d_values_ = nullptr;
-    tmp_x_values_ = nullptr;
-
-    rng_ = nullptr;
-
-    containers_allocated_ = false;
-  }
-}
-
 // ctor
 template <typename T>
 ForwardBackwardPassIOManaged<T>::ForwardBackwardPassIOManaged(
     int x_size, int d_size, std::shared_ptr<RNG<T>> rng)
     : ForwardBackwardPass<T>(x_size, d_size) {
   rng_ = rng;
-  allocateContainers();
   aux_nm_value_ = (T)-1.0;
 }
 
-// dtor
-template <typename T> ForwardBackwardPassIOManaged<T>::~ForwardBackwardPassIOManaged() {
-  freeContainers();
-}
-
-// copy construcutor
-template <typename T>
-ForwardBackwardPassIOManaged<T>::ForwardBackwardPassIOManaged(
-    const ForwardBackwardPassIOManaged<T> &other)
-    : ForwardBackwardPass<T>(other) {
-  f_io_ = other.f_io_;
-  b_io_ = other.b_io_;
-  aux_nm_value_ = other.aux_nm_value_;
-  rng_ = other.rng_;
-  checked_implemented_ = other.checked_implemented_;
-
-  if (other.containers_allocated_) {
-    allocateContainers();
-  }
-}
-
-// copy assignment
-template <typename T>
-ForwardBackwardPassIOManaged<T> &
-ForwardBackwardPassIOManaged<T>::operator=(const ForwardBackwardPassIOManaged<T> &other) {
-
-  ForwardBackwardPassIOManaged<T> tmp(other);
-  swap(*this, tmp);
-  return *this;
-}
-
-// move constructor
-template <typename T>
-ForwardBackwardPassIOManaged<T>::ForwardBackwardPassIOManaged(
-    ForwardBackwardPassIOManaged<T> &&other) {
-
-  *this = std::move(other);
-}
-
-// move assignment
-template <typename T>
-ForwardBackwardPassIOManaged<T> &
-ForwardBackwardPassIOManaged<T>::operator=(ForwardBackwardPassIOManaged<T> &&other) {
-
-  ForwardBackwardPass<T>::operator=(std::move(other));
-
-  f_io_ = other.f_io_;
-  b_io_ = other.b_io_;
-
-  // pointers
-  tmp_d_values_ = other.tmp_d_values_;
-  tmp_x_values_ = other.tmp_x_values_;
-
-  // set pointers to null
-  other.tmp_d_values_ = nullptr;
-  other.tmp_x_values_ = nullptr;
-
-  rng_ = std::move(other.rng_);
-
-  containers_allocated_ = other.containers_allocated_;
-  checked_implemented_ = other.checked_implemented_;
-  aux_nm_value_ = other.aux_nm_value_;
-  return *this;
-}
-
 template <typename T>
 void ForwardBackwardPassIOManaged<T>::setIOPar(
     const IOMetaParameter<T> &f_io, const IOMetaParameter<T> &b_io) {
@@ -232,19 +139,126 @@ void ForwardBackwardPassIOManaged<T>::setIOPar(
   checked_implemented_ = false; // need to check in forward because CUDA also shares this with CPU
 }
 
-template <typename T> void ForwardBackwardPassIOManaged<T>::ensureImplemented() {
-  if (b_io_.w_noise_type != OutputWeightNoiseType::AdditiveConstant &&
-      b_io_.w_noise_type != OutputWeightNoiseType::None) {
-    RPU_FATAL("CPU version of OutputWeightNoiseType not yet implemented ");
+template <typename T> void ForwardBackwardPassIOManaged<T>::ensureImplemented() {}
+
+template <typename T>
+void ForwardBackwardPassIOManaged<T>::applyOutputWeightNoise(
+    T **weights,
+    T *out_values,
+    const int out_size,
+    const int out_inc,
+    const T *in_values,
+    const int in_size,
+    IOMetaParameter<T> &io,
+    bool transposed) {
+
+  if (io.w_noise_type == OutputWeightNoiseType::None) {
+    return;
+  }
+
+  switch (io.w_noise_type) {
+  case OutputWeightNoiseType::AdditiveConstant:
+    if (io.w_noise > 0) {
+      T x_norm = RPU::math::nrm2<T>(in_size, in_values, 1);
+      T w_std = io.w_noise * x_norm;
+      int i_out = 0;
+      PRAGMA_SIMD
+      for (int i = 0; i < out_size; ++i) {
+        out_values[i_out] += w_std * rng_->sampleGauss();
+        i_out += out_inc;
+      }
+    }
+    break;
+  case OutputWeightNoiseType::PCMRead:
+    if (io.w_noise > 0) {
+      T w_std = io.w_noise;
+      tmp_in_values_.resize(in_size);
+
+      PRAGMA_SIMD
+      for (int j = 0; j < in_size; ++j) {
+        tmp_in_values_[j] = in_values[j] * in_values[j];
+      }
+      // likely realtively slow. Since |W|*x.^2 without GEMV...
+      int i_out = 0;
+      for (int i = 0; i < out_size; ++i) {
+        T accum = 0.0;
+        if (transposed) {
+          PRAGMA_SIMD
+          for (int j = 0; j < in_size; ++j) {
+            accum += fabs(weights[j][i]) * tmp_in_values_[j];
+          }
+        } else {
+          PRAGMA_SIMD
+          for (int j = 0; j < in_size; ++j) {
+            accum += fabs(weights[i][j]) * tmp_in_values_[j];
+          }
+        }
+        out_values[i_out] += w_std * sqrtf(accum) * rng_->sampleGauss();
+        i_out += out_inc;
+      }
+    }
+    break;
+  default:
+    RPU_FATAL("Output noise type not implemented")
+  }
+}
+
+template <typename T>
+void ForwardBackwardPassIOManaged<T>::applyIrDrop(
+    T **weights,
+    T *out_values,
+    int out_size,
+    const int out_inc,
+    const T *in_values, // inc = 1 expected
+    const int in_size,
+    IOMetaParameter<T> &io,
+    bool transposed) {
+
+  if (io.ir_drop <= 0.0) {
+    return;
+  }
+  tmp_in_values_.resize(in_size);
+  tmp_c_values_.resize(out_size);
+  tmp_out_values_.resize(out_size);
+
+  T a_scale = in_size / io.ir_drop_Gw_div_gmax;
+  // a_i = sum_j(|w_ij|*|x_j|)*n/Gw*gmax
+  for (int i = 0; i < out_size; ++i) {
+    T accum = 0.0;
+    if (transposed) {
+      PRAGMA_SIMD
+      for (int j = 0; j < in_size; ++j) {
+        accum += fabs(weights[j][i]) * fabs(in_values[j]);
+      }
+    } else {
+      PRAGMA_SIMD
+      for (int j = 0; j < in_size; ++j) {
+        accum += fabs(weights[i][j]) * fabs(in_values[j]);
+      }
+    }
+    T a = a_scale * accum;
+    // c_i = a_i*(a_i*(0.05*a_i - 0.2) + 0.5);
+    tmp_c_values_[i] = a * (a * ((T)0.05 * a - (T)0.2) + (T)0.5);
   }
 
-  if (f_io_.w_noise_type != OutputWeightNoiseType::AdditiveConstant &&
-      f_io_.w_noise_type != OutputWeightNoiseType::None) {
-    RPU_FATAL("CPU version of OutputWeightNoiseType not yet implemented ");
+  // compute x_j*(1-(1-j/n)^2)
+  PRAGMA_SIMD
+  for (int j = 0; j < in_size; ++j) {
+    T p = ((T)1 - (T)j / in_size);
+    tmp_in_values_[j] = in_values[j] * (1 - p * p);
   }
 
-  if (f_io_.bound_management == BoundManagementType::SignalChannel) {
-    RPU_FATAL("CPU version of BoundManagementType::SignalChannel not yet implemented.");
+  // y_i = y_i_ideal - ir_drop*c_i*sum_j(w_ij * x'_j)
+  RPU::math::gemv<T>(
+      CblasRowMajor, transposed ? CblasTrans : CblasNoTrans, this->d_size_, this->x_size_,
+      io.ir_drop, weights[0], this->x_size_, tmp_in_values_.data(), 1, (T)0.0,
+      tmp_out_values_.data(), 1);
+
+  int i_out = 0;
+  PRAGMA_SIMD
+  for (int i = 0; i < out_size; ++i) {
+    out_values[i_out] -= tmp_c_values_[i] * tmp_out_values_[i];
+    i_out += out_inc;
   }
 }
 
@@ -267,8 +281,9 @@ void ForwardBackwardPassIOManaged<T>::forwardVector(
     ensureImplemented();
     checked_implemented_ = true;
   }
+  tmp_x_values_.resize(this->x_size_);
 
-  T *x_value = tmp_x_values_;
+  T *x_value = tmp_x_values_.data();
   T nm_scale_value = computeNoiseManagement(
       x_input, this->x_size_, x_inc, f_io_.noise_management, aux_nm_value_, f_io_);
   bool nm = f_io_.noise_management != NoiseManagementType::None;
@@ -382,15 +397,13 @@ void ForwardBackwardPassIOManaged<T>::forwardVector(
         CblasRowMajor, CblasNoTrans, this->d_size_, this->x_size_, 1.0, weights[0], this->x_size_,
         x_value, 1, f_io_.out_noise, d_output, d_inc);
 
-    if (f_io_.w_noise > 0 && f_io_.w_noise_type == OutputWeightNoiseType::AdditiveConstant) {
-      T x_norm = RPU::math::nrm2<T>(this->x_size_, x_value, 1);
-      T w_std = f_io_.w_noise * x_norm;
-      int i_d = 0;
-      PRAGMA_SIMD
-      for (int i = 0; i < this->d_size_; ++i) {
-        d_output[i_d] += w_std * rng_->sampleGauss();
-        i_d += d_inc;
-      }
+    if (f_io_.w_noise_type != OutputWeightNoiseType::None) {
+      applyOutputWeightNoise(
+          weights, d_output, this->d_size_, d_inc, x_value, this->x_size_, f_io_, false);
+    }
+
+    if (f_io_.ir_drop != (T)0) {
+      applyIrDrop(weights, d_output, this->d_size_, d_inc, x_value, this->x_size_, f_io_, false);
     }
 
     if (sm) {
@@ -461,7 +474,9 @@ void ForwardBackwardPassIOManaged<T>::backwardVector(
   }
 
   // io managed version
-  T *d_value = tmp_d_values_;
+  tmp_d_values_.resize(this->d_size_);
+
+  T *d_value = tmp_d_values_.data();
   T nm_scale_value = computeNoiseManagement(
       d_input, this->d_size_, d_inc, b_io_.noise_management, aux_nm_value_, b_io_);
   bool nm = b_io_.noise_management != NoiseManagementType::None;
@@ -516,15 +531,13 @@ void ForwardBackwardPassIOManaged<T>::backwardVector(
       CblasRowMajor, CblasTrans, this->d_size_, this->x_size_, 1.0, weights[0], this->x_size_,
       d_value, 1, b_io_.out_noise, x_output, x_inc);
 
-  if (b_io_.w_noise > 0 && b_io_.w_noise_type == OutputWeightNoiseType::AdditiveConstant) {
-    T d_norm = RPU::math::nrm2<T>(this->d_size_, d_value, 1);
-    T w_std = b_io_.w_noise * d_norm;
-    int j_x = 0;
-    PRAGMA_SIMD
-    for (int j = 0; j < this->x_size_; j++) {
-      x_output[j_x] += w_std * rng_->sampleGauss();
-      j_x += x_inc;
-    }
+  if (b_io_.w_noise_type != OutputWeightNoiseType::None) {
+    applyOutputWeightNoise(
+        weights, x_output, this->x_size_, x_inc, d_value, this->d_size_, b_io_, true);
+  }
+
+  if (b_io_.ir_drop != (T)0) {
+    applyIrDrop(weights, x_output, this->x_size_, x_inc, d_value, this->d_size_, b_io_, true);
   }
 
   int j_x = 0;
diff --git a/src/rpucuda/rpu_forward_backward_pass.h b/src/rpucuda/rpu_forward_backward_pass.h
index 6a570682..f95c55c5 100644
--- a/src/rpucuda/rpu_forward_backward_pass.h
+++ b/src/rpucuda/rpu_forward_backward_pass.h
@@ -24,12 +24,6 @@ template <typename T> class ForwardBackwardPass {
   explicit ForwardBackwardPass(int x_size, int d_size) : x_size_(x_size), d_size_(d_size){};
   ForwardBackwardPass(){};
 
-  friend void swap(ForwardBackwardPass<T> &a, ForwardBackwardPass<T> &b) noexcept {
-    using std::swap;
-    swap(a.x_size_, b.x_size_);
-    swap(a.d_size_, b.d_size_);
-  }
-
   virtual void forwardVector(
       T **weights,
       const T *x_input,
@@ -53,26 +47,6 @@ template <typename T> class ForwardBackwardPassIOManaged : public ForwardBackwar
 public:
   explicit ForwardBackwardPassIOManaged(int x_size, int d_size, std::shared_ptr<RNG<T>> rng);
   ForwardBackwardPassIOManaged(){};
-  virtual ~ForwardBackwardPassIOManaged();
-
-  ForwardBackwardPassIOManaged(const ForwardBackwardPassIOManaged<T> &);
-  ForwardBackwardPassIOManaged<T> &operator=(const ForwardBackwardPassIOManaged<T> &);
-  ForwardBackwardPassIOManaged(ForwardBackwardPassIOManaged<T> &&);
-  ForwardBackwardPassIOManaged<T> &operator=(ForwardBackwardPassIOManaged<T> &&);
-
-  friend void
-  swap(ForwardBackwardPassIOManaged<T> &a, ForwardBackwardPassIOManaged<T> &b) noexcept {
-    using std::swap;
-    swap(static_cast<ForwardBackwardPass<T> &>(a), static_cast<ForwardBackwardPass<T> &>(b));
-    swap(a.containers_allocated_, b.containers_allocated_);
-    swap(a.f_io_, b.f_io_);
-    swap(a.b_io_, b.b_io_);
-    swap(a.tmp_x_values_, b.tmp_x_values_);
-    swap(a.tmp_d_values_, b.tmp_d_values_);
-    swap(a.aux_nm_value_, b.aux_nm_value_);
-    swap(a.rng_, b.rng_);
-    swap(a.checked_implemented_, b.checked_implemented_);
-  }
 
   void forwardVector(
       T **weights,
@@ -89,19 +63,40 @@ template <typename T> class ForwardBackwardPassIOManaged : public ForwardBackwar
 
   void setIOPar(const IOMetaParameter<T> &f_io_, const IOMetaParameter<T> &b_io_);
 
+protected:
+  void applyOutputWeightNoise(
+      T **weights,
+      T *out_values,
+      const int out_size,
+      const int out_inc,
+      const T *in_values,
+      const int in_size,
+      IOMetaParameter<T> &io,
+      bool transposed);
+
+  void applyIrDrop(
+      T **weights,
+      T *out_values,
+      int out_size,
+      const int out_inc,
+      const T *in_values,
+      const int in_size,
+      IOMetaParameter<T> &io,
+      bool transposed);
+
 private:
   void ensureImplemented();
 
-  void freeContainers();
-  void allocateContainers();
+  std::vector<T> tmp_d_values_;
+  std::vector<T> tmp_x_values_;
 
-  T *tmp_x_values_ = nullptr;
-  T *tmp_d_values_ = nullptr;
+  std::vector<T> tmp_in_values_;
+  std::vector<T> tmp_out_values_;
+  std::vector<T> tmp_c_values_;
 
   T aux_nm_value_ = -1.0;
   IOMetaParameter<T> f_io_;
   IOMetaParameter<T> b_io_;
-  bool containers_allocated_ = false;
   bool checked_implemented_ = false;
   std::shared_ptr<RNG<T>> rng_ = nullptr;
 };
diff --git a/src/rpucuda/rpu_pulsed_device.h b/src/rpucuda/rpu_pulsed_device.h
index 4a5c1d64..edc8c787 100644
--- a/src/rpucuda/rpu_pulsed_device.h
+++ b/src/rpucuda/rpu_pulsed_device.h
@@ -159,6 +159,7 @@ template <typename T> class PulsedRPUDeviceBase : public SimpleRPUDevice<T> {
       T **weights, const PulsedUpdateMetaParameter<T> &up, T current_lr, int m_batch_info){};
 
   inline T getWeightGranularity() const { return weight_granularity_; };
+  virtual T getPulseCountLearningRate(T learning_rate) { return learning_rate; };
 
 protected:
   inline void setWeightGranularity(T weight_granularity) {
diff --git a/src/rpucuda/rpu_pulsed_meta_parameter.cpp b/src/rpucuda/rpu_pulsed_meta_parameter.cpp
index 225ef892..dce73d3a 100644
--- a/src/rpucuda/rpu_pulsed_meta_parameter.cpp
+++ b/src/rpucuda/rpu_pulsed_meta_parameter.cpp
@@ -148,6 +148,7 @@ void PulsedUpdateMetaParameter<T>::calculateBlAB(
     A = (T)0.0;
     B = (T)0.0;
     BL = 0;
+    return;
   }
 
   if (fixed_BL || update_bl_management) {
@@ -158,7 +159,7 @@ void PulsedUpdateMetaParameter<T>::calculateBlAB(
     if ((weight_granularity * desired_BL) < lr) {
       A = (T)1.0;
       B = (T)1.0;
-      BL = (int)ceil(lr / weight_granularity);
+      BL = MAX((int)ceil(lr / weight_granularity), 1);
     } else {
       BL = desired_BL;
       A = sqrt(lr / (weight_granularity * BL));
diff --git a/src/rpucuda/rpu_pulsed_meta_parameter.h b/src/rpucuda/rpu_pulsed_meta_parameter.h
index 1312689e..26e91673 100644
--- a/src/rpucuda/rpu_pulsed_meta_parameter.h
+++ b/src/rpucuda/rpu_pulsed_meta_parameter.h
@@ -29,7 +29,7 @@ enum class NoiseManagementType {
   AverageAbsMaxSingleValue
 };
 
-enum class BoundManagementType { None, Iterative, IterativeWorstCase, Shift, SignalChannel };
+enum class BoundManagementType { None, Iterative, IterativeWorstCase, Shift };
 
 enum class OutputWeightNoiseType { None, AdditiveConstant, PCMRead };
 
@@ -48,6 +48,8 @@ template <typename T> struct IOMetaParameter {
   bool is_perfect = false; // short-cut to use pure floating point (only out_scale will be applied)
 
   OutputWeightNoiseType w_noise_type = OutputWeightNoiseType::None;
+  T ir_drop = (T)0.0;
+  T ir_drop_Gw_div_gmax = (T)5.7143e5; // physical ratio of wire conductance to physical gmax
 
   T inp_bound = (T)1.0;
   T inp_res = (T)1.0 / (pow((T)2.0, (T)7.0) - (T)2.0);
@@ -110,6 +112,10 @@ template <typename T> struct IOMetaParameter {
              << std::endl;
         }
       }
+      if (ir_drop > 0.0) {
+        ss << "\t ir_drop [scale]:\t" << ir_drop << "  (ir_drop_Gw_div_gmax is "
+           << ir_drop_Gw_div_gmax << ")" << std::endl;
+      }
     }
     if (out_scale != 1.0) {
       ss << "\t out_scale:\t\t" << out_scale << std::endl;
diff --git a/src/rpucuda/rpu_transfer_device.cpp b/src/rpucuda/rpu_transfer_device.cpp
index f8c577ca..c7eacb73 100644
--- a/src/rpucuda/rpu_transfer_device.cpp
+++ b/src/rpucuda/rpu_transfer_device.cpp
@@ -65,6 +65,12 @@ void TransferRPUDeviceMetaParameter<T>::printToStream(std::stringstream &ss) con
   ss << std::endl;
 
   // lr
+  if (fast_lr > 0) {
+    ss << "\tfast_lr:\t\t";
+    ss << fast_lr;
+    ss << std::endl;
+  }
+
   ss << "\ttransfer_lr: \t\t";
   if (this->_par_initialized)
     for (size_t k = 0; k < transfer_lr_vec.size(); k++) {
@@ -74,16 +80,23 @@ void TransferRPUDeviceMetaParameter<T>::printToStream(std::stringstream &ss) con
     ss << transfer_lr;
   }
   if (scale_transfer_lr) {
-    ss << " [scaled with current LR]";
+    ss << "\t[scaled with current LR]";
+  }
+  ss << std::endl;
+
+  ss << "\tn_reads_per_transfer: \t" << n_reads_per_transfer;
+  if (transfer_columns) {
+    ss << "\t[reading columns]";
+  } else {
+    ss << "\t[reading rows]";
   }
   ss << std::endl;
 
-  ss << "\tn_cols_per_transfer: \t" << n_cols_per_transfer;
   if (with_reset_prob) {
     ss << "\t[with reset p=" << with_reset_prob << "]";
   }
-  if (random_column) {
-    ss << "\t[random column]";
+  if (random_selection) {
+    ss << "\t[random selection]";
   }
   ss << std::endl;
 
@@ -112,6 +125,14 @@ void TransferRPUDeviceMetaParameter<T>::initializeWithSize(int x_size, int d_siz
     RPU_FATAL("Need at least 2 devices");
   }
 
+  if (transfer_columns) {
+    _in_size = x_size;
+    _out_size = d_size;
+  } else {
+    _in_size = d_size;
+    _out_size = x_size;
+  }
+
   this->update_policy = VectorDeviceUpdatePolicy::SingleFixed;
   this->first_update_idx = 0; // only first is updated
   this->same_context = true;
@@ -155,10 +176,10 @@ void TransferRPUDeviceMetaParameter<T>::initializeWithSize(int x_size, int d_siz
     RPU_FATAL("Expect transfer_lr_vec of size n_devices.");
   }
 
-  if (n_cols_per_transfer > x_size) {
+  if (n_reads_per_transfer > _in_size) {
     // should not be needed anyway
-    n_cols_per_transfer = x_size;
-    RPU_WARNING("too many transfers in one shot. Use x_size instead.");
+    n_reads_per_transfer = _in_size;
+    RPU_WARNING("too many transfers in one shot. Using full transfer instead.");
   }
 
   // TODO: make an default value, where the value of the transfer depends on  x_size
@@ -167,7 +188,7 @@ void TransferRPUDeviceMetaParameter<T>::initializeWithSize(int x_size, int d_siz
     T n = transfer_every;
     for (size_t i = 0; i < n_devices; i++) {
       transfer_every_vec.push_back(n);
-      n *= (T)x_size / n_cols_per_transfer;
+      n *= (T)_in_size / n_reads_per_transfer;
     }
     if (no_self_transfer) {
       transfer_every_vec[n_devices - 1] = 0;
@@ -178,8 +199,16 @@ void TransferRPUDeviceMetaParameter<T>::initializeWithSize(int x_size, int d_siz
     RPU_FATAL("Expect transfer_every_vec to be of length n_devices");
   }
 
+  if (with_reset_prob > 0 && !transfer_columns) {
+    RPU_FATAL("Reset prob is only implemented for column-transfer so far.");
+  }
+
   // IO
-  transfer_io.initializeForForward();
+  if (transfer_columns) {
+    transfer_io.initializeForForward();
+  } else {
+    transfer_io.initializeForBackward();
+  }
 
   // we turn BM off.
   if (transfer_io.bound_management != BoundManagementType::None) {
@@ -230,10 +259,11 @@ TransferRPUDevice<T>::TransferRPUDevice(const TransferRPUDevice<T> &other)
   transfer_fb_pass_ = make_unique<ForwardBackwardPassIOManaged<T>>(*other.transfer_fb_pass_);
   transfer_pwu_ = make_unique<PulsedRPUWeightUpdater<T>>(*other.transfer_pwu_);
 
-  current_col_indices_ = other.current_col_indices_;
+  current_slice_indices_ = other.current_slice_indices_;
   transfer_vecs_ = other.transfer_vecs_;
   transfer_every_ = other.transfer_every_;
   fully_hidden_ = other.fully_hidden_;
+  last_weight_ = other.last_weight_;
 }
 
 // copy assignment
@@ -255,27 +285,30 @@ template <typename T>
 TransferRPUDevice<T> &TransferRPUDevice<T>::operator=(TransferRPUDevice<T> &&other) {
   VectorRPUDevice<T>::operator=(std::move(other));
 
-  current_col_indices_ = std::move(other.current_col_indices_);
+  current_slice_indices_ = std::move(other.current_slice_indices_);
   transfer_vecs_ = std::move(other.transfer_vecs_);
   transfer_every_ = std::move(other.transfer_every_);
   transfer_fb_pass_ = std::move(other.transfer_fb_pass_);
   transfer_pwu_ = std::move(other.transfer_pwu_);
-
+  last_weight_ = std::move(other.last_weight_);
   fully_hidden_ = other.fully_hidden_;
 
   return *this;
 }
 
 template <typename T> void TransferRPUDevice<T>::setTransferVecs(const T *transfer_vecs) {
-  transfer_vecs_.resize(this->x_size_ * this->x_size_); //!!  square matrix
+  T in_size = getPar().getInSize();
+
+  transfer_vecs_.resize(in_size * in_size); //!!  square matrix
   std::fill(transfer_vecs_.begin(), transfer_vecs_.end(), (T)0.0);
 
   if (transfer_vecs == nullptr) {
     // initialize transfer vectors with unit vectors. This might be overridden
-    for (size_t i = 0; i < transfer_vecs_.size(); i += this->x_size_ + 1) {
+    for (size_t i = 0; i < transfer_vecs_.size(); i += in_size + 1) {
       transfer_vecs_[i] = 1.0;
     }
   } else {
+    // Caution: No size check!
     for (size_t i = 0; i < transfer_vecs_.size(); i++) {
       transfer_vecs_[i] = transfer_vecs[i];
     }
@@ -284,8 +317,8 @@ template <typename T> void TransferRPUDevice<T>::setTransferVecs(const T *transf
 
 template <typename T> int TransferRPUDevice<T>::resetCounters(bool force) {
 
-  current_col_indices_.resize(this->n_devices_);
-  std::fill(current_col_indices_.begin(), current_col_indices_.end(), (int)0);
+  current_slice_indices_.resize(this->n_devices_);
+  std::fill(current_slice_indices_.begin(), current_slice_indices_.end(), (int)0);
   return VectorRPUDevice<T>::resetCounters(force);
 }
 /*********************************************************************************/
@@ -304,7 +337,6 @@ void TransferRPUDevice<T>::populate(
       RPU::make_unique<ForwardBackwardPassIOManaged<T>>(this->x_size_, this->d_size_, shared_rng);
 
   transfer_fb_pass_->setIOPar(par.transfer_io, par.transfer_io);
-  // TODO: the OUT_SCALE might be different for the transfer!! How to account for that?!?
 
   transfer_pwu_ =
       RPU::make_unique<PulsedRPUWeightUpdater<T>>(this->x_size_, this->d_size_, shared_rng);
@@ -320,6 +352,22 @@ void TransferRPUDevice<T>::populate(
   fully_hidden_ = getPar().fullyHidden();   // save
 }
 
+/*********************************************************************************/
+/* getPulseCountLearningRate */
+/* Here we compute the LR for the A matrix (the SGD update). Because
+   of the device properties it is beneficial to use a constant LR
+   here, but scale the buffer with the scheduled SGD learning rate
+   later*/
+template <typename T> T TransferRPUDevice<T>::getPulseCountLearningRate(T learning_rate) {
+  const auto &par = getPar();
+
+  if (par.fast_lr > 0) {
+    return par.fast_lr;
+  } else {
+    return learning_rate;
+  }
+}
+
 /*********************************************************************************/
 /* transfer */
 template <typename T> T **TransferRPUDevice<T>::getDeviceWeights(int device_idx) const {
@@ -342,17 +390,42 @@ int TransferRPUDevice<T>::getTransferEvery(int from_device_idx, int m_batch) con
 }
 
 template <typename T>
-void TransferRPUDevice<T>::forwardUpdate(
+void TransferRPUDevice<T>::readVector(int device_idx, const T *in_vec, T *out_vec, T alpha) {
+  T **W = getDeviceWeights(device_idx);
+  if (getPar().transfer_columns) {
+    transfer_fb_pass_->forwardVector(W, in_vec, 1, out_vec, 1, alpha, false);
+  } else {
+    transfer_fb_pass_->backwardVector(W, in_vec, 1, out_vec, 1, alpha);
+  }
+}
+
+template <typename T>
+void TransferRPUDevice<T>::writeVector(
+    int device_idx, const T *in_vec, const T *out_vec, const T lr, const int m_batch_info) {
+
+  T **W = getDeviceWeights(device_idx);
+  if (getPar().transfer_columns) {
+    // in_vec is x_input
+    transfer_pwu_->updateVectorWithDevice(
+        W, in_vec, 1, out_vec, 1, lr, m_batch_info, &*this->rpu_device_vec_[device_idx]);
+  } else {
+    // in_vec is d_input
+    transfer_pwu_->updateVectorWithDevice(
+        W, out_vec, 1, in_vec, 1, lr, m_batch_info, &*this->rpu_device_vec_[device_idx]);
+  }
+}
+
+template <typename T>
+void TransferRPUDevice<T>::readAndUpdate(
     int to_device_idx,
     int from_device_idx,
     const T lr,
-    const T *x_input,
+    const T *vec, // these are the selected transfer vecs
     const int n_vec,
-    const bool trans,
     const T reset_prob,
-    const int i_col) {
+    const int i_slice) {
 
-  if (!lr) {
+  if (lr == 0.0) {
     return;
   }
 
@@ -360,68 +433,66 @@ void TransferRPUDevice<T>::forwardUpdate(
     // self update not supported per default
     return;
   }
+  const auto &par = getPar();
 
-  if (transfer_tmp_.size() < (size_t)this->d_size_) {
-    transfer_tmp_.resize(this->d_size_);
-  }
+  int in_size = par.getInSize();
+  int out_size = par.getOutSize();
+
+  transfer_tmp_.resize(out_size);
 
-  // forward / update
-  T **W;
+  // forward or backward / update
   for (int i = 0; i < n_vec; i++) {
-    const T *x = x_input + i * this->x_size_;
+    const T *v = vec + i * in_size;
 
-    transfer_fb_pass_->forwardVector(
-        this->weights_vec_[from_device_idx], x, 1, &transfer_tmp_[0], 1, (T)1.0, false);
+    readVector(from_device_idx, v, transfer_tmp_.data(), -1.0); // scale -1 for pos update
 
-    // potentially reset here (because of possible same device to-from):
-    // NOTE that with_reset_prob is COL-wise prob (elem device prob is 1)
-    if (this->rw_rng_.sampleUniform() < reset_prob) {
-      W = getDeviceWeights(from_device_idx);
-      this->rpu_device_vec_[from_device_idx]->resetCols(W, i_col, n_vec, 1, this->rw_rng_);
+    if (this->rw_rng_.sampleUniform() < reset_prob && par.transfer_columns) {
+      // potentially reset here (because of possible same device to-from):
+      // NOTE that with_reset_prob is COL-wise prob (elem device prob is 1)
+      T **W_from = getDeviceWeights(from_device_idx);
+      this->rpu_device_vec_[from_device_idx]->resetCols(W_from, i_slice, n_vec, 1, this->rw_rng_);
     }
 
     // update according to device
-    W = getDeviceWeights(to_device_idx);
-
-    transfer_pwu_->updateVectorWithDevice(
-        W, x, 1, &transfer_tmp_[0], 1,
-        -fabs(lr), // need to be negative...
-        n_vec, &*this->rpu_device_vec_[to_device_idx]);
+    writeVector(to_device_idx, v, transfer_tmp_.data(), lr, n_vec);
   }
 }
 
 template <typename T>
 void TransferRPUDevice<T>::transfer(int to_device_idx, int from_device_idx, T current_lr) {
-  int i_col = current_col_indices_[from_device_idx];
+
+  int i_slice = current_slice_indices_[from_device_idx];
   const auto &par = getPar();
-  if (par.random_column) {
-    i_col = MAX(MIN(floor(this->rw_rng_.sampleUniform() * this->x_size_), this->x_size_ - 1), 0);
+
+  int in_size = par.getInSize();
+
+  if (par.random_selection) {
+    i_slice = MAX(MIN(floor(this->rw_rng_.sampleUniform() * in_size), in_size - 1), 0);
   }
 
-  // transfer_vecs_ is always x_size-major (that is trans==false)
-  T *tvec = &transfer_vecs_[0] + i_col * this->x_size_;
-  int n_rest = this->x_size_ - i_col; // actually always x_size
+  // transfer_vecs_ is always in_size-major (that is trans==false)
+  T *tvec = &transfer_vecs_[0] + i_slice * in_size;
+  int n_rest = in_size - i_slice;
 
   T lr = par.getTransferLR(to_device_idx, from_device_idx, current_lr);
 
-  int n_transfer = MIN(par.n_cols_per_transfer, this->x_size_);
+  int n_transfer = MIN(par.n_reads_per_transfer, in_size);
 
   if (n_rest < n_transfer) {
     // rest
 
-    forwardUpdate(
-        to_device_idx, from_device_idx, lr, tvec, n_rest, false, par.with_reset_prob, i_col);
+    readAndUpdate(to_device_idx, from_device_idx, lr, tvec, n_rest, par.with_reset_prob, i_slice);
     // from beginning
-    forwardUpdate(
-        to_device_idx, from_device_idx, lr, &transfer_vecs_[0], n_transfer - n_rest, false,
+    readAndUpdate(
+        to_device_idx, from_device_idx, lr, &transfer_vecs_[0], n_transfer - n_rest,
         par.with_reset_prob, 0);
 
   } else {
-    forwardUpdate(
-        to_device_idx, from_device_idx, lr, tvec, n_transfer, false, par.with_reset_prob, i_col);
+    readAndUpdate(
+        to_device_idx, from_device_idx, lr, tvec, n_transfer, par.with_reset_prob, i_slice);
   }
 
-  current_col_indices_[from_device_idx] = (i_col + n_transfer) % this->x_size_;
+  current_slice_indices_[from_device_idx] = (i_slice + n_transfer) % in_size;
 }
 
 /*********************************************************************************/
@@ -437,6 +508,13 @@ void TransferRPUDevice<T>::doSparseUpdate(
   // we do reduce to weights in the finishUpdateCycle, because transfer is done there, too.
 }
 
+template <typename T>
+void TransferRPUDevice<T>::doDenseUpdate(T **weights, int *coincidences, RNG<T> *rng) {
+
+  this->rpu_device_vec_[0]->doDenseUpdate(this->weights_vec_[0], coincidences, rng);
+  // we do reduce to weights in the finishUpdateCycle, because transfer is done there, too.
+}
+
 template <typename T>
 void TransferRPUDevice<T>::finishUpdateCycle(
     T **weights, const PulsedUpdateMetaParameter<T> &up, T current_lr, int m_batch_info) {
diff --git a/src/rpucuda/rpu_transfer_device.h b/src/rpucuda/rpu_transfer_device.h
index 86b29bbc..6f445b2f 100644
--- a/src/rpucuda/rpu_transfer_device.h
+++ b/src/rpucuda/rpu_transfer_device.h
@@ -51,6 +51,10 @@ template <typename T> class TransferRPUDevice;
    is same for all the tranfers. However, it can be set individually
    by using the vector transfer_lr_vec (with size n_devices-1)
 
+   fast_lr: if defined, it will fix and set the first LR
+   (onto the A matrix). Otherwise the SGD learing
+   rate will be taken
+
  */
 
 template <typename T> struct TransferRPUDeviceMetaParameter : VectorRPUDeviceMetaParameter<T> {
@@ -58,14 +62,17 @@ template <typename T> struct TransferRPUDeviceMetaParameter : VectorRPUDeviceMet
   T gamma = (T)0.0; // weightening factor. [gamma vec optionally inherited from vector]
   T transfer_every = (T)1.0;
   bool units_in_mbatch = false;
-  int n_cols_per_transfer = 1;
+  int n_reads_per_transfer = 1;
   T with_reset_prob = (T)0.0;
   bool no_self_transfer = true;
-  bool random_column = false;
-
+  bool random_selection = false;
+  T fast_lr = 0.0;
   T transfer_lr = (T)1.0;
   std::vector<T> transfer_lr_vec;
   bool scale_transfer_lr = true;
+  bool transfer_columns = true; // or rows
+  int _in_size = 0;
+  int _out_size = 0;
 
   std::vector<T> transfer_every_vec;
 
@@ -86,6 +93,9 @@ template <typename T> struct TransferRPUDeviceMetaParameter : VectorRPUDeviceMet
 
   inline bool fullyHidden() const { return (!gamma && this->gamma_vec.back() == 1.0); };
 
+  inline int getInSize() const { return _in_size; };
+  inline int getOutSize() const { return _out_size; };
+
   std::string getName() const override {
     std::ostringstream ss;
     ss << "Transfer(" << this->vec_par.size() << ")";
@@ -141,8 +151,9 @@ template <typename T> class TransferRPUDevice : public VectorRPUDevice<T> {
     swap(a.transfer_fb_pass_, b.transfer_fb_pass_);
     swap(a.transfer_vecs_, b.transfer_vecs_);
     swap(a.transfer_every_, b.transfer_every_);
-    swap(a.current_col_indices_, b.current_col_indices_);
+    swap(a.current_slice_indices_, b.current_slice_indices_);
     swap(a.fully_hidden_, b.fully_hidden_);
+    swap(a.last_weight_, b.last_weight_);
   }
 
   TransferRPUDeviceMetaParameter<T> &getPar() const override {
@@ -165,25 +176,30 @@ template <typename T> class TransferRPUDevice : public VectorRPUDevice<T> {
 
   void finishUpdateCycle(
       T **weights, const PulsedUpdateMetaParameter<T> &up, T current_lr, int m_batch_info) override;
+  T getPulseCountLearningRate(T learning_rate) override;
 
   virtual int getTransferEvery(int from_device_idx, int m_batch) const;
   virtual void setTransferVecs(const T *transfer_vecs = nullptr);
   virtual void transfer(int to_device_idx, int from_device_idx, T current_lr);
-  virtual void forwardUpdate(
+  virtual void readAndUpdate(
       int to_device_idx,
       int from_device_idx,
       const T lr,
       const T *x_input,
       const int n_vec,
-      const bool trans,
       const T reset_prob,
       const int i_col);
   virtual const T *getTransferVecs() const { return &transfer_vecs_[0]; };
+  virtual void writeVector(
+      int device_idx, const T *in_vec, const T *out_vec, const T lr, const int m_batch_info);
+  virtual void readVector(int device_idx, const T *in_vec, T *out_vec, T alpha);
 
   void doSparseUpdate(
       T **weights, int i, const int *x_signed_indices, int x_count, int d_sign, RNG<T> *rng)
       override;
 
+  void doDenseUpdate(T **weights, int *coincidences, RNG<T> *rng) override;
+
 protected:
   void populate(const TransferRPUDeviceMetaParameter<T> &par, RealWorldRNG<T> *rng);
   void reduceToWeights(T **weights) const override;
@@ -194,12 +210,12 @@ template <typename T> class TransferRPUDevice : public VectorRPUDevice<T> {
   std::unique_ptr<PulsedRPUWeightUpdater<T>> transfer_pwu_ = nullptr;
   std::vector<T> transfer_vecs_;
   std::vector<T> transfer_every_;
-  std::vector<int> current_col_indices_;
+  std::vector<int> current_slice_indices_;
+  bool fully_hidden_ = false;
+  T **last_weight_ = nullptr;
 
   // no need to swap/copy.
-  T **last_weight_ = nullptr;
   std::vector<T> transfer_tmp_;
-  bool fully_hidden_ = false;
 };
 
 } // namespace RPU
diff --git a/src/rpucuda/rpu_vector_device.cpp b/src/rpucuda/rpu_vector_device.cpp
index cdba6e0d..6b608d27 100644
--- a/src/rpucuda/rpu_vector_device.cpp
+++ b/src/rpucuda/rpu_vector_device.cpp
@@ -451,6 +451,20 @@ void VectorRPUDevice<T>::doSparseUpdate(
   }
 }
 
+template <typename T>
+void VectorRPUDevice<T>::doDenseUpdate(T **weights, int *coincidences, RNG<T> *rng) {
+  if (getPar().singleDeviceUpdate()) {
+    rpu_device_vec_[current_device_idx_]->doDenseUpdate(
+        weights_vec_[current_device_idx_], coincidences, rng);
+  } else {
+
+    for (size_t k = 0; k < rpu_device_vec_.size(); k++) {
+      rpu_device_vec_[k]->doDenseUpdate(weights_vec_[k], coincidences, rng);
+    }
+  }
+  this->reduceToWeights(weights);
+}
+
 template <typename T>
 void VectorRPUDevice<T>::finishUpdateCycle(
     T **weights, const PulsedUpdateMetaParameter<T> &up, T current_lr, int m_batch_info) {
diff --git a/src/rpucuda/rpu_vector_device.h b/src/rpucuda/rpu_vector_device.h
index 0e1d02f5..e426b70a 100644
--- a/src/rpucuda/rpu_vector_device.h
+++ b/src/rpucuda/rpu_vector_device.h
@@ -199,6 +199,7 @@ template <typename T> class VectorRPUDevice : public PulsedRPUDeviceBase<T> {
   void doSparseUpdate(
       T **weights, int i, const int *x_signed_indices, int x_count, int d_sign, RNG<T> *rng)
       override;
+  void doDenseUpdate(T **weights, int *coincidences, RNG<T> *rng) override;
 
 protected:
   void populate(const VectorRPUDeviceMetaParameter<T> &par, RealWorldRNG<T> *rng);
diff --git a/src/rpucuda/rpu_weight_updater.cpp b/src/rpucuda/rpu_weight_updater.cpp
index e37f6ad6..268da418 100644
--- a/src/rpucuda/rpu_weight_updater.cpp
+++ b/src/rpucuda/rpu_weight_updater.cpp
@@ -133,12 +133,23 @@ void PulsedRPUWeightUpdater<T>::setUpPar(const PulsedUpdateMetaParameter<T> &up)
 
 template <typename T>
 bool PulsedRPUWeightUpdater<T>::checkForFPUpdate(AbstractRPUDevice<T> *rpu_device_in) {
+
+  if (rpu_device_in == nullptr) {
+    return true;
+  }
+  if (rpu_device_in->implements() == DeviceUpdateType::FloatingPoint) {
+    return true;
+  }
+  if (rpu_device_in->isPulsedDevice() && up_.pulse_type == PulseType::None) {
+    return true;
+  }
   if (rpu_device_in->hasDirectUpdate()) {
+    // also FP has direct, but that is handled above
     return false;
   }
-  return (up_.pulse_type == PulseType::None) || (rpu_device_in == nullptr) ||
-         !rpu_device_in->isPulsedDevice() ||
-         (rpu_device_in->implements() == DeviceUpdateType::FloatingPoint);
+  // omitting !isPulsedDevice
+
+  return false;
 }
 
 template <typename T>
@@ -177,14 +188,16 @@ void PulsedRPUWeightUpdater<T>::updateVectorWithDevice(
 
   // pulsed device update
   rpu_device->initUpdateCycle(weights, up_, learning_rate, m_batch_info);
+  // potentially modify the LR from the device side
+  T pc_learning_rate = rpu_device->getPulseCountLearningRate(learning_rate);
 
   if (sblm_->supports(up_.pulse_type)) {
     // envoke sparse bit line maker to get the counts and indices
     int BL = sblm_->makeCounts(
-        x_input, x_inc, d_input, d_inc, &*rng_, learning_rate < 0 ? -learning_rate : learning_rate,
-        weight_granularity, up_);
+        x_input, x_inc, d_input, d_inc, &*rng_,
+        pc_learning_rate < 0 ? -pc_learning_rate : pc_learning_rate, weight_granularity, up_);
     // positive LR actually means that positive signs *decrease* the weight (as in SGD).
-    int lr_sign = learning_rate < 0 ? -1 : 1;
+    int lr_sign = pc_learning_rate < 0 ? -1 : 1;
 
     if (BL > 0) {
 
@@ -223,10 +236,10 @@ void PulsedRPUWeightUpdater<T>::updateVectorWithDevice(
   } else {
     // use dense update
     int *coincidences = dblm_->makeCoincidences(
-        x_input, x_inc, d_input, d_inc, &*rng_, learning_rate, weight_granularity, up_);
+        x_input, x_inc, d_input, d_inc, &*rng_, pc_learning_rate, weight_granularity, up_);
     rpu_device->doDenseUpdate(weights, coincidences, &*rng_);
   }
-
+  // always the current SGD learning rate is given here
   rpu_device->finishUpdateCycle(weights, up_, learning_rate, m_batch_info);
 }
 
diff --git a/src/rpucuda/rpu_weight_updater.h b/src/rpucuda/rpu_weight_updater.h
index 8e9e0e9d..7028a1d9 100644
--- a/src/rpucuda/rpu_weight_updater.h
+++ b/src/rpucuda/rpu_weight_updater.h
@@ -94,6 +94,7 @@ template <typename T> class PulsedRPUWeightUpdater : public RPUWeightUpdater<T>
       uint32_t *d_counts);
 
   void setUpPar(const PulsedUpdateMetaParameter<T> &up);
+  inline const PulsedUpdateMetaParameter<T> &getUpPar() const { return up_; };
 
 private:
   void freeContainers();
diff --git a/src/rpucuda/weight_modifier.cpp b/src/rpucuda/weight_modifier.cpp
index 62378b50..7dddd860 100644
--- a/src/rpucuda/weight_modifier.cpp
+++ b/src/rpucuda/weight_modifier.cpp
@@ -126,6 +126,7 @@ void WeightModifier<T>::apply(
         new_weights[i] += amax * sig * rw_rng_.sampleGauss();
       }
     }
+
     break;
   }
 


Test patch:
diff --git a/cmake/dependencies_test.cmake b/cmake/dependencies_test.cmake
index 41c37e3a..c1ed8bdf 100644
--- a/cmake/dependencies_test.cmake
+++ b/cmake/dependencies_test.cmake
@@ -13,8 +13,8 @@ list(APPEND CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake/Modules)
 if(BUILD_TEST)
   include(ExternalProject)
   ExternalProject_Add(GTest
-    URL               https://github.com/google/googletest/archive/release-1.10.0.zip
-    URL_HASH          MD5=82358affdd7ab94854c8ee73a180fc53
+    URL               https://github.com/google/googletest/archive/release-1.11.0.zip
+    URL_HASH          MD5=52943a59cefce0ae0491d4d2412c120b
     CMAKE_ARGS        "-DCMAKE_CXX_FLAGS=-D_GLIBCXX_USE_CXX11_ABI\=0"
     INSTALL_COMMAND   ""
   )
diff --git a/src/rpucuda/cuda/bit_line_maker_test.cpp b/src/rpucuda/cuda/bit_line_maker_test.cpp
index 2e48c929..4c66f5c0 100644
--- a/src/rpucuda/cuda/bit_line_maker_test.cpp
+++ b/src/rpucuda/cuda/bit_line_maker_test.cpp
@@ -308,8 +308,7 @@ class BitLineMakerParTestFixture : public ::testing::TestWithParam<bool> {
     up.res = resolution;
     up.sto_round = false;
 
-    context_container = RPU::make_unique<CudaContext>(-1, false); // blocking for timings
-    context = &*context_container;
+    context = &context_container;
 
     blm = RPU::make_unique<BitLineMaker<num_t>>(context, x_size, d_size);
 
@@ -352,7 +351,7 @@ class BitLineMakerParTestFixture : public ::testing::TestWithParam<bool> {
     delete[] d_counts;
   }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
   int nK32;
   int x_size;
diff --git a/src/rpucuda/cuda/io_iterator_test.cpp b/src/rpucuda/cuda/io_iterator_test.cpp
index d3d801e5..2a25317e 100644
--- a/src/rpucuda/cuda/io_iterator_test.cpp
+++ b/src/rpucuda/cuda/io_iterator_test.cpp
@@ -64,8 +64,7 @@ template <typename T> class IteratorTestFixture : public ::testing::Test {
       unfolded_vector2[i] = unfolded_vector[i];
     }
 
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
+    context = &context_container;
     dev_orig_vector = RPU::make_unique<CudaArray<T>>(context, orig_matrix_size * N, orig_vector);
     dev_orig_vector2 = RPU::make_unique<CudaArray<T>>(context, orig_matrix_size * N, orig_vector2);
     dev_unfolded_vector =
@@ -100,7 +99,7 @@ template <typename T> class IteratorTestFixture : public ::testing::Test {
     delete[] batch_indices;
   };
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
   std::unique_ptr<CudaArray<T>> dev_orig_vector;
   std::unique_ptr<CudaArray<T>> dev_orig_vector2;
diff --git a/src/rpucuda/cuda/io_manager_test.cpp b/src/rpucuda/cuda/io_manager_test.cpp
index 92c46f44..a7b1fa78 100644
--- a/src/rpucuda/cuda/io_manager_test.cpp
+++ b/src/rpucuda/cuda/io_manager_test.cpp
@@ -93,8 +93,7 @@ class IOManagerTestFixture : public ::testing::TestWithParam<bool> {
     d_size = 490;
     m_batch = 100;
 
-    context_container = RPU::make_unique<CudaContext>(-1, false); // blocking for timings
-    context = &*context_container;
+    context = &context_container;
     iom = RPU::make_unique<InputOutputManager<num_t>>(context, x_size, d_size);
 
     num_t bound = 0.8;
@@ -170,7 +169,7 @@ class IOManagerTestFixture : public ::testing::TestWithParam<bool> {
     delete[] W;
   }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   int x_size;
diff --git a/src/rpucuda/cuda/rpucuda_expstep_test.cpp b/src/rpucuda/cuda/rpucuda_expstep_test.cpp
index e7360567..f8238c55 100644
--- a/src/rpucuda/cuda/rpucuda_expstep_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_expstep_test.cpp
@@ -38,8 +38,7 @@ template <typename Tio> class RPUCudaExpstepStateTestFixture : public ::testing:
 public:
   void SetUp() {
 
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
+    context = &context_container;
 
     x_size = 100;
     d_size = 110;
@@ -140,7 +139,7 @@ template <typename Tio> class RPUCudaExpstepStateTestFixture : public ::testing:
     d_vec2_batch.resize(m_batch * d_size);
   }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   std::unique_ptr<RPUPulsed<num_t>> layer_pulsed;
diff --git a/src/rpucuda/cuda/rpucuda_linearstep_test.cpp b/src/rpucuda/cuda/rpucuda_linearstep_test.cpp
index 5db2e443..d2075a79 100644
--- a/src/rpucuda/cuda/rpucuda_linearstep_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_linearstep_test.cpp
@@ -38,8 +38,7 @@ class RPUCudaLinearStepTestFixture : public ::testing::TestWithParam<bool> {
 public:
   void SetUp() {
 
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
+    context = &context_container;
 
     x_size = 100;
     d_size = 110;
@@ -139,7 +138,7 @@ class RPUCudaLinearStepTestFixture : public ::testing::TestWithParam<bool> {
     d_vec_batch.resize(m_batch * d_size);
   }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   std::unique_ptr<RPUPulsed<num_t>> layer_pulsed;
diff --git a/src/rpucuda/cuda/rpucuda_onesided_device_test.cpp b/src/rpucuda/cuda/rpucuda_onesided_device_test.cpp
index 8169b814..27daa139 100644
--- a/src/rpucuda/cuda/rpucuda_onesided_device_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_onesided_device_test.cpp
@@ -36,8 +36,7 @@ class RPUDeviceCudaTestFixture : public ::testing::TestWithParam<int> {
 public:
   void SetUp() {
 
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
+    context = &context_container;
 
     x_size = 5;
     d_size = 8;
@@ -119,7 +118,7 @@ class RPUDeviceCudaTestFixture : public ::testing::TestWithParam<int> {
   std::unique_ptr<AbstractRPUDevice<num_t>> rpu_device;
   std::unique_ptr<AbstractRPUDeviceCuda<num_t>> rpucuda_device;
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 };
 
diff --git a/src/rpucuda/cuda/rpucuda_powstep_test.cpp b/src/rpucuda/cuda/rpucuda_powstep_test.cpp
index e505d920..294d84b7 100644
--- a/src/rpucuda/cuda/rpucuda_powstep_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_powstep_test.cpp
@@ -38,8 +38,7 @@ class RPUCudaPowStepTestFixture : public ::testing::TestWithParam<num_t> {
 public:
   void SetUp() {
 
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
+    context = &context_container;
 
     x_size = 100;
     d_size = 110;
@@ -140,7 +139,7 @@ class RPUCudaPowStepTestFixture : public ::testing::TestWithParam<num_t> {
     d_vec_batch.resize(m_batch * d_size);
   }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   std::unique_ptr<RPUPulsed<num_t>> layer_pulsed;
diff --git a/src/rpucuda/cuda/rpucuda_pulsed_device_test.cpp b/src/rpucuda/cuda/rpucuda_pulsed_device_test.cpp
index ee194b38..954e7a45 100644
--- a/src/rpucuda/cuda/rpucuda_pulsed_device_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_pulsed_device_test.cpp
@@ -43,8 +43,7 @@ template <typename DeviceParT> class RPUDeviceTestFixture : public ::testing::Te
 public:
   void SetUp() {
 
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
+    context = &context_container;
 
     this->x_size = 53;
     this->d_size = 43;
@@ -226,7 +225,7 @@ template <typename DeviceParT> class RPUDeviceTestFixture : public ::testing::Te
 
   void TearDown() { Array_2D_Free(refweights); }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   std::unique_ptr<RPUPulsed<num_t>> layer_pulsed;
diff --git a/src/rpucuda/cuda/rpucuda_pulsed_test.cpp b/src/rpucuda/cuda/rpucuda_pulsed_test.cpp
index 96c40e83..476b15b0 100644
--- a/src/rpucuda/cuda/rpucuda_pulsed_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_pulsed_test.cpp
@@ -24,8 +24,7 @@
 #include <numeric>
 #include <random>
 
-#define TOLERANCE 1e-5
-
+#define TOLERANCE 5e-4
 namespace {
 
 using namespace RPU;
@@ -111,8 +110,7 @@ template <typename T1> class RPUCudaPulsedTestFixture : public ::testing::Test {
     layer_pulsed->disp();
 
     // copy-construct on GPU
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
+    context = &context_container;
 
     culayer_pulsed = RPU::make_unique<RPUCudaPulsed<T>>(context, *layer_pulsed);
     culayer_pulsed->disp();
@@ -174,7 +172,7 @@ template <typename T1> class RPUCudaPulsedTestFixture : public ::testing::Test {
     delete[] w_other_trans;
   }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   std::unique_ptr<RPUPulsed<T>> layer_pulsed;
@@ -212,8 +210,7 @@ template <typename T1> class RPUCudaPulsedTestFixtureNoNoise : public ::testing:
 
     is_test = false;
 
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
+    context = &context_container;
 
     x_size = 41;
     d_size = 50;
@@ -243,6 +240,8 @@ template <typename T1> class RPUCudaPulsedTestFixtureNoNoise : public ::testing:
     p_io.inp_sto_round = false;
     p_io.out_res = -1;
 
+    p_io.ir_drop = 0.3;
+
     p_io.noise_management = NoiseManagementType::AbsMax;
     p_io.bound_management = BoundManagementType::Iterative;
 
@@ -333,7 +332,7 @@ template <typename T1> class RPUCudaPulsedTestFixtureNoNoise : public ::testing:
 
   void TearDown() { delete batch_indices; }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   std::unique_ptr<RPUSimple<T>> layer;
@@ -479,10 +478,10 @@ TYPED_TEST(RPUCudaPulsedTestFixtureNoNoise, ForwardMatrixBiasNoNoise) {
 
   for (int i = 0; i < this->d_size * this->m_batch; i++) {
     // std::cout << i << ": d1 " << this->d1[i] << " vs d2 " << this->d2[i] << std::endl;
-    ASSERT_NEAR(this->d1[i], this->d2[i], 1e-5);
+    ASSERT_NEAR(this->d1[i], this->d2[i], TOLERANCE);
     ASSERT_NEAR(
         this->d2[i], this->d3[i],
-        1e-5); // because of noise/bound management: in/out bound should not matter
+        TOLERANCE); // because of noise/bound management: in/out bound should not matter
   }
 }
 
@@ -701,8 +700,8 @@ TYPED_TEST(RPUCudaPulsedTestFixtureNoNoise, BackwardMatrixBiasNoNoise) {
   this->x3 = this->x_vec_batch;
 
   for (int i = 0; i < (this->x_size - 1) * this->m_batch; i++) {
-    ASSERT_NEAR(this->x1[i], this->x2[i], 1e-5);
-    ASSERT_NEAR(this->x2[i], this->x3[i], 1e-5);
+    ASSERT_NEAR(this->x1[i], this->x2[i], TOLERANCE);
+    ASSERT_NEAR(this->x2[i], this->x3[i], TOLERANCE);
   }
 }
 
@@ -1102,7 +1101,7 @@ TYPED_TEST(RPUCudaPulsedTestFixture, BackwardMatrix) {
         cuavg[k] += cuweights[i][j] / nloop;                                                       \
         avg[k] += weights[i][j] / nloop;                                                           \
         if (no_noise) {                                                                            \
-          ASSERT_NEAR(cuweights[i][j], weights[i][j], 1e-5);                                       \
+          ASSERT_NEAR(cuweights[i][j], weights[i][j], TOLERANCE);                                  \
         }                                                                                          \
         cusig[k] += cuweights[i][j] * cuweights[i][j] / nloop;                                     \
         sig[k] += weights[i][j] * weights[i][j] / nloop;                                           \
diff --git a/src/rpucuda/cuda/rpucuda_simple_device_test.cpp b/src/rpucuda/cuda/rpucuda_simple_device_test.cpp
index 57fcfe77..40c3f3be 100644
--- a/src/rpucuda/cuda/rpucuda_simple_device_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_simple_device_test.cpp
@@ -44,8 +44,7 @@ template <typename DeviceParT> class RPUDeviceTestFixture : public ::testing::Te
 public:
   void SetUp() {
 
-    this->context_container = RPU::make_unique<CudaContext>(-1, false);
-    this->context = &*context_container;
+    context = &context_container;
 
     this->x_size = 253;
     this->d_size = 243;
@@ -229,7 +228,7 @@ template <typename DeviceParT> class RPUDeviceTestFixture : public ::testing::Te
     Array_2D_Free(refweights);
   }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   std::unique_ptr<RPUPulsed<num_t>> pulsed;
diff --git a/src/rpucuda/cuda/rpucuda_test.cpp b/src/rpucuda/cuda/rpucuda_test.cpp
index 4f5148f2..c48b162f 100644
--- a/src/rpucuda/cuda/rpucuda_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_test.cpp
@@ -35,8 +35,7 @@ template <typename T> class RPUCudaSimpleTestFixture : public ::testing::Test {
 public:
   void SetUp() {
 
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
+    context = &context_container;
 
     is_test = true;
 
@@ -83,7 +82,7 @@ template <typename T> class RPUCudaSimpleTestFixture : public ::testing::Test {
     d_vec2.resize(d_size);
   }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   std::unique_ptr<RPUSimple<T>> layer_simple;
diff --git a/src/rpucuda/cuda/rpucuda_transfer_device_test.cpp b/src/rpucuda/cuda/rpucuda_transfer_device_test.cpp
index baeb671d..898dbe57 100644
--- a/src/rpucuda/cuda/rpucuda_transfer_device_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_transfer_device_test.cpp
@@ -11,7 +11,9 @@
  */
 
 #include "cuda_util.h"
+#include "rpu_pulsed.h"
 #include "rpucuda_constantstep_device.h"
+#include "rpucuda_pulsed.h"
 #include "rpucuda_transfer_device.h"
 #include "test_helper.h"
 #include "gtest/gtest.h"
@@ -35,11 +37,9 @@ class RPUDeviceCudaTestFixture : public ::testing::TestWithParam<float> {
 public:
   void SetUp() {
 
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
-
-    x_size = 2;
+    x_size = 5;
     d_size = 3;
+    m_batch = 1;
 
     w_ref = Array_2D_Get<num_t>(d_size, x_size);
 
@@ -51,6 +51,21 @@ class RPUDeviceCudaTestFixture : public ::testing::TestWithParam<float> {
     for (int i = 0; i < d_size * x_size; i++) {
       weights[0][i] = 0;
     }
+    rx.resize(x_size * m_batch);
+    rd.resize(d_size * m_batch);
+
+    unsigned int seed = std::chrono::system_clock::now().time_since_epoch().count();
+    std::default_random_engine generator{seed};
+    std::uniform_real_distribution<num_t> udist(-1.2, 1.2);
+    auto urnd = std::bind(udist, generator);
+
+    // just assign some numbers from the weight matrix
+    for (int i = 0; i < x_size * m_batch; i++)
+      rx[i] = urnd();
+
+    for (int j = 0; j < d_size * m_batch; j++) {
+      rd[j] = urnd();
+    }
 
     up.pulse_type =
         PulseType::StochasticCompressed; // nodevice would skip the entire transfer mechanism
@@ -75,7 +90,8 @@ class RPUDeviceCudaTestFixture : public ::testing::TestWithParam<float> {
     dp->gamma = GetParam(); // meaning fully hidden
     dp->transfer_lr = 1;
     dp->transfer_every = x_size;
-    dp->n_cols_per_transfer = 1;
+    dp->n_reads_per_transfer = 1;
+    dp->transfer_columns = true;
     dp->units_in_mbatch = false;
 
     dp->transfer_io.inp_res = -1;
@@ -85,29 +101,28 @@ class RPUDeviceCudaTestFixture : public ::testing::TestWithParam<float> {
     dp->transfer_up = up;
     dp->transfer_up.pulse_type = PulseType::None; // perfect transfer
 
-    rng = new RNG<num_t>(0);
-
-    up_pwu = RPU::make_unique<PulsedWeightUpdater<num_t>>(context, x_size, d_size);
+    up_pwu = RPU::make_unique<PulsedWeightUpdater<num_t>>(&context, x_size, d_size);
 
     rpu_device = this->dp->createDeviceUnique(x_size, d_size, &rw_rng);
-    rpucuda_device = AbstractRPUDeviceCuda<num_t>::createFromUnique(context, *rpu_device);
+    rpucuda_device = AbstractRPUDeviceCuda<num_t>::createFromUnique(&context, *rpu_device);
 
-    dev_weights = RPU::make_unique<CudaArray<num_t>>(context, x_size * d_size);
+    dev_weights = RPU::make_unique<CudaArray<num_t>>(&context, x_size * d_size);
     dev_weights->assignTranspose(weights[0], d_size, x_size);
-    context->synchronize();
+    context.synchronize();
   };
 
   void TearDown() {
     Array_2D_Free<num_t>(weights);
     Array_2D_Free<num_t>(w_ref);
     delete dp;
-    delete rng;
   };
+  CudaContext context{-1, false};
 
-  int x_size, d_size, colidx;
+  int x_size, d_size, colidx, m_batch;
   num_t lifetime;
   num_t **weights;
   num_t **w_ref;
+  std::vector<num_t> rx, rd, w, w2;
   PulsedUpdateMetaParameter<num_t> up;
   TransferRPUDeviceMetaParameter<num_t> *dp;
   ConstantStepRPUDeviceMetaParameter<num_t> dp_cs;
@@ -116,11 +131,6 @@ class RPUDeviceCudaTestFixture : public ::testing::TestWithParam<float> {
   RealWorldRNG<num_t> rw_rng;
   std::unique_ptr<AbstractRPUDevice<num_t>> rpu_device;
   std::unique_ptr<AbstractRPUDeviceCuda<num_t>> rpucuda_device;
-
-  RNG<num_t> *rng;
-
-  std::unique_ptr<CudaContext> context_container;
-  CudaContext *context;
 };
 
 // define the tests
@@ -136,14 +146,14 @@ TEST_P(RPUDeviceCudaTestFixture, onSetWeights) {
   for (int i = 0; i < this->x_size * this->d_size; i++) {
     this->weights[0][i] = w_ref[0][i];
   }
-  dev_weights = RPU::make_unique<CudaArray<num_t>>(context, x_size * d_size);
+  dev_weights = RPU::make_unique<CudaArray<num_t>>(&context, x_size * d_size);
   dev_weights->assignTranspose(weights[0], d_size, x_size);
-  context->synchronize();
+  context.synchronize();
 
   if (rpu_device->onSetWeights(this->weights)) {
     rpucuda_device->populateFrom(*rpu_device); // device pars have changed (due to onSetWeights)
   }
-  context->synchronize();
+  context.synchronize();
 
   std::vector<num_t> w_vec =
       static_cast<TransferRPUDeviceCuda<num_t> *>(&*rpucuda_device)->getHiddenWeights();
@@ -167,18 +177,18 @@ TEST_P(RPUDeviceCudaTestFixture, Update) {
   dp->transfer_lr = 0; // no transfer here
   // just newly create from paramerers
   rpu_device = dp->createDeviceUnique(this->x_size, this->d_size, &this->rw_rng);
-  rpucuda_device = AbstractRPUDeviceCuda<num_t>::createFromUnique(context, *rpu_device);
+  rpucuda_device = AbstractRPUDeviceCuda<num_t>::createFromUnique(&context, *rpu_device);
 
-  CudaArray<num_t> dev_x(context, this->x_size);
+  CudaArray<num_t> dev_x(&context, this->x_size);
   dev_x.setConst(1.0);
-  CudaArray<num_t> dev_d(context, this->d_size);
+  CudaArray<num_t> dev_d(&context, this->d_size);
   dev_d.setConst(-1.0);
-  context->synchronize();
+  context.synchronize();
 
   if (rpu_device->onSetWeights(this->weights)) {
     rpucuda_device->populateFrom(*rpu_device); // device pars have changed (due to onSetWeights)
   }
-  context->synchronize();
+  context.synchronize();
 
   up_pwu->update(
       dev_x.getDataConst(), dev_d.getDataConst(), dev_weights->getData(), &*rpucuda_device,
@@ -188,7 +198,7 @@ TEST_P(RPUDeviceCudaTestFixture, Update) {
       false, // trans
       false);
   // should update all weight values of the hidden weight by -1
-  context->synchronize();
+  context.synchronize();
   auto w_vec = static_cast<TransferRPUDeviceCuda<num_t> *>(&*rpucuda_device)->getHiddenWeights();
 
   // update only on fast [nothing to transfer for first row]
@@ -200,7 +210,7 @@ TEST_P(RPUDeviceCudaTestFixture, Update) {
                                            // should actually be exactly one
     s += w_vec[i];
   }
-  std::cout << "Average weight " << s / size << " (Expected is 1.0)" << std::endl;
+  // std::cout << "Average weight " << s / size << " (Expected is 1.0)" << std::endl;
 
   // visible  weights  not
   for (int i = 0; i < size; i++) {
@@ -217,18 +227,18 @@ TEST_P(RPUDeviceCudaTestFixture, Update) {
   }
 }
 
-TEST_P(RPUDeviceCudaTestFixture, UpdateAndTransfer) {
+TEST_P(RPUDeviceCudaTestFixture, UpdateAndTransferColumns) {
 
-  CudaArray<num_t> dev_x(context, this->x_size);
+  CudaArray<num_t> dev_x(&context, this->x_size);
   dev_x.setConst(1.0);
-  CudaArray<num_t> dev_d(context, this->d_size);
+  CudaArray<num_t> dev_d(&context, this->d_size);
   dev_d.setConst(-1.0);
-  context->synchronize();
+  context.synchronize();
 
   if (rpu_device->onSetWeights(this->weights)) {
     rpucuda_device->populateFrom(*rpu_device); // device pars have changed (due to onSetWeights)
   }
-  context->synchronize();
+  context.synchronize();
 
   for (int k = 0; k < this->x_size; k++) {
     up_pwu->update(
@@ -241,7 +251,7 @@ TEST_P(RPUDeviceCudaTestFixture, UpdateAndTransfer) {
   }
   // weight values of the hidden weights should be x_size and first
   // col should be transfered once (that is set to x_size also)
-  context->synchronize();
+  context.synchronize();
   auto w_vec = static_cast<TransferRPUDeviceCuda<num_t> *>(&*rpucuda_device)->getHiddenWeights();
   dev_weights->copyTo(weights[0]);
   dev_weights->assignTranspose(weights[0], x_size, d_size);
@@ -255,10 +265,11 @@ TEST_P(RPUDeviceCudaTestFixture, UpdateAndTransfer) {
     ASSERT_EQ(w_vec[i], (num_t)this->x_size);
     s += w_vec[i];
   }
-  std::cout << "Average weight " << s / size << " (Expected is " << this->x_size << ")"
-            << std::endl;
-
   // only first col of  weights should be transferred
+  // for (int i = 0; i < size; i++) {
+  //   std::cout << "[" << i / x_size << "," << i % x_size << "]: " << w_vec[i] << ", "
+  //  	      << w_vec[i + size] << ", " << this->weights[0][i] << std::endl;
+  // }
 
   for (int i = 0; i < size; i++) {
     if (GetParam()) {
@@ -271,22 +282,217 @@ TEST_P(RPUDeviceCudaTestFixture, UpdateAndTransfer) {
   // reduce to weight.
   std::vector<num_t> rw =
       static_cast<TransferRPUDeviceCuda<num_t> *>(&*rpucuda_device)->getReduceWeightening();
-  std::cout << "RW: " << rw[0] << "," << rw[1] << std::endl;
+  // std::cout << "RW: " << rw[0] << "," << rw[1] << std::endl;
+  // for (int i = 0; i < size; i++) {
+  //   std::cout << "[" << i / x_size << "," << i % x_size << "]: " << w_vec[i] << ", "
+  // 	      << w_vec[i + size] << ", " << this->weights[0][i] << std::endl;
+  // }
+
   for (int i = 0; i < size; i++) {
-    std::cout << "[" << i / x_size << "," << i % x_size << "]: " << w_vec[i] << ", "
-              << w_vec[i + size] << ", " << this->weights[0][i] << std::endl;
+    if (GetParam()) {
+      // std::cout << "[" << i / x_size << "," << i % x_size << "]: " << w_vec[i + size] <<
+      // std::endl;
+      ASSERT_FLOAT_EQ(this->weights[0][i], rw[0] * w_vec[i] + rw[1] * w_vec[i + size]);
+    } else {
+      ASSERT_FLOAT_EQ(this->weights[0][i], i % x_size ? 0.0 : (num_t)this->x_size * rw[1]);
+    }
   }
+}
+
+TEST_P(RPUDeviceCudaTestFixture, UpdateAndTransferRows) {
+
+  dp->transfer_every = d_size;
+  dp->transfer_columns = false;
+  rpu_device = this->dp->createDeviceUnique(x_size, d_size, &rw_rng);
+  rpucuda_device = AbstractRPUDeviceCuda<num_t>::createFromUnique(&context, *rpu_device);
+
+  CudaArray<num_t> dev_x(&context, this->x_size);
+  dev_x.setConst(1.0);
+  CudaArray<num_t> dev_d(&context, this->d_size);
+  dev_d.setConst(-1.0);
+  context.synchronize();
+
+  if (rpu_device->onSetWeights(this->weights)) {
+    rpucuda_device->populateFrom(*rpu_device); // device pars have changed (due to onSetWeights)
+  }
+  context.synchronize();
+
+  for (int k = 0; k < d_size; k++) {
+    up_pwu->update(
+        dev_x.getDataConst(), dev_d.getDataConst(), dev_weights->getData(), &*rpucuda_device,
+        this->up,
+        1.0,   // lr
+        1,     // batch
+        false, // trans
+        false);
+  }
+  // weight values of the hidden weights should be d_size and first
+  // row should be transfered once thus also d_size
+  context.synchronize();
+  auto w_vec = static_cast<TransferRPUDeviceCuda<num_t> *>(&*rpucuda_device)->getHiddenWeights();
+  dev_weights->copyTo(weights[0]);
+  dev_weights->assignTranspose(weights[0], x_size, d_size);
+  dev_weights->copyTo(weights[0]);
+
+  // update only on fast [nothing to transfer for first row]
+  int size = this->d_size * this->x_size;
+  // hidden weights updated
+  num_t s = 0;
+  for (int i = 0; i < size; i++) {
+    ASSERT_EQ(w_vec[i], (num_t)this->d_size);
+    s += w_vec[i];
+  }
+  // only first row of  weights should be transferred
+  // for (int i = 0; i < size; i++) {
+  //   std::cout << "[" << i / x_size << "," << i % x_size << "]: " << w_vec[i] << ", "
+  //  	      << w_vec[i + size] << ", " << this->weights[0][i] << std::endl;
+  // }
+
+  for (int i = 0; i < size; i++) {
+    // wvec should be x_size major
+    if (GetParam()) {
+      ASSERT_FLOAT_EQ(w_vec[i + size], i >= x_size ? 0.0 : (num_t)this->d_size);
+    } else {
+      ASSERT_FLOAT_EQ(w_vec[i + size], 0.0); // should not be used
+    }
+  }
+
+  // reduce to weight.
+  std::vector<num_t> rw =
+      static_cast<TransferRPUDeviceCuda<num_t> *>(&*rpucuda_device)->getReduceWeightening();
 
   for (int i = 0; i < size; i++) {
     if (GetParam()) {
-      std::cout << "[" << i / x_size << "," << i % x_size << "]: " << w_vec[i + size] << std::endl;
       ASSERT_FLOAT_EQ(this->weights[0][i], rw[0] * w_vec[i] + rw[1] * w_vec[i + size]);
     } else {
-      ASSERT_FLOAT_EQ(this->weights[0][i], i % x_size ? 0.0 : (num_t)this->x_size * rw[1]);
+      ASSERT_FLOAT_EQ(this->weights[0][i], i >= x_size ? 0.0 : (num_t)this->d_size * rw[1]);
     }
   }
 }
 
+TEST_P(RPUDeviceCudaTestFixture, CUDAvsCPU) {
+
+  PulsedMetaParameter<num_t> p;
+  p.up = up;
+  p.up.x_res_implicit = 0.01;
+  p.up.d_res_implicit = 0.01;
+  p.up.pulse_type = PulseType::DeterministicImplicit;
+
+  dp->transfer_up.pulse_type = PulseType::DeterministicImplicit;
+  dp->transfer_up.x_res_implicit = 0.01;
+  dp->transfer_up.d_res_implicit = 0.01;
+  dp->transfer_up.desired_BL = 10;
+
+  dp->transfer_io.is_perfect = true;
+
+  CudaArray<num_t> dev_x(&context, x_size * m_batch, rx.data());
+  CudaArray<num_t> dev_d(&context, d_size * m_batch, rd.data());
+  context.synchronize();
+
+  auto *rpu = new RPUPulsed<num_t>(x_size, d_size);
+  rpu->populateParameter(&p, dp);
+  rpu->setWeights(this->weights[0]);
+
+  rpu->setLearningRate(1.0);
+  auto *rpucuda = new RPUCudaPulsed<num_t>(context.getStream(), *rpu);
+  rpucuda->setLearningRate(1.0);
+  context.synchronize();
+
+  int size = this->d_size * this->x_size;
+  // double check whether weights are correct
+  w.resize(x_size * d_size);
+  rpu->getWeights(w.data());
+  w2.resize(x_size * d_size);
+  rpucuda->getWeights(w2.data());
+
+  for (int i = 0; i < size; i++) {
+    ASSERT_NEAR(w[i], w2[i], 1.0e-4);
+  }
+
+  context.synchronize();
+  for (int k = 0; k < size; k++) {
+    rpu->update(rx.data(), rd.data(), false, m_batch, false, false);
+    rpucuda->update(dev_x.getData(), dev_d.getData(), false, m_batch, false, false);
+  }
+  // weight values of the hidden weights should be x_size and first
+  // col should be transfered once (that is set to dw_min)
+  context.synchronize();
+  w.resize(x_size * d_size);
+  rpu->getWeights(w.data());
+  w2.resize(x_size * d_size);
+  rpucuda->getWeights(w2.data());
+
+  for (int i = 0; i < size; i++) {
+    ASSERT_NEAR(w[i], w2[i], 1.0e-4);
+  }
+}
+
+TEST_P(RPUDeviceCudaTestFixture, CUDAvsCPURows) {
+
+  PulsedMetaParameter<num_t> p;
+  p.up = up;
+  p.up.x_res_implicit = 0.01;
+  p.up.d_res_implicit = 0.01;
+  p.up.pulse_type = PulseType::DeterministicImplicit;
+
+  dp->transfer_columns = false;
+  dp->transfer_up.pulse_type = PulseType::DeterministicImplicit;
+  dp->transfer_up.x_res_implicit = 0.01;
+  dp->transfer_up.d_res_implicit = 0.01;
+  dp->transfer_up.desired_BL = 10;
+
+  dp->transfer_io.is_perfect = true;
+
+  CudaArray<num_t> dev_x(&context, x_size * m_batch, rx.data());
+  CudaArray<num_t> dev_d(&context, d_size * m_batch, rd.data());
+  context.synchronize();
+
+  auto *rpu = new RPUPulsed<num_t>(x_size, d_size);
+  rpu->populateParameter(&p, dp);
+  rpu->setWeights(this->weights[0]);
+  // rpu->disp();
+  // rpu->dispParameter();
+  rpu->setLearningRate(1.0);
+  auto *rpucuda = new RPUCudaPulsed<num_t>(context.getStream(), *rpu);
+  rpucuda->setLearningRate(1.0);
+  context.synchronize();
+  // rpucuda->disp();
+  // rpucuda->dispParameter();
+
+  int size = this->d_size * this->x_size;
+  // double check whether weights are correct
+  w.resize(x_size * d_size);
+  rpu->getWeights(w.data());
+  w2.resize(x_size * d_size);
+  rpucuda->getWeights(w2.data());
+
+  for (int i = 0; i < size; i++) {
+    ASSERT_NEAR(w[i], w2[i], 1.0e-4);
+  }
+
+  context.synchronize();
+  for (int k = 0; k < size; k++) {
+    rpu->update(rx.data(), rd.data(), false, m_batch, false, false);
+    rpucuda->update(dev_x.getData(), dev_d.getData(), false, m_batch, false, false);
+  }
+  // weight values of the hidden weights should be x_size and first
+  // col should be transfered once (that is set to dw_min)
+  context.synchronize();
+  w.resize(x_size * d_size);
+  rpu->getWeights(w.data());
+  w2.resize(x_size * d_size);
+  rpucuda->getWeights(w2.data());
+
+  // std::cout << "CPU vs CUDA" << std::endl;
+  // for (int i = 0; i < size; i++) {
+  //   std::cout << "[" << i / x_size << "," << i % x_size << "]: " << w[i] << " vs. " << w2[i] <<
+  //   std::endl;
+  // }
+  for (int i = 0; i < size; i++) {
+    ASSERT_NEAR(w[i], w2[i], 1.0e-4);
+  }
+}
+
 } // namespace
 
 int main(int argc, char **argv) {
diff --git a/src/rpucuda/cuda/rpucuda_vector_device_test.cpp b/src/rpucuda/cuda/rpucuda_vector_device_test.cpp
index 9e1e8d89..0c75dca8 100644
--- a/src/rpucuda/cuda/rpucuda_vector_device_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_vector_device_test.cpp
@@ -54,7 +54,7 @@ template <> void setAdditionalValues(TransferRPUDeviceMetaParameter<num_t> *vp,
   vp->transfer_io.out_res = -1;
   vp->transfer_every = 1;
   vp->transfer_lr = 0.5;
-  vp->n_cols_per_transfer = 1;
+  vp->n_reads_per_transfer = 1;
   vp->gamma = value;
 }
 
@@ -62,8 +62,7 @@ template <typename VectorDeviceParT> class RPUDeviceTestFixture : public ::testi
 public:
   void SetUp() {
 
-    this->context_container = RPU::make_unique<CudaContext>(-1, false);
-    this->context = &*context_container;
+    context = &context_container;
 
     this->x_size = 4;
     this->d_size = 5;
@@ -261,7 +260,7 @@ template <typename VectorDeviceParT> class RPUDeviceTestFixture : public ::testi
   };
   void TearDown() { Array_2D_Free(refweights); }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   std::unique_ptr<RPUPulsed<num_t>> layer_pulsed;
diff --git a/src/rpucuda/cuda/weight_clipper_cuda_test.cpp b/src/rpucuda/cuda/weight_clipper_cuda_test.cpp
index 1709b8f4..a94a1410 100644
--- a/src/rpucuda/cuda/weight_clipper_cuda_test.cpp
+++ b/src/rpucuda/cuda/weight_clipper_cuda_test.cpp
@@ -36,8 +36,7 @@ class WeightClipperTestFixture : public ::testing::TestWithParam<bool> {
 public:
   void SetUp() {
 
-    context_container = RPU::make_unique<CudaContext>(-1, false);
-    context = &*context_container;
+    context = &context_container;
 
     x_size = 100;
     d_size = 130;
@@ -69,7 +68,7 @@ class WeightClipperTestFixture : public ::testing::TestWithParam<bool> {
     delete[] v;
   }
 
-  std::unique_ptr<CudaContext> context_container;
+  CudaContext context_container{-1, false};
   CudaContext *context;
 
   int x_size, d_size;
diff --git a/src/rpucuda/rpu_transfer_device_test.cpp b/src/rpucuda/rpu_transfer_device_test.cpp
index 5d183f5d..92693617 100644
--- a/src/rpucuda/rpu_transfer_device_test.cpp
+++ b/src/rpucuda/rpu_transfer_device_test.cpp
@@ -243,6 +243,61 @@ TEST_P(RPUDeviceTestFixture, doSparseUpdateWithTransfer) {
   delete rpu_device;
 }
 
+TEST_P(RPUDeviceTestFixture, doSparseUpdateWithTransferRows) {
+  this->dp->transfer_columns = false;
+  rpu_device = this->dp->createDevice(this->x_size, this->d_size, &this->rw_rng);
+  rpu_device->onSetWeights(this->weights); // all zero
+  rpu_device->initUpdateCycle(this->weights, this->up, 1, 1);
+
+  // last row update, net 1 (n_pos-n_neg)
+  float dx = rpu_device->getWeightGranularity() * (n_pos - n_neg);
+  int rowidx = this->d_size - 1;
+  rpu_device->doSparseUpdate(
+      this->weights, rowidx, this->x_indices, this->n_neg + this->n_pos, (num_t)-1.0, this->rng);
+  rpu_device->finishUpdateCycle(this->weights, this->up, 1, 1); // to signal the end of the update
+
+  for (int i = 0; i < this->d_size; i++) {
+    rpu_device->initUpdateCycle(this->weights, this->up, 1, 1);   // to signal the end of the update
+    rpu_device->finishUpdateCycle(this->weights, this->up, 1, 1); // to signal the end of the update
+  }
+  // should have transfered all cols once.
+
+  num_t ***w_vec = rpu_device->getWeightVec();
+
+  // update only on fast [nothing to transfer for first row]
+  for (size_t m = 0; m < this->dp->vec_par.size(); m++) {
+    for (int j = 0; j < this->d_size; j++) {
+      for (int i = 0; i < this->x_size; i++) {
+        if (j == rowidx && i == this->colidx) {
+          if (m == 1 && GetParam() == 0) {
+            // if fullyHidden (GetParam==0) then actual weight is updated directly
+            ASSERT_FLOAT_EQ(w_vec[m][j][i], 0);
+          } else {
+            // should be fully transferred, since transfer_lr = 1
+            ASSERT_FLOAT_EQ(w_vec[m][j][i], dx);
+          }
+        } else {
+          ASSERT_FLOAT_EQ(w_vec[m][j][i], 0);
+        }
+      }
+    }
+  }
+
+  // reduce to weight. Weightening does not necessarily sums to 1
+  const num_t *reduce_weightening = rpu_device->getReduceWeightening();
+  num_t sc = reduce_weightening[0] + reduce_weightening[1];
+  for (int j = 0; j < this->d_size; j++) {
+    for (int i = 0; i < this->x_size; i++) {
+      if (j == rowidx && i == this->colidx) {
+        ASSERT_FLOAT_EQ(this->weights[j][i], sc * dx);
+      } else {
+        ASSERT_FLOAT_EQ(this->weights[j][i], 0);
+      }
+    }
+  }
+  delete rpu_device;
+}
+
 TEST_P(RPUDeviceTestFixture, Decay) {
 
   for (int i = 0; i < this->x_size * this->d_size; i++) {
diff --git a/tests/helpers/testcases.py b/tests/helpers/testcases.py
index a26f4921..b519f73f 100644
--- a/tests/helpers/testcases.py
+++ b/tests/helpers/testcases.py
@@ -23,17 +23,17 @@
 class AihwkitTestCase(TestCase):
     """Test case that contains common asserts and functions for aihwkit."""
 
-    def assertTensorAlmostEqual(self, tensor_a, tensor_b):
+    def assertTensorAlmostEqual(self, tensor_a, tensor_b, decimal=6):
         """Assert that two tensors are almost equal."""
         # pylint: disable=invalid-name
         array_a = tensor_a.detach().cpu().numpy()
         array_b = tensor_b.detach().cpu().numpy()
-        assert_array_almost_equal(array_a, array_b)
+        assert_array_almost_equal(array_a, array_b, decimal=decimal)
 
-    def assertNotAlmostEqualTensor(self, tensor_a, tensor_b):
+    def assertNotAlmostEqualTensor(self, tensor_a, tensor_b, decimal=6):
         """Assert that two tensors are not equal."""
         # pylint: disable=invalid-name
-        assert_raises(AssertionError, self.assertTensorAlmostEqual, tensor_a, tensor_b)
+        assert_raises(AssertionError, self.assertTensorAlmostEqual, tensor_a, tensor_b, decimal)
 
 
 class ParametrizedTestCase(AihwkitTestCase):