深度学习是一种基于人工神经网络解释数据的机器学习方法。具体来说，我们提供机器可以理解的数据，并构建从数据中学习表示的神经网络模型。我们可以使用这种技术来构建识别语音、从图像中分类对象、理解文本、翻译语言、转换数据域等模型。基本神经网络包括全连接层(FCL)卷积神经网络 ( CNN )和递归神经网络 ( RNN )。这些架构在数据分类、区域理解和顺序关系方面表现出很高的准确性。
深度学习需要大量的计算，以便广泛使用。然而，这个问题得到了解决，因为我们可以通过使用 GPU 计算能力来显著减少训练时间。这是因为神经网络的基本架构是基于矩阵运算的，而 GPU 是为此而优化的硬件平台。具体来说，深度学习的创新是通过英伟达 CUDA 加速来解决的，因为深度学习中的许多算法都可以加速。
在本章中，我们将简要回顾神经网络操作，并讨论如何在图形处理器上加速这些操作。作为实践，我们将使用 cuDNN 和 cuBLAS CUDA 库实现一个卷积网络。cudn 库是英伟达的 CUDA 库，专门优化深度学习操作。我们将分三个部分介绍它的实现。我们还将介绍图形处理器如何优化所需的操作。然后，我们将通过比较长短时记忆 ( LSTM )网络的性能来讲述使用 cuDNN 库是如何有效的。然后，我们将使用英伟达工具扩展 ( 英伟达)来介绍深度学习中的评测方法。这将测量图形处理器上的网络操作，以便我们可以分析时间线中的操作并了解它们的性能。
在本章中，我们将涵盖以下主题:
- 利用 CUBLAS 实现全连接层加速
- 基于 cuDNN 的元素层
- cuDNN/CUDA 中的 Softmax 和损失函数
- 带有 cuDNN 的卷积神经网络
- 带有 CUDA 的递归神经网络
- 剖析深度学习框架
本章要求安装 cudn 库和 CUDA 工具包。我们还需要支持 CUDA 的图形处理器。本章将涵盖深度学习的基础及其性能，因此不需要新的 GPU 功能。换句话说，如果你覆盖了前面章节的大部分内容，你将有一个合适的图形处理器可以使用。
要安装 cuDNN 库，需要从https://developer.nvidia.com/cudnn下载软件包。您需要登录 NVIDIA 开发人员网站才能访问下载页面。如果您还没有 NVIDIA 开发人员帐户，您需要注册该帐户。确保 cudn 是用您安装的 CUDA 版本编译的。
全连接层是深度学习的基本架构。让我们回顾一下它的操作，看看 CUDA 是如何在前向和反向传播过程中加速神经网络的。然后，我们将把它们应用于图形处理器。
神经网络的基本操作是在输入数据和参数之间执行点操作。我们称之为感知。在深度学习中，神经网络以分层的方式连接多个感知。我们称之为前馈神经网络。下图显示了感知器和基本神经网络:
感知器的基本操作是用输入数据和适当的权重创建点积。然后，它执行具有激活功能的非线性操作，例如 sigmoid 或整流器线性单元 ( ReLU )。在前馈神经网络中，操作只是仿射变换，然后应用激活函数。一个向量将作为输入输入到神经网络，并将其与两层中每个节点之间的权重参数相乘。
为了训练神经网络，我们执行前向传播、损耗计算和梯度反向传播，然后使用更新参数。让我们简单介绍一下。然后，我们将使用 cuBLAS 和其他 CUDA 操作来匹配每个步骤。
正向操作可由以下等式表示:
这里，
是给定输入向量的预测结果，
，
是权重参数矩阵，
是激活函数。我们可以看到，全连通层的基本运算是矩阵运算。因此，我们需要对输入和激活函数实现矩阵乘法运算。因为我们承担分类任务，所以我们使用 softmax 函数来标准化输出，并在下一层获得概率分布结果。
为了获得真值之间的损失，我们对标签应用单向编码，并通过获得每个元素的熵来获得交叉熵损失，如下所示:
我们可以通过每个交叉熵损失的和得到总损失值。然后，我们可以从前面的等式中获得梯度。这看起来是一个复杂的操作，但可以简化，如下所示:
现在，我们将梯度传播到上一层，这被称为反向传播。在这个任务中，我们使用链式规则来获得每个权重和偏差参数的梯度。然后，我们可以更新权重参数的设置和偏差。例如，我们可以通过以下等式获得权重和偏差的梯度:
我们可以通过下面的等式获得传播到前一层的梯度:
这里，
是激活函数的梯度。因此，我们需要从第二层获取
作为第一层。然后，第一层的权重和偏差梯度可以通过以下等式获得:
现在，我们可以根据梯度后代规则更新权重和偏差，如下所示:
、
这里，
是迭代步骤。
激活函数
的梯度和类型可以不同。下一节将介绍这个激活层的实现。激活函数的推导可由以下等式表示:
、
因此，神经网络运算是一组线性代数运算，可以被 cuBLAS 库覆盖。实现的代码可以在01_ann中找到。我们将在实现全连接层、T3】实现层操作和实现 softmax 层部分介绍这些实现细节。
在我们编写代码之前，让我们介绍一下如何将操作打包到层配置中:
- 首先，我们执行正向操作。
- 然后，我们执行反向操作。
- 然后我们从梯度得到一个权重更新。
- 最后，输出层将获得损失。
以这种方式，该层可以被配置如下:
根据工作流程，它有标准化的输入和输出以及两种类型的输入。左侧数据路径将使用输入来命名，而右侧将使用输出来命名。数据分两个阶段输入(向前和向后)。我们将使用 blobs 来管理参数和输入/输出数据。blob 是跨层处理的数据的包装，有助于管理内存空间。我们将使用这种设计来简化网络的每一层配置。每个层都有每个斑点的描述符和前向/后向处理操作。
现在，让我们创建一个图层类，它将是所有图层的基类。下面的代码展示了class公共函数是如何堆叠的。并且，你可以在layer.h找到它的实现，在01_ann/src/ directory找到layer.cu。这不仅包括向前和向后操作，还包括重量更新控制和损失计算:
class Layer
{
public:
Layer();
~Layer();
std::string get_name() { return name_; }
virtual Blob<float> *forward(Blob<float> *input) = 0;
virtual Blob<float> *backward(Blob<float> *grad_input) = 0;
virtual float get_loss(Blob<float> *target);
virtual int get_accuracy(Blob<float> *target);
void set_cuda_context(CudaContext *context) { cuda_ = context; }
/* weights update control */
void freeze() { freeze_ = true; }
void unfreeze() { freeze_ = false;}
void set_load_pretrain() { load_pretrain_ = true; }
void set_gradient_stop() { gradient_stop_ = true; }为了支持这些操作，层类维护了几个 cuDNN 描述符、blob 指针和权重更新控制器。当我们介绍网络实施时，将会介绍详细的实施:
protected:
std::string name_;
// Tensor descriptor for the input/output tensor
cudnnTensorDescriptor_t input_desc_;
cudnnTensorDescriptor_t output_desc_;
// filter and bias descriptor for weights and biases
cudnnFilterDescriptor_t filter_desc_;
cudnnTensorDescriptor_t bias_desc_;
// output memory
Blob<float> *input_ = nullptr; /* x */
Blob<float> *output_ = nullptr; /* y */
Blob<float> *grad_input_ = nullptr; /* dx */
Blob<float> *grad_output_ = nullptr; /* dy */
// master weights & bias
bool freeze_ = false; /* control parameter updates */
Blob<float> *weights_ = nullptr; /* w */
Blob<float> *biases_ = nullptr; /* b */
Blob<float> *grad_weights_ = nullptr; /* dw */
Blob<float> *grad_biases_ = nullptr; /* db */
int batch_size_ = 0; // mini-batch size
// cuda handle container
CudaContext *cuda_ = nullptr;
// initialize weights along with the input size
void init_weight_bias(unsigned int seed = 0);
void update_weights_biases(float learning_rate);
// pretrain parameters
bool load_pretrain_ = false;
int load_parameter();
int save_parameter();
// gradient stop tagging
bool gradient_stop_ = false;
friend class Network;
}该层类将在其他部分的深度学习网络实施中使用。因此，它有用于 cuDNN 操作的cudnnTensorDescriptor_t变量，以及get_loss()和get_accuracy()函数。
在我们的实现中，我们将使用名为Blob的数据容器。它的名字是从 Caffe 借来的。这允许我们存储张量或网络参数及其尺寸信息和记忆点。我们将用这个连接每一层。这有助于每个层根据输入张量的大小信息初始化其权重。此外，每一层都可以基于Blob的信息验证其结果。
这个斑点需要神经网络中的尺寸信息，如下面一行代码所示。然后，它的构造函数将根据大小信息创建一个主机端缓冲区:
Blob<T>(int n, int c, int h, int w)Blob还可以处理主机和设备中的记忆，并可以帮助我们访问这些记忆。Blob具有以下内存访问助手功能:
// get specified memory pointer
ftype *ptr() { return h_ptr_; }
// get cuda memory
ftype *cuda()
{
if (d_ptr_ == nullptr)
cudaMalloc((void**)&d_ptr_, sizeof(ftype) * len());
return d_ptr_;
}
// transfer data between memory
ftype *to(DeviceType target) {
ftype *ptr = nullptr;
if (target == host)
{
cudaMemcpy(h_ptr_, cuda(), sizeof(ftype) * len(),
cudaMemcpyDeviceToHost);
ptr = h_ptr_;
}
else // DeviceType::cuda
{
cudaMemcpy(cuda(), h_ptr_, sizeof(ftype) * len(),
cudaMemcpyHostToDevice);
ptr = d_ptr_;
}
return ptr;
}正如我们前面讨论的，Blob可以存储张量，我们还需要提供张量形状信息作为 cuDNN APIs 所需的描述符。因此，Blob可以使用以下代码创建和设置张量描述符:
/* Tensor Control */
bool is_tensor_ = false;
cudnnTensorDescriptor_t tensor_desc_;
cudnnTensorDescriptor_t tensor()
{
if (is_tensor_)
return tensor_desc_;
cudnnCreateTensorDescriptor(&tensor_desc_);
cudnnSetTensor4dDescriptor(tensor_desc_,
CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT,
n_, c_, h_, w_);
is_tensor_ = true;
return tensor_desc_;
}现在，让我们使用Blob实现一个完全连接的层。
在本节中，我们将使用 cuBLAS 编写一个完全连接的网络。对于这一层，我们将创建一个从Layer类派生的Dense类。类构造函数将接收默认的层配置信息，如下所示:
Dense::Dense(std::string name, int output_size)
{
name_ = name;
output_size_ = output_size;
}但这不足以配置整个层。缺失的信息将从输入中提供，因为输入大小将由前一层决定。现在，让我们讨论正向传播。
在前向传播中，我们可以将前向过程分为两个步骤，如下所示:
由于重量大小不必受批次大小的影响，我们只考虑输入重量和输出重量的数量。另一方面，数据馈送斑点，如输入和输出，受到批次大小的影响。因此，我们使用滤波器和输入数据的 GEMM 运算可以设计如下:
隐藏的输出将与偏置值相加。输入数据不限于来自数据加载器的数据。当我们堆叠层时，前一层的输出将是当前层的输入数据。正向操作可以如下实现:
Blob<float> *Dense::forward(Blob<float> *input) {
.. { blob initialization } ..
// output = weights^T * input (without biases)
cublasSgemm(cuda_->cublas(),
CUBLAS_OP_T, CUBLAS_OP_N, output_size_,
batch_size_, input_size_,
&cuda_->one, weights_->cuda(), input_size_,
input_->cuda(), input_size_,
&cuda_->zero, output_->cuda(), output_size_);
// output += biases * one_vec^T
cublasSgemm(cuda_->cublas(),
CUBLAS_OP_N, CUBLAS_OP_N, output_size_, batch_size_, 1,
&cuda_->one, biases_->cuda(), output_size_, one_vec, 1,
&cuda_->one, output_->cuda(), output_size_);
return output_;
}在第一次迭代中，每一层都需要初始化它的权重和偏差。例如，这个Dense层可以初始化它的权重、偏差和输出张量元素。我们可以将这个初始化任务分成两个阶段。首先是权重和偏差，如下所示:
// initialize weights and biases
if (weights_ == nullptr)
{
// setup parameter size information
input_size_ = input->c() * input->h() * input->w();
// initialize weight, bias, and output
weights_ = new Blob<float>(1, 1, input_size_, output_size_);
biases_ = new Blob<float>(1, 1, output_size_);
}接下来的阶段是关于更新输入信息和初始化输出 blob。当它是新的或需要重新配置时，我们需要执行以下操作。在这个任务中，我们还需要创建一个充满我们批量大小的向量。这将用于偏差添加:
// initilaize input and output
if (input_ == nullptr || batch_size_ != input->n())
{
input_ = input;
batch_size_ = input->n();
if (output_ == nullptr)
output_ = new Blob<float>(batch_size_, output_size_);
else
output_->reset(batch_size_, output_size_);
output_->tensor();
if (d_one_vec != nullptr)
cudaFree(d_one_vec);
checkCudaErrors(cudaMalloc((void**)&d_one_vec, sizeof(float) * batch_size_));
init_one_vec<<< (batch_size_+BLOCK_DIM_1D-1)/BLOCK_DIM_1D, BLOCK_DIM_1D >>>(d_one_vec, batch_size_);
if (!freeze_)
init_weight_bias();
}这个初始化任务不仅触发了第一次迭代，还触发了批量大小的改变。在培训阶段不需要检查批次大小，但在测试阶段会很有用。这是因为训练和推理中的批量大小不同。在这种情况下，我们需要按照新的批处理大小创建一个输出 blob。输出张量的大小被确定为通道大小。输出 blob 的创建代码如下，创建一个大小为(batch_size_、output_size_、1、1)的 blob:
output_ = new Blob<float>(batch_size_, output_size_);这就产生了扁平张量。然后，我们喂养这些张量，这要求它们在通道中对齐。这种对齐在 softmax 层中是特别需要的。我们将在 softmax 层的实现中介绍这一点。
这个阶段的另一个重要任务是初始化权重和偏差。在我们的实现中，我们将使用 ReLU 作为激活器。我们将使用普通初始化器(https://arxiv.org/abs/1502.01852)技术来使网络可训练。遵循上一篇文章中的指导原则，可以使用以下等式生成所需的权重值:
是前一层输入的数量。为此，我们可以在更新输入张量信息后初始化参数。此外，偏置值将被初始化为0。下面的代码显示了这个的实现:
void Layer::init_weight_bias(unsigned int seed)
{
// Create random network
std::random_device rd;
std::mt19937 gen(seed == 0 ? rd() : static_cast<unsigned int>
(seed));
// He normal distribution
float range = sqrt(6.f / input_->size());
std::uniform_real_distribution<> dis(-range, range);
for (int i = 0; i < weights_->len(); i++)
weights_->ptr()[i] = static_cast<float>(dis(gen));
for (int i = 0; i < biases_->len(); i++)
biases_->ptr()[i] = 0.f;
// copy initialized value to the device
weights_->to(DeviceType::cuda);
biases_->to(DeviceType::cuda);
}现在，让我们讨论反向传播。
正如我们前面讨论的，从下一层开始的梯度会传播到这一层。基于传播梯度，我们需要获得权重、偏差和数据的三个梯度(输入梯度)。我们需要创建可以存储它们的 blobs。它们的大小不取决于批次大小，所以我们只需要确保我们创建了它们。下面的代码展示了我们如何为此目的创建 blobs:
if (grad_weights_ == nullptr) {
grad_output_ = grad_output;
grad_weights_ = new Blob<float>(weights_->shape());
grad_biases_ = new Blob<float>(biases_->shape());
grad_input_ = new Blob<float>(input_->shape());
}在前面的代码中，grad_output_表示从下一层传播的输出数据的梯度，grad_input_表示将传播到上一层的输入数据的梯度。因此，我们不需要创建一个grad_output_斑点。如果你觉得这些命名惯例很混乱，那么如果你把grad_input_看作
而把grad_input_看作
可能会更容易。
下面的代码展示了我们如何实现这一点:
Blob<float> *Dense::backward(Blob<float> *grad_output) {
.. { blob initialization } ..
// db = (dy) * one_vec
cublasSgemv(cuda_->cublas(),
CUBLAS_OP_N,
output_size_, batch_size_,
&cuda_->one,
grad_output_->cuda(), output_size_,
one_vec, 1,
&cuda_->zero,
grad_biases_->cuda(), 1);
// dw = x * (dy)^T
cublasSgemm(cuda_->cublas(),
CUBLAS_OP_N, CUBLAS_OP_T,
input_size_, output_size_, batch_size_,
&cuda_->one,
input_->cuda(), input_size_,
grad_output_->cuda(), output_size_,
&cuda_->zero,
grad_weights_->cuda(), input_size_);
// dx = W * dy
if (!gradients_stop_)
cublasSgemm(cuda_->cublas(),
CUBLAS_OP_N, CUBLAS_OP_N,
input_size_, batch_size_, output_size_,
&cuda_->one,
weights_->cuda(), input_size_,
grad_output_->cuda(), output_size_,
&cuda_->zero,
grad_input_->cuda(), input_size_);
return grad_input_;
}如果该层是模型中的第一层，我们也可以跳过计算输入数据的梯度，因为我们不必对它做任何事情。
权重和偏差更新将在我们想要更新权重时进行。在本节中，我们将为此使用随机梯度下降 ( SGD )。该操作也可以用于其他层。在这里，我们将这个函数放在Layer类中。重量更新也可以通过cublas功能完成，如下所示:
void Layer::update_weights_biases(float learning_rate)
{
float eps = -1.f * learning_rate;
if (weights_ != nullptr && grad_weights_ != nullptr) {
// w = w + eps * dw
cublasSaxpy(cuda_->cublas(),
weights_->len(),
&eps,
grad_weights_->cuda(), 1,
weights_->cuda(), 1);
}
if (biases_ != nullptr && grad_biases_ != nullptr)
{
// b = b + eps * db
cublasSaxpy(cuda_->cublas(),
biases_->b(),
&eps,
grad_biases_->cuda(), 1,
biases_->cuda(), 1);
}
}如您所见，我们可以用学习速率更新权重和偏差。当然，您也可以更改eps操作来应用其他优化算法。
在 C/C++ 编程中，程序员应该讲述如何在终止类实例时返回已使用的资源。按照我们的设计，如果图层有权重参数，并且可以从渐变中更新，那么图层最多会创建六个斑点。下面的代码显示了层终止代码，它终止内部创建的斑点:
Layer::~Layer()
{
if (output_ != nullptr) delete output_;
if (grad_input_ != nullptr) delete grad_input_;
if (weights_ != nullptr) delete weights_;
if (biases_ != nullptr) delete biases_;
if (grad_weights_ != nullptr) delete grad_weights_;
if (grad_biases_ != nullptr) delete grad_biases_;
}输入斑点或张量描述符将由其他层或斑点终端处理。图层类是其他图层的基类。因此，我们可以专注于终止自定义创建的资源，因为当我们终止任何派生的层时，这个终止代码将被一起调用。
即使我们已经构建了网络和层，我们也应该开发一些额外的层来完成网络。例如，我们没有实现激活、softmax 和损失计算层。我们将在接下来的章节中介绍这些层。
在神经网络层中有许多元素操作。激活功能是这些操作之一。cuDNN 库提供了六个激活函数:sigmoid、ReLU、tanh、clipped ReLU、eLU 和 identity。在 cuDNN 库中，cudnnActivationForward()做正向运算，cudnnActivationBackward()做反向运算。
我们来看看cuddnnActivationForward()函数的界面，如下图:
cudnnStatus_t cudnnActivationForward( cudnnHandle_t handle,
cudnnActivationDescriptor_t activationDesc,
const void *alpha, const cudnnTensorDescriptor_t xDesc,
const void *x, const void *beta,
const cudnnTensorDescriptor_t yDesc, void *y)使用cudnnActivationDescriptor_t，我们可以确定激活功能的类型。Alpha 和 beta 是标量值，决定了要添加的输入速率。xDesc和yDesc保存张量的形状信息。可以使用cudnnCreateTensorDescriptor()创建它们。
当你查看cudnnActivationBackward()函数时，dy是下一层的渐变输入，dx是上一层的渐变输出。在这种情况下，y成为输入。以这种方式，dyDesc提供梯度输入形状信息，而dxDesc提供梯度输出形状信息:
cudnnStatus_t cudnnActivationBackward( cudnnHandle_t handle,
cudnnActivationDescriptor_t activationDesc,
const void *alpha, const cudnnTensorDescriptor_t yDesc,
const void *y,
const cudnnTensorDescriptor_t dyDesc, const void *dy,
const cudnnTensorDescriptor_t xDesc, const void *x,
const void *beta, const cudnnTensorDescriptor_t dxDesc, void *dx)一般来说，我们可以预期层与层之间的张量形状不会改变。因此，我们可以对x和dx使用相同的张量描述符。和使用y和dy一样。
现在，让我们使用 cuDNN 应用编程接口实现启用 cuDNN 的激活功能。为了使用 cuDNN API，我们需要提供一个张量描述符来指定 cudn 函数的输入和输出张量维度。我们还需要指定激活操作。
虽然我们的示例实现没有使用层接口，但是我们需要将示例集成到层接口中。在我们的层设计中，激活层可以这样实现:
class Activation: public Layer
{
public:
Activation(std::string name, cudnnActivationMode_t mode,
float coef = 0.f);
~Activation();
Blob<float> *forward(Blob<float> *input);
Blob<float> *backward(Blob<float> *grad_input);
private:
cudnnActivationDescriptor_t act_desc_;
cudnnActivationMode_t mode_;
float coef_;
};在初始化步骤，我们需要创建几个张量描述符和一个激活描述符。cuDNN 库要求开发人员提供张量大小或对应于 API 的任何其他操作句柄:
Activation::Activation(std::string name, cudnnActivationMode_t mode, float coef)
{
name_ = name;
mode_ = mode;
coef_ = coef;
cudnnCreateActivationDescriptor(&act_desc_);
cudnnSetActivationDescriptor(act_desc_, mode, CUDNN_PROPAGATE_NAN, coef);
}在 cuDNN 中，我们使用激活描述符来指定激活函数操作。我们通过cudnnSetActivationDescriptor()功能来实现。然后，可以确定cudnnActivationForward/Backward()功能的运行。我们将在下一节讨论这个问题。然而，在此之前，我们需要实现类析构函数，以便它销毁激活描述符，如下所示:
cudnnDestroyActivationDescriptor(activation_desc);现在，让我们来介绍一下激活层的向前和向后操作。
这也称为谨慎操作。这一层不需要我们处理权重和偏差，因此它比密集层更容易实现。
在第一次迭代中，我们需要初始化输入描述符、输出描述符和输出 blob。当批处理大小改变时，我们将更新输出 blob。然而，我们不必初始化权重和偏差，因为它没有这些。下面的代码展示了它的实现:
if (input_ == nullptr || batch_size_ != input->n())
{
input_ = input;
input_desc_ = input->tensor();
batch_size_ = input->n();
if (output_ == nullptr)
output_ = new Blob<float>(input->shape());
else
output_->reset(input->shape());
output_desc_ = output_->tensor();
}初始化后，我们使用 cuDNN 中的cudnnActivationForward()函数进行激活过程，如下所示:
cudnnActivationForward(cudnnHandle, act_desc_,
&one, input_desc_, d_input, &zero, output_desc_, d_output);这个激活函数的操作是在我们初始化这个层时决定的，正如我们前面讨论的。
下一步是实现反向传播。我们将重用已经有的输入/输出张量描述符。现在，我们必须初始化我们想要反向传播的梯度:
if (grad_input_ != grad_output_)
{
grad_output_ = grad_output;
grad_input_ = new Blob<float>(input_->shape());
grad_input_->reset(input_->shape());
}初始化后，我们可以调用cudnnActivationBackward()函数，如下所示:
cudnnActivationBackward(cudnnHandle, activation_desc,
&one, output_desc_, output_->cuda(), output_desc_,
d_grad_output, input_desc_, input_->cuda(),
&zero, input_desc_, grad_input_->cuda());请注意，我们重用了在正向传递中创建的输入张量描述符和输出张量描述符。我们可以这样做，因为激活操作不会改变张量的大小。我们可以通过使用 cuDNN 应用编程接口激活反向传播来简化我们的实现。
cudnnActivationBackward()功能的输出为d_grad_input。正如我们在上一节中所描述的，这个渐变将被传递到下层。
现在，我们将实现 softmax 层，并将我们的层实现集成为一个网络。然后，我们将讨论全连通层在图像分类任务中的准确性。
对于 MNIST 数据集分类，我们将使用 softmax 分类器。softmax 函数对输入进行归一化，并生成
概率的概率分布。softmax 操作可表示如下:
cuDNN 的 softmax 转发功能支持该操作，以及通道和所有实例。之前，我们将密集层的输出与通道对齐。因此，我们将对通道应用 softmax 操作。
为了确认我们的训练是否有效，我们需要计算损失函数。软最大损失函数称为交叉熵损失，因为它的损失函数用于获得跨
概率的损失。损失函数如下:
我们需要获得这个软最大损失的梯度来更新神经网络。幸运的是，推导后 softmax 损失的梯度很简单，如下所示:
对于正向操作，我们将使用 cuDNN 函数来获取 softmax 的输出。要获得渐变，自定义操作更加直观和简单。
现在，让我们看看如何使用 cudn 和 CUDA 代码实现 softmax 层。
我们可以从 cuDNN 库中使用cudnnSoftmaxForward()获得软最大成本函数的输出:
cudnnSoftmaxForward(cudnnHandle, CUDNN_SOFTMAX_ACCURATE,
CUDNN_SOFTMAX_MODE_CHANNEL,
&one, input_desc, d_input, &zero, output_desc, d_output);这种情况下最重要的参数设置之一是CUDNN_SOFTMAX_MODE_CHANNEL。该选项根据输入张量描述符信息启用通道级 softmax 操作。通过这样做，我们可以提供通过密集层的小批量输入的通道对齐的张量。
softmax 层中的反向传递不同于其他层实现。该操作将输入数据的标签作为输入，并获得适当的梯度。如前所述，软最大损耗的梯度可以通过以下公式获得:
我们可以使用cublasSaxpy()来实现这个操作，如下所示:
// set grad_input_ as predict
cudaMemcpyAsync(grad_input_->cuda(), output_->cuda(),
output_->buf_size(), cudaMemcpyDeviceToDevice));
// set grad_input_ = predict - target
cublasSaxpy(cuda_->cublas(), target->len(), &cuda_->minus_one,
target->cuda(), 1, grad_input_->cuda(), 1));在前面的代码中，目标斑点包含一个热编码的目标向量，因此将负目标向量添加到预测值会产生适当的梯度。之后，我们需要在传播到上一层之前对批处理梯度进行标准化，如下所示:
int grad_output_size = target->n() * target->c() * target->h() * target->w();
float scale = 1.0f / static_cast<float>(target->n());
cublasSscal(cuda_->cublas(), grad_output_size, &scale, grad_input_->cuda(), 1);由于这引入了加权和的平均值，我们可以预期每个批次的梯度被归一化。
计算 softmax 的损失值是可选的。这意味着它的价值在训练和推理中没有被考虑。但是，我们可以将此作为培训的指标。
正如我们之前讨论的，softmax 损耗函数应实现以下等式:
我们可以从每个样本的输出中获得损失，并使用核函数累积它们，如下所示:
__global__ void
softmax_loss_kernel(float *reduced_loss, float *predict,
float *target, int size)
{
int batch_idx = blockDim.x * blockIdx.x + threadIdx.x;
extern __shared__ float s_data[];
float loss = 0.f;
// each thread calculate entropy for each data
// and accumulate to shared memory
if (batch_idx > 0)
return;
for (int c = 0; c < num_outputs; c++)
loss += target[batch_idx * num_outputs + c] * \
logf(predict[batch_idx * num_outputs + c]);
workspace[batch_idx] = -loss;
// Then, we do reduction the result to calculate loss
// Using 1 thread block
if (blockIdx.x > 0) return;
// Cumulate workspace data
s_data[threadIdx.x] = 0.f;
for (int i = 0; i < batch_size; i += blockDim.x)
s_data[threadIdx.x] += workspace[threadIdx.x + i];
__syncthreads();
// Reduction
for (unsigned int stride = blockDim.x / 2; stride > 0; stride >>= 1)
{
if (threadIdx.x + stride < batch_size)
s_data[threadIdx.x] += s_data[threadIdx.x + stride];
__syncthreads();
}
if (threadIdx.x == 0)
reduced_loss[blockIdx.x] = s_data[0];
}该操作使用并行约简，我们在第 3 章、 CUDA 线程编程中介绍过，以批量方式获取累积损失值。由于我们将只使用这个减少的损失值来确认训练，我们将简单地监控它的输出，而不是取它的平均值。
现在，让我们用 MNIST 数据集加载器集成我们已经实现的所有层。
整个过程的一个重要部分是为特定数据集提供数据加载器。在本实验中，我们将使用包含 60，000 个样本的 MNIST 数据集。当初始化时，我们告诉数据加载器它应该加载火车还是测试集。之后，数据加载器将在数据集中加载一些神奇的数字，以及所有样本及其标签。加载的数据将存储在向量中，并用相同的随机种子进行混洗。由于数据加载器构建并混洗样本向量，训练循环或测试循环可以为每次迭代获得随机化的输入数据。完全实现的代码可以在本书 GitHub 存储库中的src/mnist.cpp文件中找到。
当我们有多个层时，我们需要一个可以用神经网络操作管理这些层的对象，即前向/后向传播和权重更新。在本实验中，我们将有一个层数组，并迭代该数组进行正向处理。例如，可以使用以下代码执行正向操作:
Blob<float> *Network::forward(Blob<float> *input) {
output_ = input;
for (auto layer : layers_)
output_ = layer->forward(output_);
return output_;
}反向传播也可以通过以相反的顺序迭代数组来完成:
void Network::backward(Blob<float> *target) {
Blob<float> *gradient = target;
// back propagation.. update weights internally.....
for (auto layer = layers_.rbegin(); layer != layers_.rend(); layer++) {
// getting back propagation status with gradient size
gradient = (*layer)->backward(gradient);
}
}如您所见，我们管理矢量中的图层，并拥有每个图层的操作。向网络中添加一个新层更简单，如下面的代码所示:
void Network::add_layer(Layer *layer) {
layers_.push_back(layer);
}通过使用Network类，我们可以使用各种模型管理功能，比如参数更新、图层注册、图层初始化等等。此外，我们可以构建一个类似现代深度学习框架的神经网络。例如，我们可以创建如下模型:
// step 1\. loading dataset
MNIST data_loader = MNIST("./dataset");
// create training dataset loader and shuffling the data
data_loader.train(batch_size, true);
// step 2\. model initialization
Network model;
model.add_layer(new Dense("dense1", 500)); // 1st layer
model.add_layer(new Dense("dense2", 10)); // 2nd layer
model.cuda(); // set cuda context for each layer我们还可以进行以下训练循环:
// get data sample's shared buffer
Blob<float> *train_data = data_loader.get_data();
// get target's shared buffer
Blob<float> *train_target = data_loader.get_target();
// load data and targets with the batch size
data_loader.get_batch();
tp_count = 0; step = 0;
while (step < num_steps)
{
// transfer loaded data to the GPU
train_data->to(cuda);
train_target->to(cuda);
model.forward(train_data); // forward
model.backward(train_target); // backward
learning_rate *= 1.f / (1.f + lr_decay * step);
model.update(learning_rate); // update
step = data_loader.next(true); // load next data
... monitoring logic ...
}对于测试阶段，我们为测试数据集创建另一个数据集加载器，并且只使用正向传递进行迭代。下面的代码显示了它的实现:
test_data_loader.test(batch_size_test); // create test dataset loader
Blob<float> *test_data = test_data_loader.get_data(); // get sample data shared buffer
Blob<float> *test_target = test_data_loader.get_target(); // get target shared buffer
test_data_loader.get_batch(); // load samples and targets with the batch size
tp_count = 0; step = 0;
while (step < num_steps_test) {
// transfer loaded data to the GPU
test_data->to(cuda);
test_target->to(cuda);
model.forward(test_data); // forward
tp_count += model.get_accuracy(test_target);
step = test_data_loader.next(); // load next data
}
float accuracy = 100.f * tp_count / num_steps_test / batch_size_test;在测试阶段，我们将在完成测试数据集中所有样本的测试后获得准确性。现在，我们需要在测试循环后获得精度。
现在，让我们运行我们实现的代码，看看它的结果。对于训练阶段，我们将迭代 2400 步，批量为 256。MNIST 数据集在训练集中有 60，000 个样本。2400 步意味着我们将经历大约 10 个时代的迭代。可以使用以下命令编译示例代码:
$ nvcc -run -m64 -std=c++ 11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lcublas -lcudnn -lnvToolsExt -o train ./train.cpp ./src/layer.cu ./src/loss.cu ./src/mnist.cpp ./src/network.cpp下面的截图显示了我们实施的培训和测试输出:
在训练迭代中，网络从训练数据集获得了 92%的准确率。然而，测试准确率只有 77%，相对于训练结果来说，这是一个相对较低的分数。推理在训练和推理之间显示出很大的准确性差距的原因有很多。一个可能的原因是，完全连接的层没有考虑前面截图中显示的区域信息。在深度学习中，我们使用卷积层使网络学习空间信息。
现在，让我们用 cuDNN 实现卷积层，将其添加到网络中，并比较模型的性能。
cuDNN 库为卷积运算提供了优化的性能。通过创建一个卷积层，我们将覆盖应用编程接口的前向和后向操作配置。
卷积网络层使用其权重对输入数据执行卷积。当你想建立一个能感知区域信息的神经网络时，这种网络架构非常有用。回想一下第 7 章、CUDA中的并行编程模式中的卷积实现，它需要相当大的内存带宽，需要进一步优化才能获得最佳性能。然而，使用 cuDNN 库，我们也可以获得最佳性能，因为我们不必重新发明轮子。
卷积层的实现类似于全连接层的实现。然而，有两个不同之处，这要归功于 cuDNN 库:我们不必像以前那样完全实现那么多细节，我们需要为操作分配一个工作空间大小。对于每一个卷积操作——正向、滤波器反向和输入反向——都需要额外的存储空间，这取决于它们的算法。该算法可以根据给定的输入/输出/滤波器张量维数而变化。详细的 API 调用将在后面处理。
像其他层一样，它有三个工作阶段。对于推理阶段，我们称之为cudnnConvolutionForward()和cudnnAddTensor()。对于落后阶段，我们称之为cudnnConvolutionBackwardData()、cudnnConvolutionBackwardFilter()、cudnnConvolutionBackwardBias()。最后，在更新阶段，我们可以重用完全连接的层中的代码。该层的配置概述如下:
在深度学习神经网络中，通常使用汇集层和卷积网络。池化层只是按照一个简单的规则选择要输出的输入数据。下图显示了最大池化的示例:
使用 cuDNN 库，我们将实现这两个卷积运算。
像一个完全连接的层，这个卷积层有权重和偏差参数。在完全连接层，我们使用了 cuBLAS，它不需要 cuDNN 相关的描述符。然而，我们将使用 cuDNN 卷积函数，因此我们需要使用滤波器描述符和卷积运算描述符。下面的代码显示了在构建层时我们应该初始化哪些资源:
Conv2D::Conv2D(std::string name,
int out_channels, kernel_size, stride, padding, dilation):
out_channels_(out_channels), kernel_size_(kernel_size),
stride_(stride), padding_(padding), dilation_(dilation) {
name_ = name;
cudnnCreateFilterDescriptor(&filter_desc_);
cudnnCreateConvolutionDescriptor(&conv_desc_);
cudnnSetConvolution2dDescriptor(conv_desc_,
padding_, padding_, stride_, stride_, dilation_,dilation_,
CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT);
}由于我们在构建模型时提供卷积运算信息，因此我们可以指定卷积描述符。然而，滤波器的操作可以在推断时指定，因为我们可以在那时学习输入张量的大小。现在，让我们在卷积层实现正向传递。
如前所述，我们可以用输入张量大小初始化卷积层。这个输入张量的大小对输出张量的大小有影响。下面的代码显示了正向传递中的参数初始化步骤:
// initialize weights and bias
if (weights_ == nullptr) {
// initialize containers handles
cudnnSetFilter4dDescriptor(filter_desc_,
CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW,
out_channels_, input->c(), kernel_size_, kernel_size_);
weights_ = new Blob<float>(out_channels_, input->c(), kernel_size_, kernel_size_);
biases_ = new Blob<float>(1, out_channels_); // bias size
bias_desc_ = biases_->tensor();
}然后，我们需要更新输入资源，初始化输出 blob，创建 cuDNN 工作空间，并初始化权重参数，如下所示:
// initilaize input and output
if (input_ == nullptr || batch_size_ != input->n()) {
// initialize input
input_ = input;
input_desc_ = input->tensor();
batch_size_ = input->n();
// getting output tensor size
cudnnGetConvolution2dForwardOutputDim(
conv_desc_, input_desc_, filter_desc_,
&output_size_[0], &output_size_[1],
&output_size_[2], &output_size_[3]);
// initialize output blob
if (output_ == nullptr)
output_ = new Blob<float>(output_size_);
else
output_->reset(output_size_);
output_desc_ = output_->tensor();
// initialize weights
if (!freeze_)
init_weight_bias();
// initialize workspace for cudnn
set_workspace();
}为了获得输出张量大小，我们使用cudnnGetConvolution2dForwardOutputDim()函数。该函数基于输入张量大小、卷积运算和滤波器大小输出尺寸信息。然后，我们重用在完全连接层中使用的相同参数初始化代码。
要调用 cuDNN 的卷积 API，需要提供它的工作算法和工作空间内存。我们这样做是因为 cuDNN 根据卷积大小选择最优卷积算法，它的测量需要立即进行。算法确定后，cuDNN 可以确定工作空间大小。卷积层需要对前向通道、输入数据梯度和权重梯度进行卷积运算。我们需要单独处理每个算法，但是我们只能分配一个工作空间，因为该工作空间专门用于每个卷积操作。
因此，我们在所需的每个卷积算法工作空间大小中创建具有最大大小的工作空间。下面的代码展示了我们如何使用它们和管理工作空间:
Conv2d::set_workspace() {
size_t temp_size = 0;
// fwd
cudnnGetConvolutionForwardAlgorithm(cuda_->cudnn(),
input_desc_, filter_desc_, conv_desc_, output_desc_,
CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &conv_fwd_algo_);
cudnnGetConvolutionForwardWorkspaceSize(cuda_->cudnn(),
input_desc_, filter_desc_, conv_desc_, output_desc_,
conv_fwd_algo_, &temp_size);
workspace_size = std::max(workspace_size, temp_size);
// bwd - data
cudnnGetConvolutionBackwardDataAlgorithm(cuda_->cudnn(),
filter_desc_, output_desc_, conv_desc_, input_desc_,
CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST, 0,
&conv_bwd_data_algo_);
cudnnGetConvolutionBackwardDataWorkspaceSize(cuda_->cudnn(),
filter_desc_, output_desc_, conv_desc_, input_desc_,
conv_bwd_data_algo_, &temp_size);
workspace_size = std::max(workspace_size, temp_size);
// bwd - filter
cudnnGetConvolutionBackwardFilterAlgorithm(cuda_->cudnn(),
input_desc_, output_desc_, conv_desc_, filter_desc_,
CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST, 0,
&conv_bwd_filter_algo_);
cudnnGetConvolutionBackwardFilterWorkspaceSize(cuda_->cudnn(),
input_desc_, output_desc_, conv_desc_, filter_desc_,
conv_bwd_filter_algo_, &temp_size);
workspace_size = std::max(workspace_size, temp_size);
if (workspace_size > 0) {
if (d_workspace != nullptr)
cudaFree(d_workspace);
cudaMalloc((void**)&d_workspace, workspace_size);
}
}每个卷积算法都用单独的类型指定，即cudnnConvolutionFwdAlgo_t、cudnnConvolutionBwdDataAlgo_t和cudnnConvolutionBwdFilterAlgo_t。我们可以通过将它们声明为类成员变量来使用它们，即conv_fwd_algo_、conv_bwd_data_algo_和conv_bwd_filter_algo_。
现在，我们编写初始化后的前向处理代码。我们与滤波器进行卷积，并添加一个偏差。以下代码显示了 cuDNN 卷积正向实现:
cudnnConvolutionForward(cuda_->cudnn(), &cuda_->one, input_desc_, input_->cuda(), \
filter_desc_, weights_->cuda(), conv_desc_, conv_fwd_algo_, d_workspace, workspace_size, \
&cuda_->zero, output_desc_, output_->cuda());
cudnnAddTensor(cuda_->cudnn(), &cuda_->one, bias_desc_, biases_->cuda(), \
&cuda_->one, output_desc_, output_->cuda());卷积的结果将使用输出斑点传递到下一层。
在反向传播中，我们应该计算偏差的梯度、权重的梯度和输入数据的梯度。为此，我们需要在第一次迭代时创建 blobs，以便存储它们。它们的大小不取决于批次大小，所以我们只需要确保它们被创建。初始化步骤可以如下实现:
// initialize grad_output back-propagation space
if (grad_weights_ == nullptr) {
grad_output_ = grad_output;
grad_weights_ = new Blob<float>(weights_->shape());
grad_biases_ = new Blob<float>(1, biases_->c());
grad_input_ = new Blob<float>(input_->shape());
}然后，我们称之为 cuDNN 向后卷积 API，如下所示:
Blob<float> *Conv2D::backward(Blob<float> *grad_output) {
... { initialization step } ...
// gradients of biases
cudnnConvolutionBackwardBias(cuda_->cudnn(),
&cuda_->one,
output_desc_, grad_output->cuda(),
&cuda_->zero,
bias_desc_, grad_biases_->cuda());
// gradients of weights
cudnnConvolutionBackwardFilter(cuda_->cudnn(),
&cuda_->one,
input_desc_, input_->cuda(),
output_desc_, grad_output_->cuda(),
conv_desc_, conv_bwd_filter_algo_, d_workspace, workspace_size,
&cuda_->zero,
filter_desc_, grad_weights_->cuda());
// gradients of input data
if (!gradient_stop_)
cudnnConvolutionBackwardData(cuda_->cudnn(),
&cuda_->one,
filter_desc_, weights_->cuda(),
output_desc_, grad_output->cuda(),
conv_desc_, conv_bwd_data_algo_, d_workspace, workspace_size,
&cuda_->zero,
input_desc_, grad_input_->cuda());然后，我们将输入数据的梯度传递到前一层，以传播梯度。我们将在更新步骤中使用基类的梯度更新代码来更新权重和偏差的梯度。当我们在完全连接的层中实现后向传播时，我们讨论了这一点。如果这是第一层，我们也可以跳过计算输入数据的梯度。
池层有两个特征。首先，与卷积层相比，它的输出大小不同，cuDNN 为此提供了相应的 API。第二，它没有任何内部重量。
要指定池操作，我们可以使用 cuDNN 的cudnnPoolingDescriptor_t函数，在类构造函数中创建并指定 cuDNN 的池描述符，如下所示:
cudnnCreatePoolingDescriptor(&pool_desc_);
cudnnSetPooling2dDescriptor(pool_desc_, mode_, CUDNN_PROPAGATE_NAN,
kernel_size_, kernel_size_, padding_, padding_, stride_, stride_);现在，让我们实现池层的前向和后向操作。
汇集层有助于减小张量的大小。因此，我们需要计算输出大小。我们可以使用cudnnGetPooling2dForwardOutputDim()函数计算大小，就像我们在卷积层实现中所做的那样。此外，张量大小取决于批次大小。这意味着如果批量改变，我们需要更新张量大小。下面的代码展示了如何初始化输入和输出斑点:
if (input_ == nullptr || batch_size_ != input->n()) {
input_ = input;
// resource initialize
input_desc_ = input_->tensor();
batch_size_ = input->n();
// setting output
cudnnGetPooling2dForwardOutputDim(pool_desc_, input_desc_,
&output_size_[0], &output_size_[1], &output_size_[2],
&output_size_[3]);
if (output_ == nullptr)
output_ = new Blob<float>(output_size_);
else
output_->reset(output_size_);
output_desc_ = output_->tensor();
}对于正向传递，我们调用cudnnPoolingForward()函数，如下所示:
Blob<float> *Pooling::forward(Blob<float> *input) {
... { initialization step } ...
cudnnPoolingForward(cudnnHandle, pool_desc_, &one,
input_desc_, input_->cuda(),
&zero, output_desc_, output_->cuda());
}对于反向传播步骤，我们调用cudnnPoolingBackward()函数，如下所示:
Blob<float> *Pooling::backward(Blob<float> *grad_output) {
if (grad_input_ == nullptr)
grad_input_ = new Blob<float>(input_->shape());
cudnnPoolingBackward(cudnnHandle, pool_desc_,
&one, output_desc_, output_->cuda(),
output_desc_, grad_output->cuda(),
input_desc_, input_->cuda(),
&zero, input_desc_, grad_input_->cuda());
}汇集层的输入张量形状和输入梯度相同，输出形状和输出梯度相同。因此，我们可以分别重用输入和输出的张量描述符。
现在，让我们将这些集成到单个卷积层实现中。
现在，我们将更新我们以前的网络，LeNet。网络代码可以编写如下:
Network model;
model.add_layer(new Conv2D("conv1", 20, 5));
model.add_layer(new Pooling("pool", 2, 0, 2, CUDNN_POOLING_MAX));
model.add_layer(new Conv2D("conv2", 50, 5));
model.add_layer(new Pooling("pool", 2, 0, 2, CUDNN_POOLING_MAX));
model.add_layer(new Dense("dense1", 500));
model.add_layer(new Activation("relu", CUDNN_ACTIVATION_RELU));
model.add_layer(new Dense("dense2", 10));
model.add_layer(new Softmax("softmax"));
model.cuda();现在，我们可以开始训练和推理阶段，因为我们已经配置了我们的层，使它们相互连接。让我们用下面的命令编译代码:
$ nvcc -run -m64 -std=c++ 11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lcublas -lcudnn -lnvToolsExt -o train ./train.cpp ./src/layer.cu ./src/loss.cu ./src/mnist.cpp ./src/network.cpp然后，我们可以看到如下的训练和测试结果:
如您所见，与仅使用完全连接的网络相比，该网络实现了更高的训练精度和推理。我们还可以通过查看 NVIDIA 配置文件来确认其操作，如下所示:
最新的 NVIDIA GPUs 支持混合精密运算深度学习。我们不会在这本书里讨论这个问题，因为它超出了我们的范围。但是，如果您想了解更多信息，可以访问英伟达在/usr/src/cudnn_samples_v7/conv_sample提供的示例。要访问这个示例，您需要从 cuDNN 网页下载示例。这个示例代码展示了如何使用 cuDNN 库使用混合精度操作。
为了让 cuDNN APIs 与张量核一起工作，我们需要设置数学类型，如下所示:
cudnnSetConvolutionMathType(cudnnConvDesc, CUDNN_TENSOR_OP_MATH);然后，我们需要使用cudnnSetTensorNdDescriptor()初始化输入和输出张量的张量描述符。这为张量提供了填充，以便我们获得优化的张量核心性能。
一个很好的基于 cuDNN 的实现是cudnn-training:https://github.com/tbennun/cudnn-training。它将 LeNet 实现为一系列 cuDNN 函数。您可以沿着每一行来查看 CUDNN 函数是如何工作的。
如果您有兴趣使用 cuDNN 部署您的网络，请查看以下关于 GTC-CNN 使用 cuDNN 进行推理的视频(https://developer.nvidia.com/gtc/2019/video/S9644/video)。这篇演讲介绍了使用 cuDNN 进行 CNN 推理的有用的性能优化技巧。
在深度学习训练中使用半精度需要的不仅仅是 FP16 操作的利用率。我们需要在 FP16 中计算张量，同时在 FP32 中保持权重。此外，一些操作需要 FP32。我们称之为混合精度。cuDNN 库提供了一个名为 mnistCUDNN 的混合精度推理示例。此示例显示了输入和图层数据类型的转换。如果你想了解更多深度学习和训练中的混合精度运算，请阅读以下文章:https://devblogs . NVIDIA . com/video-mixed-precision-technologies-tensor-cores-deep-learning/。
现在，我们将从性能方面介绍深度学习中的其他 GPU 使用注意事项。
rrn 允许您在深度学习中分析顺序数据。虽然这个网络有顺序依赖性，但仍有很大的优化空间。在这一节中，我们将介绍它的算法以及 cuDNN 如何提供优化的性能。
RNN 有很多种，但是 cuDNN 只支持四种，分别是 RNN 带 ReLU，RNN 带 tanh，LSTM，GRU，它们有两个输入:来自前一个网络的隐藏参数和来自源的输入。根据它们的类型，它们有不同的操作。在本实验中，我们将讲述 LSTM 行动。下图显示了 LSTM 的正向操作:
从计算的角度来看，有八个矩阵矩阵乘法和许多元素操作。根据这个估计，我们可以预期 LSTM 可能是内存受限的，因为每个操作都是内存受限的。另一方面，CUDNN 提供了cudnnRNNForwardInference()和cudnnRNNFowardTraining() RNN 功能。我们将通过测量这个函数的性能和模拟 LSTM 来介绍使用这个函数的好处。为此，我们将实现一个虚拟 LSTM 层，并将其性能与 cudn LSTM 函数进行比较。
出于测试目的，我们将像这样设置超参数:
int mode = 2; // LSTM in CUDNN
int seq_length = 512;
int num_layers = 4;
int hidden_size = 512;
int input_size = hidden_size;
int batch_size = 32;
float dropout_rate = 0;
bool bidirectional = 0;
int persistent = 0;序列长度或隐藏大小可能会有所不同，具体取决于问题。在这个测试中，我们将使用512作为长度，这在序列研究中被大量使用。CUDNN API 需要更多的选项才能工作，例如辍学率、双向或单向以及持久 rnn。我们将只在这部分测试香草 LSTM。
让我们编写一些代码，作为 LSTM 层执行cudnnRNNForwardTraining()函数:
- 我们需要初始化输入和输出内存空间。要执行 cuDNN 的 RNN 应用编程接口，我们需要使用以下变量:
// hx, cx, hy, cy, dhy, dcy, dhx, and dcs can be null.
void *x; // input
void *hx = nullptr; // input of initial hidden state
void *cx = nullptr; // input of cell state (LSTM)
void *y; // output
void *hy = nullptr; // output of final hidden state
void *cy = nullptr; // output of final cell state (LSTM)
void *dy; // input of gradient
void *dhy = nullptr; // input of final hidden state
void *dcy = nullptr; // input of final cell state (LSTM)
void *dx; // output of gradient at the input of rnn
void *dhx = nullptr; // output of gradient at the initial hidden state
void *dcx = nullptr; // output of gradient at the initial cell state这些变量是 LSTM 的输入和输出。为了提供输入和获得输出，我们需要分配适当的内存空间。遵循 LSTM 定义，我们需要考虑输入、输出和隐藏层的长度。这些尺寸可以确定如下:
int input_length = seq_length * input_size * batch_size;
int output_length = seq_length * hidden_size * batch_size;
int hidden_length = hidden_size * batch_size * num_layers;然后，我们可以为每个项目分配内存。
- 现在，我们需要为库登 RNN 应用编程接口设置张量描述符。下面的代码显示了我们应该设置的所需张量描述符:
cudnnTensorDescriptor_t x_desc[seq_length], y_desc[seq_length], \
dx_desc[seq_length], dy_desc[seq_length];
cudnnTensorDescriptor_t hx_desc, cx_desc;
cudnnTensorDescriptor_t dhx_desc, dcx_desc;
cudnnTensorDescriptor_t hy_desc, cy_desc;
cudnnTensorDescriptor_t dhy_desc, dcy_desc;对于输入和输出描述符，我们需要初始化每个元素，即批处理大小及其输入大小。其他隐藏张量描述符用层数、批次大小和隐藏大小初始化。本节将不介绍如何编写初始化代码。但是，如果您想了解更多信息，可以查看10_deep_learning/03_rnn文件中的代码。
- 我们还必须为 RNN 操作提供一个工作空间，就像我们为卷积操作所做的那样:
void *workspace;
cudnnFilterDescriptor_t w_desc, dw_desc;
cudnnSetRNNDescriptor_v6(cudnnHandle, rnn_desc,
hidden_size, num_layers, dropout_desc, CUDNN_LINEAR_INPUT,
bidirectional ? CUDNN_BIDIRECTIONAL : CUDNN_UNIDIRECTIONAL,
CUDNN_LSTM, CUDNN_RNN_ALGO_STANDARD, CUDNN_DATA_FLOAT));
size_t weight_size;
cudnnGetRNNParamsSize(cudnnHandle, rnn_desc, x_desc[0], &weight_size, CUDNN_DATA_FLOAT);
cudaMalloc((void**)&workspace, weight_size);然后，我们可以根据工作空间的大小设置过滤器描述符，如下所示:
dimW = {weight_size / sizeof(float), 1, 1}
cudnnCreateFilterDescriptor(&w_desc);
cudnnCreateFilterDescriptor(&dw_desc);
cudnnSetFilterNdDescriptor(w_desc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, 3, dimW);
cudnnSetFilterNdDescriptor(dw_desc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, 3, dimW);
cudnnRNNForwardTraining(cudnnHandle, rnn_desc, seq_length,
x_desc, x, hx_desc, hx, cx_desc, cx,
w_desc, w,
y_desc, y, hy_desc, hy, cy_desc, cy,
workspace, workspace_size, reserved_space,
reserved_size);我们可以使用cudaEvnetRecoard()和 flops 计算来衡量它们的性能。例如，正向操作可以用以下等式配置:
然后，我们将通过将批量从 32 增加到 256 来测试我们的实现。适用的测试范围可能不同，GPU 的内存大小也可能不同。
在本节中，我们实现了基于 LSTM 的模拟和cudnnRNNForwardTraining()调用。我们的部分模拟版本只有 GEMM 操作，这是最计算密集型的。现在，让我们比较一下这些实现的性能。
在我们的实施中，我们将重点模拟 LSTM 的主要行动，而不是完全实施。
让我们确定 LSTM 网络的超参数。一般来说，输入序列长度范围从 512 到 2，048。层数不同。但是由于 tanh 操作，不能很大。对于输入大小，我们将使用 512。通常，就 RNN 使用率而言，批次大小在 32 到 256 之间。CUDNN 需要更多关于辍学率的输入，双向还是单向，以及我们是否在使用持久 RNN。我们只是现在不用它们。我们的 LSTM 配置信息如下:
现在，我们将有一个部分实现的 LSTM 操作来测量计算强度。正如我们前面讨论的，LSTM 有两个矩阵乘法，我们需要计算。LSTM 操作将为输入序列的每个元素以及每个层计算该值。然后，操作可以配置如下:
for (int layer = 0; layer < num_layers; layer++) {
for (int linear_layer = 0; linear_layer < 4; linear_layer++) {
for (int sequence = 0; sequence < seq_length; sequence++) {
cublasSgemm(cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N,
hidden_size, input_size, batch_size,
&alpha, input_weight, input_size, x, input_size,
&beta, h, hidden_size);
cublasSgemm(cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N,
hidden_size, hidden_size, batch_size,
&alpha, recurrent_weight, hidden_size,
h, hidden_size,
&beta, y, hidden_size);
}
}
}我们可以使用更多的元素操作，但它只是近似计算强度，所以我们现在将省略它们。
让我们比较它们的性能以及不同的批次大小，如下所示在main()函数中实现的代码:
for (int step = 1; step <= 8; step++)
{
batch_size = 32 * step;
printf("Batch Size: %3d\n", batch_size);
rnn_operation(seq_length, num_layers, hidden_size, input_size,
batch_size, dropout_rate, bidirectional, mode, persistent);
cublas_operation(mode, 2ull, input_size, hidden_size, seq_length, batch_size, num_layers);
}并且，我们可以使用以下命令编译并执行示例源代码:
$ nvcc -run -m64 -std=c++ 11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lcublas -lcudnn -lcurand -o rnn ./rnn.cpp下图显示了特斯拉 V100 卡上 cuBLAS 和 cuDNN 的实测性能:
在上图中，这两种实现在性能上表现出巨大的差异。cuDNN 的 LSTM 性能比使用 cuBLAS 的模拟 LSTM 好得多。此外，LSTM 操作的表现遵循特斯拉 V100 图形处理器的屋顶线。另一方面，两个 SGEMM 操作没有显示出这种性能，因为矩阵大小不足以获得完全的性能。要从特斯拉 V100 获得 10 个 TFlops，矩阵大小应该与 1，024 的平方相似或更大。然而，正如我们所看到的，我们的矩阵大小大约是 512 的平方。
LSTM optimization is explained in the following NVIDIA article: https://devblogs.nvidia.com/optimizing-recurrent-neural-networks-cudnn-5. It combines matrix-matrix multiplications, fusing element-wise operations, multiple streams, and multi-layer parallelization.
One of the optimization versions of the RNN is the persistent RNN (https://svail.github.io/persistent_rnns), which was introduced by Greg Diamos. Although his implementation does not include LSTM and GRU, you can learn how the RNN can be optimized.
一般来说，我们使用深度学习框架(如 TensorFlow、PyTorch 和 MxNet)来开发和研究神经网络。由于这些框架，我们可以有效地开发复杂的模型。然而，当涉及到性能工程时，由于剖析工具的能力，理解框架下的 GPU 操作是一条陡峭的学习曲线。例如，当模型简单时，使用 chrome tracing 进行概要分析是有用的，但是当模型复杂时就没用了。
在第 5 章、 CUDA 应用分析和调试中，我们介绍了 NVIDIA 工具扩展 ( NVTX )，它允许我们在 GPU 应用中进行自定义注释，并使用 NVIDIA Nsight Systems 查看时间线。对于复杂的应用，程序员分析它们的性能并找到瓶颈是很有用的。
在本节中，我们将通过修改 ResNet-50 示例代码来介绍如何在 PyTorch 和 TensorFlow 中使用 NVTX。示例代码可以在本书 GitHub 存储库中的10_deep_learining/05_framework_profile文件夹中找到。您可以从https://github.com/nvidia/DeepLearningExamples获得原始源代码。
为了进行轻松的工作环境配置，我们将为 PyTorch 和 TensorFlow 使用 NVIDIA GPU 云 ( NGC )深度学习容器。如果您需要了解 NGC 或集装箱的基本用法，请访问本书中的 NGC 附录。
现在，让我们先从 PyTorch 开始。
在 PyTorch 中，我们可以使用torch.cuda.nvtx.range_push("foo")和torch.cuda.nvtx.range_pop()放置自定义标签。这保持了原有的 CUDA NVTX APIs，即nvtxRangePush()和nvtxRangePop()。让我们看看 NVTX 注释如何帮助我们理解时间轴中的深度学习操作。在以下步骤中，我们将使用05_framework_profile/pytorch/RN50v1.5文件中的 ResNet-50 示例代码:
- 我们将在
train()函数的训练循环中放置 NVTX 注释来注释step值。该功能可以在image_classificaiton/training.py文件中找到。下面的屏幕截图分别显示了第 234 行和第 260 行的训练循环和 NVTX 注释:
在前面的代码中，训练操作在step函数中实现，该函数由get_train_step()函数定义。因此，我们需要在该函数中放置 NVTX 注释来了解更多信息。
- 让我们在第 164 行给
get_train_step()函数添加一些 NVTX 注释。该函数返回_step()函数，包括训练操作。因此，我们将在这个函数中放置 NVTX 注释。训练过程是前向和后向传播、全约简和优化(更新权重)。以下屏幕截图显示了第 166 行和第 171 行的前向传播注释:
这样，我们可以在剩余的操作上放置其他注释。
- 我们也可以有模型层的 NVTX 注释。在本例中，ResNet-50 模型在
image_classification/resnet.py文件中实现。以下屏幕截图显示了网络注释示例:
正如我们所看到的，我们可以按照 ResNet 架构放置 NVTX 注释。如果我们在每个构建块中放置注释，我们可以获得更多信息。
- 现在，让我们分析一下模型。如前所述，我们将使用名为 PyTorch 的 NGC 深度学习容器。
imagenet数据集位于/raid/datasets/imagenet/raw-data文件夹中。为了限制分析时间范围，我们将使用延迟选项(-y)和持续时间选项(-d)。下面的代码显示了一个 bash shell 脚本，它执行容器并分析网络:
#/bin/bash
CODE_PATH="RN50v1.5"
DATASET_PATH="/raid/datasets/imagenet/raw-data/"
OUTPUT_NAME="resnet50_pyt"
# default profile
docker run --rm -ti --runtime=nvidia \
-v $(pwd)/${CODE_PATH}:/workspace \
-v ${DATASET_PATH}:/imagenet \
nvcr.io/nvidia/pytorch:19.08-py3 \
nsys profile -t cuda,nvtx,cudnn,cublas -o ${OUTPUT_NAME}
-f true -w true -y 60 -d 20 \
python /workspace/main.py --arch resnet50 -b 64
--fp16 /imagenet执行后，前面的代码在 RN50v1.5 目录中生成分析结果，即resnet50_pyt.qdrep。
- 最后，使用 NVIDIA Nsight 系统打开分析输出
resnet50_pyt.qdrep，并查看操作。下面的截图显示了带有 NVTX 注释的测量步骤:
在这里，我们可以看到后向操作花费的时间是前向操作的两倍。此外，PyTorch 将训练循环和反向传播的宿主线程分开。查看内核概要分析，最耗时的点是按元素执行内核。让我们放大转发通道来查看层的执行时间，如下图所示:
在这里，我们可以看到第二个卷积块需要最长的时间来完成。如果这一层有低效点，我们可以进一步挖掘。如果某个特定的内核函数被确定为瓶颈，需要进行优化，我们也可以使用 NVIDIA Nsight Compute 对其进行分析。比较主机应用编程接口跟踪和图形处理器，我们可以看到持续时间是不同的。这是因为主机和 GPU 的操作是异步的。所以，我们在从主机测量 GPU 执行时间时需要谨慎。现在，让我们看看优化步骤，如下图所示:
我们可以看到，主机和 GPU 的测量执行时间存在巨大差异。主机的测量执行时间为 25.367 毫秒，而图形处理器的时间为 4.048 毫秒。其操作主要是元素操作，其执行被延迟到反向传播完成。我们也可以找到异步执行。之后可以看到cudaDeviceSynchronize()功能，防止当前步骤被下一步更新。
我们还可以通过设置一个环境，即CUDA_LAUNCH_BLOCKING=1，来禁用这些异步操作。我们可以使用环境选项(-e)将它传递给系统的配置文件选项。然后，我们可以用主机和内核函数分析应用的align操作。
PyTorch 在其 CUDA 对象中有几个 NVTX 特色的 API。PyTorch 文档可在https://py torch . org/docs/stable/_ modules/torch/cuda/nvtx . html找到。通过直接调用 PyTorch 中的 NVTX API，调用 CUDA NVTX APIs。这意味着我们可以在概要时间线中获得定制标记的 NVTX 标记。
分析张量流图需要我们有一个支持 NVTX 注释的 NVTX 插件。要在 TensorFlow 中使用 NVTX 注释，我们需要使用以下命令安装nvtx-plugins-tf Python 插件:
$ pip install nvtx-plugins-tf然而，如果我们使用的 NGC 张量流容器晚于 19.08 版本，我们就不必这样做了
TensorFlow 图 API 是符号 API，所以需要特定的编程方法。NVTX 插件为此提供了两个选项:装饰器和 Python 函数。
下面是一个 NVTX 装饰器的示例:
import nvtx.plugins.tf as nvtx_tf
ENABLE_NVTX=true
@nvtx_tf.ops.trace(message='Dense Block', domain_name='Forward',
grad_domain_name='Gradient', enabled=ENABLE_NVTX,
trainable=True)
def dense_layer(x):
x = tf.layers.dense(x, 1000, activation=tf.nn.relu, name='dense_1')
x = tf.layers.dense(x, 1000, activation=tf.nn.relu, name='dense_2’)
return x下面是一个 NVTX Python 函数的示例:
import nvtx.plugins.tf as nvtx_tf
ENABLE_NVTX=true
x, nvtx_context = nvtx_tf.ops.start(x, message='Dense Block', \
domain_name='Forward’, grad_domain_name='Gradient’,
enabled=ENABLE_NVTX, trainable=True)
x = tf.layers.dense(x, 1000, activation=tf.nn.relu, name='dense_1')
x = tf.layers.dense(x, 1000, activation=tf.nn.relu, name='dense_2’)
x = nvtx_tf.ops.end(x, nvtx_context)NVTX 插件提供了 NVTXHook，它允许我们分析 TF 估计器和会话。例如，我们可以如下使用钩子:
from nvtx.plugins.tf.estimator import NVTXHook
nvtx_callback = NVTXHook(skip_n_steps=1, name='Train’)
training_hooks=[]
training_hooks.append(nvtx_callback)然后，我们可以使用以下代码将此应用于任一选项:
with tf.train.MonitoredSession(hooks=training_hooks) as sess:或者，我们可以使用以下代码:
tf.estimator.Estimator(hooks=training_hooks, ...)现在，让我们将此应用于示例 ResNet-50 代码，并查看操作。示例代码可以在05_framework_profile/tensorflow/RN50v1.5文件夹中找到:
- 让我们从将
NVTXHook应用于估计器开始。训练图的定义可以在第 312 行的runtime/runner.py文件中找到。在构建图表之前，我们将把NVTXHook添加到钩子列表中，如下面的代码块所示:
- 然后，我们将把 NVTX 注释应用到模型构建函数中。
model_build()功能可以在model/resnet_v1_5.py文件的ResnetModel类中找到。下面的代码显示了一个使用 Python 函数在model_build()函数中的conv1层放置 NVTX 注释的示例:
在前面的代码中，当使用nvtx_tf.ops.start()和nvtx_tf.ops.end()功能时，我们需要谨慎使用正确的输入和输出。仅将 NVTX 注释放置在其他层中。确保最终完全连接的层的输出是网络的输出。
我们还必须禁用代码来检查它拥有的可训练变量的数量。如果 NVTX 的trainable参数值为True，则尺寸会发生变化。在resnet_v1_5.py文件的第 174 行，有一个断言代码块，用于检查该变量的编号。简单评论一下，如下:
- 我们还使用 NVTX 装饰器作为 ResNet 构建模块。在
model/blocks目录中，我们可以找到conv2d_blocks.py和resnet_bottleneck_block.py中的conv2d和 ResNet 瓶颈块实现。在conv2d_blocks.py文件中，我们可以修饰conv2d_block()函数来标注 NVTX 概要文件，如下所示:
同样，我们可以对resnet_bottleneck_block.py文件做同样的操作:
- 现在，让我们分析一下模型。就像我们使用 PyTorch 容器一样，我们将使用 TensorFlow 的 NGC 容器。我们将假设
imagenet数据集的tfrecord文件位于/raid/datasets/imagenet/tfrecord目录中。下面的代码显示了一个 bash shell 脚本，它执行容器并分析网络:
#/bin/bash
CODE_PATH="RN50v1.5"
DATASET_PATH="/raid/datasets/imagenet/tfrecord"
OUTPUT_NAME="resnet50_tf"
# default profile
docker run --rm -ti --runtime=nvidia \
-v $(pwd):/result \
-v $(pwd)/${CODE_PATH}:/workspace \
-v ${DATASET_PATH}:/imagenet \
nvcr.io/nvidia/tensorflow:19.08-py3 \
nsys profile -t cuda,nvtx,cudnn,cublas -o ${OUTPUT_NAME}
-f true -w true -y 40 -d 20 \
python /workspace/main.py --mode=training_benchmark
--warmup_steps 200 \
--num_iter 500 --iter_unit batch
--results_dir=results --batch_size 64当我们执行这个函数时，我们会在RN50v1.5目录中得到resnet50_tf.qdrep文件。
- 最后，让我们回顾一下使用英伟达系统分析的输出:
在这里，我们可以确认反向传播的时间是正向传播的两倍。这个示例代码没有与中央处理器和图形处理器同步。正因为如此，我们可以看到主机和 GPU 之间的时间差更大。当我们在构建块中放置附加注释时，我们将能够在层中看到子块注释。
使用 NVIDIA Nsight 系统进行性能分析在监控多图形处理器训练中的全减少执行时间方面提供了额外的好处。以下屏幕截图显示了使用两个图形处理器进行训练的图形处理器的分析结果:
在高亮显示的行中，我们可以看到ncclAllRecude()函数，它同时调用反向传播。这样做，我们就不会得到全归约运算的延迟。这个示例代码使用 Horovod 来训练多个 GPU。如果你想了解更多，请访问 Horovod 的 GitHub 页面:https://github.com/horovod/horovod。您可以从这里获得文档和示例代码。
在本章中，我们学习了如何使用 CUDA 库来获得深度学习和性能优势。当我们回顾它们的用途时，我们将它们与每一步的深度学习机制相匹配。多亏了我们可以使用的深度学习库，我们可以实现一个简单的有线电视新闻网，而不需要实现算法。然后，我们使用 NVTX 注释在 PyTorch 和 TensorFlow 中描述了 ResNet-50 模型。
对于一些深度学习的工程师和研究人员来说，实现基本算法可能是不切实际的。但是，了解性能因素和基本操作可以帮助您构建高效且有效的深度学习产品。如今，我们看到许多基于深度学习的产品化服务。工程师花费大量资源生产他们训练好的模型，以及训练他们的模型，以便他们获得尽可能低的错误率。希望您能够深入了解如何在深度学习应用中使用 NVTX 概要分析。利用这些知识，你可以从你的图形处理器中获得更多。祝你好运！