C---深度学习-二-

C++ 深度学习（二）

原文：Deep Learning with C++

译者：飞龙

协议：CC BY-NC-SA 4.0

}

sb_stemmer_delete(st);

return out;

}

#### 基于嵌入的向量化

让我们看看如何使用预训练的 GloVe 和 Eigen 将词元转换为向量——这与手动部分中关于独热/频率/序数编码和基本词元到向量映射的向量化概念相同——现在通过简洁的向量化库操作来实现，以提高速度和清晰度：

- `glove_eigen_demo.cpp`：

```cpp

#include <Eigen/Dense>

#include <unordered_map>

#include <fstream>

#include <sstream>

#include <vector>

#include <string>

using EmbMap = std::unordered_map<std::string, Eigen::VectorXd>;

EmbMap loadGlove(const std::string &path, int dim)

{

    EmbMap E;

    std::ifstream f(path);

    std::string line, w;

    while (std::getline(f, line))

    {

        std::istringstream ss(line);

        if (!(ss >> w))

        {

            continue;

        }

        Eigen::VectorXd v(dim);

        for (int i = 0; i < dim; ++i)

        {

            ss >> v[i];

        }

第2章：C++中的数据准备与预处理

}

E.emplace(std::move(w), std::move(v));

}

return E;

}

Eigen::MatrixXd embedTokens(

    const std::vector<std::string> &toks,

    const EmbMap &E,

    int dim

)

{

    Eigen::MatrixXd M = Eigen::MatrixXd::Zero(toks.size(), dim);

    for (size_t i = 0; i < toks.size(); ++i)

    {

        auto it = E.find(toks[i]);

        if (it != E.end())

        {

            M.row(i) = it->second.transpose();

        }

    }

    return M;

}

上下文嵌入（通过ONNX Runtime的BERT）

要使用ONNX Runtime生成上下文嵌入，请从你的分词器中传递令牌ID和掩码（例如，通过SentencePiece或Hugging Face分词器的C API）：

• bert_onnx_demo.cpp：

#include <onnxruntime/onnxruntime_cxx_api.h>

#include <vector>

#include <array>

std::vector<float> runBertCLS(

    const std::vector<int64_t> &ids,

    const std::vector<int64_t> &mask,

    const char *model_path

)

{

    Ort::Env env{ORT_LOGGING_LEVEL_ERROR, "bert"};

    Ort::Session session{env, model_path, Ort::SessionOptions{}};

    std::array<int64_t, 2> shape{1, (int64_t)ids.size()};

    auto mem = Ort::MemoryInfo::CreateCpu(

        OrtArenaAllocator,

        OrtMemTypeDefault

    );

    auto t_ids = Ort::Value::CreateTensor<int64_t>(

        mem,

        const_cast<int64_t *>(ids.data()),

        ids.size(),

        shape.data(),

        2

    );

    auto t_mask = Ort::Value::CreateTensor<int64_t>(

        mem,

        const_cast<int64_t *>(mask.data()),

        mask.size(),

        shape.data(),

        2

    );

    const char *in_names[] = {"input_ids", "attention_mask"};

    const char *out_names[] = {"last_hidden_state"}; // or "logits", etc.

    Ort::Value in_vals[] = {

        std::move(t_ids),

        std::move(t_mask)

    };

    auto outs = session.Run(

        Ort::RunOptions{},

        in_names,

        in_vals,

        2,

        out_names,

        1

    );

    float *p = outs[0].GetTensorMutableData<float>();

    return std::vector<float>(p, p + 768); // adjust hidden size if needed

}

序列填充与截断（仅头文件的辅助工具）

让我们看看如何将分词后的ID序列填充/截断到固定长度（为Transformer做好准备）：

• pad_truncate_demo.cpp：

#include <vector>

#include <algorithm>

// Pad/truncate to fixed length (e.g., Transformer input)

template <class T>

std::vector<T> padOrTruncate(const std::vector<T>& ids, size_t max_len, T pad) {

    if (ids.size() >= max_len) return {ids.begin(), ids.begin() + max_len};

    auto out = ids; out.resize(max_len, pad); return out;

}

// 1 for real tokens, 0 for padding

inline std::vector<int64_t> makeMask(size_t len, size_t max_len) {

    std::vector<int64_t> m(std::min(len, max_len), 1); m.resize(max_len, 0); return m;

}

高级技术 – 命名实体识别/情感分析/词形还原（ONNX Runtime）

NER、情感分析和词形还原是核心的NLP任务：NER定位并标注文本中的真实世界实体；情感分析对作者的态度进行分类（例如正面/负面）；词形还原则将单词还原为词典中的基本形式，以规范化语言。现代Transformer模型为NER和情感分析学习丰富的上下文，而词形还原依赖于语言学规则和词典，共同为下游应用形成干净、结构化的信号。

让我们看看如何用C++运行这些经典任务：

• nlp_advanced_demo.cpp：

// NER (ONNX Runtime, dtype-aware INT32/INT64).

void run_ner(const std::string& onnx_path) {

    Ort::Env env(ORT_LOGGING_LEVEL_WARNING, "ner");

    Ort::Session sess(env, onnx_path.c_str(), Ort::SessionOptions{});

    // Toy example: [CLS] John lives in New York . [SEP]

    std::vector<i64> ids{101,2198,3268,1999,2047,2259,1012,102};

    std::vector<i64> ones(ids.size(),1), zeros(ids.size(),0);

    Shape2 shp{1,(i64)ids.size()};

    auto mem = Ort::MemoryInfo::CreateCpu(OrtArenaAllocator, OrtMemTypeDefault);

    // Match model dtypes (some models want INT32, others INT64).

    auto d_ids = dtype_of(sess, "input_ids");

    auto d_mask = dtype_of(sess, "attention_mask");

    auto d_seg = dtype_of(sess, "token_type_ids"); // may be absent

    std::vector<const char*> names; names.reserve(3);

    std::vector<Ort::Value> vals; vals.reserve(3);

    names.push_back("input_ids"); vals.emplace_back(make_tensor(mem, d_ids, ids, shp));

    names.push_back("attention_mask"); vals.emplace_back(make_tensor(mem, d_mask, ones, shp));

    if (d_seg != ONNX_TENSOR_ELEMENT_DATA_TYPE_UNDEFINED) {

        names.push_back("token_type_ids"); vals.emplace_back(make_tensor(mem, d_seg, zeros, shp));

    }

    const char* out = "logits"; // typical for token-classification exports

    auto outs = sess.Run(Ort::RunOptions{}, names.data(), vals.data(), vals.size(), &out, 1);

    // Argmax over label dimension per token.

    const float* y = outs.front().GetTensorData<float>();

    auto info = outs.front().GetTensorTypeAndShapeInfo();

    auto s = info.GetShape(); // [1, T, C]

    i64 T = s[1], C = s[2];

    for (i64 t = 0; t < T; ++t) {

        const float* row = y + t*C;

        auto idx = std::max_element(row, row + C) - row;

        std::cout << idx << (t + 1 < T ? ' ' : '\n');

    }

}

上述函数加载一个词元分类ONNX模型，并为单个句子构建最小输入，自动将你的整数向量转换为模型所需的数据类型（INT32或INT64）。它执行一次前向传播以获得每个词元的logits，并为每个词元打印一个argmax标签索引——这正是NER背后的核心步骤：

// Sentiment (ONNX Runtime, dtype-aware INT32/INT64)

void run_sentiment(const std::string& onnx_path) {

    Ort::Env env(ORT_LOGGING_LEVEL_WARNING, "sent");

    Ort::SessionOptions so; Ort::Session sess(env, onnx_path.c_str(), so);

    std::vector<int64_t> ids = {101, 1996, 3185, 2001, 2307, 999, 102}; // [CLS] the movie was great ! [SEP]

    std::vector<int64_t> ones(ids.size(), 1), zeros(ids.size(), 0);

    std::array<int64_t,2> shape{1, (int64_t)ids.size()};

    Ort::MemoryInfo mem = Ort::MemoryInfo::CreateCpu(OrtArenaAllocator, OrtMemTypeDefault);

    auto dids = dtype_for(sess, "input_ids");

    auto dmask = dtype_for(sess, "attention_mask");

    auto dseg = dtype_for(sess, "token_type_ids");

    std::vector<const char*> in_names = {"input_ids","attention_mask"};

    std::vector<Ort::Value> in_vals;

    in_vals.emplace_back(make_int_tensor(mem, dids, ids, shape));

    in_vals.emplace_back(make_int_tensor(mem, dmask, ones, shape));

    if (dseg != ONNX_TENSOR_ELEMENT_DATA_TYPE_UNDEFINED) {

        in_names.push_back("token_type_ids");

        in_vals.emplace_back(make_int_tensor(mem, dseg, zeros, shape));

    }

    const char* out_name = "logits";

    auto outs = sess.Run(Ort::RunOptions{}, in_names.data(), in_vals.data(), in_vals.size(), &out_name, 1);

    const float* y = outs.front().GetTensorData<float>();

    auto info = outs.front().GetTensorTypeAndShapeInfo(); auto shp = info.GetShape(); // [1,C]

    int64_t C = shp[1];

    std::cout << "sentiment=" << (int)std::distance(y, std::max_element(y, y + C)) << "\n";

}

此函数加载一个序列分类（情感）ONNX模型，为单个句子构建最小输入，并针对input_ids、attention_mask以及可选的token_type_ids适配模型的整数数据类型（INT32或INT64）。它运行模型，读取[1, C]形状的logits，并打印argmax类别索引——你之后可以将索引映射到标签名称（例如，negative/positive）。

在介绍了高级NLP工具之后，我们现在转向其他数据类型预处理——图像、视频和音频。图像通常需要调整大小/裁剪、RGB转换、按模型的均值/标准差进行缩放/归一化，并排列通道（例如NCHW）。视频则增加了时间维度：均匀或关键帧采样、时间维度上的尺寸调整、逐帧变换，有时还需要光流或视频片段归一化。音频通常涉及重采样（例如16 kHz）、混音为单声道、响度归一化，以及诸如log-mel频谱图之类的时频变换。在所有这些模态中，匹配模型原始的训练流程和张量形状对于稳定精度至关重要。

第 2 章

由于篇幅限制，我们只展示几个小而高效的功能——足以演示如何导入正确的库并组合一个最小的预处理流水线。目标是清晰而非面面俱到：调整大小/裁剪和归一化图像，对音频进行重采样和加窗，以及为短视频片段采样帧，同时匹配模型期望的形状和数据类型。每个片段都倾向于惯用的、依赖较少的模式，你可以将其直接引入自己的代码库并按需扩展。

图像数据

要预处理图像数据，可以应用以下技术：

• 调整大小、裁剪、翻转、旋转和透视变换
• 直方图均衡化、颜色归一化和对比度调整
• 使用 Sobel 或 Canny 滤波器进行边缘检测和特征提取

请看下面的示例：

image_preproc_demo.cpp：

#include <opencv2/opencv.hpp>

#include <iostream>

int main(int argc, char** argv) {

    if (argc < 2) { std::cerr << "usage: " << argv[0] << " input.jpg\n"; return 1; }

    cv::Mat img = cv::imread(argv[1]);

    if (img.empty()) { std::cerr << "load fail\n"; return 2; }

    // 1) Geometry: resize -> center-crop 224x224

    cv::Mat r;

    cv::resize(img, r, {256, 256});

    int s = 224, x = (r.cols - s) / 2, y = (r.rows - s) / 2;

    cv::Mat crop = r(cv::Rect(x, y, s, s));

    // 2) Contrast: histogram equalization on Y (luminance)

    cv::Mat ycc, eq;

    cv::cvtColor(crop, ycc, cv::COLOR_BGR2YCrCb);

    std::vector<cv::Mat> ch;

    cv::split(ycc, ch);

    cv::equalizeHist(ch[0], ch[0]);

    cv::merge(ch, ycc);

    cv::cvtColor(ycc, eq, cv::COLOR_YCrCb2BGR);

    // 3) Edges: Canny

    cv::Mat gray, edges;

    cv::cvtColor(eq, gray, cv::COLOR_BGR2GRAY);

    cv::Canny(gray, edges, 100, 200);

    cv::imwrite("resized.jpg", r);

    cv::imwrite("crop.jpg", crop);

    cv::imwrite("equalized.jpg", eq);

    cv::imwrite("edges.jpg", edges);

    return 0;

}

音频/视频数据

音频/视频数据预处理技术包括以下内容：

• 音频的 MFCC 等特征提取技术
• 视频帧和音轨的时间对齐与同步
• 面向音乐和语音处理的语谱图与色度特征生成

来看一个生成语谱图的示例：

spectrogram_demo.cpp：

// Core spectrogram loop (after reading mono samples into

// std::vector<double> x)

const int N = 512, H = 160, B = N / 2 + 1;

std::vector<double> win(N);

for (int n = 0; n < N; ++n)

    win[n] = 0.5 * (1 - cos(2 * M_PI * n / (N - 1)));

double* in = (double*)fftw_malloc(sizeof(double) * N);

fftw_complex* out = (fftw_complex*)fftw_malloc(sizeof(fftw_complex) * B);

fftw_plan plan = fftw_plan_dft_r2c_1d(N, in, out, FFTW_ESTIMATE);

int T = (int)((x.size() + H - 1) / H);

std::vector<double> S((size_t)T * B);  // log-magnitude spectrogram [time t, bin k]

for (int t = 0; t < T; ++t) {

    size_t start = (size_t)t * H;

    for (int n = 0; n < N; ++n) {

        double s = (start + n < x.size()) ? x[start + n] : 0.0;

        in[n] = s * win[n];  // Hann window

    }

    fftw_execute(plan);  // FFT frame

    for (int k = 0; k < B; ++k) {

        double re = out[k][0], im = out[k][1];

        S[(size_t)t * B + k] = 20.0 * log10(std::sqrt(re * re + im * im) + 1e-12);  // dB

    }

}

接下来，我们将深入探讨归一化（缩放/标准化）和数据增强（噪声、平移、时间/频率掩蔽），以提升模型的鲁棒性。

探索归一化与数据增强技术

在本节中，所有代码都将汇集到 multimodal_preproc_demo.cpp 文件中。

归一化策略

在多模态流水线中，归一化可确保各特征范围的一致性，从而促进模型训练并改善收敛。它还能缓解多模态输入中数据尺度不一致所带来的问题。实践中，请按如下方式对各模态进行归一化：

• 文本数据：使用预训练模型（例如 BERT 和 GloVe）对嵌入进行归一化。实施缩放以保持维度一致性。
• 图像数据：将像素强度缩放到 [0, 1] 或 [-1, 1] 等范围，并对深度学习模型执行按通道的标准化。
• 数值数据：使用对数缩放、Box-Cox 或 Yeo-Johnson 等变换来处理偏态分布。

以下是一个图像归一化示例：

#include <opencv2/opencv.hpp>

cv::Mat normalizeImage(const cv::Mat& image) {

    cv::Mat normalized;

    image.convertTo(normalized, CV_32F, 1.0 / 255.0);

    return normalized;

}

• 文本归一化（嵌入）：对向量进行 L2 缩放，使余弦相似度有意义，且向量模长不会占主导地位：

#include <vector>

#include <cmath>

void l2_normalize(std::vector<float> &x)

{

    double s = 0;

    for (float v : x)

        s += v * v;

    s = std::sqrt(s);

    if (s)

        for (float &v : x)

            v /= (float)s;

}

在深入建模之前，最后一个手段是强化数据集本身。数据增强通过生成保留标签且让模型接触自然变化的逼真变体，拓宽有效训练分布——在样本稀缺、类别不平衡或出现过拟合时至关重要。

数据增强技术

数据增强通过对现有数据施加变换来生成多样化的数据集。这种方法可以提高模型的泛化能力并防止过拟合。实践中，可以按如下方式对各模态进行增强：

• 图像增强：应用随机旋转、翻转、高斯噪声、颜色抖动和仿射变换
• 文本增强：使用释义、回译、同义词替换，以及随机插入或删除词元
• 数值数据增强：添加高斯噪声、通过插值生成合成样本，或实现自助采样

来看一个向图像添加高斯噪声的示例：

#include <opencv2/opencv.hpp>

#include <random>

cv::Mat addGaussianNoise(const cv::Mat& image, double mean, double stddev) {

    cv::Mat noise(image.size(), CV_32F);

    std::default_random_engine generator;

    std::normal_distribution<float> distribution(mean, stddev);

    for (int i = 0; i < noise.rows; ++i) {

        for (int j = 0; j < noise.cols; ++j) {

            noise.at<float>(i, j) = distribution(generator);

        }

    }

    cv::Mat noisyImage;

    image.convertTo(noisyImage, CV_32F);

    noisyImage += noise;

    return noisyImage;

}

• 增强（词元丢弃）：随机丢弃词元，以提升对缺失/含噪单词的鲁棒性：

#include <vector>

#include <string>

#include <random>

std::vector<std::string> token_dropout(const std::vector<std::string> &toks, float p)

{

    std::mt19937 rng(42);

    std::bernoulli_distribution keep(1.0f - p);

    std::vector<std::string> out;

    out.reserve(toks.size());

    for (auto &t : toks)

        if (keep(rng))

            out.push_back(t);

    return out;

}

现在，我们从单样本变换转向端到端数据流水线，展示逐项预处理如何通过 PyTorch Dataset API（用于加载器/迭代器）接入训练和生产流程。

集成 PyTorch Dataset API 构建数据流水线

PyTorch Dataset API 支持灵活的自定义实现，用于高效的数据加载和预处理。模块化设计可无缝集成到复杂的数据工作流和分布式环境中。

以下是一个自定义 PyTorch 实现：

custom_dataset_demo.cpp：

#include <torch/torch.h>

struct CustomDataset : torch::data::datasets::Dataset<CustomDataset> {

    CustomDataset(torch::Tensor X, torch::Tensor y)

        : data_(X.contiguous().to(torch::kFloat32)),

          labels_(y.contiguous().to(torch::kLong)) {}

    torch::data::Example<> get(size_t index) override {

        return {data_[index], labels_[index]};

    }

    torch::optional<size_t> size() const override {

        return data_.size(0);

    }

private:

    torch::Tensor data_, labels_;

};

到目前为止，我们优化了数据的内容；接下来，我们要优化数据的移动方式。当数据集超出 RAM 容量时，内存映射、分布式文件系统和批量 I/O 等技术可以让流水线保持高速度和资源高效——在扩展规模之前，让我们先从一个最小的 mmap 示例开始。

第2章

管理大规模数据集

当数据集超出RAM容量时，有三种模式可以派上用场：内存映射、批量（流式）I/O和分布式文件系统。这里我们重点介绍前两种，因为它们可移植性强，并且易于在独立C++中演示；分布式文件系统需要集群基础设施，超出了本书范围。有关代码片段来源的可运行示例，请参阅bigdata_demo.cpp。

内存映射

内存映射I/O允许你将文件视为连续的内存块，避免了显式读取循环和额外复制——非常适合对大型数据集进行快速随机访问。接下来展示的资源获取即初始化（RAII）包装器以只读方式打开文件，将其映射到进程地址空间，并确保在销毁时干净地取消映射/关闭（可移植操作系统接口（POSIX）mmap；在Windows上，使用CreateFileMapping/MapViewOfFile）：

struct MappedFile {

    void* data=nullptr; size_t size=0; int fd=-1;

    explicit MappedFile(const std::string& path) {

        fd = ::open(path.c_str(), O_RDONLY);

        if (fd < 0) throw std::runtime_error("open failed");

        struct stat st{}; if (fstat(fd, &st) != 0) throw std::runtime_error("fstat failed");

        size = static_cast<size_t>(st.st_size);

        data = ::mmap(nullptr, size, PROT_READ, MAP_PRIVATE, fd, 0);

        if (data == MAP_FAILED) throw std::runtime_error("mmap failed");

    }

    ~MappedFile(){ if (data && data!=MAP_FAILED) ::munmap(data,size); if (fd>=0) ::close(fd); }

};

当不应映射整个文件时——因为它太大、位于远程，或者你希望严格控制内存边界——可以改用分块读取；在下一节中，我们将以固定大小的批次流式处理数据。

批量处理

对于无法轻松放入RAM的文件——或者当你希望内存使用可预测时——批量（分块）处理会读取固定大小的块，处理后再前进到下一个。以下辅助函数以B字节块流式读取文件，以便你可以增量处理每个部分：

// Stream large files without loading all at once.

#include <fstream>

#include <vector>

void process_in_batches(const std::string& path, size_t B=1<<20) {

    std::ifstream f(path, std::ios::binary);

    std::vector<char> buf(B);

    while (f) {

        f.read(buf.data(), buf.size());

        std::streamsize n = f.gcount();

        // process buf[0..n)

    }

}

总而言之，这些预处理模式——涵盖文本、图像、音频和大规模存储——构成了一个实用的工具包，可用于准备任意规模的数据。有了这些基础组件，我们现在进入本章总结。

总结

在本章中，你学习了如何将杂乱、异构的数据转换为可供模型使用的张量。我们诊断了常见问题（尺度不匹配、缺失值、异常值、数据泄漏），并应用了有原则的修复方法：前向/后向填充和统计插补器；从独热和频率到学习嵌入的分类编码器；通过最小-最大、Z-score和稳健变换进行缩放；以及使用PCA进行降维。对于时间序列数据，你构建了滚动统计量、差分和频谱特征；对于表格数据，你构建了交互项和多项式项。我们还强调了数据增强，以在不改变标签的情况下拓宽有效的训练分布。

在工程方面，你练习了用C++实现这些想法：用于插补和编码的小型可测试实用程序；用于向量化变换的Armadillo/Eigen；以及用于生产速度和数值稳定性的库驱动版本（例如PCA、k-NN插补、FFT后端）。你学会了使用内存映射和批量处理来流式处理大型数据集，并通过将拟合（在训练窗口/队列上）与变换（在验证/测试/生产上）分离来防范数据泄漏。

这些技能立竿见影：更干净的输入意味着更快的收敛、更好的泛化以及更可靠的推理。稳健的预处理还能减少运维意外——定义良好的模式、一致的缩放和无泄漏的流水线是任何高质量深度学习系统的基础。

在下一章，我们将从数据转向计算。你将设置计算统一设备架构（CUDA），编写并启动你的第一个内核，管理设备内存，并分析性能——以便本章准备好的张量能够高效地流经GPU加速的训练和推理。

问题

在继续之前，请使用以下提示检查你的理解：

1. 为什么对输入进行归一化？最小-最大缩放和标准化有何不同？
2. 对于输出为[1, C]的序列分类，如何获得预测类别？
3. 列举两种保持标签不变的图像增强方法，以及一种在卷积神经网络之前常用的归一化技术。
4. 短时傅里叶变换（STFT）频谱图捕捉到什么？窗口大小与时间/频率分辨率之间的权衡是什么？

延伸阅读

以下是一些推荐的资源，可加深你的理解（我们会在每章中使用相同的标识）：

• ONNX Runtime C++ API（会话、张量、I/O绑定）
• Hugging Face Optimum：将Transformers导出为ONNX
• OpenCV文档：调整大小、色彩空间、直方图均衡化
• PyTorch C++（LibTorch）数据集/DataLoader指南
• STFT/频谱图入门（信号处理文本）
• POSIX mmap 和内存映射I/O参考资料
• 数据密集型应用系统设计：构建可靠、可扩展、可维护系统的核心思想，Martin Kleppmann 著，O'Reilly Media（大规模数据模式）

答案

以下是我们建议的答案——请对照检查你的回答：

1. 归一化可以稳定训练并加速收敛。最小-最大缩放将数据缩放到固定范围（例如[0,1]）；标准化使数据以零均值和单位方差为中心和缩放。
2. 在类别维度C上取argmax来选取预测标签。
3. 增强方法：随机翻转、随机旋转（也可以裁剪/亮度扰动）。归一化：将像素缩放到[0,1]，并应用逐通道的均值/标准差标准化。
4. 频谱图展示能量随时间和频率的变化。较大的窗口提高频率分辨率，但会降低时间分辨率（反之亦然）。

获取本书的PDF版本及更多信息

扫描二维码（或访问packtpub.com/unlock）。按书名搜索本书，确认版本，然后按照页面上的步骤操作。

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_108_0.png
https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_108_1.png

注意：请保留发票。直接从Packt购买不需要发票。

第3章：CUDA在C++深度学习中的GPU加速

计算统一设备架构（CUDA）是NVIDIA的并行计算平台和编程模型，用于图形处理单元（GPU）。通过暴露数千个轻量级线程，CUDA可以加速数值模拟、深度学习（DL）和图像处理等工作负载，远超典型CPU性能。在本章中，你将设置工具链，学习CUDA的执行和内存模型（内核、网格/块/线程、共享/全局内存），并使用核心库构建第一个内核。到本章结束时，你将能够理解GPU/CPU数据移动，编写和启动内核，分析和优化代码，并避免常见的性能陷阱。

在本章中，我们将涵盖以下主题：

• GPU加速应用与纯CPU应用
• 安装和设置CUDA
• 使用CUDA优化C++代码
• 如何在CUDA中进行调试

GPU加速应用与纯CPU应用

GPU是一种面向吞吐量的处理器，设计用于并行运行大量轻量级线程，而CPU则针对少量低延迟的复杂线程进行了优化。这种差异改变了你组织工作、管理内存和分析性能的方式。GPU加速应用与纯CPU应用在几个重要方面有所不同：

• 大规模并行：现代 GPU 可包含数千个核心，允许多个操作同时发生。与一个或少数几个 CPU 核心迭代处理数据不同，成百上千个线程会并行处理数据块。

• 编程模型：借助 CUDA，你可以决定应用程序中哪些部分适合并行化，并将它们转换为内核（kernels）。然后，通过执行配置（块和线程的数量）启动这些内核，该配置决定了 GPU 资源的使用方式。

• 内存考量：由于 GPU 是带有自己内存的独立设备，你需要处理数据传输，或使用 CPU 和 GPU 都能访问的统一内存（Unified Memory）（cudaMallocManaged）。

优秀的应用程序越来越多地同时使用这两种设备；这也被称为异构并行计算。不过，在本章中，我们将重点放在 GPU 上。

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_110_0.png

图 3.1：CPU（汽车）与 GPU（自行车队）——延迟与吞吐量

可以把 CPU 想象成一辆快车，把 GPU 想象成一支自行车队。对于蜿蜒路线上的单个乘客来说，汽车（CPU）非常出色：核心速度高、缓存大，旨在最小化延迟。自行车队（GPU）则同时将许多小包裹运送到附近的地址：成千上万个更简单的骑手（线程）并行投递，从而最大化吞吐量。

• GPU（自行车）胜出的场景：大规模、规整、数据并行的负载，可以干净地拆分到许多线程上——矩阵乘法、卷积、图像滤波器、蒙特卡洛——尤其是在大批次和合并内存访问的情况下。

• CPU（汽车）胜出的场景：具有紧密串行依赖、分支密集逻辑、频繁同步的任务，或者内核启动/主机-设备传输开销和线程发散会吞没收益的情况。

异构并行计算将延迟优化的 CPU 与面向吞吐量的 GPU 配对，让各自运行最擅长的工作。用 CUDA 的术语来说，CPU（主机）负责编排控制密集型逻辑、I/O 和批处理，而 GPU（设备）执行数据并行内核。良好的设计会尽量减少 PCIe/NVLink 上的主机-设备通信量，让热数据驻留在 GPU 上，使用页锁定内存，并通过流和事件将传输与计算重叠。统一内存可以加快开发速度；若要达到峰值性能，则应优先使用显式拷贝，并在需要时使用预取（prefetch）和建议（advice）。构建并发地分阶段处理输入、计算和输出的流水线，以充分利用两种引擎；在能够减少启动开销或全局内存遍历时，融合内核。使用 Nsight Systems/Compute 进行性能分析，以平衡利用率、暴露带宽瓶颈，并决定哪些部分留在 CPU 上、哪些放在 GPU 上。

在编写第一个内核之前，我们需要一个可用的 CUDA 工具链。本节将引导你在 Linux/WSL 和 Windows 上安装 NVIDIA 驱动和 CUDA 工具包、设置路径并验证环境，同时介绍版本兼容性相关注意事项。

安装并配置 CUDA

在开始编写和运行 CUDA 程序之前，你需要搭建合适的开发环境。这包括确保你拥有支持 CUDA 的 GPU、安装合适的驱动、下载并配置 CUDA 工具包，以及可选地设置集成开发环境（IDE）或代码编辑器。本节将指导你在 Windows、Linux 和 macOS 系统上完成必要步骤，并提供备选的云方案。

安装 NVIDIA 驱动

如果你没有合适的本地系统，可以从 亚马逊云服务（AWS）、Azure、谷歌云平台（GCP） 或 IBM Cloud 等云服务商租用 GPU 加速的虚拟机（VM）或实例。这些服务可让你无需安装任何本地硬件即可快速搭建 CUDA 环境。

CUDA 需要较新且兼容的 NVIDIA 驱动。在大多数系统上，安装 CUDA 工具包时也会安装合适的驱动。不过，你也可以按以下方式独立安装驱动：

• Windows：

  • 从 NVIDIA 的下载驱动页面（https://www.nvidia.com/en-us/drivers/）下载最新的 NVIDIA 驱动。
  • 运行安装程序并按照提示操作。可能需要重启。
  • 安装完成后，打开 NVIDIA 控制面板或运行 nvidia-smi 以验证驱动已生效。

• Linux（基于 Ubuntu/Debian）：

  • 使用包管理器或 NVIDIA 网站提供的 .run 文件安装驱动。例如，在 Ubuntu 上，你可以运行以下命令：
    sudo apt-get update
sudo apt-get install nvidia-driver-xxx
  • 将 xxx 替换为推荐的驱动版本号。安装完成后，重启系统并使用 nvidia-smi 进行验证。

• Linux（基于 RHEL/CentOS/SUSE）：

  • 遵循特定发行版的说明，或使用 NVIDIA CUDA 仓库提供的软件包。更详细的指南可在 NVIDIA Linux 安装指南（https://docs.nvidia.com/cuda/）中找到。

• macOS：

  • 官方对较新 macOS 版本的 NVIDIA 驱动支持有限。如果你使用的是仍支持 NVIDIA Web Driver 的旧版 macOS（10.14 之前），请从 NVIDIA 网站下载并安装最新的兼容驱动。
  • 由于原生支持有限，许多 macOS 用户现在选择外置 GPU（eGPU）扩展坞或云方案。

下载并安装 CUDA 工具包

CUDA 工具包包含 CUDA 编译器（nvcc）、库、示例代码和必要的开发工具。安装步骤如下：

1. 下载 CUDA 工具包：访问 NVIDIA CUDA 工具包下载页面（https://developer.nvidia.com/cuda-downloads），选择你的操作系统、架构、发行版（Linux）和安装程序类型。

2. 运行安装程序，具体操作如下：

   • Windows：双击 .exe 安装程序并按照屏幕上的说明操作。默认情况下，CUDA 工具包将安装在 C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\vXX.X。
   • Linux：使用 .run 文件，或按照下载页面上的说明使用你的包管理器。例如，在 Ubuntu 上，你可以执行以下操作：
     sudo dpkg -i cuda-repo-ubuntuXXXX.deb
sudo apt-get update
sudo apt-get install cuda
     安装后，通过编辑 ~/.bashrc，将 CUDA 的 bin 目录添加到 PATH 环境变量中：
     echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc
echo 'export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc
source ~/.bashrc
   • macOS：官方已逐步停止在较新 macOS 版本上支持 CUDA 工具包。对于较旧的 macOS 版本，请运行 .pkg 安装程序并按照 NVIDIA 的说明操作。

3. 验证：安装完成后，运行以下命令，验证 nvcc 是否已加入 PATH 环境变量：

   nvcc --version

   你应该会看到输出，其中详细列出了所安装的 CUDA 版本。

安装 CUDA 有时可能颇具挑战，尤其是在处理驱动兼容性、编译器版本或缺少运行时库（如 Windows 上的 DLL 文件或 Linux 上的共享库）时。开发人员还可能遇到与 PATH 或 LD_LIBRARY_PATH 等环境变量相关的问题；这些变量必须包含 CUDA 的二进制目录和库目录，nvcc 等工具才能正常运行。

如果在安装过程中遇到问题，以下外部资源提供了详细的分步指南和故障排查步骤：

• NVIDIA CUDA 工具包下载页面：https://developer.nvidia.com/cuda-downloads

• NVIDIA CUDA 快速入门指南：https://docs.nvidia.com/cuda/cuda-quick-start-guide/

• CUDA 安装指南（Linux）：https://docs.nvidia.com/cuda/cuda-installation-guide-linux/

• CUDA 安装指南（Windows）：https://docs.nvidia.com/cuda/cuda-installation-guide-microsoft-windows/

• CUDA 和 cuDNN 安装实用教程（含故障排除）：https://www.digitalocean.com/community/tutorials/install-cuda-cudnn-for-gpu

这些指南包含分步说明、兼容性注意事项，以及针对环境配置和依赖问题的常见修复方法。在解决安装问题或版本冲突时，NVIDIA 官方文档通常是最可靠的参考资料。

安装其他依赖项

根据你的项目需求，可能需要以下工具：

• C++ 编译器：
  • 在 Linux 上，gcc 和 g++ 通常可通过包管理器获取。
  • 在 Windows 上，可以使用 Microsoft Visual Studio（2019 或 2022），或依赖 Visual Studio Build Tools。
  • 在 macOS 上，如果你的 CUDA 版本兼容，可以使用 clang++（属于 Xcode 命令行工具的一部分）。
• IDE（集成开发环境）：
  • Windows：Microsoft Visual Studio 与 CUDA 项目集成良好。
  • Linux 和 macOS：常见选择包括 Visual Studio Code、带 CUDA 插件的 Eclipse 或 JetBrains CLion。这些 IDE 通常为 CUDA 代码提供语法高亮、代码补全和集成调试功能。
• 库和框架：
  • NVIDIA 提供了众多加速库（cuBLAS、cuFFT、cuDNN 等）。如果你的项目涉及线性代数、信号处理或深度学习，请安装并链接这些库。
  • 有关库安装和使用的详细信息，请查阅 NVIDIA CUDA Toolkit 文档（https://docs.nvidia.com/cuda/）。

使用基于云的 CUDA 环境

如果硬件兼容性或环境搭建的复杂性成为问题，可以考虑使用基于云的开发环境：

• AWS：提供配备 NVIDIA GPU 的 EC2 实例（P3、G4 和 P4 实例）
• Microsoft Azure：为计算密集型任务提供各种支持 GPU 的 VM 实例
• GCP：提供配备 NVIDIA GPU 的计算实例，用于训练和推理工作负载
• NVIDIA LaunchPad 或 NGC 容器：提供预配置的容器，其中包含 CUDA、各类库和 AI 框架，可简化环境搭建流程

你可以通过安全外壳（SSH）连接到这些环境，并像在本地机器上一样运行 CUDA 程序。

在本节中，我们将通过 Google Colab 使用 CUDA 来走一条捷径：

1. 在 Colab 中使用 CUDA 有两种方式——一种是将 CUDA 代码写入 .cpp 文件，另一种是使用 nvcc4jupyter Colab 扩展：

   https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_115_0.png

   图 3.2：要更改运行时，请转到“运行时”选项卡并选择“更改运行时类型”

2. 接下来，在对话框中，将硬件加速器设置为 T4 GPU，然后点击保存以启用 CUDA，如下所示。

   https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_116_0.png

   图 3.3：将运行时类型更改为 GPU 可启用 GPU 编程

启用 GPU 运行时后，让我们验证 CUDA 工具链并加载 Colab 扩展，以便直接从 Notebook 单元格编译 CUDA 代码。

3. 使用 Google Colab，你可以通过以下命令检查 CUDA 的安装情况：

   
   !nvcc --version
   
   

   输出的结果应如下所示：

   https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_116_1.png

   图 3.4：!nvcc --version 命令的输出

4. 接下来，我们需要安装 nvcc4jupyter 扩展，以便从 Notebook 单元格运行 nvcc：

   
   !pip install nvcc4jupyter
   
   %load_ext nvcc4jupyter
   
   

   图 3.5 显示了输出应该是什么样子。

   https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_117_0.png

   图 3.5：安装 nvcc4jupyter 并加载其扩展

之后，你可以使用以下任一方式处理代码：

5. 在 .cpp 文件中编写代码。Colab 是为编写 Python 代码而设计的；但是，借助新安装的扩展，我们实际上可以编写出 .cpp 文件并在本地编译。图 3.6 展示了具体做法。

   https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_117_1.png

   图 3.6：如何使用 Colab 编写 .cpp 文件并编译

6. 直接使用 nvcc4jupyter 扩展编写代码。使用该扩展时，你可能会注意到 Python 静态检查器给出的一些警告。请忽略这些警告。

   https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_118_0.png

   图 3.7：如何在 Colab 中编译 .cpp 文件

当你拥有支持 CUDA 的 GPU 时，请安装兼容的 NVIDIA 驱动程序和 CUDA Toolkit（包含 nvcc 和核心库），然后进行端到端验证：nvidia-smi 命令应列出你的 GPU/驱动程序，nvcc --version 命令应报告预期的工具包版本，并且像 deviceQuery（或 bandwidthTest）这样的示例程序应能成功编译和运行。你可以在 Windows/Linux/WSL 上进行本地开发，或使用云选项（例如，通过 nvcc4jupyter 或编写/编译 .cpp 文件的 Colab）。请注意驱动程序与工具包之间可能不匹配、Linux 上可能缺失 PATH/LD_LIBRARY_PATH、WSL2 的要求以及 Colab 的临时会话限制。确认兼容性并成功运行测试后，你就可以编写、启动和分析你的第一个 CUDA 内核了。

使用 CUDA 优化 C++ 代码

本节将展示如何利用 CUDA 将性能关键的 C++ 代码转变为快速、GPU 加速的实现。我们将从识别热点和建立 CPU 基线开始，然后将数据并行任务映射到内核，选择有效的网格/块大小，并管理主机与设备间的内存以尽量减少数据传输。你将学习实用技术——合并访问、共享内存分块、占用率调整和基于流的重叠操作——以及何时应该使用 cuBLAS/cuFFT/cuDNN 而不是手写代码。我们将使用 Nsight Systems/Compute 进行测量，分析加速比（Amdahl 定律/Roofline 模型），并通过正确性测试安全地迭代，以确保优化带来真实、可复现的性能提升。

CUDA 编程模型

CUDA 提供了一种异构编程模型，其中主机代码在 CPU 上运行，而设备代码（内核）在 GPU 上运行。CPU 和 GPU 各自拥有独立的内存空间和执行上下文，通常需要显式的数据拷贝（除非使用统一内存）。通常，你需要识别算法中数据并行的部分，并以内核的形式将其卸载到 GPU 上。

CUDA 内核本质上是一个在 GPU 上运行的 C/C++ 函数（带有一些语言扩展和限制）。现代 NVIDIA GPU 实际上可以并发运行多个内核，但为了简单起见，在示例或入门章节中，我们通常假设某一时刻只有一个内核处于活动状态。内核一经启动，便由 GPU 上的众多线程执行。

CUDA 线程与 CPU 线程在几个重要方面存在差异：

• 大规模并行性：每次内核启动可涉及数千到数百万个线程，所有线程在不同数据元素或数据区域上执行相同的代码
• 轻量级创建：GPU 线程的创建开销极小，线程间的上下文切换主要由硬件管理，与 CPU 线程相比速度极快
• 数据并行聚焦：每个线程通常独立处理整个数据集的一部分（或子集），使得 GPU 能够实现大规模并行

主机负责启动内核，并在设备内核启动之间执行串行代码；也就是说，管理内存、与设备之间的数据交换以及错误处理。

CPU 基线：C++ 中的数组加法

让我们从一个简单的基于 CPU 的程序开始，该程序将两个浮点数数组相加：

#include <iostream>

#include <algorithm>

#include <cmath>

// 对两个向量执行逐元素累加

void vectorAddCPU(int length, float* inputA, float* inputB)

{

    for (int idx = 0; idx < length; ++idx)

    {

        inputB[idx] += inputA[idx];

    }

}

int main()

{

    const int SIZE = 1 << 20; // 1,048,576 个元素

    // 在主机（CPU）上分配内存

    float* vecA = new float[SIZE];

    float* vecB = new float[SIZE];

    // 用示例值填充数组

第 3 章

for (int i = 0; i < SIZE; ++i)

{

    vecA[i] = 0.5f;

    vecB[i] = 2.5f;

}

// Execute vector addition

vectorAddCPU(SIZE, vecA, vecB);

// Verify correctness

float largestDeviation = 0.0f;

for (int i = 0; i < SIZE; ++i)

{

    float expected = 3.0f;

    largestDeviation = std::max(largestDeviation, std::fabs(vecB[i] - expected));

}

std::cout << "Largest deviation: " << largestDeviation << std::endl;

// Release allocated memory

delete[] vecA;

delete[] vecB;

return 0;

}

编译并运行该程序：

$ g++ vector_add_cpu.cpp -o vector_add_cpu

$ ./vector_add_cpu

输出如下：

Max error: 0.000000

在代码中，*x 表示指向浮点数组的指针。在 C++（以及 CUDA）中，星号（*）表示该变量是一个指针，也就是说它存储的是某个值的内存地址，而不是值本身（例如，float *x 表示 x 是一个指向浮点数的指针）。如果为其分配了内存（例如 x = new float[N]; // In CPU memory 或 cudaMallocManaged(&x, N * sizeof(float)); // In CUDA Unified Memory），那么 x 就保存着这块已分配内存的地址，我们可以使用数组下标来访问其中的元素：x[0] = 1.0f; 将 1.0f 赋给第一个元素。或者，我们也可以使用指针运算：*(x + 0) = 1.0f;，这与 x[0] = 1.0f 等价。

转换为 CUDA 代码

为了在 GPU 上并行化这种计算，我们将 add() 函数修改为作为 CUDA 内核运行。__global__ 关键字用于指定一个在 GPU 上运行的函数：

__global__

void add(int n, float *x, float *y) {

    int index = threadIdx.x;

    int stride = blockDim.x;

    for (int i = index; i < n; i += stride)

        y[i] = x[i] + y[i];

}

CUDA 提供了可以从 CPU 和 GPU 两方面访问的统一内存（Unified Memory）。使用 cudaMallocManaged() 分配内存，使用 cudaFree() 释放内存：

float *x, *y;

cudaMallocManaged(&x, N*sizeof(float));

cudaMallocManaged(&y, N*sizeof(float));

这个 CPU 基线程序既为我们提供了正确性检查，也提供了性能参考：我们分配并初始化两个数组，运行一个直接的 add() 循环，然后验证 Max error: 0.0 以确认逻辑无误。请保留该版本（及其运行时间），作为后续所有工作的基准。接下来，我们将把这同一段计算移植到 CUDA——把 add 变成内核，将迭代映射到线程/线程块，并管理内存（统一内存或显式拷贝）——然后重新验证结果，并与这个基线对比加速比。

启动 CUDA 内核

我们不直接调用 add()，而是使用三尖括号语法在 GPU 上启动它：

add<<<1, 1>>>(N, x, y);

由于 CUDA 内核是异步执行的，我们需要在访问结果之前同步设备：

cudaDeviceSynchronize();

让我们仔细看看这段代码：

• __global__ void GPUFunction()：__global__ 修饰符标记一个内核——这是一个在 GPU 上执行、通常由 CPU（主机）启动的函数。内核的返回类型必须是 void。对于动态并行（计算能力 ≥ 3.5，-rdc=true），内核实际上可以从设备端启动其他内核。

• GPUFunction<<<1, 1>>>();：三尖括号启动配置指定了线程层次结构：<<<gridDim, blockDim, [sharedMemBytes], [stream]>>>。<<<1,1>>> 表示 一个线程块，每个块中一个线程——这对 “hello world” 演示很有用；实际工作负载会使用许多线程块和线程。在内核内部，blockIdx、blockDim、threadIdx 等内建变量让每个线程可以计算自己的工作索引。

• cudaDeviceSynchronize();：内核启动是异步的——主机会立即继续执行。cudaDeviceSynchronize() 函数会阻塞主机，直到默认流中此前所有工作完成为止（有助于计时和错误检查）。提示：在启动后立即使用 cudaGetLastError() 或 cudaPeekAtLastError() 检查启动错误。

完整的 CUDA 实现

我们已经介绍了编写 CUDA 所需的样板内容；现在，让我们把它组合成一个完整且可运行的示例：

#include <iostream>

#include <math.h>

__global__

void add(int n, float *x, float *y) {

    int index = threadIdx.x;

    int stride = blockDim.x;

    for (int i = index; i < n; i += stride)

        y[i] = x[i] + y[i];

}

int main(void) {

    int N = 1<<20;

    float *x, *y;

    cudaMallocManaged(&x, N*sizeof(float));

    cudaMallocManaged(&y, N*sizeof(float));

    for (int i = 0; i < N; i++) {

        x[i] = 1.0f;

        y[i] = 2.0f;

    }

    add<<<1, 256>>>(N, x, y);

    cudaDeviceSynchronize();

    float maxError = 0.0f;

    for (int i = 0; i < N; i++)

        maxError = fmax(maxError, fabs(y[i]-3.0f));

    std::cout << "Max error: " << maxError << std::endl;

    cudaFree(x);

    cudaFree(y);

    return 0;

}

现在，你可以使用 nvcc（NVIDIA 的 CUDA C/C++ 编译器）来构建并可选地运行你的程序，命令如下：

$ nvcc -arch=sm_70 -o out some-CUDA.cu -run

我们来分解一下上述命令：

• nvcc：NVIDIA CUDA 编译器驱动程序。
• -arch=sm_70：指定为计算能力 7.0 的 GPU（例如 V100）编译。请根据你的系统按需调整。
• -o out：将编译生成的二进制文件写入名为 out 的可执行文件。
• -run：在成功编译后自动运行该可执行文件。

如果省略 -run，你只需编译：

nvcc -arch=sm_70 -o out some-CUDA.cu

./out

在 Colab notebook 中，你可以执行以下内容：

%%shell

nvcc add.cu -o add_cuda

./add_cuda

预期输出如下：

Max error: 0.000000

正确性和编译都已验证之后，下一个考虑因素是时间花在了哪里。在下一节中，我们将对该程序进行性能分析——使用 CUDA 事件对主机-设备传输和内核计时，然后使用 Nsight Systems/Compute 检查实际占用率、内存吞吐量以及线程束/分支效率——从而定位瓶颈并指导优化。

分析 CUDA 性能

使用 nvprof 对执行过程进行性能分析：

$ nvprof ./add_cuda

该命令的输出打印结果如图 3.8 所示：

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_124_0.png

图 3.8：使用 CUDA 进行数组加法的性能分析结果

你的计时结果可能因 GPU 不同而有所差异。

CUDA 线程层次结构

一个 CUDA 内核可能由数千甚至数百万个线程执行。这些线程被组织成线程块，线程块又被组织成网格。理解这种线程层次结构对于扩展你的应用程序至关重要：

• 线程：执行内核代码的最小工作单元
• 线程块：一组线程（一维上最多 1,024 个），它们可以共享内存并彼此同步
• 网格：一个或多个线程块的集合

图 3.9 展示了 CUDA 的执行层次结构（线程 → 线程块 → 网格）以及每一层的内存空间（每线程的寄存器/局部内存、每线程块的共享内存、设备全局内存）。在选择启动维度和规划数据移动时，请牢记这一模型。

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_125_0.png

图 3.9：CUDA 线程层次结构与内存空间（线程 ↔ 局部内存，线程块 ↔ 共享内存，网格 ↔ 全局内存）

有了这些信息，你就可以通过 CUDA 的三尖括号语法来指定执行配置——线程块数量与每线程块线程数——从而启动内核。执行配置可以如下所示：

someKernel<<<NUMBER_OF_BLOCKS, THREADS_PER_BLOCK>>>(parameters);

接下来，我们将看到如何利用刚才学到的知识，使用多个线程来加速这个 for 循环。

使用多个线程加速

现在你已经用单个线程运行了执行加法操作的内核；要充分利用 GPU，我们需要让它并行化。关键就在于 CUDA 的 <<<1, 1>>> 语法。<<< >>> 中的值是执行配置——它们设置启动时的网格和线程块维度，这决定了有多少个 GPU 线程来运行该内核。具体来说，我们有如下内容：

• 第一个参数表示要启动的线程块数量
• 第二个参数表示每个线程块中的线程数量

add<<<1, 1>>> 这一配置意味着：

• 创建了一个线程块
• 在该线程块中创建了一个线程

前面的示例中，内核add()函数只启动了一个线程；因为它逐个遍历整个数组，所以它可以正确地完成整个数组的加法。然而，如果你启动多个线程，它们在没有协调的情况下读写同一个y[i]元素，就会产生竞争条件——最终值取决于哪个线程先读或先写。例如，一个线程可能写入y[i] = x[i] + y_old，而另一个线程看到这个更新后的值，可能计算出y[i] = x[i] + (x[i] + y_old)，从而得到错误结果。

为了利用GPU的并行性，需要将工作分配到多个线程。由于NVIDIA硬件以32线程的线程束（warp）为单位调度线程，因此请选择32的倍数作为块大小（例如256），并启动足够多的块来覆盖数据。相应地更新代码：

add<<<1, 256>>>(N, x, y);

如果按原样运行，并行版本仍然会让每个线程重复同样的工作，而不是将工作划分开。为了在多个线程之间划分数组，可以使用CUDA的内建变量为每个线程计算唯一索引。这里，threadIdx.x是线程在块内的位置，blockDim.x是块中的线程数，blockIdx.x标识块。一种常见模式是int i = blockIdx.x * blockDim.x + threadIdx.x；这样，每个线程只更新自己的元素，从而实现真正的并行执行。

我们需要修改循环，让并行线程跨步遍历数组：

int blockSize = 256;

int numBlocks = (N + blockSize - 1) / blockSize;

add<<<numBlocks, blockSize>>>(N, x, y);

修改内核，使其在所有线程上计算：

__global__

void add(int n, float *x, float *y)

{

    int index = threadIdx.x;

    int stride = blockDim.x;

    for (int i = index; i < n; i += stride)

        y[i] = x[i] + y[i];

}

以下是我们加法函数或内核的修改后代码。我们将其保存到add_block.cu文件中。代码分解如下：

• 我们创建一个使数组元素相加的内核：

__global__

void add(int n, float *x, float *y) {

    int index = threadIdx.x;

    int stride = blockDim.x;

    for (int i = index; i < n; i += stride)

        y[i] = x[i] + y[i];

}

• 我们启动main，设置元素总数（N = 1<<20 = 1,048,576），并声明输入/输出数组的指针：

int main(void) {

    int N = 1 << 20; // 1M elements

    float *x, *y;

• 我们分配统一内存（Unified Memory），使CPU和GPU都能访问x和y，而无需显式调用cudaMemcpy：

cudaMallocManaged(&x, N * sizeof(float));

cudaMallocManaged(&y, N * sizeof(float));

• 我们以1个块 × 256个线程启动内核（网格跨步循环覆盖所有N个元素）。cudaDeviceSynchronize()函数会在主机读取结果之前等待GPU工作完成：

add<<<1, 256>>>(N, x, y);

cudaDeviceSynchronize();

• 我们释放GPU托管内存并从main返回：

cudaFree(x);

cudaFree(y);

return 0;

然后，我们编译并对其运行性能分析器：

$ nvcc add_block.cu -o add_block

$ nvprof ./add_block

与之前的性能分析相比，速度有了很大的提升。

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_128_0.png

图3.10：对内核进行并行化（add<<<blocks,256>>>）后的nvprof摘要

在图3.10中，我们可以看到与单线程基线相比，时间显著下降；现在大多数开销来自分配（cudaMallocManaged）和同步（cudaDeviceSynchronize），内核通过cudaLaunchKernel启动。统一内存显示有24次CPU缺页错误。

我们能否进一步提高速度以获得更高性能？让我们开始吧。

使用多个块进行扩展

CUDA GPU由多个并行处理器组成，称为流式多处理器（SM），每个处理器能够并发运行多个线程块。例如，基于Ampere架构的NVIDIA A100 GPU拥有108个SM，每个SM最多支持2048个活动线程。为了充分利用GPU的并行处理能力，内核应该以多个线程块启动。

执行配置（<<<numBlocks, blockSize>>>）中的第一个参数定义线程块数量，第二个参数指定每块线程数。这些线程块共同构成了所谓的网格。给定需要处理的N个元素和每块256个线程，所需的块数可按如下方式计算：

numBlocks = (N + blockSize - 1) / blockSize

这确保所有元素都被处理，即使N不是blockSize的整数倍。让我们使用以下设置：

int blockSize = 256;

int numBlocks = (N + blockSize - 1) / blockSize;

add<<<numBlocks, blockSize>>>(N, x, y);

使用扁平线程索引为每个线程分配唯一的工作：i = blockIdx.x * blockDim.x + threadIdx.x。图3.11展示了gridDim.x = 4096和blockDim.x = 256时的映射关系；例如，blockIdx.x = 2, threadIdx.x = 3 → i = 2*256 + 3 = 515。

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_129_0.png

图3.11：计算全局线程索引——对于gridDim.x = 4096和blockDim.x = 256，每个块贡献256个连续索引；全局索引为 i = blockIdx.x * blockDim.x + threadIdx.x

上图展示了在CUDA中使用blockDim.x、gridDim.x和threadIdx.x对一维数组进行索引的方法。每个线程首先计算其块的偏移量（blockIdx.x * blockDim.x），然后加上其在块内的位置（threadIdx.x），从而确定其唯一索引。由此得到以下公式：

index = blockIdx.x * blockDim.x + threadIdx.x

这种索引策略是CUDA编程中的基本模式，确保每个线程都能高效地处理数据的不同部分。

使用块和线程更新后的内核代码如下：

__global__

void add(int n, float *x, float *y)

{

    int index = blockIdx.x * blockDim.x + threadIdx.x;

    int stride = blockDim.x * gridDim.x;

    for (int i = index; i < n; i += stride)

        y[i] = x[i] + y[i];

}

完整代码概述如下：

1. 包含CUDA内核。

该内核使用网格跨步循环模式，因此任意数量的线程/块都可以覆盖n个元素：

#include <iostream>

#include <math.h>

// Kernel function to add the elements of two arrays

__global__

void add(int n, float *x, float *y)

{

    int index = blockIdx.x * blockDim.x + threadIdx.x;

    int stride = blockDim.x * gridDim.x;

    for (int i = index; i < n; i += stride)

        y[i] = x[i] + y[i];

}

2. 接下来，编写main()，设置问题规模N，并在统一内存中分配数组。

分配CPU和GPU都可以访问的数组（CUDA统一内存）：

int main(void)

{

    int N = 1<<20;

    float *x, *y;

    // Allocate Unified Memory – accessible from CPU or GPU

    cudaMallocManaged(&x, N*sizeof(float));

    cudaMallocManaged(&y, N*sizeof(float));

3. 我们在主机上初始化输入。

只需将x填充为1，将y填充为2，这样所有位置的正确答案都是3：

    // initialize x and y arrays on the host

    for (int i = 0; i < N; i++) {

        x[i] = 1.0f;

        y[i] = 2.0f;

    }

4. 选择启动配置并运行内核。

选择一个合理的块大小（32的倍数）和足够覆盖N的块数：

    // Run kernel on 1M elements on the GPU

    int blockSize = 256;

    int numBlocks = (N + blockSize - 1) / blockSize;

    add<<<numBlocks, blockSize>>>(N, x, y);

5. 同步、验证并清理。

等待GPU完成，检查结果，然后释放内存：

    // Wait for GPU to finish before accessing on host

    cudaDeviceSynchronize();

    // Check for errors (all values should be 3.0f)

    float maxError = 0.0f;

    for (int i = 0; i < N; i++)

        maxError = fmax(maxError, fabs(y[i]-3.0f));

    std::cout << "Max error: " << maxError << std::endl;

    // Free memory

    cudaFree(x);

    cudaFree(y);

    return 0;

}

为了编译和执行，让我们运行以下命令：

nvcc add_grid.cu -o add_grid

nvprof ./add_grid

对其执行进行性能分析后，结果如图 3.12所示：

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_132_0.png

图 3.12：add_grid 的 nvprof 摘要——grid-stride 内核运行正确；大部分时间花在 cudaMallocManaged 以及内核启动/同步开销上，报告了 24 次统一内存 CPU 缺页

这是运行多个块带来的又一次巨大提速！或者你可能还不觉得这有什么了不起；这是因为该示例足够小，能够快速完成计算，所以你没有看到明显的收益，不过让我们看看接下来的例子。

如何在 CUDA 中进行调试

在 CUDA 中，错误可能出现在三个地方——API 调用、内核启动本身，以及异步执行期间。首先检查每个运行时 API 的返回代码。

检查 CUDA API 调用

大多数 CUDA 函数会返回一个 cudaError_t 错误码。你应该检查返回值：

cudaError_t err = cudaMallocManaged(&a, size);

if (err != cudaSuccess) {

    printf("Error: %s\n", cudaGetErrorString(err));

}

即使所有 API 调用都成功，启动本身仍可能无效（例如，维度错误）。由于内核返回 void，可以使用 cudaGetLastError() 捕获启动时的错误。

检查内核启动错误

内核的返回类型是 void，因此启动本身不会产生 cudaError_t 错误。要检测启动时发生的问题——例如非法的网格/块大小或资源限制——可以使用 cudaGetLastError()（或 cudaPeekAtLastError()）查询运行时，它会为最近一次启动返回一个 cudaError_t 错误；请注意，cudaGetLastError() 同时也会清除错误状态：

someKernel<<<1, -1>>>(); // -1 is not a valid number of threads.

cudaError_t err;

err = cudaGetLastError(); // cudaGetLastError will return the error from above.

if (err != cudaSuccess)

{

    printf("Error: %s\n", cudaGetErrorString(err));

}

成功的启动并不保证执行一定成功。由于内核是异步运行的，需要进行同步（例如 cudaDeviceSynchronize()）并检查结果，才能发现设备端的失败。

检查异步执行错误

由于内核是异步运行的，你必须调用一个同步函数（例如 cudaDeviceSynchronize）并检查其结果，以捕获内核执行期间发生的任何错误：

err = cudaDeviceSynchronize();

if (err != cudaSuccess) {

    printf("Kernel Execution Error: %s\n", cudaGetErrorString(err));

}

为避免重复的样板代码，可以将 CUDA 调用封装在一个检查状态并打印可读错误信息的小型辅助函数/宏中——请在整个练习中都使用这种模式。

CUDA 错误处理函数

为了减少样板代码，可以定义一个宏来封装 CUDA 调用并检查其返回代码——下面是可供后续练习复用的示例：

#include <stdio.h>

#include <assert.h>

inline cudaError_t checkCuda(cudaError_t result)

{

    if (result != cudaSuccess) {

        fprintf(stderr, "CUDA Runtime Error: %s\n",

            cudaGetErrorString(result));

        assert(result == cudaSuccess);

    }

    return result;

}

int main()

{

    checkCuda(cudaDeviceSynchronize());

}

在许多深度学习和高性能计算应用中，数据并不严格是一维的。相反，你可能会处理多维张量——例如具有高度和宽度的图像，或者具有批量大小和通道维度的高阶张量。一维执行配置（例如 <<<gridSize, blockSize>>>）仍然可以通过将数据索引展平到单个轴上来工作，但如果需要跟踪多个维度（例如卷积层中的批量、通道、高度和宽度），这种方法可能会变得笨拙且容易出错。

幸运的是，CUDA 通过 dim3 类型支持二维（2D）和三维（3D）线程块和网格。要定义二维或三维的网格或块，可以这样使用 CUDA 的 dim3 类型：

dim3 threads_per_block(16, 16, 1);

dim3 number_of_blocks(16, 16, 1);

someKernel<<<number_of_blocks, threads_per_block>>>();

<<<>>> 实例中的顺序很重要。每个块的最大线程总数为 1024。以下是每个维度的最大值表。

维度	最大值
网格大小 dGrid.x	2^31 - 1
网格大小 dGrid.y	65535
网格大小 dGrid.z	65535
块大小 dBlock.x	1024
块大小 dBlock.y	1024
块大小 dBlock.z	64

表 3.1：CUDA 最大启动维度和每块线程限制

CUDA 错误会在三个地方出现：API 调用、内核启动和异步执行。始终检查每个运行时调用的返回值；在 <<<grid, block>>> 之后，调用 cudaGetLastError() 来捕获启动问题，然后调用 cudaDeviceSynchronize() 并再次检查以发现设备端故障。使用一个小型错误检查辅助函数/宏来保持一致性。对于索引和性能，在 2D/3D 网格/块中优先使用 dim3；请记住顺序是 <<<grid, block>>>，并遵守每块 1,024 个线程的限制（参见表 3.1）。

摘要

本章介绍了 CUDA 作为 NVIDIA 的并行模型，用于在 GPU 上运行海量轻量级线程，将其与仅使用 CPU 的执行进行了对比，并展示了各自何时更胜一筹。你搭建了工作环境（驱动程序 + CUDA 工具包），对其进行了验证（使用 nvidia-smi、nvcc --version 和示例运行），并了解了在 Colab 中使用 nvcc4jupyter 的快速路径。我们建立了一个 CPU 基线（通过数组相加内核），然后将其移植到 CUDA：编写内核、选择启动配置 (<<<grid, block>>>)，并使用每线程索引和 grid-stride 循环将迭代映射到线程。你学习了 CUDA 执行层次结构（线程 | 块 | 网格）和内存作用域（线程本地、共享、全局），了解了为什么块大小应为 32（线程束）的倍数，以及如何从一个块扩展到多个块以充分利用 SM。我们使用 nvprof 对运行进行了性能分析（并提到 Nsight Systems/Compute 可进行更深入的分析），解释了统一内存缺页，并识别了常见瓶颈（启动开销、数据传输、同步）。最后，你为 API 调用、启动错误和异步故障添加了健壮的错误处理，并了解了当数据是多维时如何使用 2D/3D dim3 网格/块。

在下一章中，我们将从使用 Eigen 手动实现神经网络（NN）转向使用 LibTorch（PyTorch 官方 C++ API），它为张量、NN 模块、优化器和数据集提供了高级抽象。你将学习如何使用 torch::nn 层定义模型，使用内置优化器训练模型，并利用自动微分而不是手动计算梯度。我们还将通过 LibTorch 引入 CUDA 的 GPU 加速，使相同的模型能够高效扩展到更大的数据集和更深的架构。

问题

检查你的理解——在继续之前回答这些问题：

1. <<<grid, block>>> 中的执行配置指定了什么？
2. 如果多个线程在没有协调的情况下更新同一 i 实例的 y[i]，会出现什么问题？在数组相加内核中我们如何避免这个问题？
3. 说出两个端到端验证 CUDA 安装的命令（一个用于驱动程序，一个用于工具包）。

扩展阅读

以下是一些推荐参考，可帮助你加深 CUDA 调试和优化技能：

• CUDA C 编程指南（NVIDIA）：核心模型、内存、执行层次结构
• CUDA 最佳实践指南（NVIDIA）：占用率、内存合并、流、统一内存技巧
• Nsight Systems 和 Nsight Compute 用户指南：性能分析方法和指标
• CUDA 示例：grid-stride 循环、归约等的参考实现
• Mark Harris，《如何优化 CUDA 中的数据传递》：页锁定内存、重叠
• 关于 Roofline 和 Amdahl 定律的论文/文章，用于推理可实现的加速比
• cuBLAS/cuDNN/cuFFT 文档：何时优先使用库而不是自定义内核

第 3 章

答案

以下是前面问题的答案（请先尝试回答问题后再查看）：

1. 它定义了内核的并行形状：网格中有多少个块以及每块线程数（可选地指定共享内存字节数和流）。

2. 数据竞争；多个线程不可预测地读取/写入同一位置。避免方法是给每个线程分配一个唯一索引（并且只更新它自己的元素）；例如，i = blockIdx.x * blockDim.x + threadIdx.x（或者当多个线程需要合并结果时，使用原子操作/归约）。

3. nvidia-smi（驱动/GPU 可见性）和 nvcc --version（工具包）；编译/运行如 deviceQuery 这样的示例即可完成检查。

获取本书的 PDF 版本及更多内容

扫描二维码（或访问 packtpub.com/unlock）。按名称搜索这本书，确认版本，然后按照页面上的步骤操作。

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_137_0.pnghttps://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_137_1.png

注意：请保管好你的发票。从 Packt 直接购买无需发票。

第 2 部分：在 C++ 中构建和训练神经网络

在本书的这一部分中，你将从核心概念转向在 C++ 中进行动手模型开发。从一个基础神经网络开始，你将逐步构建多层感知机、卷积神经网络、循环神经网络、长短期记忆（LSTM）网络、生成模型以及基础大语言模型组件。在此过程中，你将学习实用的训练工作流程、正则化方法、优化策略，以及通过分布式训练、并行化和模型压缩来扩展工作负载的技术。到本部分结束时，你将能够使用 C++ 设计、训练和优化各种深度学习模型，以应对真实世界的任务。

本书的这一部分包含以下章节：

• 第 4 章，在 C++ 中构建基础神经网络
• 第 5 章，C++ 中的多层感知机
• 第 6 章，C++ 中的卷积神经网络
• 第 7 章，C++ 中的循环神经网络与长短期记忆网络
• 第 8 章，C++ 中的生成网络、自编码器与大语言模型
• 第 9 章，C++ 中的 Transformer 与大语言模型微调

第 4 章：在 C++ 中构建基础神经网络

机器学习（ML）使算法能够发现数据中的模式，并利用这些模式对新出现的未见示例进行预测，而无需为每种情况手工编码规则。神经网络是受大脑（粗略地）启发的一类 ML 模型。在 图 4.1 中，生物神经元通过树突接收信号，在细胞体中整合信号，并沿轴突发出输出。感知机——神经网络的基本构建块——与此对应：输入 x_i 通过具有强度（权重）w_i 的 突触 到达，组合成加权和 z = Σ w_i x_i + b（偏置 b 的作用类似于阈值），然后通过激活函数 y = σ(z)（例如 ReLU、sigmoid、tanh）产生输出。你可以回看本书的 第 1 章 来复习这些内容。

突触前神经元通过突触沿轴突将信号发送到突触后神经元的树突。具有突触强度 w_i 的输入 x_i 在胞体中被整合为 Σ w_i x_i。如果足够强，就会产生一个输出脉冲——这正是感知机加权和与激活函数的灵感来源。改编自 Nature：https://www.nature.com/articles/s41427-018-0101-y。

将许多感知机堆叠起来形成层（输入 → 隐藏 → 输出），使网络能够逼近复杂的非线性映射：从像素到标签、从特征到概率、从文本到嵌入向量，等等。我们通过损失函数度量误差，并通过梯度下降和反向传播调整 {w_i, b}，从而使网络未来的预测不断改进。

在接下来的篇幅中，我们将把这个示意图变为代码：从单个感知机开始，添加常见的激活函数和损失函数，在 C++ 中组装一个最小的多层感知机（MLP），并在一个玩具数据集上训练它——为后续更深的网络和 GPU 加速打下基础。

在本章中，我们将介绍以下主题：

• 理解基础知识
• 构建前向传播与损失函数（例如用于分类的 BCE/交叉熵）
• 使用 LibTorch 在 C++ 中实现神经元和激活函数
• 推导/反向传播梯度，并使用随机梯度下降更新权重

到本章结束时，你将拥有一个自包含的 C++ 程序，用于训练基本的感知机——包括前向传播、反向传播和优化循环——并且有信心扩展它（更多层/激活函数、正则化、更好的初始化），之后继续迈向 GPU 加速。

理解基础知识

机器学习的核心是学习：通过调整内部参数——权重和偏置——来捕获训练数据中的模式，并能泛化到新的未见示例。我们从数据集开始，它是一组示例的集合；每个示例包含特征（可测量的输入），在监督学习中还包含标签或目标（期望的输出）。在训练期间，算法迭代地调整参数，以减小衡量训练集预测误差的损失。之后，我们在独立的测试（或验证）集上进行评估，以估计学到的关系在模型训练期间未见过的数据上的泛化能力。

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_142_0.png

图 4.1：线性回归与逻辑回归对比。左：对带噪声的一维数据进行线性回归，最小二乘拟合为 ŷ = 0.77x + 0.68。右：对两个二维类别进行逻辑回归，决策边界为 p = 0.5（σ(w^T x + b) = 0.5）。线性模型最小化 MSE；逻辑模型通过 sigmoid 输出概率，并使用二元交叉熵进行训练。

单个神经元本质上是一个仿射模型加上可选的非线性。在线性回归（图 4.1）中，我们用 ŷ = w_0 + Σ_j w_j x_j 预测一个标量，并通过使用诸如MSE（均方误差）的损失函数最小化预测值（ŷ）与正确标签（y）之间的差异来学习 {w_j}。当然，对于较小的问题，你可以使用正规方程求得闭式解，但随着数据规模和特征数量的增加，正规方程不再可行；相反，我们使用梯度下降——这也是我们将在深度神经网络中依赖的优化器。

逻辑回归（图 4.1）通过将 sigmoid 应用于线性得分 ŷ = σ(z) = 1 / (1 + e^{-z})（其中 z = w_0 + Σ_j w_j x_j）将输出转换为概率，并使用二元交叉熵损失函数训练神经元。二元逻辑回归的决策边界是超平面 w^T x + b = 0（等价于 y = 0）；非线性边界来自特征变换或堆叠层。

在深入神经网络之前，让我们先说明如何在 C++ 中实现线性和逻辑回归。本节将在更简单的上下文中介绍权重更新和代价函数的概念。

构建前向传播和损失函数（例如用于分类的 BCE/交叉熵）

单个神经元计算 z = w x + b，然后应用激活函数 ŷ = σ(z)。使用恒等激活函数（σ(z) = z），得到线性回归，通过最小化MSE来训练。使用sigmoid（σ(z) = 1 / (1 + e^{-z})），得到逻辑回归，使用二元交叉熵进行训练。你可以将这两种情况视为神经网络的最简形式或特例：一个神经元、一个输出。掌握它们后，过渡到多层网络就会很自然——在各层之间堆叠相同的构建块，并对所有层进行反向传播。

让我们使用梯度下降在 C++ 中实现线性回归。

1. 导入合成数据集：我们将对从 y = 2.5x + 0.9 + ε 采样得到的带噪声的点拟合一条直线。

int main() {

    mt19937 rng(42);

    uniform_real_distribution<double> U(0.0, 1.0);

    normal_distribution<double> N(0.0, 0.2);

    // Synthetic data

    const int n = 200;

    vector<double> x(n), y(n);

在C++中构建基础神经网络

for (int i = 0; i < n; ++i) {

    x[i] = U(rng);

    y[i] = 2.5 * x[i] + 0.7 + N(rng);

}

• 2. 初始化参数和超参数： 初始化权重w和偏置b（初始化为0即可），并设置学习率和迭代周期。

double w = 0.0; // weight

double b = 0.0; // bias

double lr = 0.1; // learning rate

const int epochs = 1000;

• 3. 对于每个迭代周期，我们将数据遍历一次。 对于每个数据点 (x, y)：
  • 预测： y_hat = w * x + b。
  • 单样本MSE： mse = (y_hat - y)^2。
  • 梯度： ∂mse/∂w = 2 * (y_hat - y) * x，∂mse/∂b = 2 * (y_hat - y)。
  • 更新参数。

for (int i = 0; i < n; ++i) {

    double y_hat = w * x[i] + b; // prediction

    double resid = y_hat - y[i]; // residual

    double mse_i = resid * resid; // per-sample MSE (for clarity/logging if desired)

    w -= lr * (2.0 * resid * x[i]); // gradient step from MSE

    b -= lr * (2.0 * resid);

}

既然我们已经完成了线性回归，接下来转向它的概率近亲——逻辑回归——它通过 sigmoid 函数将线性得分映射为类别概率。我们将用一个 sigmoid 神经元对两个二维高斯簇进行分类来演示它。

• 1. sigmoid函数可以如下创建：

static inline double sigmoid(double z) {

    return 1.0 / (1.0 + std::exp(-z));

}

• 2. 创建简单的线性可分数据集： 两个簇：类别0在(-2, -2)附近，类别1在(2, 2)附近。

// Synthetic 2D dataset (two Gaussian blobs)

std::mt19937 rng(123);

std::normal_distribution<double> N0x(-2.0, 1.0), N0y(-2.0, 1.0);

std::normal_distribution<double> N1x(2.0, 1.0), N1y(2.0, 1.0);

const int n_per = 200;

std::vector<std::array<double,2>> X;

std::vector<int> y;

X.reserve(2 * n_per);

y.reserve(2 * n_per);

for (int i = 0; i < n_per; ++i) {

    X.push_back({N0x(rng), N0y(rng)}); y.push_back(0);

    X.push_back({N1x(rng), N1y(rng)}); y.push_back(1);

}

const int n = static_cast<int>(X.size());

• 3. 初始化参数和超参数： 一个逻辑回归神经元具有权重w1、w2和偏置b。我们将使用固定的学习率。

double w1 = 0.0, w2 = 0.0, b = 0.0;

double lr = 0.1;

const int epochs = 3000;

• 4. 梯度下降循环（交叉熵损失）： 对于每个迭代周期，计算二分类交叉熵的梯度。对于样本 (x, y)，其中 z = w^T x + b，p = σ(z)，

• 5. ∂BCE/∂z = p - y。链式法则得到：∂BCE/∂w = (p - y) * x，∂BCE/∂b = (p - y)。我们对整个数据集的梯度求平均，然后更新参数。

for (int epoch = 1; epoch <= epochs; ++epoch) {

    double gw1 = 0.0, gw2 = 0.0, gb = 0.0;

    for (int i = 0; i < n; ++i) {

        double z = w1 * X[i][0] + w2 * X[i][1] + b; // linear score

        double p = sigmoid(z); // probability (ŷ)

        double diff = (p - y[i]); // d(BCE)/dz

        gw1 += diff * X[i][0] / n;

        gw2 += diff * X[i][1] / n;

        gb += diff / n;

    }

    w1 -= lr * gw1;

    w2 -= lr * gw2;

    b -= lr * gb;

}

• 6. 用BCE监控训练： 计算平均二分类交叉熵来跟踪训练进度。

if (epoch % 500 == 0) {

    double loss = 0.0;

    for (int i = 0; i < n; ++i) {

        double p = sigmoid(w1*X[i][0] + w2*X[i][1] + b);

        p = std::clamp(p, 1e-12, 1.0 - 1e-12); // numeric safety

        loss += -( y[i]*std::log(p) + (1 - y[i])*std::log(1 - p) );

    }

    loss /= n;

    std::cout << "Epoch " << epoch << " BCE=" << loss

              << " w=[" << w1 << "," << w2 << "] b=" << b << "\n";

}

• 7. 评估准确率并展示决策边界： 我们使用0.5阈值进行分类；决策边界为 w1 * x1 + w2 * x2 + b = 0。

int correct = 0;

for (int i = 0; i < n; ++i) {

    double p = sigmoid(w1*X[i][0] + w2*X[i][1] + b);

    int pred = (p >= 0.5);

    if (pred == y[i]) correct++;

}

std::cout << "\nAccuracy: " << (100.0 * correct / n) << "%\n";

std::cout << "Decision boundary (p=0.5): w1*x1 + w2*x2 + b = 0\n";

线性回归和逻辑回归是很好的基线——快速、简单且可解释——但它们只能建模线性关系。除非你手工设计特征（多项式、交互项、核函数），否则它们的决策边界仍然是仿射的，在复杂数据上很快就会碰壁。特征工程可以有所帮助，但它是手动的、脆弱的，而且随着模式变得高阶或非平稳，很难扩展。

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_147_0.png

图4.2：深度神经网络中的层次化特征学习

从原始像素出发，早期层检测明暗模式和边缘；中层将边缘组合成图案和部件；后续层组装出高级特征。输出层将这些学习到的表示映射为类别标签。

图4.2改编自Fraija-Fernández N.等人发表于《Frontiers in Environmental Science》3:80 (2016)的《用于监测海洋生态系统的环境DNA宏条形码：现状与未来展望综述》中的图2。 https://doi.org/10.3389/fenvs.2015.00080 ——采用知识共享署名许可（CC BY 4.0）。

如图4.2所示，深度网络通过端到端学习特征来避免手动特征工程：堆叠带非线性激活的仿射层，可以产生丰富的决策边界和层次化表示。这种能力支撑了计算机视觉、自然语言处理和强化学习领域的最先进成果。在本章剩余部分，我们将用C++和LibTorch实现一个MLP，让你的模型可以直接从数据中学习这些变换。

使用LibTorch在C++中实现神经元和激活函数

一个神经元只是一个仿射变换加上（可选的）非线性：z = W x + b, y = σ(z)。在LibTorch中，你可以使用内置的torch::nn::Linear层来实现W x + b，并通过torch::relu、torch::sigmoid、torch::tanh或torch::nn::ReLU等函数/模块应用激活。在neuron_demo.cpp中，我们提供了一个包含一个仿射层和激活（ReLU）的简单架构。

// neuron_demo.cpp

#include <torch/torch.h>

#include <iostream>

int main()

{

    torch::manual_seed(42);

    // Pick device (CUDA if available)

    torch::Device device = torch::cuda::is_available() ? torch::kCUDA : torch::kCPU;

    std::cout << "Using device: " << (device.is_cuda() ? "CUDA" : "CPU") << "\n";

    const int64_t batch = 4, in_f = 5, out_f = 3;

    auto x = torch::randn({batch, in_f}).to(device);

    torch::nn::Linear fc(in_f, out_f);

    fc->to(device);

    torch::nn::init::kaiming_uniform_(fc->weight, /*a=*/std::sqrt(5.0));

    torch::nn::init::zeros_(fc->bias);

    auto z = fc->forward(x); // [batch, out_f]

    auto y = torch::relu(z); // activation

    std::cout << "Input shape : " << x.sizes() << "\n";

    std::cout << "Logits shape: " << z.sizes() << "\n";

    std::cout << "Output shape: " << y.sizes() << "\n";

}

第 4 章：在 C++ 中构建基础神经网络

std::cout << "Output sample:\n"

          << y.slice(/*dim=*/0, 0, 2) << "\n";

return 0;

}

你刚刚构建了一个带有非线性激活（如 ReLU）的单一仿射层。很自然的下一步是将这些层堆叠起来：第一层学习简单模式，下一层将它们组合成更丰富的特征，最后一层将这些特征映射到预测。这就是 MLP——最小的深度网络——也是我们在用 LibTorch/CUDA 扩展之前所使用的骨架。

为了让数学和代码更透明，我们将从一个极小的 MLP 和一个具有 2 个特征的玩具数据集开始：

• 架构：

  • 输入层：2 个单元
  • 隐藏层：3 个单元
  • 输出层：1 个单元（二分类或简单回归）

  https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_149_0.png

  图 4.3：我们将实现的 MLP 架构

• 数据流：输入向量与输入到隐藏层的权重相乘，再加上偏置；我们应用 sigmoid 得到隐藏激活。这些激活再与隐藏层到输出层的权重相乘，再次加上偏置，并通过另一个 sigmoid 产生输出。

• 训练与损失：为简单起见，我们将使用 MSE 通过 梯度下降 进行训练（后续分类任务可换成交叉熵）。梯度通过 反向传播 计算，并用于更新两组权重和偏置。

我们将使用 Eigen 进行矩阵运算，以从头展示其机制。一旦你明白了这一点，就可以用同样的逻辑将张量和模块替换为 LibTorch。

我们先创建 NeuralNetwork 类（头文件）。这个头文件列出了模型的参数和API。我们存储两个权重矩阵（[hidden, input] 和 [output, hidden]）和两个偏置向量（每层一个）。激活函数是 sigmoid，通过 Eigen 逐元素实现。公有方法包括：构造函数用于设置大小和学习率，feedforward 用于预测，train 用于在单个样本上执行一步学习。（所有向量都是列向量；Eigen 的 VectorXd 默认是列向量。）

#ifndef NEURAL_NETWORK_H

#define NEURAL_NETWORK_H

#include <Eigen/Dense>

#include <cmath>

using namespace Eigen;

class NeuralNetwork {

private:

    MatrixXd weights_input_hidden; // [hidden, input]

    MatrixXd weights_hidden_output; // [output, hidden]

    VectorXd bias_hidden; // [hidden]

    VectorXd bias_output; // [output]

    double learning_rate;

    // Sigmoid activation applied elementwise

    VectorXd sigmoid(const VectorXd& x) {

        return 1.0 / (1.0 + (-x.array()).exp());

    }

    // Derivative of sigmoid given its output (y = sigmoid(x))

    VectorXd sigmoid_derivative(const VectorXd& y) {

        return y.array() * (1.0 - y.array());

    }

public:

    NeuralNetwork(int input_size, int hidden_size, int output_size, double lr);

    // Forward pass for a single example

    VectorXd feedforward(const VectorXd& input);

    // One training step on a single (input, target) pair

    void train(const VectorXd& input, const VectorXd& target);

};

#endif

接下来，我们将实现构造函数。我们随机初始化权重和偏置，这样隐藏单元就不会全都学到相同的东西（对称破缺）。learning_rate（double）控制每个梯度步长的大小。出于教学目的，我们使用 Random()；对于更深的网络，请考虑 Xavier/He 初始化。

NeuralNetwork::NeuralNetwork(int input_size, int hidden_size, int output_size, double lr) {

    learning_rate = lr;

    // Random init (for teaching; in practice consider Xavier/He)

    weights_input_hidden = MatrixXd::Random(hidden_size, input_size);

    weights_hidden_output = MatrixXd::Random(output_size, hidden_size);

    bias_hidden = VectorXd::Random(hidden_size);

    bias_output = VectorXd::Random(output_size);

}

对于前向传播，给定输入向量 x，我们计算隐藏层预激活 (z_1 = W_1 x + b_1)，用 sigmoid 压缩得到 (a_1 = \sigma(z_1))，然后计算输出层预激活 (z_2 = W_2 a_1 + b_2)，再次应用 sigmoid 得到最终预测。这与数学图示一致，也正好为反向传播做好了准备。

VectorXd NeuralNetwork::feedforward(const VectorXd& input) {

    // Hidden layer: z1 = W1*x + b1; a1 = sigmoid(z1)

    VectorXd hidden_input = weights_input_hidden * input + bias_hidden;

    VectorXd hidden_output = sigmoid(hidden_input);

    // Output layer: z2 = W2*a1 + b2; a2 = sigmoid(z2)

    VectorXd final_input = weights_hidden_output * hidden_output + bias_output;

    VectorXd final_output = sigmoid(final_input);

    return final_output;

}

最后，反向传播使我们能够把误差转化为参数更新。训练要做两件事：一次前向传播获得预测，然后一次反向传播把损失信号转化为梯度。对于 MSE 和 sigmoid 输出，输出层 delta (\delta_{\text{out}}) 为 ((y - \hat{y}) \odot \sigma'(\hat{y}))；隐藏层 delta (\delta_{\text{hid}}) 为 ((W_{\text{ho}}^\top \delta_{\text{out}}) \odot \sigma'(a_{\text{hid}}))。然后我们执行简单的随机梯度下降（SGD）更新：在每个权重矩阵上加上 (\eta \times \nabla_W)，在每个偏置上加上 (\eta \times \nabla_b)。如果你不太理解下面的内容，别担心，我们将在下一节更详细地解释反向传播。

void NeuralNetwork::train(const VectorXd& input, const VectorXd& target) {

    // 1) Forward

    VectorXd hidden_input = weights_input_hidden * input + bias_hidden;

    VectorXd hidden_output = sigmoid(hidden_input);

    VectorXd final_input = weights_hidden_output * hidden_output + bias_output;

    VectorXd final_output = sigmoid(final_input);

    // 2) Error (MSE-style)

    VectorXd output_error = target - final_output; // size: [output]

    // 3) Backprop deltas (use deltas, not raw error)

    VectorXd output_delta = sigmoid_derivative(final_output).cwiseProduct(output_error);

    VectorXd hidden_error = weights_hidden_output.transpose() * output_delta;

    VectorXd hidden_delta = sigmoid_derivative(hidden_output).cwiseProduct(hidden_error);

    // 4) SGD updates (+= because delta uses (target - output))

    weights_hidden_output += learning_rate * (output_delta * hidden_output.transpose());

    bias_output += learning_rate * output_delta;

    weights_input_hidden += learning_rate * (hidden_delta * input.transpose());

    bias_hidden += learning_rate * hidden_delta;

}

有了前向传播和合适的损失函数（BCE/交叉熵），我们现在就可以反向传播这些误差并搭建训练循环——将原始预测转变为一个真正会学习的模型。

推导/反向传播梯度并用随机梯度下降更新权重

你已经见识过反向传播的实际应用。其核心就是链式法则在网络中反向运行：前向传播计算出预测值 (\hat{y}) 和损失 (L = (y - \hat{y})^2) 之后，一次反向传播会复用中间值，高效地获得每一层的梯度，如 图 4.4 所示。具体来说，从输出端开始，形成一个误差信号（即“delta”），然后通过权重的转置和激活函数的导数将其向后传递，从而获得更早层的 delta。

https://github.com/OpenDocCN/freelearn-dl-zh/tree/master/docs/deep-learning-with-c/img/f45df257dfb11cd9b58ac2c5f6935210_153_0.png

图 4.4：围绕输出预激活 (z_2) 的局部反向传播

对于 MSE 和 sigmoid 输出，两层 MLP 的 delta 和梯度如下：

• 输出层 delta：(\delta_{\text{out}} = (y - \hat{y}) \odot \sigma'(\hat{y}))

• 隐藏层 delta：(\delta_{\text{hid}} = (W_{\text{ho}}^\top \delta_{\text{out}}) \odot \sigma'(a_{\text{hid}}))

• 梯度（外积）：

[
\frac{\partial L}{\partial W_{\text{ho}}} = \delta_{\text{out}} a_{\text{hid}}^\top, \quad
\frac{\partial L}{\partial b_{\text{o}}} = \delta_{\text{out}}
]

[
\frac{\partial L}{\partial W_{\text{ih}}} = \delta_{\text{hid}} x^\top, \quad
\frac{\partial L}{\partial b_{\text{h}}} = \delta_{\text{hid}}
]

梯度下降随后会以小步长（即学习率 lr）沿梯度的反方向微调参数。根据我们这里的 δ 定义（使用 y = sigmoid(z)），这些更新表现为加法形式：

W ← W + lr × (δ × a^⊤), b ← b + lr × δ.

下面是一个单样本反向传播 + SGD（MSE + sigmoid）示例：

static VectorXd sigmoid(const VectorXd& v) {

    return 1.0 / (1.0 + (-v.array()).exp());

}

static VectorXd sigmoid_deriv_from_output(const VectorXd& y) { // y = sigmoid(z)

    return y.array() * (1.0 - y.array());

}

double train_one_mse(const VectorXd& x, const VectorXd& t) {

    // ----- Forward -----