Add backward functions for ReLU, Sigmoid and Tanh

include/aidge/backend/cuda/operator/ReLUImpl.hpp

@@ -36,6 +36,7 @@ private:
         cudnnActivationMode_t mReLUDesc = nullptr;
     #endif
     std::shared_ptr<Tensor> mInputFallback;
+	std::shared_ptr<Tensor> mOutputGradFallback;
 
 public:
     ReLUImpl_cuda(const ReLU_Op &op) : OperatorImpl(op, "cuda") {}
@@ -46,10 +47,12 @@ public:
 
 public:
     void forward();
+	void backward();
     ~ReLUImpl_cuda();
 
 private:
     template <class T> void forward_(const Tensor& input);
+	template <class T> void backward_(const Tensor& output_grad);
 };
 
 namespace {

include/aidge/backend/cuda/operator/SigmoidImpl.hpp

@@ -36,6 +36,7 @@ private:
         cudnnActivationMode_t mSigmoidDesc = nullptr;
     #endif
     std::shared_ptr<Tensor> mInputFallback;
+	std::shared_ptr<Tensor> mOutputGradFallback;	
 
 public:
     SigmoidImpl_cuda(const Sigmoid_Op &op) : OperatorImpl(op, "cuda") {}
@@ -46,10 +47,12 @@ public:
 
 public:
     void forward();
+	void backward();
     ~SigmoidImpl_cuda();
 
 private:
     template <class T> void forward_(const Tensor& input);
+	template <class T> void backward_(const Tensor& output_grad);	
 };
 
 namespace {

include/aidge/backend/cuda/operator/TanhImpl.hpp

@@ -36,6 +36,7 @@ private:
         cudnnActivationMode_t mTanhDesc = nullptr;
     #endif
     std::shared_ptr<Tensor> mInputFallback;
+	std::shared_ptr<Tensor> mOutputGradFallback;
 
 public:
     TanhImpl_cuda(const Tanh_Op &op) : OperatorImpl(op, "cuda") {}
@@ -46,10 +47,12 @@ public:
 
 public:
     void forward();
+	void backward();	
     ~TanhImpl_cuda();
 
 private:
     template <class T> void forward_(const Tensor& input);
+	template <class T> void backward_(const Tensor& output_grad);
 };
 
 namespace {

src/operator/ReLUImpl.cpp

@@ -64,6 +64,55 @@ void Aidge::ReLUImpl_cuda::forward_(const Tensor& input) {
                                std::static_pointer_cast<Tensor>(op.getRawOutput(0))->getImpl()->rawPtr()));
 }
 
