Repository: IBM/aihwkit (language: python)

Issue:
WeightModifierType.DROP_CONNECT not implemented
## Description

While training with [InferenceRPUConfig](https://aihwkit.readthedocs.io/en/latest/api/aihwkit.simulator.configs.configs.html#aihwkit.simulator.configs.configs.InferenceRPUConfig) and the modifier type `WeightModifierType.DROP_CONNECT`, I get a "not implemented" error.

However this option is [documented](https://aihwkit.readthedocs.io/en/latest/api/aihwkit.simulator.parameters.enums.html#aihwkit.simulator.parameters.enums.WeightModifierType.DROP_CONNECT). It should also work with another modifier type if `pdrop>0`, which doesn't seem true. If I create a modifier as `WeightModifierParameter(type=WeightModifierType.ADD_NORMAL, std_dev=0, pdrop=0.2)` there is no error message and no drop-connect.

## How to reproduce

Based on [this example](https://github.com/IBM/aihwkit/blob/master/examples/05_simple_layer_hardware_aware.py).

```python
# pylint: disable=invalid-name

# Imports from PyTorch.
from torch import Tensor
from torch.nn.functional import mse_loss

# Imports from aihwkit.
from aihwkit.nn import AnalogLinear
from aihwkit.optim import AnalogSGD
from aihwkit.simulator.configs import (
    InferenceRPUConfig,
    WeightNoiseType,
    WeightClipType,
    WeightModifierType,
)
from aihwkit.inference import PCMLikeNoiseModel, GlobalDriftCompensation

if __name__ == '__main__':
    # Prepare the datasets (input and expected output).
    x = Tensor([[0.1, 0.2, 0.4, 0.3], [0.2, 0.1, 0.1, 0.3]])
    y = Tensor([[1.0, 0.5], [0.7, 0.3]])

    # Define a single-layer network, using inference/hardware-aware training tile
    rpu_config = InferenceRPUConfig()
    rpu_config.forward.out_res = -1.0  # Turn off (output) ADC discretization.
    rpu_config.forward.w_noise_type = WeightNoiseType.ADDITIVE_CONSTANT
    rpu_config.forward.w_noise = 0.02  # Short-term w-noise.

    rpu_config.clip.type = WeightClipType.FIXED_VALUE
    rpu_config.clip.fixed_value = 1.0
    rpu_config.modifier.pdrop = 0.1  # Drop connect
    rpu_config.modifier.type = WeightModifierType.DROP_CONNECT  # <-- Modifier type changed from the example ----
    rpu_config.modifier.rel_to_actual_wmax = True

    # Inference noise model.
    rpu_config.noise_model = PCMLikeNoiseModel(g_max=25.0)

    # drift compensation
    rpu_config.drift_compensation = GlobalDriftCompensation()

    print(rpu_config)

    model = AnalogLinear(4, 2, bias=True, rpu_config=rpu_config)

    # Define an analog-aware optimizer, preparing it for using the layers.
    opt = AnalogSGD(model.parameters(), lr=0.1)
    opt.regroup_param_groups(model)

    print(next(model.analog_tiles()).tile)

    for epoch in range(100):
        # Delete old gradient
        opt.zero_grad()
        # Add the training Tensor to the model (input).
        pred = model(x)
        # Add the expected output Tensor.
        loss = mse_loss(pred, y)
        # Run training (backward propagation).
        loss.backward()

        opt.step()

        print("Loss error: {:.16f}".format(loss))

    model.eval()

    # Do inference with drift.
    pred_before = model(x)

    print("Correct value:\t {}".format(y.detach().cpu().numpy().flatten()))
    print("Prediction after training:\t {}".format(pred_before.detach().cpu().numpy().flatten()))

    for t_inference in [0.0, 1.0, 20.0, 1000.0, 1e5]:
        model.drift_analog_weights(t_inference)
        pred_drift = model(x)
        print(
            "Prediction after drift (t={}, correction={:1.3f}):\t {}".format(
                t_inference,
                next(model.analog_tiles()).alpha.cpu().numpy(),
                pred_drift.detach().cpu().numpy().flatten(),
            )
        )
```

It gives this error message:

```
Traceback (most recent call last):
  File ".../src/debug.py", line 61, in <module>
    opt.step()
  File ".../lib/python3.9/site-packages/torch/optim/optimizer.py", line 280, in wrapper
    out = func(*args, **kwargs)
  File ".../lib/python3.9/site-packages/torch/utils/_contextlib.py", line 115, in decorate_context
    return func(*args, **kwargs)
  File ".../lib/python3.9/site-packages/aihwkit/optim/analog_optimizer.py", line 145, in step
    param.analog_tile.post_update_step()
  File ".../lib/python3.9/site-packages/torch/utils/_contextlib.py", line 115, in decorate_context
    return func(*args, **kwargs)
  File ".../lib/python3.9/site-packages/aihwkit/simulator/tiles/inference.py", line 307, in post_update_step
    self.tile.modify_weights(self._tmp["weight_modify_params"])  # type: ignore
RuntimeError: Error in weight_modifier.cpp:204  Requested WeightModifierType not implemented in CPU version.
```

The message seems to imply that it is implemented with CUDA, but it is not. I get this error message:
```
...
RuntimeError: Error in weight_modifier_cuda.cu:593  Requested WeightModifierType not implemented.
```

The default exemple include `WeightModifierType.ADD_NORMAL` with a drop connect. Which doesn't apply the drop connect, according to my tests.

## Expected behavior

The `WeightModifierType.DROP_CONNECT` modifier and the `pdrop>0` with a different modifier should actually apply drop connect. 

## Other information

I might be wrong, but I think the drop connect was working with an old version of aihwkit.

* Pytorch version: 2.2.1
* Package version: 0.8 and 0.9
* OS: Linux
* Python version: 3.9.2 and 3.10.12
* Install method: with pip for CPU version and pip + [CUDA.Dockerfile](https://github.com/IBM/aihwkit/blob/master/CUDA.Dockerfile) for CUDA version



Discussion / hints:

Fix patch:
diff --git a/src/rpucuda/cuda/bit_line_maker.cu b/src/rpucuda/cuda/bit_line_maker.cu
index 04d81960..6d6255ee 100644
--- a/src/rpucuda/cuda/bit_line_maker.cu
+++ b/src/rpucuda/cuda/bit_line_maker.cu
@@ -162,7 +162,7 @@ namespace RPU {
                                                                                                    \
   int batch_idx = m_batch - 1;                                                                     \
   for (int j = 0; j < size * nK32; j++) {                                                          \
-    counts[j] = tmp32[(batch_idx)*size * nK32 + j];                                                \
+    counts[j] = tmp32[(batch_idx) * size * nK32 + j];                                              \
   }                                                                                                \
                                                                                                    \
   CUDA_CALL(cudaDeviceSynchronize());                                                              \
diff --git a/src/rpucuda/cuda/cuda_util.cu b/src/rpucuda/cuda/cuda_util.cu
index dc184fdd..9d051aeb 100644
--- a/src/rpucuda/cuda/cuda_util.cu
+++ b/src/rpucuda/cuda/cuda_util.cu
@@ -20,7 +20,7 @@
 
 #define DISABLE_SHARED_MUTEX 1
 
-#define IDX2F(i, j, ld) ((((j)-1) * (ld)) + ((i)-1))
+#define IDX2F(i, j, ld) ((((j) - 1) * (ld)) + ((i) - 1))
 
 // this should be not necesary, because device id is set individually
 // per thread. However, if one would want to use 2 GPUs within one
diff --git a/src/rpucuda/cuda/pwu_kernel.h b/src/rpucuda/cuda/pwu_kernel.h
index 4dc9c0db..2e231106 100644
--- a/src/rpucuda/cuda/pwu_kernel.h
+++ b/src/rpucuda/cuda/pwu_kernel.h
@@ -103,7 +103,7 @@ getIdxToLoad<false>(int batch_index, int count_index, int sz, int m_batch, int c
                                                                                                    \
     uint32_t x_and_d = x & d;                                                                      \
     n = __popc(x_and_d);                                                                           \
-    n -= ((x_and_d)&1);                                                                            \
+    n -= ((x_and_d) & 1);                                                                          \
                                                                                                    \
     if (nK32 > 1) {                                                                                \
       int i_d = 0;                                                                                 \
@@ -451,8 +451,8 @@ __global__ void kernelUpdateWFunctor(
       } sum_n = 0;                                                                                 \
       last_negative = 0;                                                                           \
                                                                                                    \
-      int pos_n = __popc((~negative) & n); int neg_n = __popc((negative)&n); T dw_pos = (T)pos_n;  \
-      T dw_neg = (T)neg_n;                                                                         \
+      int pos_n = __popc((~negative) & n); int neg_n = __popc((negative) & n);                     \
+      T dw_pos = (T)pos_n; T dw_neg = (T)neg_n;                                                    \
                                                                                                    \
       if (noise_std_dw > (T)0.0) {                                                                 \
         if (pos_n > 0) {                                                                           \
@@ -937,7 +937,7 @@ __global__ void kernelUpdateWBatchSharedSum(
       }
 
     } // within range
-  }   // batch strides
+  } // batch strides
   if (within_range) {
     weights[idx] = w;
   }
@@ -1010,7 +1010,7 @@ __global__ void kernelUpdateWBatchSharedSumBoundCheck(
       }
 
     } // within range
-  }   // batch strides
+  } // batch strides
 
   if (within_range) {
     weights[idx] = w;
@@ -1111,7 +1111,7 @@ __global__ void kernelUpdateWBatchSharedFunctor(
       } // batch
 
     } // within range
-  }   // batch strides
+  } // batch strides
 
   if (within_range) {
     weights[idx] = w;
@@ -1505,7 +1505,7 @@ __global__ void kernelUpdateWBatchSharedWeightOutputFunctor(
       } // batch
 
     } // within range
-  }   // batch strides
+  } // batch strides
 
   if (within_range) {
     weights[idx] = w;
diff --git a/src/rpucuda/cuda/pwu_kernel_parameter_base.h b/src/rpucuda/cuda/pwu_kernel_parameter_base.h
index 390f684a..09ef1fcd 100644
--- a/src/rpucuda/cuda/pwu_kernel_parameter_base.h
+++ b/src/rpucuda/cuda/pwu_kernel_parameter_base.h
@@ -127,7 +127,7 @@ template <typename T> class PWUKernelParameterBase {
     if (this->use_bo64 == 1) {
       this->use_bo64 = 2;
     }
-  };                                                        // debug hack
+  }; // debug hack
   inline void forceNonTrans() { this->out_trans = false; }; // debug hack
   inline void force32() { this->use_bo64 = 0; };            // debug hack
 
