From cc99347560ab8ff659c3e22cd1efe1c19df94d5c Mon Sep 17 00:00:00 2001
From: Francesco Gatti <gattifrancesco@hotmail.it>
Date: Thu, 29 Jun 2017 10:06:49 +0000
Subject: [PATCH] MulAdd implemented

---
 CMakeLists.txt  |  9 ++++-----
 include/Layer.h | 18 ++++++++++++++++++
 include/utils.h |  3 +++
 src/MulAdd.cpp  | 45 +++++++++++++++++++++++++++++++++++++++++++++
 src/utils.cpp   | 10 ++++++++++
 tests/test.cpp  |  2 ++
 6 files changed, 82 insertions(+), 5 deletions(-)
 create mode 100644 src/MulAdd.cpp

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 4722ded..c79fc8c 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -8,12 +8,11 @@ cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS
 cuda_add_library(kernels SHARED src/kernels/activation_elu.cu)
 
 include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS})
-add_library(tkDNN SHARED src/Layer.cpp src/LayerWgs.cpp 
-    src/Dense.cpp src/Activation.cpp src/Conv2d.cpp src/Conv3d.cpp src/Flatten.cpp
+add_library(tkDNN SHARED src/Layer.cpp src/LayerWgs.cpp
+    src/Dense.cpp src/Activation.cpp src/Conv2d.cpp src/Conv3d.cpp src/Flatten.cpp src/MulAdd.cpp
     src/Network.cpp src/utils.cpp)
-target_link_libraries(tkDNN kernels)
+target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDA_TOOLKIT_ROOT_DIR}/lib/libcudnn.so)
 
 add_executable(tkDNNtest tests/test.cpp)
 message(${CUDA_LIBRARIES})
-target_link_libraries(tkDNNtest tkDNN 
-    ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDA_TOOLKIT_ROOT_DIR}/lib/libcudnn.so)
\ No newline at end of file
+target_link_libraries(tkDNNtest tkDNN)
\ No newline at end of file
diff --git a/include/Layer.h b/include/Layer.h
index d015540..927154e 100644
--- a/include/Layer.h
+++ b/include/Layer.h
@@ -191,5 +191,23 @@ protected:
     value_type *dstData;  //where results will be putted
 };
 
+
+/**
+    MulAdd layer
+    apply a multiplication and then an addition for each data
+*/
+class MulAdd : public Layer {
+
+public:
+    MulAdd(Network *net, dataDim_t input_dim, value_type mul, value_type add); 
+    virtual ~MulAdd();
+
+    value_type* infer(dataDim_t &dim, value_type* srcData);
+
+protected:
+    value_type mul, add;
+    value_type *dstData, *add_vector;  //where results will be putted
+};
+
 }
 #endif //LAYER_H
\ No newline at end of file
diff --git a/include/utils.h b/include/utils.h
index f1108e5..ed50073 100644
--- a/include/utils.h
+++ b/include/utils.h
@@ -65,6 +65,9 @@
 void readBinaryFile(const char* fname, int size, value_type** data_h, value_type** data_d);
 void printDeviceVector(int size, value_type* vec_d);
 void resize(int size, value_type **data);
+
 void matrixTranspose(cublasHandle_t handle, value_type* srcData, value_type* dstData, int rows, int cols);
 
+void matrixMulAdd(  cublasHandle_t handle, value_type* srcData, value_type* dstData, 
+                    value_type* add_vector, int dim, value_type mul);
 #endif //UTILS_H
\ No newline at end of file
diff --git a/src/MulAdd.cpp b/src/MulAdd.cpp
new file mode 100644
index 0000000..dd7cfd2
--- /dev/null
+++ b/src/MulAdd.cpp
@@ -0,0 +1,45 @@
+#include <iostream>
+
+#include "Layer.h"
+#include "kernels.h"
+
+namespace tkDNN {
+
+MulAdd::MulAdd(Network *net, dataDim_t input_dim, value_type mul, value_type add) : 
+    Layer(net, input_dim) {
+
+    this->mul = mul;
+    this->add = add;
+
+    int size = input_dim.tot();
+
+    // create a vector with all value setted to add 
+    value_type *add_vector_h = new value_type[size];
+    for(int i=0; i<size; i++)
+        add_vector_h[i] = add;
+
+    checkCuda( cudaMalloc(&add_vector, size*sizeof(value_type)));
+    checkCuda( cudaMemcpy(add_vector, add_vector_h, size*sizeof(value_type), cudaMemcpyHostToDevice));
+    delete [] add_vector_h;
+
+
+    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) );
+}
+
+MulAdd::~MulAdd() {
+
+    checkCuda( cudaFree(add_vector) );
+    checkCuda( cudaFree(dstData) );
+}
+
+value_type* MulAdd::infer(dataDim_t &dim, value_type* srcData) {
+
+    matrixMulAdd(net->cublasHandle, srcData, dstData, add_vector, input_dim.tot(), mul);
+    
+    //update data dimensions    
+    dim = output_dim;
+
+    return dstData;
+}
+
+}
\ No newline at end of file
diff --git a/src/utils.cpp b/src/utils.cpp
index 66c2e97..9414395 100644
--- a/src/utils.cpp
+++ b/src/utils.cpp
@@ -53,4 +53,14 @@ void matrixTranspose(cublasHandle_t handle, value_type* srcData, value_type* dst
     float const alpha(1.0);
     float const beta(0.0);
     checkERROR( cublasSgeam( handle, CUBLAS_OP_T, CUBLAS_OP_N, m, n, &alpha, A, n, &beta, A, m, clone, m ));
+}
+
+void matrixMulAdd(  cublasHandle_t handle, value_type* srcData, value_type* dstData, 
+                    value_type* add_vector, int dim, value_type mul) {
+
+    checkCuda( cudaMemcpy(dstData, add_vector, dim*sizeof(value_type), cudaMemcpyDeviceToDevice));
+        
+    value_type alpha = mul;
+    checkERROR( cublasSaxpy(handle, dim, &alpha, srcData, 1, dstData, 1));
+
 }
\ No newline at end of file
diff --git a/tests/test.cpp b/tests/test.cpp
index 3f2c727..9b54887 100644
--- a/tests/test.cpp
+++ b/tests/test.cpp
@@ -19,6 +19,7 @@ int main() {
     tkDNN::Conv3d     c1 (&net, a0.output_dim, 4, 2, 2, 2, 1, 1, 1, c1_bin, c1_bias_bin);
     tkDNN::Activation a1 (&net, c1.output_dim, tkDNN::ACTIVATION_ELU);
     tkDNN::Flatten    f1 (&net, a1.output_dim);
+    tkDNN::MulAdd     m1 (&net, f1.output_dim, 2, 1);
 
     // Load input
     value_type *data;
@@ -35,6 +36,7 @@ int main() {
     data = c1.infer(dim, data); dim.print();
     data = a1.infer(dim, data); dim.print();
     data = f1.infer(dim, data); dim.print();
+    data = m1.infer(dim, data); dim.print();
 
     TIMER_STOP
     // Print result