+void Aidge::ReLUImpl_cuda::backward() {
+    const OperatorTensor& op = static_cast<const OperatorTensor&>(mOp);
+
+    assert(op.getOutput(0)->grad() && "missing output #0");
+
+    const auto& output_grad = op.getOutput(0)->grad()->refCastFrom(mOutputGradFallback, *op.getOutput(0)->grad());
+
+    // Lazy-initialize CuDNN ReLU descriptor
+    if (mReLUDesc == nullptr) {
+		#if CUDNN_VERSION >= 5000
+			CHECK_CUDNN_STATUS(cudnnCreateActivationDescriptor(&mReLUDesc));
+			CHECK_CUDNN_STATUS(cudnnSetActivationDescriptor(
+				mReLUDesc, CUDNN_ACTIVATION_RELU, CUDNN_NOT_PROPAGATE_NAN, 0.0));
+		#else
+			mReLUDesc = CUDNN_ACTIVATION_RELU;
+		#endif
+    }
+
+    // Do the actual backward computation
+    // Template is only for scaling parameters, which are always in float
+    // excepted when the convolution is performed in double precision.
+    if (op.getInput(0)->grad()->dataType() == DataType::Float64) {
+        backward_<double>(output_grad);
+    }
+    else {
+        backward_<float>(output_grad);
+    }
+}
+
+template <class T>
+void Aidge::ReLUImpl_cuda::backward_(const Tensor& output_grad) {
+    const OperatorTensor& op = static_cast<const OperatorTensor&>(mOp);
+    const typename Cuda::cudnn_scaling_type<T>::type alpha = 1.0f;
+    const typename Cuda::cudnn_scaling_type<T>::type beta = 0.0f;
+    CHECK_CUDNN_STATUS(
+        cudnnActivationBackward(CudaContext::cudnnHandle(),
+                               mReLUDesc,
+                               &alpha,
+							   std::dynamic_pointer_cast<TensorImpl_cuda_>(op.getOutput(0)->getImpl())->getCudnnTensorDesc(*op.getOutput(0)),
+                               std::static_pointer_cast<Tensor>(op.getRawOutput(0))->getImpl()->rawPtr(),
+							   std::dynamic_pointer_cast<TensorImpl_cuda_>(output_grad.getImpl())->getCudnnTensorDesc(output_grad),
+                               output_grad.getImpl()->rawPtr(),
+                               std::dynamic_pointer_cast<TensorImpl_cuda_>(op.getInput(0)->getImpl())->getCudnnTensorDesc(*op.getInput(0)),
+                               std::static_pointer_cast<Tensor>(op.getRawInput(0))->getImpl()->rawPtr(),							   
+                               &beta,
+							   std::dynamic_pointer_cast<TensorImpl_cuda_>(op.getInput(0)->grad()->getImpl())->getCudnnTensorDesc(*op.getInput(0)->grad()),
+                               op.getInput(0)->grad()->getImpl()->rawPtr()));
+}
+
 Aidge::ReLUImpl_cuda::~ReLUImpl_cuda() {
     if (mReLUDesc != nullptr) {
 		#if CUDNN_VERSION >= 5000

src/operator/SigmoidImpl.cpp

@@ -64,6 +64,55 @@ void Aidge::SigmoidImpl_cuda::forward_(const Tensor& input) {
                                std::static_pointer_cast<Tensor>(op.getRawOutput(0))->getImpl()->rawPtr()));
 }
 
+void Aidge::SigmoidImpl_cuda::backward() {
+    const OperatorTensor& op = static_cast<const OperatorTensor&>(mOp);
+
+    assert(op.getOutput(0)->grad() && "missing output #0");
+
+    const auto& output_grad = op.getOutput(0)->grad()->refCastFrom(mOutputGradFallback, *op.getOutput(0)->grad());
+
+    // Lazy-initialize CuDNN Sigmoid descriptor
+    if (mSigmoidDesc == nullptr) {
+		#if CUDNN_VERSION >= 5000
+			CHECK_CUDNN_STATUS(cudnnCreateActivationDescriptor(&mSigmoidDesc));
+			CHECK_CUDNN_STATUS(cudnnSetActivationDescriptor(
+				mSigmoidDesc, CUDNN_ACTIVATION_SIGMOID, CUDNN_NOT_PROPAGATE_NAN, 0.0));
+		#else
+			mSigmoidDesc = CUDNN_ACTIVATION_SIGMOID;
+		#endif
+    }
+
+    // Do the actual backward computation
+    // Template is only for scaling parameters, which are always in float
+    // excepted when the convolution is performed in double precision.
+    if (op.getInput(0)->grad()->dataType() == DataType::Float64) {
+        backward_<double>(output_grad);
+    }
+    else {
+        backward_<float>(output_grad);
+    }
+}
+
+template <class T>
+void Aidge::SigmoidImpl_cuda::backward_(const Tensor& output_grad) {
+    const OperatorTensor& op = static_cast<const OperatorTensor&>(mOp);
+    const typename Cuda::cudnn_scaling_type<T>::type alpha = 1.0f;
+    const typename Cuda::cudnn_scaling_type<T>::type beta = 0.0f;
+    CHECK_CUDNN_STATUS(
+        cudnnActivationBackward(CudaContext::cudnnHandle(),
+                               mSigmoidDesc,
+                               &alpha,
+							   std::dynamic_pointer_cast<TensorImpl_cuda_>(op.getOutput(0)->getImpl())->getCudnnTensorDesc(*op.getOutput(0)),
+                               std::static_pointer_cast<Tensor>(op.getRawOutput(0))->getImpl()->rawPtr(),
+							   std::dynamic_pointer_cast<TensorImpl_cuda_>(output_grad.getImpl())->getCudnnTensorDesc(output_grad),
+                               output_grad.getImpl()->rawPtr(),
+                               std::dynamic_pointer_cast<TensorImpl_cuda_>(op.getInput(0)->getImpl())->getCudnnTensorDesc(*op.getInput(0)),
+                               std::static_pointer_cast<Tensor>(op.getRawInput(0))->getImpl()->rawPtr(),							   
+                               &beta,
+							   std::dynamic_pointer_cast<TensorImpl_cuda_>(op.getInput(0)->grad()->getImpl())->getCudnnTensorDesc(*op.getInput(0)->grad()),
+                               op.getInput(0)->grad()->getImpl()->rawPtr()));
+}
+
 Aidge::SigmoidImpl_cuda::~SigmoidImpl_cuda() {
     if (mSigmoidDesc != nullptr) {
 		#if CUDNN_VERSION >= 5000

src/operator/TanhImpl.cpp

@@ -64,6 +64,55 @@ void Aidge::TanhImpl_cuda::forward_(const Tensor& input) {
                                std::static_pointer_cast<Tensor>(op.getRawOutput(0))->getImpl()->rawPtr()));
 }
 
