MulAdd implemented

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)
+target_link_libraries(tkDNNtest tkDNN)

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

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

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;
+}
+
+}

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));
+
 }

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