diff --git a/src/rpucuda/cuda/rpucuda_hidden_device.cu b/src/rpucuda/cuda/rpucuda_hidden_device.cu
index a0fd3e46..362721bd 100644
--- a/src/rpucuda/cuda/rpucuda_hidden_device.cu
+++ b/src/rpucuda/cuda/rpucuda_hidden_device.cu
@@ -25,9 +25,9 @@ namespace RPU {
   } else {                                                                                         \
     hw += hs_dw;                                                                                   \
   }                                                                                                \
-  if (hw > (T)1.0 || hw < (T)-1.0) {                                                               \
+  if (hw > (T)1.0 || hw < (T) - 1.0) {                                                             \
                                                                                                    \
-    T dw = (hw > (T)1) ? ((T)par_4.w) : ((T)-par_4.y);                                             \
+    T dw = (hw > (T)1) ? ((T)par_4.w) : ((T) - par_4.y);                                           \
     hw = (T)0.0;                                                                                   \
     if (noise_std_dw > (T)0.0) {                                                                   \
       T stoch_value = curand_normal(&local_state);                                                 \
diff --git a/src/rpucuda/cuda/rpucuda_onesided_device.h b/src/rpucuda/cuda/rpucuda_onesided_device.h
index 3a9fc0fe..878729c1 100644
--- a/src/rpucuda/cuda/rpucuda_onesided_device.h
+++ b/src/rpucuda/cuda/rpucuda_onesided_device.h
@@ -53,7 +53,7 @@ template <typename T> class OneSidedRPUDeviceCuda : public VectorRPUDeviceCuda<T
   };
   OneSidedRPUDeviceCuda<T> *clone() const override { return new OneSidedRPUDeviceCuda<T>(*this); };
 
-  void setHiddenUpdateIdx(int idx) override{};
+  void setHiddenUpdateIdx(int idx) override {};
 
   void runUpdateKernel(
       pwukp_t<T> kpars,
diff --git a/src/rpucuda/cuda/rpucuda_simple_device.h b/src/rpucuda/cuda/rpucuda_simple_device.h
index 4143fbb3..e0d8f206 100644
--- a/src/rpucuda/cuda/rpucuda_simple_device.h
+++ b/src/rpucuda/cuda/rpucuda_simple_device.h
@@ -49,7 +49,7 @@ template <typename T> class AbstractRPUDeviceCuda {
       T *d_buffer) = 0;
   virtual bool hasDirectUpdate() const = 0;
   virtual int getHiddenUpdateIdx() const { return 0; };
-  virtual void setHiddenUpdateIdx(int idx){};
+  virtual void setHiddenUpdateIdx(int idx) {};
   virtual void dumpExtra(RPU::state_t &extra, const std::string prefix) = 0;
   virtual void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict) = 0;
   virtual void populateFrom(const AbstractRPUDevice<T> &rpu_device) = 0;
diff --git a/src/rpucuda/cuda/rpucuda_transfer_device.h b/src/rpucuda/cuda/rpucuda_transfer_device.h
index 635443c8..c241e00a 100644
--- a/src/rpucuda/cuda/rpucuda_transfer_device.h
+++ b/src/rpucuda/cuda/rpucuda_transfer_device.h
@@ -52,7 +52,7 @@ template <typename T> class TransferRPUDeviceCuda : public VectorRPUDeviceCuda<T
   };
   TransferRPUDeviceCuda<T> *clone() const override { return new TransferRPUDeviceCuda<T>(*this); };
 
-  void setHiddenUpdateIdx(int idx) override{};
+  void setHiddenUpdateIdx(int idx) override {};
   void dumpExtra(RPU::state_t &extra, const std::string prefix) override;
   void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict) override;
 