+void Aidge::TanhImpl_cuda::backward() {
+    const OperatorTensor& op = static_cast<const OperatorTensor&>(mOp);
+
+    assert(op.getOutput(0)->grad() && "missing output #0");
+
+    const auto& output_grad = op.getOutput(0)->grad()->refCastFrom(mOutputGradFallback, *op.getOutput(0)->grad());
+
+    // Lazy-initialize CuDNN Tanh descriptor
+    if (mTanhDesc == nullptr) {
+		#if CUDNN_VERSION >= 5000
+			CHECK_CUDNN_STATUS(cudnnCreateActivationDescriptor(&mTanhDesc));
+			CHECK_CUDNN_STATUS(cudnnSetActivationDescriptor(
+				mTanhDesc, CUDNN_ACTIVATION_SIGMOID, CUDNN_NOT_PROPAGATE_NAN, 0.0));
+		#else
+			mTanhDesc = CUDNN_ACTIVATION_SIGMOID;
+		#endif
+    }
+
+    // Do the actual backward computation
+    // Template is only for scaling parameters, which are always in float
+    // excepted when the convolution is performed in double precision.
+    if (op.getInput(0)->grad()->dataType() == DataType::Float64) {
+        backward_<double>(output_grad);
+    }
+    else {
+        backward_<float>(output_grad);
+    }
+}
+
+template <class T>
+void Aidge::TanhImpl_cuda::backward_(const Tensor& output_grad) {
+    const OperatorTensor& op = static_cast<const OperatorTensor&>(mOp);
+    const typename Cuda::cudnn_scaling_type<T>::type alpha = 1.0f;
+    const typename Cuda::cudnn_scaling_type<T>::type beta = 0.0f;
+    CHECK_CUDNN_STATUS(
+        cudnnActivationBackward(CudaContext::cudnnHandle(),
+                               mTanhDesc,
+                               &alpha,
+							   std::dynamic_pointer_cast<TensorImpl_cuda_>(op.getOutput(0)->getImpl())->getCudnnTensorDesc(*op.getOutput(0)),
+                               std::static_pointer_cast<Tensor>(op.getRawOutput(0))->getImpl()->rawPtr(),
+							   std::dynamic_pointer_cast<TensorImpl_cuda_>(output_grad.getImpl())->getCudnnTensorDesc(output_grad),
+                               output_grad.getImpl()->rawPtr(),
+                               std::dynamic_pointer_cast<TensorImpl_cuda_>(op.getInput(0)->getImpl())->getCudnnTensorDesc(*op.getInput(0)),
+                               std::static_pointer_cast<Tensor>(op.getRawInput(0))->getImpl()->rawPtr(),							   
+                               &beta,
+							   std::dynamic_pointer_cast<TensorImpl_cuda_>(op.getInput(0)->grad()->getImpl())->getCudnnTensorDesc(*op.getInput(0)->grad()),
+                               op.getInput(0)->grad()->getImpl()->rawPtr()));
+}
+
 Aidge::TanhImpl_cuda::~TanhImpl_cuda() {
     if (mTanhDesc != nullptr) {
 		#if CUDNN_VERSION >= 5000