diff --git a/CMakeLists.txt b/CMakeLists.txt
index d843668..deb6371 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -9,9 +9,12 @@ 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/Flatten.cpp src/MulAdd.cpp src/Pooling.cpp
+    src/Dense.cpp src/Activation.cpp src/Conv2d.cpp src/Flatten.cpp src/MulAdd.cpp src/Pooling.cpp src/Softmax.cpp
     src/Network.cpp src/utils.cpp)
 target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn)
 
 add_executable(tkDNNtest tests/test.cpp)
 target_link_libraries(tkDNNtest tkDNN)
+
+add_executable(mnist tests/mnist/test.cpp)
+target_link_libraries(mnist tkDNN)
diff --git a/include/Layer.h b/include/Layer.h
index 3ee566a..98aead8 100644
--- a/include/Layer.h
+++ b/include/Layer.h
@@ -207,5 +207,20 @@ protected:
     bool poolOn3d;
 };
 
+/**
+    Softmax layer
+*/
+class Softmax : public Layer {
+
+public:
+    Softmax(Network *net, dataDim_t input_dim); 
+    virtual ~Softmax();
+
+    virtual value_type* infer(dataDim_t &dim, value_type* srcData);
+
+protected:
+    value_type *dstData;  //where results will be putted
+};
+
 }
 #endif //LAYER_H
diff --git a/src/Activation.cpp b/src/Activation.cpp
index caf8aff..23007ae 100644
--- a/src/Activation.cpp
+++ b/src/Activation.cpp
@@ -35,6 +35,8 @@ Activation::Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t
 Activation::~Activation() {
 
     checkCuda( cudaFree(dstData) );
+
+    checkCUDNN( cudnnDestroyActivationDescriptor(activDesc) );
 }
 
 value_type* Activation::infer(dataDim_t &dim, value_type* srcData) {
diff --git a/src/Softmax.cpp b/src/Softmax.cpp
new file mode 100644
index 0000000..bd59ea7
--- /dev/null
+++ b/src/Softmax.cpp
@@ -0,0 +1,48 @@
+#include <iostream>
+
+#include "Layer.h"
+#include "kernels.h"
+
+namespace tkDNN {
+
+Softmax::Softmax(Network *net, dataDim_t input_dim) : 
+    Layer(net, input_dim) {
+
+    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) );
+
+    checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc,
+                                        net->tensorFormat,
+                                        net->dataType,
+                                        input_dim.n*input_dim.l, 
+                                        input_dim.c,
+                                        input_dim.h, input_dim.w) );
+    checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
+                                            net->tensorFormat,
+                                            net->dataType,
+                                            input_dim.n*input_dim.l, 
+                                            input_dim.c,
+                                            input_dim.h, input_dim.w) );
+}
+
+Softmax::~Softmax() {
+
+    checkCuda( cudaFree(dstData) );
+}
+
+value_type* Softmax::infer(dataDim_t &dim, value_type* srcData) {
+
+    value_type alpha = value_type(1);
+    value_type beta  = value_type(0);
+    checkCUDNN( cudnnSoftmaxForward(net->cudnnHandle,
+                                    CUDNN_SOFTMAX_ACCURATE ,
+                                    CUDNN_SOFTMAX_MODE_CHANNEL,
+                                    &alpha,
+                                    srcTensorDesc,
+                                    srcData,
+                                    &beta,
+                                    dstTensorDesc,
+                                    dstData) );
+    return dstData;
+}
+
+}