diff --git a/src/rpucuda/cuda/update_management_helper.cu b/src/rpucuda/cuda/update_management_helper.cu
index 01c24b0d..919a9078 100644
--- a/src/rpucuda/cuda/update_management_helper.cu
+++ b/src/rpucuda/cuda/update_management_helper.cu
@@ -363,7 +363,7 @@ int debugKernelTranslateTransFormatToBatchOrder64Format(
       for (int i = 0; i < k; i++) { // k is smaller than 32 because nK32==1
         kagg_t current_cK = Kc + i;
         kagg_t iB = (current_cK) >> 5;
-        int ibit = (current_cK)&0x1f;
+        int ibit = (current_cK) & 0x1f;
         if ((c & (one << i)) > 0) {
           counts_out_ref[iB + idx * nBref] |= ((uint64_t)1) << ibit;
         }
diff --git a/src/rpucuda/cuda/weight_clipper_cuda.h b/src/rpucuda/cuda/weight_clipper_cuda.h
index cb8d9305..3d12b45d 100644
--- a/src/rpucuda/cuda/weight_clipper_cuda.h
+++ b/src/rpucuda/cuda/weight_clipper_cuda.h
@@ -26,8 +26,8 @@ template <typename T> class WeightClipperCuda {
 
   void apply(T *weights, const WeightClipParameter &wclpar);
 
-  void dumpExtra(RPU::state_t &extra, const std::string prefix){};
-  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict){};
+  void dumpExtra(RPU::state_t &extra, const std::string prefix) {};
+  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict) {};
 
 private:
   CudaContextPtr context_ = nullptr;
diff --git a/src/rpucuda/cuda/weight_modifier_cuda.cu b/src/rpucuda/cuda/weight_modifier_cuda.cu
index c51fbd43..0380a739 100644
--- a/src/rpucuda/cuda/weight_modifier_cuda.cu
+++ b/src/rpucuda/cuda/weight_modifier_cuda.cu
@@ -589,6 +589,11 @@ void WeightModifierCuda<T>::apply(
     break;
   }
 
+  case WeightModifierType::DropConnect: {
+    // to not get into the default branch. Drop connect is handled below
+    break;
+  }
+
   default:
     RPU_FATAL("Requested WeightModifierType not implemented.");
   }
