dlib C++ Library - cudnn_dlibapi.cpp

// Copyright (C) 2015 Davis E. King (davis@dlib.net)
// License: Boost Software License See LICENSE.txt for the full license.
#ifndef DLIB_DNN_CuDNN_CPP_
#define DLIB_DNN_CuDNN_CPP_
#ifdef DLIB_USE_CUDA
#include "cudnn_dlibapi.h"
#include "tensor.h"
#include <cudnn.h>
#include <iostream>
#include <string>
#include <vector>
#include "cuda_utils.h"
#include "cpu_dlib.h"
#include "cuda_dlib.h"
#include "tensor_tools.h"
static const char* cudnn_get_error_string(cudnnStatus_t s)
{
 switch(s)
 {
 case CUDNN_STATUS_NOT_INITIALIZED: 
 return "CUDA Runtime API initialization failed.";
 case CUDNN_STATUS_ALLOC_FAILED: 
 return "CUDA Resources could not be allocated.";
 case CUDNN_STATUS_BAD_PARAM:
 return "CUDNN_STATUS_BAD_PARAM";
 case CUDNN_STATUS_EXECUTION_FAILED:
 return "CUDNN_STATUS_EXECUTION_FAILED";
 case CUDNN_STATUS_NOT_SUPPORTED:
 return "CUDNN_STATUS_NOT_SUPPORTED";
 case CUDNN_STATUS_ARCH_MISMATCH:
 return "CUDNN_STATUS_ARCH_MISMATCH: Your GPU is too old and not supported by cuDNN";
 default:
 return "A call to cuDNN failed";
 }
}
// Check the return value of a call to the cuDNN runtime for an error condition.
#define CHECK_CUDNN(call) \
do{ \
 const cudnnStatus_t error = call; \
 if (error != CUDNN_STATUS_SUCCESS) \
 { \
 std::ostringstream sout; \
 sout << "Error while calling " << #call << " in file " << __FILE__ << ":" << __LINE__ << ". ";\
 sout << "code: " << error << ", reason: " << cudnn_get_error_string(error);\
 throw dlib::cudnn_error(sout.str()); \
 } \
}while(false)
namespace dlib
{
 namespace cuda 
 {
 // ------------------------------------------------------------------------------------
 static cudnnTensorDescriptor_t descriptor(const tensor& t) 
 {
 return (const cudnnTensorDescriptor_t)t.get_cudnn_tensor_descriptor().get_handle();
 }
 static cudnnTensorDescriptor_t descriptor(const tensor_descriptor& t) 
 {
 return (const cudnnTensorDescriptor_t)t.get_handle();
 }
 // ------------------------------------------------------------------------------------
 class cudnn_context
 {
 public:
 // not copyable
 cudnn_context(const cudnn_context&) = delete;
 cudnn_context& operator=(const cudnn_context&) = delete;
 cudnn_context()
 {
 handles.resize(16);
 }
 ~cudnn_context()
 {
 for (auto h : handles)
 {
 if (h)
 cudnnDestroy(h);
 }
 }
 cudnnHandle_t get_handle (
 ) 
 { 
 int new_device_id;
 CHECK_CUDA(cudaGetDevice(&new_device_id));
 // make room for more devices if needed
 if (new_device_id >= (long)handles.size())
 handles.resize(new_device_id+16);
 // If we don't have a handle already for this device then make one
 if (!handles[new_device_id])
 CHECK_CUDNN(cudnnCreate(&handles[new_device_id]));
 // Finally, return the handle for the current device
 return handles[new_device_id];
 }
 private:
 std::vector<cudnnHandle_t> handles;
 };
 static cudnnHandle_t context()
 {
 thread_local cudnn_context c;
 return c.get_handle();
 }
 // ------------------------------------------------------------------------------------
 class cudnn_device_buffer
 {
 public:
 // not copyable
 cudnn_device_buffer(const cudnn_device_buffer&) = delete;
 cudnn_device_buffer& operator=(const cudnn_device_buffer&) = delete;
 cudnn_device_buffer()
 {
 buffers.resize(16);
 }
 ~cudnn_device_buffer()
 {
 }
 std::shared_ptr<resizable_cuda_buffer> get_buffer (
 )
 {
 int new_device_id;
 CHECK_CUDA(cudaGetDevice(&new_device_id));
 // make room for more devices if needed
 if (new_device_id >= (long)buffers.size())
 buffers.resize(new_device_id+16);
 // If we don't have a buffer already for this device then make one
 std::shared_ptr<resizable_cuda_buffer> buff = buffers[new_device_id].lock();
 if (!buff)
 {
 buff = std::make_shared<resizable_cuda_buffer>();
 buffers[new_device_id] = buff;
 }
 // Finally, return the buffer for the current device
 return buff;
 }
 private:
 std::vector<std::weak_ptr<resizable_cuda_buffer>> buffers;
 };
 static std::shared_ptr<resizable_cuda_buffer> device_global_buffer()
 {
 thread_local cudnn_device_buffer buffer;
 return buffer.get_buffer();
 }
 // ------------------------------------------------------------------------------------
 class cudnn_activation_descriptor
 {
 public:
 // not copyable 
 cudnn_activation_descriptor(const cudnn_activation_descriptor&) = delete;
 cudnn_activation_descriptor& operator=(const cudnn_activation_descriptor&) = delete;
 cudnn_activation_descriptor(
 cudnnActivationMode_t mode,
 cudnnNanPropagation_t reluNanOpt,
 double reluCeiling
 )
 {
 CHECK_CUDNN(cudnnCreateActivationDescriptor(&handle));
 CHECK_CUDNN(cudnnSetActivationDescriptor(handle, mode, reluNanOpt, reluCeiling));
 }
 ~cudnn_activation_descriptor()
 {
 cudnnDestroyActivationDescriptor(handle);
 }
 cudnnActivationDescriptor_t get_handle (
 ) 
 { 
 return handle; 
 }
 private:
 cudnnActivationDescriptor_t handle;
 };
 static cudnnActivationDescriptor_t relu_activation_descriptor()
 {
 thread_local cudnn_activation_descriptor des(CUDNN_ACTIVATION_RELU, CUDNN_PROPAGATE_NAN,0);
 return des.get_handle();
 }
 static cudnnActivationDescriptor_t sigmoid_activation_descriptor()
 {
 thread_local cudnn_activation_descriptor des(CUDNN_ACTIVATION_SIGMOID, CUDNN_PROPAGATE_NAN,0);
 return des.get_handle();
 }
 static cudnnActivationDescriptor_t tanh_activation_descriptor()
 {
 thread_local cudnn_activation_descriptor des(CUDNN_ACTIVATION_TANH, CUDNN_PROPAGATE_NAN,0);
 return des.get_handle();
 }
 // ------------------------------------------------------------------------------------
 tensor_descriptor::
 tensor_descriptor(
 ) : handle(nullptr)
 {
 }
 tensor_descriptor::
 ~tensor_descriptor()
 {
 set_size(0,0,0,0);
 }
 void tensor_descriptor::
 set_size(
 int n, 
 int k,
 int nr, 
 int nc 
 )
 {
 if (handle)
 {
 cudnnDestroyTensorDescriptor((cudnnTensorDescriptor_t)handle);
 handle = nullptr;
 }
 if (n != 0 && nr != 0 && nc != 0 && k != 0)
 {
 cudnnTensorDescriptor_t h;
 CHECK_CUDNN(cudnnCreateTensorDescriptor(&h));
 handle = h;
 CHECK_CUDNN(cudnnSetTensor4dDescriptor((cudnnTensorDescriptor_t)handle,
 CUDNN_TENSOR_NCHW,
 CUDNN_DATA_FLOAT,
 n,
 k,
 nr,
 nc));
 }
 }
 void tensor_descriptor::
 get_size (
 int& n, 
 int& k,
 int& nr,
 int& nc
 ) const
 {
 if (handle)
 {
 int nStride, cStride, hStride, wStride;
 cudnnDataType_t datatype;
 CHECK_CUDNN(cudnnGetTensor4dDescriptor((cudnnTensorDescriptor_t)handle,
 &datatype,
 &n,
 &k,
 &nr,
 &nc,
 &nStride,
 &cStride,
 &hStride,
 &wStride));
 }
 else
 {
 n = 0;
 k = 0;
 nr = 0;
 nc = 0;
 }
 }
 // ------------------------------------------------------------------------------------
 void add(
 float beta,
 tensor& dest,
 float alpha,
 const tensor& src
 )
 {
 DLIB_CASSERT(
 (have_same_dimensions(src, dest) ||
 (src.num_samples()==1 && src.k()==dest.k() && src.nr()==1 && src.nc()==1) ||
 (src.num_samples()==1 && src.k()==dest.k() && src.nr()==dest.nr() && src.nc()==dest.nc()) ||
 (src.num_samples()==1 && src.k()==1 && src.nr()==dest.nr() && src.nc()==dest.nc()) ||
 (src.num_samples()==dest.num_samples() && src.k()==1 && src.nr()==1 && src.nc()==1)) &&
 is_same_object(src,dest) == false , 
 "\n\t dest.num_samples(): " << dest.num_samples()
 <<"\n\t dest.k(): " << dest.k()
 <<"\n\t dest.nr(): " << dest.nr()
 <<"\n\t dest.nc(): " << dest.nc()
 <<"\n\t src.num_samples(): " << src.num_samples()
 <<"\n\t src.k(): " << src.k()
 <<"\n\t src.nr(): " << src.nr()
 <<"\n\t src.nc(): " << src.nc()
 );
 if (dest.size() == src.size() && beta == 1)
 {
 // Call the dlib function in this case since it's faster than the one that
 // comes with cuDNN (at least as of cuDNN v4).
 add_scaled(dest, alpha, src);
 return;
 }
 else if (src.num_samples()==dest.num_samples() && src.k()==1 && src.nr()==1 && src.nc()==1)
 {
 add_cv_to_all_columns(beta, dest, alpha, src);
 return;
 }
 CHECK_CUDNN(cudnnAddTensor(context(),
 &alpha,
 descriptor(src),
 src.device(),
 &beta,
 descriptor(dest),
 dest.device()));
 }
 void assign_conv_bias_gradient (
 tensor& grad,
 const tensor& gradient_input
 )
 {
 DLIB_CASSERT(
 grad.num_samples() == 1 &&
 grad.k() >= 1 &&
 grad.nr() == 1 &&
 grad.nc() == 1 &&
 gradient_input.k() == grad.k() &&
 gradient_input.size() > 0 &&
 is_same_object(grad,gradient_input) == false
 );
 const float alpha = 1;
 const float beta = 0;
 CHECK_CUDNN(cudnnConvolutionBackwardBias(context(),
 &alpha,
 descriptor(gradient_input),
 gradient_input.device(),
 &beta,
 descriptor(grad),
 grad.device()));
 }
 // ------------------------------------------------------------------------------------
 void batch_normalize_inference (
 const double eps,
 resizable_tensor& dest,
 const tensor& src,
 const tensor& gamma, 
 const tensor& beta,
 const tensor& running_means,
 const tensor& running_variances
 )
 {
 DLIB_CASSERT(
 gamma.num_samples() == 1 && 
 gamma.nr() == src.nr() &&
 gamma.nc() == src.nc() &&
 gamma.k() == src.k() &&
 have_same_dimensions(gamma, beta) &&
 have_same_dimensions(gamma, running_means) &&
 have_same_dimensions(gamma, running_variances) && 
 eps > 0, 
 "\ngamma.num_samples(): " << gamma.num_samples() << 
 "\ngamma.k(): " << gamma.k() << 
 "\ngamma.nr(): " << gamma.nr() << 
 "\ngamma.nc(): " << gamma.nc() << 
 "\nbeta.num_samples(): " << beta.num_samples() << 
 "\nbeta.k(): " << beta.k() << 
 "\nbeta.nr(): " << beta.nr() << 
 "\nbeta.nc(): " << beta.nc() << 
 "\nrunning_means.num_samples(): " << running_means.num_samples() << 
 "\nrunning_means.k(): " << running_means.k() << 
 "\nrunning_means.nr(): " << running_means.nr() << 
 "\nrunning_means.nc(): " << running_means.nc() << 
 "\nrunning_variances.num_samples(): " << running_variances.num_samples() << 
 "\nrunning_variances.k(): " << running_variances.k() << 
 "\nrunning_variances.nr(): " << running_variances.nr() << 
 "\nrunning_variances.nc(): " << running_variances.nc() << 
 "\nsrc.k(): " << src.k() << 
 "\nsrc.nr(): " << src.nr() << 
 "\nsrc.nc(): " << src.nc() <<
 "\neps: " << eps 
 );
 const float in_scale = 1;
 const float out_scale = 0;
 dest.copy_size(src);
 CHECK_CUDNN(cudnnBatchNormalizationForwardInference(
 context(),
 CUDNN_BATCHNORM_PER_ACTIVATION,
 &in_scale,
 &out_scale,
 descriptor(src),
 src.device(),
 descriptor(dest),
 dest.device(),
 descriptor(gamma),
 gamma.device(),
 beta.device(),
 running_means.device(),
 running_variances.device(),
 eps));
 }
 void batch_normalize (
 const double eps,
 resizable_tensor& dest,
 resizable_tensor& means,
 resizable_tensor& invstds,
 const double averaging_factor,
 resizable_tensor& running_means,
 resizable_tensor& running_variances,
 const tensor& src,
 const tensor& gamma, 
 const tensor& beta 
 )
 {
 DLIB_CASSERT(0 <= averaging_factor && averaging_factor <= 1, "averaging_factor: " << averaging_factor);
 DLIB_CASSERT(averaging_factor==1 || have_same_dimensions(running_means,means));
 DLIB_CASSERT(averaging_factor==1 || have_same_dimensions(running_variances,invstds));
 DLIB_CASSERT(
 src.num_samples() > 1 &&
 gamma.num_samples() == 1 && 
 beta.num_samples() == 1 && 
 gamma.nr() == beta.nr() && beta.nr() == src.nr() &&
 gamma.nc() == beta.nc() && beta.nc() == src.nc() &&
 gamma.k() == beta.k() && beta.k() == src.k() &&
 eps > 0, 
 "\ngamma.num_samples(): " << gamma.num_samples() << 
 "\ngamma.k(): " << gamma.k() << 
 "\ngamma.nr(): " << gamma.nr() << 
 "\ngamma.nc(): " << gamma.nc() << 
 "\nbeta.num_samples(): " << beta.num_samples() << 
 "\nbeta.k(): " << beta.k() << 
 "\nbeta.nr(): " << beta.nr() << 
 "\nbeta.nc(): " << beta.nc() << 
 "\nsrc.k(): " << src.k() << 
 "\nsrc.nr(): " << src.nr() << 
 "\nsrc.nc(): " << src.nc() <<
 "\neps: " << eps 
 );
 const float in_scale = 1;
 const float out_scale = 0;
 dest.copy_size(src);
 means.set_size(1, src.k(), src.nr(), src.nc());
 invstds.copy_size(means);
 running_means.copy_size(means);
 running_variances.copy_size(means);
 // cuDNN requires that running_means and running_variances be initialized to
 // some valid float values even if the averaging factor would have ignored
 // them. 
 if (averaging_factor == 1)
 {
 running_means = 0;
 running_variances = 1;
 }
 CHECK_CUDNN(cudnnBatchNormalizationForwardTraining(
 context(),
 CUDNN_BATCHNORM_PER_ACTIVATION,
 &in_scale,
 &out_scale,
 descriptor(src),
 src.device(),
 descriptor(dest),
 dest.device(),
 descriptor(gamma),
 gamma.device(),
 beta.device(),
 averaging_factor,
 running_means.device(),
 running_variances.device(),
 eps,
 means.device(),
 invstds.device()));
 }
 void batch_normalize_gradient(
 const double eps,
 const tensor& gradient_input,
 const tensor& means,
 const tensor& invstds,
 const tensor& src,
 const tensor& gamma,
 tensor& src_grad,
 tensor& gamma_grad, 
 tensor& beta_grad 
 )
 {
 const long num = src.k()*src.nr()*src.nc();
 DLIB_CASSERT(src.num_samples() > 1);
 DLIB_CASSERT(num == (long)means.size());
 DLIB_CASSERT(num == (long)invstds.size());
 DLIB_CASSERT(num == (long)gamma.size());
 DLIB_CASSERT(num == (long)gamma_grad.size());
 DLIB_CASSERT(num == (long)beta_grad.size());
 DLIB_CASSERT(have_same_dimensions(gradient_input, src));
 DLIB_CASSERT(have_same_dimensions(gradient_input, src_grad));
 DLIB_CASSERT(eps > 0);
 const float in_scale = 1;
 const float out_scale = 1;
 const float in_scale_params = 1;
 const float out_scale_params = 0;
 CHECK_CUDNN(cudnnBatchNormalizationBackward(
 context(),
 CUDNN_BATCHNORM_PER_ACTIVATION,
 &in_scale,
 &out_scale,
 &in_scale_params,
 &out_scale_params,
 descriptor(src),
 src.device(),
 descriptor(gradient_input),
 gradient_input.device(),
 descriptor(src_grad),
 src_grad.device(),
 descriptor(gamma),
 gamma.device(),
 gamma_grad.device(),
 beta_grad.device(),
 eps,
 means.device(),
 invstds.device()));
 }
 // ------------------------------------------------------------------------------------
 void batch_normalize_conv_inference (
 const double eps,
 resizable_tensor& dest,
 const tensor& src,
 const tensor& gamma, 
 const tensor& beta,
 const tensor& running_means,
 const tensor& running_variances
 )
 {
 DLIB_CASSERT(
 gamma.num_samples() == 1 && 
 gamma.nr() == 1 &&
 gamma.nc() == 1 &&
 gamma.k() == src.k() &&
 have_same_dimensions(gamma, beta) &&
 have_same_dimensions(gamma, running_means) &&
 have_same_dimensions(gamma, running_variances) &&
 eps > 0, 
 "\ngamma.num_samples(): " << gamma.num_samples() << 
 "\ngamma.k(): " << gamma.k() << 
 "\ngamma.nr(): " << gamma.nr() << 
 "\ngamma.nc(): " << gamma.nc() << 
 "\nbeta.num_samples(): " << beta.num_samples() << 
 "\nbeta.k(): " << beta.k() << 
 "\nbeta.nr(): " << beta.nr() << 
 "\nbeta.nc(): " << beta.nc() << 
 "\nrunning_means.num_samples(): " << running_means.num_samples() << 
 "\nrunning_means.k(): " << running_means.k() << 
 "\nrunning_means.nr(): " << running_means.nr() << 
 "\nrunning_means.nc(): " << running_means.nc() << 
 "\nrunning_variances.num_samples(): " << running_variances.num_samples() << 
 "\nrunning_variances.k(): " << running_variances.k() << 
 "\nrunning_variances.nr(): " << running_variances.nr() << 
 "\nrunning_variances.nc(): " << running_variances.nc() << 
 "\nsrc.k(): " << src.k() << 
 "\nsrc.nr(): " << src.nr() << 
 "\nsrc.nc(): " << src.nc() <<
 "\neps: " << eps 
 );
 const float in_scale = 1;
 const float out_scale = 0;
 dest.copy_size(src);
 CHECK_CUDNN(cudnnBatchNormalizationForwardInference(
 context(),
 CUDNN_BATCHNORM_SPATIAL,
 &in_scale,
 &out_scale,
 descriptor(src),
 src.device(),
 descriptor(dest),
 dest.device(),
 descriptor(gamma),
 gamma.device(),
 beta.device(),
 running_means.device(),
 running_variances.device(),
 eps));
 }
 void batch_normalize_conv (
 const double eps,
 resizable_tensor& dest,
 resizable_tensor& means,
 resizable_tensor& invstds,
 const double averaging_factor,
 resizable_tensor& running_means,
 resizable_tensor& running_variances,
 const tensor& src,
 const tensor& gamma, 
 const tensor& beta 
 )
 {
 DLIB_CASSERT(0 <= averaging_factor && averaging_factor <= 1, "averaging_factor: " << averaging_factor);
 DLIB_CASSERT(averaging_factor==1 || have_same_dimensions(running_means,means));
 DLIB_CASSERT(averaging_factor==1 || have_same_dimensions(running_variances,invstds));
 DLIB_CASSERT(
 src.num_samples() > 1 &&
 gamma.num_samples() == 1 && 
 beta.num_samples() == 1 && 
 gamma.nr() == 1 && 
 beta.nr() == 1 && 
 gamma.nc() == 1 && 
 beta.nc() == 1 && 
 gamma.k() == beta.k() && beta.k() == src.k() &&
 eps > 0, 
 "\ngamma.num_samples(): " << gamma.num_samples() << 
 "\ngamma.k(): " << gamma.k() << 
 "\ngamma.nr(): " << gamma.nr() << 
 "\ngamma.nc(): " << gamma.nc() << 
 "\nbeta.num_samples(): " << beta.num_samples() << 
 "\nbeta.k(): " << beta.k() << 
 "\nbeta.nr(): " << beta.nr() << 
 "\nbeta.nc(): " << beta.nc() << 
 "\nsrc.k(): " << src.k() << 
 "\nsrc.nr(): " << src.nr() << 
 "\nsrc.nc(): " << src.nc() <<
 "\neps: " << eps 
 );
 const float in_scale = 1;
 const float out_scale = 0;
 dest.copy_size(src);
 means.set_size(1, src.k());
 invstds.copy_size(means);
 running_means.copy_size(means);
 running_variances.copy_size(means);
 // cuDNN requires that running_means and running_variances be initialized to
 // some valid float values even if the averaging factor would have ignored
 // them. 
 if (averaging_factor == 1)
 {
 running_means = 0;
 running_variances = 1;
 }
 CHECK_CUDNN(cudnnBatchNormalizationForwardTraining(
 context(),
 CUDNN_BATCHNORM_SPATIAL,
 &in_scale,
 &out_scale,
 descriptor(src),
 src.device(),
 descriptor(dest),
 dest.device(),
 descriptor(gamma),
 gamma.device(),
 beta.device(),
 averaging_factor,
 running_means.device(),
 running_variances.device(),
 eps,
 means.device(),
 invstds.device()));
 }
 void batch_normalize_conv_gradient(
 const double eps,
 const tensor& gradient_input,
 const tensor& means,
 const tensor& invstds,
 const tensor& src,
 const tensor& gamma,
 tensor& src_grad,
 tensor& gamma_grad, 
 tensor& beta_grad 
 )
 {
 DLIB_CASSERT(src.k() == (long)means.size());
 DLIB_CASSERT(src.k() == (long)invstds.size());
 DLIB_CASSERT(src.k() == (long)gamma.size());
 DLIB_CASSERT(src.k() == (long)gamma_grad.size());
 DLIB_CASSERT(src.k() == (long)beta_grad.size());
 DLIB_CASSERT(have_same_dimensions(gradient_input, src));
 DLIB_CASSERT(have_same_dimensions(gradient_input, src_grad));
 DLIB_CASSERT(eps > 0);
 const float in_scale = 1;
 const float out_scale = 1;
 const float in_scale_params = 1;
 const float out_scale_params = 0;
 CHECK_CUDNN(cudnnBatchNormalizationBackward(
 context(),
 CUDNN_BATCHNORM_SPATIAL,
 &in_scale,
 &out_scale,
 &in_scale_params,
 &out_scale_params,
 descriptor(src),
 src.device(),
 descriptor(gradient_input),
 gradient_input.device(),
 descriptor(src_grad),
 src_grad.device(),
 descriptor(gamma),
 gamma.device(),
 gamma_grad.device(),
 beta_grad.device(),
 eps,
 means.device(),
 invstds.device()));
 }
 // ------------------------------------------------------------------------------------
 // ------------------------------------------------------------------------------------
 tensor_conv::
 tensor_conv(
 ) : 
 filter_handle(nullptr),
 conv_handle(nullptr),
 forward_algo(0),
 backward_data_algo(0),
 backward_filters_algo(0)
 {
 clear();
 }
 void tensor_conv::
 clear (
 )
 {
 if (filter_handle) 
 cudnnDestroyFilterDescriptor((cudnnFilterDescriptor_t)filter_handle);
 if (conv_handle) 
 cudnnDestroyConvolutionDescriptor((cudnnConvolutionDescriptor_t)conv_handle);
 filter_handle = nullptr;
 conv_handle = nullptr;
 out_num_samples = 0;
 out_k = 0;
 out_nr = 0;
 out_nc = 0;
 stride_y = 0;
 stride_x = 0;
 padding_y = 0;
 padding_x = 0;
 data_num_samples = 0;
 data_k = 0;
 data_nr = 0;
 data_nc = 0;
 filters_num_samples = 0;
 filters_k = 0;
 filters_nr = 0;
 filters_nc = 0;
 forward_algo = 0;
 backward_data_algo = 0;
 backward_filters_algo = 0;
 forward_workspace_size_in_bytes = 0;
 backward_data_workspace_size_in_bytes = 0;
 backward_filters_workspace_size_in_bytes = 0;
 forward_workspace.reset();
 backward_data_workspace.reset();
 backward_filters_workspace.reset();
 workspace.reset();
 }
 void tensor_conv::
 setup(
 const tensor& data,
 const tensor& filters,
 int stride_y_,
 int stride_x_,
 int padding_y_,
 int padding_x_
 ) 
 {
 DLIB_CASSERT(data.k() == filters.k());
 // if the last call to setup gave the same exact settings then don't do
 // anything.
 if (stride_y_ == stride_y && 
 stride_x_ == stride_x &&
 padding_y_ == padding_y && 
 padding_x_ == padding_x &&
 data_num_samples == data.num_samples() &&
 data_k == data.k() &&
 data_nr == data.nr() &&
 data_nc == data.nc() &&
 filters_num_samples == filters.num_samples() &&
 filters_k == filters.k() &&
 filters_nr == filters.nr() &&
 filters_nc == filters.nc())
 {
 return;
 }
 clear();
 try
 {
 stride_y = stride_y_;
 stride_x = stride_x_;
 padding_y = padding_y_;
 padding_x = padding_x_;
 data_num_samples = data.num_samples();
 data_k = data.k();
 data_nr = data.nr();
 data_nc = data.nc();
 filters_num_samples = filters.num_samples();
 filters_k = filters.k();
 filters_nr = filters.nr();
 filters_nc = filters.nc();
 CHECK_CUDNN(cudnnCreateFilterDescriptor((cudnnFilterDescriptor_t*)&filter_handle));
 CHECK_CUDNN(cudnnSetFilter4dDescriptor((cudnnFilterDescriptor_t)filter_handle, 
 CUDNN_DATA_FLOAT, 
 CUDNN_TENSOR_NCHW,
 filters.num_samples(),
 filters.k(),
 filters.nr(),
 filters.nc()));
 CHECK_CUDNN(cudnnCreateConvolutionDescriptor((cudnnConvolutionDescriptor_t*)&conv_handle));
#if CUDNN_MAJOR >= 6
 CHECK_CUDNN(cudnnSetConvolution2dDescriptor((cudnnConvolutionDescriptor_t)conv_handle,
 padding_y, // vertical padding
 padding_x, // horizontal padding
 stride_y,
 stride_x,
 1, 1, // must be 1,1
 CUDNN_CROSS_CORRELATION,
 CUDNN_DATA_FLOAT)); // could also be CUDNN_CONVOLUTION
#else
 CHECK_CUDNN(cudnnSetConvolution2dDescriptor((cudnnConvolutionDescriptor_t)conv_handle,
 padding_y, // vertical padding
 padding_x, // horizontal padding
 stride_y,
 stride_x,
 1, 1, // must be 1,1
 CUDNN_CROSS_CORRELATION)); // could also be CUDNN_CONVOLUTION
#endif
 CHECK_CUDNN(cudnnGetConvolution2dForwardOutputDim(
 (const cudnnConvolutionDescriptor_t)conv_handle,
 descriptor(data),
 (const cudnnFilterDescriptor_t)filter_handle,
 &out_num_samples,
 &out_k,
 &out_nr,
 &out_nc));
 tensor_descriptor dest_desc;
 dest_desc.set_size(out_num_samples,out_k,out_nr,out_nc);
 // Pick which forward algorithm we will use and allocate the necessary
 // workspace buffer.
 cudnnConvolutionFwdAlgo_t forward_best_algo;
 CHECK_CUDNN(cudnnGetConvolutionForwardAlgorithm(
 context(), 
 descriptor(data),
 (const cudnnFilterDescriptor_t)filter_handle,
 (const cudnnConvolutionDescriptor_t)conv_handle,
 descriptor(dest_desc),
 dnn_prefer_fastest_algorithms()?CUDNN_CONVOLUTION_FWD_PREFER_FASTEST:CUDNN_CONVOLUTION_FWD_NO_WORKSPACE,
 std::numeric_limits<size_t>::max(),
 &forward_best_algo));
 forward_algo = forward_best_algo;
 CHECK_CUDNN(cudnnGetConvolutionForwardWorkspaceSize( 
 context(),
 descriptor(data),
 (const cudnnFilterDescriptor_t)filter_handle,
 (const cudnnConvolutionDescriptor_t)conv_handle,
 descriptor(dest_desc),
 forward_best_algo,
 &forward_workspace_size_in_bytes));
 // Pick which backward data algorithm we will use and allocate the
 // necessary workspace buffer.
 cudnnConvolutionBwdDataAlgo_t backward_data_best_algo;
 CHECK_CUDNN(cudnnGetConvolutionBackwardDataAlgorithm(
 context(),
 (const cudnnFilterDescriptor_t)filter_handle,
 descriptor(dest_desc),
 (const cudnnConvolutionDescriptor_t)conv_handle,
 descriptor(data),
 dnn_prefer_fastest_algorithms()?CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST:CUDNN_CONVOLUTION_BWD_DATA_NO_WORKSPACE,
 std::numeric_limits<size_t>::max(),
 &backward_data_best_algo));
 backward_data_algo = backward_data_best_algo;
 CHECK_CUDNN(cudnnGetConvolutionBackwardDataWorkspaceSize(
 context(),
 (const cudnnFilterDescriptor_t)filter_handle,
 descriptor(dest_desc),
 (const cudnnConvolutionDescriptor_t)conv_handle,
 descriptor(data),
 backward_data_best_algo,
 &backward_data_workspace_size_in_bytes));
 // Pick which backward filters algorithm we will use and allocate the
 // necessary workspace buffer.
 cudnnConvolutionBwdFilterAlgo_t backward_filters_best_algo;
 CHECK_CUDNN(cudnnGetConvolutionBackwardFilterAlgorithm(
 context(),
 descriptor(data),
 descriptor(dest_desc),
 (const cudnnConvolutionDescriptor_t)conv_handle,
 (const cudnnFilterDescriptor_t)filter_handle,
 dnn_prefer_fastest_algorithms()?CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST:CUDNN_CONVOLUTION_BWD_FILTER_NO_WORKSPACE,
 std::numeric_limits<size_t>::max(),
 &backward_filters_best_algo));
 // cuDNN 5.1 has a bug that causes
 // cudnnGetConvolutionBackwardFilterAlgorithm() to pick the winograd
 // algorithm even for cases where cuDNN doesn't support it, leading to
 // incorrect outputs. So here we check if we are in a case where winograd
 // isn't supported and manually overrule
 // cudnnGetConvolutionBackwardFilterAlgorithm() by picking a safe
 // algorithm.
 if (dnn_prefer_fastest_algorithms() && 
 !(stride_x == 1 && stride_y == 1 && ((filters_nr==3&&filters_nc==3) || (filters_nr==5&&filters_nc==5)))
 )
 {
 backward_filters_best_algo = CUDNN_CONVOLUTION_BWD_FILTER_ALGO_0;
 }
 backward_filters_algo = backward_filters_best_algo;
 CHECK_CUDNN(cudnnGetConvolutionBackwardFilterWorkspaceSize( 
 context(),
 descriptor(data),
 descriptor(dest_desc),
 (const cudnnConvolutionDescriptor_t)conv_handle,
 (const cudnnFilterDescriptor_t)filter_handle,
 backward_filters_best_algo,
 &backward_filters_workspace_size_in_bytes));
 workspace = device_global_buffer();
 }
 catch(...)
 {
 clear();
 throw;
 }
 }
 tensor_conv::
 ~tensor_conv (
 )
 {
 clear();
 }
 void tensor_conv::operator() (
 const bool add_to_output,
 resizable_tensor& output,
 const tensor& data,
 const tensor& filters
 )
 {
 DLIB_CASSERT(stride_y > 0 && stride_x > 0, "You must call setup() before calling this function");
 output.set_size(out_num_samples, out_k, out_nr, out_nc);
 (*this)(add_to_output, static_cast<tensor&>(output), data, filters);
 }
 void tensor_conv::operator() (
 const bool add_to_output,
 tensor& output,
 const tensor& data,
 const tensor& filters
 )
 {
 DLIB_CASSERT(is_same_object(output,data) == false);
 DLIB_CASSERT(is_same_object(output,filters) == false);
 DLIB_CASSERT(filters.k() == data.k());
 DLIB_CASSERT(stride_y > 0 && stride_x > 0, "You must call setup() before calling this function");
 DLIB_CASSERT(filters.nc() <= data.nc() + 2*padding_x,
 "Filter windows must be small enough to fit into the padded image."
 << "\n\t filters.nc(): " << filters.nc() 
 << "\n\t data.nc(): " << data.nc() 
 << "\n\t padding_x: " << padding_x 
 );
 DLIB_CASSERT(filters.nr() <= data.nr() + 2*padding_y,
 "Filter windows must be small enough to fit into the padded image."
 << "\n\t filters.nr(): " << filters.nr() 
 << "\n\t data.nr(): " << data.nr() 
 << "\n\t padding_y: " << padding_y 
 );
 DLIB_CASSERT(output.num_samples() == data.num_samples(),out_num_samples << " " << data.num_samples());
 DLIB_CASSERT(output.k() == filters.num_samples());
 DLIB_CASSERT(output.nr() == 1+(data.nr()+2*padding_y-filters.nr())/stride_y);
 DLIB_CASSERT(output.nc() == 1+(data.nc()+2*padding_x-filters.nc())/stride_x);
 const float alpha = 1;
 const float beta = add_to_output ? 1 : 0;
 // Since cudnnConvolutionForward() is an asynchronous call, we need to hold a
 // reference to the workspace buffer so we can be sure it isn't reallocated
 // while the function is still executing on the device. But each time we come
 // here, we make sure to grab the latest workspace buffer so that, globally, we
 // minimize the number of such buffers.
 forward_workspace = workspace->get(forward_workspace_size_in_bytes);
 CHECK_CUDNN(cudnnConvolutionForward(
 context(),
 &alpha,
 descriptor(data),
 data.device(),
 (const cudnnFilterDescriptor_t)filter_handle,
 filters.device(),
 (const cudnnConvolutionDescriptor_t)conv_handle,
 (cudnnConvolutionFwdAlgo_t)forward_algo,
 forward_workspace,
 forward_workspace_size_in_bytes,
 &beta,
 descriptor(output),
 output.device()));
 }
 void tensor_conv::get_gradient_for_data (
 const bool add_to_output,
 const tensor& gradient_input, 
 const tensor& filters,
 tensor& data_gradient
 )
 {
 const float alpha = 1;
 const float beta = add_to_output ? 1 : 0;
 // Since cudnnConvolutionBackwardData() is an asynchronous call, we need to hold a
 // reference to the workspace buffer so we can be sure it isn't reallocated
 // while the function is still executing on the device. But each time we come
 // here, we make sure to grab the latest workspace buffer so that, globally, we
 // minimize the number of such buffers.
 backward_data_workspace = workspace->get(backward_data_workspace_size_in_bytes);
 CHECK_CUDNN(cudnnConvolutionBackwardData(context(),
 &alpha,
 (const cudnnFilterDescriptor_t)filter_handle,
 filters.device(),
 descriptor(gradient_input),
 gradient_input.device(),
 (const cudnnConvolutionDescriptor_t)conv_handle,
 (cudnnConvolutionBwdDataAlgo_t)backward_data_algo,
 backward_data_workspace,
 backward_data_workspace_size_in_bytes,
 &beta,
 descriptor(data_gradient),
 data_gradient.device()));
 }
 void tensor_conv::
 get_gradient_for_filters (
 const bool add_to_output,
 const tensor& gradient_input, 
 const tensor& data,
 tensor& filters_gradient
 )
 {
 const float alpha = 1;
 const float beta = add_to_output ? 1 : 0;
 // Since cudnnConvolutionBackwardFilter() is an asynchronous call, we need to hold a
 // reference to the workspace buffer so we can be sure it isn't reallocated
 // while the function is still executing on the device. But each time we come
 // here, we make sure to grab the latest workspace buffer so that, globally, we
 // minimize the number of such buffers.
 backward_filters_workspace = workspace->get(backward_filters_workspace_size_in_bytes);
 CHECK_CUDNN(cudnnConvolutionBackwardFilter(context(),
 &alpha,
 descriptor(data),
 data.device(),
 descriptor(gradient_input),
 gradient_input.device(),
 (const cudnnConvolutionDescriptor_t)conv_handle,
 (cudnnConvolutionBwdFilterAlgo_t)backward_filters_algo,
 backward_filters_workspace,
 backward_filters_workspace_size_in_bytes,
 &beta,
 (const cudnnFilterDescriptor_t)filter_handle,
 filters_gradient.device()));
 }
 // ------------------------------------------------------------------------------------
 // ------------------------------------------------------------------------------------
 pooling::pooling (
 ) : handle(nullptr),window_height(0),window_width(0),stride_y(0),stride_x(0),padding_y(0), padding_x(0)
 {
 }
 pooling::~pooling(
 )
 {
 clear();
 }
 void pooling::
 clear(
 )
 {
 if (handle)
 cudnnDestroyPoolingDescriptor((cudnnPoolingDescriptor_t)handle);
 handle = nullptr;
 window_height = 0;
 window_width = 0;
 stride_y = 0;
 stride_x = 0;
 padding_y = 0;
 padding_x = 0;
 }
 void pooling::
 setup_max_pooling(
 int window_height_,
 int window_width_,
 int stride_y_,
 int stride_x_,
 int padding_y_,
 int padding_x_ 
 )
 {
 setup(window_height_, window_width_, stride_y_, stride_x_, padding_y_, padding_x_, CUDNN_POOLING_MAX);
 do_max_pooling = true;
 }
 void pooling::
 setup_avg_pooling(
 int window_height_,
 int window_width_,
 int stride_y_,
 int stride_x_,
 int padding_y_,
 int padding_x_
 )
 {
 setup(window_height_, window_width_, stride_y_, stride_x_, padding_y_, padding_x_, CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING);
 do_max_pooling = false;
 }
 void pooling::
 setup(
 int window_height_,
 int window_width_,
 int stride_y_,
 int stride_x_,
 int padding_y_,
 int padding_x_,
 int pooling_mode
 )
 {
 DLIB_CASSERT (window_height_ > 0 && window_width_ > 0 && 
 stride_y_ > 0 && stride_x_ > 0 , 
 "window_height_: " << window_height_ 
 << "\t\n window_width_: " << window_width_ 
 << "\t\n stride_y_: " << stride_y_ 
 << "\t\n stride_x_: " << stride_x_ );
 DLIB_CASSERT( 0 <= padding_y_ && padding_y_ < window_height_ && 
 0 <= padding_x_ && padding_x_ < window_width_,
 "window_height_: " << window_height_ 
 << "\t\n window_width_: " << window_width_ 
 << "\t\n padding_y_: " << padding_y_ 
 << "\t\n padding_x_: " << padding_x_ );
 if (window_height == window_height_ &&
 window_width == window_width_ &&
 stride_y == stride_y_ &&
 stride_x == stride_x_ && 
 padding_y == padding_y_ &&
 padding_x == padding_x_
 )
 {
 return;
 }
 clear();
 try
 {
 window_height = window_height_;
 window_width = window_width_;
 stride_x = stride_x_;
 stride_y = stride_y_;
 padding_y = padding_y_;
 padding_x = padding_x_;
 cudnnPoolingDescriptor_t poolingDesc;
 CHECK_CUDNN(cudnnCreatePoolingDescriptor(&poolingDesc));
 handle = poolingDesc;
 CHECK_CUDNN(cudnnSetPooling2dDescriptor(poolingDesc,
 (cudnnPoolingMode_t)pooling_mode,
 CUDNN_PROPAGATE_NAN,
 window_height,
 window_width,
 padding_y,
 padding_x, 
 stride_y,
 stride_x));
 }
 catch(...)
 {
 clear();
 throw;
 }
 }
 void pooling::
 operator() (
 resizable_tensor& dest,
 const tensor& src
 )
 {
 DLIB_CASSERT(window_width <= src.nc() + 2*padding_x,
 "Pooling windows must be small enough to fit into the padded image."
 << "\n\t window_width: " << window_width 
 << "\n\t src.nc(): " << src.nc() 
 << "\n\t padding_x: " << padding_x 
 );
 DLIB_CASSERT(window_height <= src.nr() + 2*padding_y,
 "Pooling windows must be small enough to fit into the padded image."
 << "\n\t window_height: " << window_height 
 << "\n\t src.nr(): " << src.nr() 
 << "\n\t padding_y: " << padding_y 
 );
 const float alpha = 1;
 const float beta = 0;
 int outN;
 int outC;
 int outH;
 int outW;
 CHECK_CUDNN(cudnnGetPooling2dForwardOutputDim((const cudnnPoolingDescriptor_t)handle,
 descriptor(src),
 &outN,
 &outC,
 &outH,
 &outW));
 dest.set_size(outN,outC,outH,outW);
 DLIB_CASSERT(dest.num_samples() == src.num_samples());
 DLIB_CASSERT(dest.k() == src.k());
 DLIB_CASSERT(dest.nr() == 1 + (src.nr() + 2*padding_y - window_height)/stride_y, 
 "\n stride_y: " << stride_y <<
 "\n padding_y: " << padding_y <<
 "\n window_height: " << window_height <<
 "\n src.nr(): " << src.nr() <<
 "\n dest.nr(): " << dest.nr() <<
 "\n src.nr()/stride_y: " << src.nr()/stride_y); 
 DLIB_CASSERT(dest.nc() == 1 + (src.nc() + 2*padding_x - window_width)/stride_x, 
 "\n stride_x: " << stride_x <<
 "\n padding_x: " << padding_x <<
 "\n window_width: " << window_width <<
 "\n src.nc(): " << src.nc() <<
 "\n dest.nc(): " << dest.nc() <<
 "\n src.nc()/stride_x: " << src.nc()/stride_x); 
 CHECK_CUDNN(cudnnPoolingForward(context(),
 (const cudnnPoolingDescriptor_t)handle,
 &alpha,
 descriptor(src),
 src.device(),
 &beta,
 descriptor(dest),
 dest.device()));
 }
 void pooling::get_gradient(
 const tensor& gradient_input, 
 const tensor& dest,
 const tensor& src,
 tensor& grad 
 )
 {
 DLIB_CASSERT(have_same_dimensions(gradient_input,dest));
 DLIB_CASSERT(have_same_dimensions(src,grad));
 const float alpha = 1;
 const float beta = 1;
 CHECK_CUDNN(cudnnPoolingBackward(context(),
 (const cudnnPoolingDescriptor_t)handle,
 &alpha,
 descriptor(dest),
 dest.device(),
 descriptor(gradient_input),
 gradient_input.device(),
 descriptor(src),
 src.device(),
 &beta,
 descriptor(grad),
 grad.device()));
 }
 // ------------------------------------------------------------------------------------
 // ------------------------------------------------------------------------------------
 void softmax (
 tensor& dest,
 const tensor& src
 )
 {
 DLIB_CASSERT(have_same_dimensions(dest,src));
 if (src.size() == 0)
 return;
 const float alpha = 1;
 const float beta = 0;
 CHECK_CUDNN(cudnnSoftmaxForward(context(),
 CUDNN_SOFTMAX_ACCURATE,
 CUDNN_SOFTMAX_MODE_CHANNEL,
 &alpha,
 descriptor(src),
 src.device(),
 &beta,
 descriptor(dest),
 dest.device()));
 }
 void softmax_gradient (
 tensor& grad,
 const tensor& dest,
 const tensor& gradient_input
 )
 {
 DLIB_CASSERT(
 have_same_dimensions(dest,gradient_input) == true &&
 have_same_dimensions(dest,grad) == true );
 if (dest.size() == 0)
 return;
 const float alpha = 1;
 const float beta = is_same_object(grad,gradient_input) ? 0 : 1;
 CHECK_CUDNN(cudnnSoftmaxBackward(context(),
 CUDNN_SOFTMAX_ACCURATE,
 CUDNN_SOFTMAX_MODE_CHANNEL,
 &alpha,
 descriptor(dest),
 dest.device(),
 descriptor(gradient_input),
 gradient_input.device(),
 &beta,
 descriptor(grad),
 grad.device()));
 }
 // ------------------------------------------------------------------------------------
 // ------------------------------------------------------------------------------------
 void softmax_all (
 tensor& dest,
 const tensor& src
 )
 {
 DLIB_CASSERT(have_same_dimensions(dest,src));
 if (src.size() == 0)
 return;
 const float alpha = 1;
 const float beta = 0;
 CHECK_CUDNN(cudnnSoftmaxForward(context(),
 CUDNN_SOFTMAX_ACCURATE,
 CUDNN_SOFTMAX_MODE_INSTANCE,
 &alpha,
 descriptor(src),
 src.device(),
 &beta,
 descriptor(dest),
 dest.device()));
 }
 void softmax_all_gradient (
 tensor& grad,
 const tensor& dest,
 const tensor& gradient_input
 )
 {
 DLIB_CASSERT(
 have_same_dimensions(dest,gradient_input) == true &&
 have_same_dimensions(dest,grad) == true );
 if (dest.size() == 0)
 return;
 const float alpha = 1;
 const float beta = is_same_object(grad,gradient_input) ? 0 : 1;
 CHECK_CUDNN(cudnnSoftmaxBackward(context(),
 CUDNN_SOFTMAX_ACCURATE,
 CUDNN_SOFTMAX_MODE_INSTANCE,
 &alpha,
 descriptor(dest),
 dest.device(),
 descriptor(gradient_input),
 gradient_input.device(),
 &beta,
 descriptor(grad),
 grad.device()));
 }
 // ------------------------------------------------------------------------------------
 // ------------------------------------------------------------------------------------
 void sigmoid (
 tensor& dest,
 const tensor& src
 )
 {
 DLIB_CASSERT(have_same_dimensions(dest,src));
 if (src.size() == 0)
 return;
 const float alpha = 1;
 const float beta = 0;
 CHECK_CUDNN(cudnnActivationForward(context(),
 sigmoid_activation_descriptor(),
 &alpha,
 descriptor(src),
 src.device(),
 &beta,
 descriptor(dest),
 dest.device()));
 }
 void sigmoid_gradient (
 tensor& grad,
 const tensor& dest,
 const tensor& gradient_input
 )
 {
 DLIB_CASSERT(
 have_same_dimensions(dest,gradient_input) == true &&
 have_same_dimensions(dest,grad) == true );
 if (dest.size() == 0)
 return;
 const float alpha = 1;
 const float beta = is_same_object(grad,gradient_input) ? 0 : 1;
 CHECK_CUDNN(cudnnActivationBackward(context(),
 sigmoid_activation_descriptor(),
 &alpha,
 descriptor(dest),
 dest.device(),
 descriptor(gradient_input),
 gradient_input.device(),
 descriptor(dest),
 dest.device(),
 &beta,
 descriptor(grad),
 grad.device()));
 }
 // ------------------------------------------------------------------------------------
 void relu (
 tensor& dest,
 const tensor& src
 )
 {
 DLIB_CASSERT(have_same_dimensions(dest,src));
 if (src.size() == 0)
 return;
 const float alpha = 1;
 const float beta = 0;
 CHECK_CUDNN(cudnnActivationForward(context(),
 relu_activation_descriptor(),
 &alpha,
 descriptor(src),
 src.device(),
 &beta,
 descriptor(dest),
 dest.device()));
 }
 void relu_gradient (
 tensor& grad,
 const tensor& dest,
 const tensor& gradient_input
 )
 {
 DLIB_CASSERT(
 have_same_dimensions(dest,gradient_input) == true &&
 have_same_dimensions(dest,grad) == true );
 if (dest.size() == 0)
 return;
 const float alpha = 1;
 const float beta = is_same_object(grad,gradient_input) ? 0 : 1;
 CHECK_CUDNN(cudnnActivationBackward(context(),
 relu_activation_descriptor(),
 &alpha,
 descriptor(dest),
 dest.device(),
 descriptor(gradient_input),
 gradient_input.device(),
 descriptor(dest),
 dest.device(),
 &beta,
 descriptor(grad),
 grad.device()));
 }
 // ------------------------------------------------------------------------------------
 void tanh (
 tensor& dest,
 const tensor& src
 )
 {
 DLIB_CASSERT(have_same_dimensions(dest,src));
 if (src.size() == 0)
 return;
 const float alpha = 1;
 const float beta = 0;
 CHECK_CUDNN(cudnnActivationForward(context(),
 tanh_activation_descriptor(),
 &alpha,
 descriptor(src),
 src.device(),
 &beta,
 descriptor(dest),
 dest.device()));
 }
 void tanh_gradient (
 tensor& grad,
 const tensor& dest,
 const tensor& gradient_input
 )
 {
 DLIB_CASSERT(
 have_same_dimensions(dest,gradient_input) == true &&
 have_same_dimensions(dest,grad) == true);
 if (dest.size() == 0)
 return;
 const float alpha = 1;
 const float beta = is_same_object(grad,gradient_input) ? 0 : 1;
 CHECK_CUDNN(cudnnActivationBackward(context(),
 tanh_activation_descriptor(),
 &alpha,
 descriptor(dest),
 dest.device(),
 descriptor(gradient_input),
 gradient_input.device(),
 descriptor(dest),
 dest.device(),
 &beta,
 descriptor(grad),
 grad.device()));
 }
 // ------------------------------------------------------------------------------------
 }
}
#endif // DLIB_USE_CUDA
#endif // DLIB_DNN_CuDNN_CPP_