quantization - bugfixes

6 files changed, 71 insertions, 20 deletions

diff --git a/include/caffe/quantized_layer.hpp b/include/caffe/quantized_layer.hpp
index 9ce4c00d..18873e87 100644
--- a/include/caffe/quantized_layer.hpp
+++ b/include/caffe/quantized_layer.hpp
@@ -38,9 +38,9 @@ public:
    * @param fl The number of bits in the fractional part.
    */
   void Trim2FixedPoint_cpu(Ftype* data, const int cnt, bool power2_range, const int bitwidth,
-      const int rounding, int fracbits, float scale, float offset, bool unsigned_quant, bool clip);
+      const int rounding, int fracbits, float scale, float offset, bool unsigned_quant, bool clip, bool roundup);
   void Trim2FixedPoint_gpu(Ftype* data, const int cnt, bool power2_range, const int bitwidth,
-      const int rounding, int fracbits, float scale, float offset, bool unsigned_quant, bool clip);
+      const int rounding, int fracbits, float scale, float offset, bool unsigned_quant, bool clip, bool roundup);
           
   /**
    * @brief Generate random number in [0,1) range.
diff --git a/src/caffe/layers/conv_layer.cu b/src/caffe/layers/conv_layer.cu
index 50ef113e..791e5d97 100644
--- a/src/caffe/layers/conv_layer.cu
+++ b/src/caffe/layers/conv_layer.cu
@@ -7,7 +7,7 @@ namespace caffe {
 template <typename Ftype, typename Btype>
 void ConvolutionLayer<Ftype, Btype>::Forward_gpu(const vector<Blob*>& bottom,
       const vector<Blob*>& top) {
-  //this->Quantize_gpu(bottom, top);
+  this->Quantize_gpu(bottom, top);
   const Ftype* weight = this->blobs_[0]->template gpu_data<Ftype>();
   for (int i = 0; i < bottom.size(); ++i) {
     const Ftype* bottom_data = bottom[i]->gpu_data<Ftype>();
diff --git a/src/caffe/layers/cudnn_conv_layer.cu b/src/caffe/layers/cudnn_conv_layer.cu
index f42c210d..0c907b30 100644
--- a/src/caffe/layers/cudnn_conv_layer.cu
+++ b/src/caffe/layers/cudnn_conv_layer.cu
@@ -11,7 +11,7 @@ namespace caffe {
 template<typename Ftype, typename Btype>
 void CuDNNConvolutionLayer<Ftype, Btype>::Forward_gpu(const vector<Blob*>& bottom,
     const vector<Blob*>& top) {
-  //this->Quantize_gpu(bottom, top);
+  this->Quantize_gpu(bottom, top);
   const Ftype* weight = this->blobs_[0]->template gpu_data<Ftype>();
   shared_ptr<GPUMemory::Workspace>& ws = GPUMemory::workspace_[Caffe::current_device()];
   if (use_v7grouping()) {
diff --git a/src/caffe/net.cpp b/src/caffe/net.cpp
index a4bab339..4edba6db 100644
--- a/src/caffe/net.cpp
+++ b/src/caffe/net.cpp
@@ -2161,7 +2161,10 @@ void Net::EnableQuantizationForSelectedLayers() {
       std::transform(layer_type_lower.begin(), layer_type_lower.end(), layer_type_lower.begin(),
               [](unsigned char c) {return std::tolower(c);}
       );
-      if(layer_type_lower.find("data") != string::npos) {
+      if(layer_type_lower.find("Data") != string::npos) {
+          max_blobs_to_quantize = 1;
+      }
+      if(layer_type_lower.find("Input") != string::npos) {
           max_blobs_to_quantize = 1;
       }
 
@@ -2172,13 +2175,16 @@ void Net::EnableQuantizationForSelectedLayers() {
               layer_type == "Bias" || layer_type == "Pooling") {
           is_quantized_layer_type = true;
       }
-      if(layer_type_lower.find("data") != string::npos) {
+      if(layer_type_lower.find("Data") != string::npos) {
+          is_quantized_layer_type = true;
+      }
+      if(layer_type_lower.find("Input") != string::npos) {
           is_quantized_layer_type = true;
       }
 
       //quantize weights
       if(net_qparam.quantize_weights()) {
-          if(is_quantized_layer_type && (!is_merged_layer) && (!is_ignored_layer_name)) {
+          if(is_quantized_layer_type /*&& (!is_merged_layer)*/ && (!is_ignored_layer_name)) {
               if(layer_type == "Convolution" || layer_type == "InnerProduct" || layer_type == "Deconvolution") {
                   QuantizationParameter& quantization_param = *layers_[layer_id]->mutable_layer_param().mutable_quantization_param();
                   for(int blob_id=0; blob_id<layers_[layer_id]->blobs().size(); blob_id++) {
diff --git a/src/caffe/quantized_layer.cpp b/src/caffe/quantized_layer.cpp
index e686d76c..8407b43e 100644
--- a/src/caffe/quantized_layer.cpp
+++ b/src/caffe/quantized_layer.cpp
@@ -51,7 +51,7 @@ void QuantizedLayer<Ftype, Btype>::QuantizeWeights_cpu(Ftype* data, const int bl
   case QuantizationParameter_Precision_DYNAMIC_FIXED_POINT:
     Trim2FixedPoint_cpu(data, count, param.power2_scale_weights(), qparam_w.bitwidth(),
         param.rounding_scheme(), qparam_w.fracbits(), qparam_w.scale_target(),
-        qparam_w.offset(), qparam_w.unsigned_quant(), clip);
+        qparam_w.offset(), qparam_w.unsigned_quant(), clip, false);
     break;
   case QuantizationParameter_Precision_FLOAT:
         break;
@@ -72,7 +72,7 @@ void QuantizedLayer<Ftype, Btype>::QuantizeLayerInputs_cpu(Ftype* data, const in
     case QuantizationParameter_Precision_DYNAMIC_FIXED_POINT:
       Trim2FixedPoint_cpu(data, count, param.power2_scale_activations(), qparam_in.bitwidth(),
           param.rounding_scheme(), qparam_in.fracbits(), qparam_in.scale_target(),
-          qparam_in.offset(), qparam_in.unsigned_quant(), true);
+          qparam_in.offset(), qparam_in.unsigned_quant(), true, true);
       break;
     case QuantizationParameter_Precision_FLOAT:
           break;
@@ -91,7 +91,7 @@ void QuantizedLayer<Ftype, Btype>::QuantizeLayerOutputs_cpu(
     case QuantizationParameter_Precision_DYNAMIC_FIXED_POINT:
       Trim2FixedPoint_cpu(data, count, param.power2_scale_activations(), qparam_out.bitwidth(),
           param.rounding_scheme(), qparam_out.fracbits(), qparam_out.scale_target(),
-          qparam_out.offset(), qparam_out.unsigned_quant(), true);
+          qparam_out.offset(), qparam_out.unsigned_quant(), true, true);
       break;
     case QuantizationParameter_Precision_FLOAT:
           break;
@@ -103,7 +103,7 @@ void QuantizedLayer<Ftype, Btype>::QuantizeLayerOutputs_cpu(
 
 template<typename Ftype, typename Btype>
 void QuantizedLayer<Ftype, Btype>::Trim2FixedPoint_cpu(Ftype* data, const int cnt, bool power2_range, const int bitwidth,
-    const int rounding, int fracbits, float scale, float offset, bool unsigned_quant, bool clip) {
+    const int rounding, int fracbits, float scale, float offset, bool unsigned_quant, bool clip, bool roundup) {
   float inv_scale = 1.0f/scale;
 
   int qrange = unsigned_quant? bitwidth :  (bitwidth - 1);
@@ -121,7 +121,13 @@ void QuantizedLayer<Ftype, Btype>::Trim2FixedPoint_cpu(Ftype* data, const int cn
     // Round data
     switch (rounding) {
     case QuantizationParameter_Rounding_NEAREST:
-      data[index] = round(data[index]);
+        //data[index] = round(data[index]);
+        if(roundup) {
+            //data[index] = int(data[index]+0.5);
+            data[index] = (int(data[index] * 4096) + 2048)>>12;
+        } else {
+            data[index] = int(data[index] >= 0? (data[index]+0.5) : (data[index]-0.5));
+        }
       break;
     case QuantizationParameter_Rounding_STOCHASTIC:
       data[index] = floor(data[index] + RandUniform_cpu());
diff --git a/src/caffe/quantized_layer.cu b/src/caffe/quantized_layer.cu
index 3a9f3f70..69c2e974 100644
--- a/src/caffe/quantized_layer.cu
+++ b/src/caffe/quantized_layer.cu
@@ -46,7 +46,7 @@ void QuantizedLayer<Ftype, Btype>::QuantizeWeights_gpu(Ftype* data, const int bl
   case QuantizationParameter_Precision_DYNAMIC_FIXED_POINT:
     Trim2FixedPoint_gpu(data, count, param.power2_scale_weights(), qparam_w.bitwidth(),
         param.rounding_scheme(), qparam_w.fracbits(), qparam_w.scale_target(),
-        qparam_w.offset(), qparam_w.unsigned_quant(), clip);
+        qparam_w.offset(), qparam_w.unsigned_quant(), clip, false);
     break;
   case QuantizationParameter_Precision_FLOAT:
           break;
@@ -66,7 +66,7 @@ void QuantizedLayer<Ftype, Btype>::QuantizeLayerInputs_gpu(
     case QuantizationParameter_Precision_DYNAMIC_FIXED_POINT:
       Trim2FixedPoint_gpu(data, count, param.power2_scale_activations(), qparam_in.bitwidth(),
           param.rounding_scheme(), qparam_in.fracbits(), qparam_in.scale_target(),
-          qparam_in.offset(), qparam_in.unsigned_quant(), true);
+          qparam_in.offset(), qparam_in.unsigned_quant(), true, true);
       break;
     case QuantizationParameter_Precision_FLOAT:
           break;
@@ -85,7 +85,7 @@ void QuantizedLayer<Ftype, Btype>::QuantizeLayerOutputs_gpu(Ftype* data,
     case QuantizationParameter_Precision_DYNAMIC_FIXED_POINT:
       Trim2FixedPoint_gpu(data, count, param.power2_scale_activations(), qparam_out.bitwidth(),
           param.rounding_scheme(), qparam_out.fracbits(), qparam_out.scale_target(),
-          qparam_out.offset(), qparam_out.unsigned_quant(), true);
+          qparam_out.offset(), qparam_out.unsigned_quant(), true, true);
       break;
     case QuantizationParameter_Precision_FLOAT:
           break;
@@ -96,7 +96,7 @@ void QuantizedLayer<Ftype, Btype>::QuantizeLayerOutputs_gpu(Ftype* data,
 }
 
 template <typename Dtype>
-__global__ void Trim2FixedPoint_kernel(Dtype* data, const int cnt,
+__global__ void Trim2FixedPoint_roundassymetric_kernel(Dtype* data, const int cnt,
     const int bitwidth, const int rounding, float scale, float inv_scale, float offset, float min_data, float max_data, bool clip) {
     CUDA_KERNEL_LOOP(index, cnt) {
 
@@ -105,7 +105,8 @@ __global__ void Trim2FixedPoint_kernel(Dtype* data, const int cnt,
     // Round data
     switch (rounding) {
     case QuantizationParameter_Rounding_NEAREST:
-      data[index] = rint(data[index]);
+      //data[index] = rint(data[index]);
+      data[index] = int(data[index] >= 0? (data[index]+0.5) : (data[index]-0.5));
       break;
     case QuantizationParameter_Rounding_STOCHASTIC:
       data[index] = __float2int_rd(data[index] + RandUniform_device(index));
@@ -124,20 +125,58 @@ __global__ void Trim2FixedPoint_kernel(Dtype* data, const int cnt,
   }
 }
 
+template <typename Dtype>
+__global__ void Trim2FixedPoint_roundup_kernel(Dtype* data, const int cnt,
+    const int bitwidth, const int rounding, float scale, float inv_scale, float offset, float min_data, float max_data, bool clip) {
+    CUDA_KERNEL_LOOP(index, cnt) {
+
+    data[index] = (data[index] * scale) + offset;
+
+    // Round data
+    switch (rounding) {
+    case QuantizationParameter_Rounding_NEAREST:
+      //data[index] = rint(data[index]);
+      //data[index] = (data[index]+0.5);
+      data[index] = (int(data[index] * 4096) + 2048)>>12;
+      break;
+    case QuantizationParameter_Rounding_STOCHASTIC:
+      //data[index] = __float2int_rd(data[index] + RandUniform_device(index));
+      break;
+    default:
+      break;
+    }
+
+    // Saturate data
+    if(clip) {
+      data[index] = (data[index]>(Dtype)max_data? (Dtype)max_data:
+        (data[index]<(Dtype)min_data?(Dtype)min_data:data[index]));
+    }
+
+    data[index] = (data[index] - offset) * inv_scale;
+  }
+}
+
 
 template<typename Ftype, typename Btype>
 void QuantizedLayer<Ftype, Btype>::Trim2FixedPoint_gpu(Ftype* data, const int cnt, bool power2_range,
-      const int bitwidth, const int rounding, int fracbits, float scale, float offset, bool unsigned_quant, bool clip) {
+      const int bitwidth, const int rounding, int fracbits, float scale, float offset, bool unsigned_quant, bool clip,
+      bool roundup) {
   float inv_scale = 1.0f/scale;
 
   int qrange = unsigned_quant? bitwidth :  (bitwidth - 1);
   float min_data = unsigned_quant? 0 : -(powf(2, qrange));
   float max_data = +(powf(2, qrange) - 1);
 
-  Trim2FixedPoint_kernel<<<CAFFE_GET_BLOCKS(cnt), CAFFE_CUDA_NUM_THREADS>>>(
-      data, cnt, bitwidth, rounding, scale, inv_scale, offset, min_data, max_data, clip);
+  if(roundup) {
+      Trim2FixedPoint_roundup_kernel<<<CAFFE_GET_BLOCKS(cnt), CAFFE_CUDA_NUM_THREADS>>>(
+          data, cnt, bitwidth, rounding, scale, inv_scale, offset, min_data, max_data, clip);
+  } else {
+      Trim2FixedPoint_roundassymetric_kernel<<<CAFFE_GET_BLOCKS(cnt), CAFFE_CUDA_NUM_THREADS>>>(
+          data, cnt, bitwidth, rounding, scale, inv_scale, offset, min_data, max_data, clip);
+  }
 }
 
+
 template void QuantizedLayer<double, double>::Quantize_gpu(const vector<Blob*>& bottom,const vector<Blob*>& top);
 template void QuantizedLayer<double, float>::Quantize_gpu(const vector<Blob*>& bottom,const vector<Blob*>& top);
 template void QuantizedLayer<double, float16>::Quantize_gpu(const vector<Blob*>& bottom,const vector<Blob*>& top);