diff --git a/src/rpucuda/cuda/weight_remapper_cuda.h b/src/rpucuda/cuda/weight_remapper_cuda.h
index e1e90606..522dc872 100644
--- a/src/rpucuda/cuda/weight_remapper_cuda.h
+++ b/src/rpucuda/cuda/weight_remapper_cuda.h
@@ -42,8 +42,8 @@ template <typename T> class WeightRemapperCuda {
       T *biases = nullptr,
       int *channel_exceded = nullptr);
 
-  void dumpExtra(RPU::state_t &extra, const std::string prefix){};
-  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict){};
+  void dumpExtra(RPU::state_t &extra, const std::string prefix) {};
+  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict) {};
 
 private:
   CudaContextPtr context_ = nullptr;
diff --git a/src/rpucuda/dense_bit_line_maker.h b/src/rpucuda/dense_bit_line_maker.h
index 1f87e417..ed9ac2f4 100644
--- a/src/rpucuda/dense_bit_line_maker.h
+++ b/src/rpucuda/dense_bit_line_maker.h
@@ -58,8 +58,8 @@ template <typename T> class DenseBitLineMaker {
   bool supports(RPU::PulseType pulse_type) const;
 
   /* Ignore the buffer / counts, as they will be generated anew each sample.*/
-  void dumpExtra(RPU::state_t &extra, const std::string prefix){};
-  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict){};
+  void dumpExtra(RPU::state_t &extra, const std::string prefix) {};
+  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict) {};
 
 private:
   void freeContainers();
diff --git a/src/rpucuda/math_util.cpp b/src/rpucuda/math_util.cpp
index 2a6f8b23..bdd677e8 100644
--- a/src/rpucuda/math_util.cpp
+++ b/src/rpucuda/math_util.cpp
@@ -73,14 +73,14 @@ void gemm<half_t>(
     const int ldb,
     const half_t beta,
     half_t *C,
-    const int ldc){
-    // TODO: DOES HGEMM JUST NOT work for some reasons? MKL FP16 different from half_t ?
-    // RPU_INFO("A: " << (float)A[0] << ", B[0] " << B[0] << ", C[0] " << C[0]);
-    // cblas_hgemm(
-    //     Order, TransA, TransB, M, N, K, alpha, (const unsigned short *) A, lda,
-    //     (const unsigned short *) B, ldb, beta, (unsigned short *) C, ldc);
-
-    // just use sgemm for now (quite slow to copy)
+    const int ldc) {
+  // TODO: DOES HGEMM JUST NOT work for some reasons? MKL FP16 different from half_t ?
+  // RPU_INFO("A: " << (float)A[0] << ", B[0] " << B[0] << ", C[0] " << C[0]);
+  // cblas_hgemm(
+  //     Order, TransA, TransB, M, N, K, alpha, (const unsigned short *) A, lda,
+  //     (const unsigned short *) B, ldb, beta, (unsigned short *) C, ldc);
+
+  // just use sgemm for now (quite slow to copy)
 };
 #endif
 
diff --git a/src/rpucuda/rpu.h b/src/rpucuda/rpu.h
index 787059a8..b2212118 100644
--- a/src/rpucuda/rpu.h
+++ b/src/rpucuda/rpu.h
@@ -138,9 +138,9 @@ template <typename T> class RPUAbstract {
 
   T getLearningRate() const { return learning_rate_; };
 
-  virtual void finishUpdateCalculations(){};
-  virtual void finishAllCalculations(){};
-  virtual void makeUpdateAsync(){};
+  virtual void finishUpdateCalculations() {};
+  virtual void finishAllCalculations() {};
+  virtual void makeUpdateAsync() {};
 
 protected:
   int x_size_ = 0;
@@ -324,8 +324,8 @@ template <typename T> class RPUSimple : public RPUAbstract<T> {
      which usually are drawn during instantiation of the RPU object
      based on parameters defining their probabilty distributions. */
   virtual void getDeviceParameterNames(std::vector<std::string> &names) const { names.clear(); };
-  virtual void getDeviceParameter(std::vector<T *> &data_ptrs){};
-  virtual void setDeviceParameter(const std::vector<T *> &data_ptrs){};
+  virtual void getDeviceParameter(std::vector<T *> &data_ptrs) {};
+  virtual void setDeviceParameter(const std::vector<T *> &data_ptrs) {};
 
   /* These dumps extra state vectors that are not returned by
      getDeviuceParameters or getWeights*/
@@ -333,7 +333,7 @@ template <typename T> class RPUSimple : public RPUAbstract<T> {
   virtual void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict);
 
   virtual int getHiddenUpdateIdx() const { return 0; };
-  virtual void setHiddenUpdateIdx(int idx){};
+  virtual void setHiddenUpdateIdx(int idx) {};
 
   /* Decaying the weights once. Alpha can be a factor additionally
      applied to the current decay rate*/
@@ -416,7 +416,7 @@ template <typename T> class RPUSimple : public RPUAbstract<T> {
   virtual void setDeltaWeights(T *dw_extern);
   virtual T *getDeltaWeights() const { return delta_weights_extern_[0]; };
 
-  virtual void setVerbosityLevel(int verbose){};
+  virtual void setVerbosityLevel(int verbose) {};
 
   /* public interfaces for forward/backward/update. Format is
      expected in x-major order. However, the batch dimension comes
@@ -553,7 +553,7 @@ template <typename T> class RPUSimple : public RPUAbstract<T> {
   virtual T *copyToMatrixBiasBuffer(const T *X_input_without_bias, int m_batch, bool x_trans);
   virtual void
   copyFromMatrixBiasBuffer(T *X_input_without_bias, int m_batch, bool x_trans, T *bias_buffer);
-  virtual void releaseMatrixBiasBuffer(){};
+  virtual void releaseMatrixBiasBuffer() {};
   virtual T *getMatrixBiasBuffer(int m_batch);
   void forwardMatrixBias(
       const T *X_input_without_bias,
diff --git a/src/rpucuda/rpu_dynamic_transfer_device.cpp b/src/rpucuda/rpu_dynamic_transfer_device.cpp
index e03c9c57..0fc51f86 100644
--- a/src/rpucuda/rpu_dynamic_transfer_device.cpp
+++ b/src/rpucuda/rpu_dynamic_transfer_device.cpp
@@ -34,9 +34,8 @@ void DynamicTransferRPUDeviceMetaParameter<T>::printToStream(std::stringstream &
      << std::endl;
 
   if (experimental_fast_lr_feedback) {
-    ss << "\t Fast LR Feedback"
-       << ", target: " << experimental_feedback_target << ", mod: " << experimental_feedback_mod
-       << "]" << std::endl;
+    ss << "\t Fast LR Feedback" << ", target: " << experimental_feedback_target
+       << ", mod: " << experimental_feedback_mod << "]" << std::endl;
   }
 
   ChoppedTransferRPUDeviceMetaParameter<T>::printToStream(ss);
diff --git a/src/rpucuda/rpu_onesided_device.h b/src/rpucuda/rpu_onesided_device.h
index cb7d7ec5..d7bf9b6c 100644
--- a/src/rpucuda/rpu_onesided_device.h
+++ b/src/rpucuda/rpu_onesided_device.h
@@ -109,7 +109,7 @@ template <typename T> class OneSidedRPUDevice : public VectorRPUDevice<T> {
   void finishUpdateCycle(
       T **weights, const PulsedUpdateMetaParameter<T> &up, T current_lr, int m_batch_info) override;
 
-  void setHiddenUpdateIdx(int idx) override{};
+  void setHiddenUpdateIdx(int idx) override {};
 
   void doSparseUpdate(
       T **weights, int i, const int *x_signed_indices, int x_count, int d_sign, RNG<T> *rng)
diff --git a/src/rpucuda/rpu_pulsed_device.h b/src/rpucuda/rpu_pulsed_device.h
index 97f83127..09fddb72 100644
--- a/src/rpucuda/rpu_pulsed_device.h
+++ b/src/rpucuda/rpu_pulsed_device.h
@@ -167,10 +167,10 @@ template <typename T> class PulsedRPUDeviceBase : public SimpleRPUDevice<T> {
       const T *x_input = nullptr,
       const int x_inc = 1,
       const T *d_input = nullptr,
-      const int d_inc = 1){};
+      const int d_inc = 1) {};
   // called when update completed
   virtual void finishUpdateCycle(
-      T **weights, const PulsedUpdateMetaParameter<T> &up, T current_lr, int m_batch_info){};
+      T **weights, const PulsedUpdateMetaParameter<T> &up, T current_lr, int m_batch_info) {};
 
   inline T getWeightGranularity() const { return weight_granularity_; };
   inline T getNumStates() const { return num_states_; };
@@ -185,7 +185,7 @@ template <typename T> class PulsedRPUDeviceBase : public SimpleRPUDevice<T> {
      initUpdateCycle. Can be used to do some additional computation on
      the input */
   virtual void
-  initWithUpdateInput(const T *x_input, const int x_inc, const T *d_input, const int d_inc){};
+  initWithUpdateInput(const T *x_input, const int x_inc, const T *d_input, const int d_inc) {};
 
   void dumpExtra(RPU::state_t &extra, const std::string prefix) override {
     SimpleRPUDevice<T>::dumpExtra(extra, prefix);
@@ -440,7 +440,7 @@ public:
       return new DEVICENAME##RPUDeviceMetaParameter<T>(*this);                                     \
     };                                                                                             \
                                                                                                    \
-    T calcWeightGranularity() const override{GRANULARITY_BODY};                                    \
+    T calcWeightGranularity() const override { GRANULARITY_BODY };                                 \
     T calcNumStates() const override {                                                             \
       return (this->w_max - this->w_min) / calcWeightGranularity();                                \
     };                                                                                             \
diff --git a/src/rpucuda/rpu_simple_device.h b/src/rpucuda/rpu_simple_device.h
index 23e6ba59..57faf637 100644
--- a/src/rpucuda/rpu_simple_device.h
+++ b/src/rpucuda/rpu_simple_device.h
@@ -136,7 +136,7 @@ template <typename T> class AbstractRPUDevice {
   virtual int getHiddenWeightsCount() const = 0;
   virtual void setHiddenWeights(const std::vector<T> &data) = 0;
   virtual int getHiddenUpdateIdx() const { return 0; };
-  virtual void setHiddenUpdateIdx(int idx){};
+  virtual void setHiddenUpdateIdx(int idx) {};
   virtual void dumpExtra(RPU::state_t &extra, const std::string prefix) = 0;
   virtual void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict) = 0;
 
@@ -220,11 +220,11 @@ template <typename T> class SimpleRPUDevice : public AbstractRPUDevice<T> {
   SimpleRPUDevice<T> *clone() const override { return new SimpleRPUDevice<T>(*this); }
 
   void getDPNames(std::vector<std::string> &names) const override { names.clear(); };
-  void getDeviceParameter(T **weights, std::vector<T *> &data_ptrs) override{};
-  void setDeviceParameter(T **out_weights, const std::vector<T *> &data_ptrs) override{};
+  void getDeviceParameter(T **weights, std::vector<T *> &data_ptrs) override {};
+  void setDeviceParameter(T **out_weights, const std::vector<T *> &data_ptrs) override {};
   int getHiddenWeightsCount() const override { return 0; };
-  void setHiddenWeights(const std::vector<T> &data) override{};
-  void printDP(int x_count, int d_count) const override{};
+  void setHiddenWeights(const std::vector<T> &data) override {};
+  void printDP(int x_count, int d_count) const override {};
   void printToStream(std::stringstream &ss) const override { this->getPar().printToStream(ss); };
   void disp(std::stringstream &ss) const override {
     ss << "Device " << this->getPar().getName() << " [" << this->x_size_ << "," << this->d_size_
diff --git a/src/rpucuda/rpu_transfer_device.h b/src/rpucuda/rpu_transfer_device.h
index 93940ff8..8b1a8dd9 100644
--- a/src/rpucuda/rpu_transfer_device.h
+++ b/src/rpucuda/rpu_transfer_device.h
@@ -89,7 +89,7 @@ template <typename T> struct TransferRPUDeviceMetaParameter : VectorRPUDeviceMet
       int n_total_devices);
 
   virtual void initializeWithSize(int x_size, int d_size);
-  void initialize() override{/* do nothing */};
+  void initialize() override { /* do nothing */ };
 
   inline bool fullyHidden() const { return (!gamma && this->gamma_vec.back() == (T)1.0); };
 
@@ -181,7 +181,7 @@ template <typename T> class TransferRPUDevice : public VectorRPUDevice<T> {
 
   void getDeviceParameter(T **weights, std::vector<T *> &data_ptrs) override;
   void setDeviceParameter(T **out_weights, const std::vector<T *> &data_ptrs) override;
-  void setHiddenUpdateIdx(int idx) override{};
+  void setHiddenUpdateIdx(int idx) override {};
 
   void finishUpdateCycle(
       T **weights, const PulsedUpdateMetaParameter<T> &up, T current_lr, int m_batch_info) override;
diff --git a/src/rpucuda/rpu_weight_updater.h b/src/rpucuda/rpu_weight_updater.h
index 0f26148c..d57692f5 100644
--- a/src/rpucuda/rpu_weight_updater.h
+++ b/src/rpucuda/rpu_weight_updater.h
@@ -41,8 +41,8 @@ template <typename T> class RPUWeightUpdater {
       const int d_inc,
       const T learning_rate);
 
-  virtual void dumpExtra(RPU::state_t &extra, const std::string prefix){};
-  virtual void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict){};
+  virtual void dumpExtra(RPU::state_t &extra, const std::string prefix) {};
+  virtual void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict) {};
 
 protected:
   int x_size_ = 0;
diff --git a/src/rpucuda/sparse_bit_line_maker.h b/src/rpucuda/sparse_bit_line_maker.h
index b706a977..2d5290a4 100644
--- a/src/rpucuda/sparse_bit_line_maker.h
+++ b/src/rpucuda/sparse_bit_line_maker.h
@@ -71,8 +71,8 @@ template <typename T> class SparseBitLineMaker {
   bool supports(RPU::PulseType pulse_type) const;
 
   /* Ignore the buffer / indices, as they will be generated anew each sample.*/
-  void dumpExtra(RPU::state_t &extra, const std::string prefix){};
-  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict){};
+  void dumpExtra(RPU::state_t &extra, const std::string prefix) {};
+  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict) {};
 
 private:
   void freeContainers();
diff --git a/src/rpucuda/weight_clipper.h b/src/rpucuda/weight_clipper.h
index 8aa72a5b..a0507cf8 100644
--- a/src/rpucuda/weight_clipper.h
+++ b/src/rpucuda/weight_clipper.h
@@ -76,8 +76,8 @@ template <typename T> class WeightClipper {
   /* in-place clipping of weights */
   void apply(T *weights, const WeightClipParameter &wclpar);
 
-  void dumpExtra(RPU::state_t &extra, const std::string prefix){};
-  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict){};
+  void dumpExtra(RPU::state_t &extra, const std::string prefix) {};
+  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict) {};
 
 private:
   void clip(T *weights, T clip_value);
diff --git a/src/rpucuda/weight_modifier.cpp b/src/rpucuda/weight_modifier.cpp
index fcc453f0..380156e7 100644
--- a/src/rpucuda/weight_modifier.cpp
+++ b/src/rpucuda/weight_modifier.cpp
@@ -200,6 +200,11 @@ void WeightModifier<T>::apply(
     break;
   }
 
+  case WeightModifierType::DropConnect: {
+    // will be done below
+    break;
+  }
+
   default:
     RPU_FATAL("Requested WeightModifierType not implemented in CPU version.");
   }
diff --git a/src/rpucuda/weight_remapper.h b/src/rpucuda/weight_remapper.h
index 3b73fd60..321d2009 100644
--- a/src/rpucuda/weight_remapper.h
+++ b/src/rpucuda/weight_remapper.h
@@ -118,8 +118,8 @@ template <typename T> class WeightRemapper {
       T *scales = nullptr,
       T *biases = nullptr);
 
-  void dumpExtra(RPU::state_t &extra, const std::string prefix){};
-  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict){};
+  void dumpExtra(RPU::state_t &extra, const std::string prefix) {};
+  void loadExtra(const RPU::state_t &extra, const std::string prefix, bool strict) {};
 
 private:
   std::vector<T> max_values_;


Test patch:
diff --git a/src/rpucuda/cuda/rpucuda_mixedprec_device_test.cpp b/src/rpucuda/cuda/rpucuda_mixedprec_device_test.cpp
index d84bb22b..1d78da21 100644
--- a/src/rpucuda/cuda/rpucuda_mixedprec_device_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_mixedprec_device_test.cpp
@@ -44,7 +44,7 @@ class MixedPrecRPUDeviceCudaTestFixtureSmall : public ::testing::Test {
     dp.setDevicePar(dp_cs);
   };
 
-  void TearDown(){};
+  void TearDown() {};
 
   int x_size, d_size;
   MixedPrecRPUDeviceMetaParameter<num_t> dp;
diff --git a/src/rpucuda/cuda/rpucuda_mixedprec_int_device_test.cpp b/src/rpucuda/cuda/rpucuda_mixedprec_int_device_test.cpp
index ca6d2d3c..41f07361 100644
--- a/src/rpucuda/cuda/rpucuda_mixedprec_int_device_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_mixedprec_int_device_test.cpp
@@ -34,7 +34,7 @@ class MixedPrecIntRPUDeviceCudaTestFixtureSmall : public ::testing::Test {
     dp.setDevicePar(dp_cs);
   };
 
-  void TearDown(){};
+  void TearDown() {};
 
   int x_size, d_size;
   MixedPrecIntRPUDeviceMetaParameter<num_t> dp;
diff --git a/src/rpucuda/cuda/rpucuda_pulsed_device_test.cpp b/src/rpucuda/cuda/rpucuda_pulsed_device_test.cpp
index 77496c78..a66b2202 100644
--- a/src/rpucuda/cuda/rpucuda_pulsed_device_test.cpp
+++ b/src/rpucuda/cuda/rpucuda_pulsed_device_test.cpp
@@ -37,7 +37,7 @@ namespace {
 
 using namespace RPU;
 
-template <typename DeviceParT> void specific_settings(DeviceParT &par){};
+template <typename DeviceParT> void specific_settings(DeviceParT &par) {};
 
 template <> void specific_settings(PiecewiseStepRPUDeviceMetaParameter<num_t> &par) {
   par.piecewise_up_vec = std::vector<num_t>{0.1, 0.5, 1.0, 0.3, 0.1};