© 2026 async. fork自：GitHub
主要介绍并行计算与 CUDA 编程的基础概念。文章首先区分串行、并发与并行三种程序执行方式，并说明 CPU 多核并行和 GPU 大规模线程并行的差异。随后以向量加法为例，讲解 CUDA 程序从数据准备、显存申请、数据传输、核函数执行到结果拷回的完整流程，并进一步介绍 Grid、Block、Thread 线程层级、SIMT 执行模型、GPU 硬件单元以及 nvcc 编译流程。最后结合 Nsight Systems 分析程序性能，引出主机与设备间数据传输优化的重要性。
CPU 的并行主要通过多核来执行。一台现代的 CPU 通常有多个核心，每个核心都能独立执行指令，这样可以把不同的任务分配到不同的核心上同时运行。如操作系统、Web 服务等。
GPU 的并行是通过将一个大任务拆分成海量相似的小任务，让许多计算单元同时做同样模式的工作。GPU 的内部拥有大量更轻量的执行单元，通过牺牲单个线程的灵活性换来极强的批量吞吐能力。
GPU 的硬件单元指的是流式多处理器（Streaming Multiprocessor， SM），是 GPU 内部的核心工作单元，GPU 会将所有的线程分配到一个个 SM 上执行。其中单个 SM 中又主要有 CUDA Core 和 Tensor Core：
CUDA Core 通常处理的是一个线程的一条普通算术指令，如整数运算、浮点运算、加减乘除、逻辑判断、地址计算以及一些通用指令等，而 Tensor Core 主要用于专门加速矩阵乘加操作的，其最典型的操作即：
也就是矩阵乘法再加累加。这类操作在深度学习里极其常见，因为全连接层本质上是矩阵乘法，卷积在很多实现里也会转成矩阵乘法，Transformer 里的 attention、MLP 也高度依赖 GEMM。
CPU 首先进行程序初始化、数据准备等工作，而后将数据从主机内存拷贝至设备全局内存。而后 SM 读取全局内存中的数据，并行执行计算任务，再将结果写回全局内存，最后将结果从设备全局内存拷贝回主机内存。
CPU 加法的过程主要如下：准备和初始化数据，定义加法函数，依靠循环来进行所有的元素加法，调用函数后验证结果。
CPU先准备和初始化数据：
然后将数据传输到 GPU，类似于 C 语言中的申请堆内存，GPU 中的内存也需要先申请再使用：
内存申请成功后传输数据：
而后 GPU 从全局内存中读取并计算后写回。这里的计算是要写在 GPU 上运行的加法函数并进行调用，在 CPU 的加法运算中使用了循环，而在 GPU 中要使用的是 SIMT，SIMT 指挥每个线程，需要组织结构和编号。在 CUDA 中主要是 Grid→Block→Thread 的层级关系。
__syncthreads() 做块内同步，还会被分配到同一个 SM 上执行，SM 上可以同时驻留多个 Block。不同的 Block 之间一般不直接同步。Block 还有自己的索引：blockIdx.x,blockIdx.y,blockIdx.z，每个 Block 内部的线程也有自己的局部编号：threadIdx.x,threadIdx.y,threadIdx.z。Block 的尺寸在进行 kernel lauch 时指定，如 dim3 blockDim(256) 和 dim3 blockDim(16,16) 都是表示每个 block 有 256 个线程。kernel<<gridDim, blockDim>>(...); 里面的 gridDim 表示 Block 的数量，blockDim 表示每个 Block 中有多少个 Thread。因此就需要定义 Block 的数量和大小来指挥线程并行计算，还要定义 GPU 上的加法函数，结合定义的信息调用 GPU 加法函数。
其中核函数定义为：
其中的 dim3 是 CUDA 表示线程层级结构的类型，<<>> 中的内容传递线程层级信息给核函数，核函数（kernel）即设备侧的入口函数，__global__ 表示这是一个核函数，最后将结果拷回主机：
完整的代码如下：
最后通过以下命令编译运行即可：
nvcc 是 CUDA 的编译器，是 CUDA Toolkit 的一部分。其中编译的流程大致如下：
nvcc 将同一份 .cu 源文件拆分成两条路径来处理：一条是 host 路径，用于生成 CPU 侧的目标代码，另一条是 device 路径，用于生成 GPU 侧的目标代码。具体而言，nvcc 会先对源文件中的 device 部分进行处理，生成 PTX 或 cubin 这类设备端中间结果，并将其封装进 fatbinary；随后再对 host 部分重新预处理，把 CUDA 的扩展语法转换成宿主编译器能够理解的标准 C/C++ 形式，同时把前面生成的 fatbinary 嵌入进去，最后再交给 g++、clang++ 或 MSVC 这样的 host compiler 生成 host object。
英伟达提供了一个系统级性能分析工具 Nsight Systems，便于对编写的 CUDA 程序进行性能评估。
通过以上命令就能够对运行的程序进行性能评估。观察得到的结果是 HtoD 和 DtoH 的耗时远远大于核函数的时间。设定的指标为：
要进一步优化，就需要从 HtoD 和 DtoH 这两个时间入手。
一般而言，启用越多的线程，就能获得越大的并行度，越大的并行度也就能获取越好的性能。一个设备最多能够启用的线程数量可以通过网络/官网查询获取，也可以通过代码获取，NVIDIA 提供了 cudaDeviceProp 来通过代码查询。
#include <stdio.h>
#include <stdlib.h>
#include <vector>
using namespace std;
void add_cpu(vector<float>& c, const vector<float>& a, const vector<float>& b){
for(size_t i = 0; i < c.size(); i++){
c[i] = a[i] + b[i];
}
}
int main(){
const size_t SIZE = 1 << 20;
std::vector<float> a(SIZE, 1.0f);
std::vector<float> b(SIZE, 2.0f);
std::vector<float> c(SIZE, 0.0f);
add_cpu(c, a, b);
return 0;
}
const size_t SIZE = 1 << 20; // 1M elements
size_t bytes = SIZE * sizeof(float);
std::vector<float> h_a(SIZE, 1.0f);
std::vector<float> h_b(SIZE, 2.0f);
std::vector<float> h_c(SIZE, 0.0f);
float *d_a, *d_b, *d_c;
CUDA_CHECK(cudaMalloc((void**)&d_a, bytes));
CUDA_CHECK(cudaMalloc((void**)&d_b, bytes));
CUDA_CHECK(cudaMalloc((void**)&d_c, bytes));CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyHostToDevice)); int threadsPerBlock = 256;
int blocksPerGrid = (SIZE + threadsPerBlock - 1) / threadsPerBlock;
// 启动 kernel
vec_add<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, SIZE);
// 检查 kernel 启动错误
CUDA_CHECK(cudaGetLastError());
// 等待 GPU 执行完成
CUDA_CHECK(cudaDeviceSynchronize());
__global__ void vec_add(const float* a, const float* b, float* c, size_t n) {
size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}// 把结果拷回主机
CUDA_CHECK(cudaMemcpy(h_c.data(), d_c, bytes, cudaMemcpyDeviceToHost));
// 简单验证
for (size_t i = 0; i < 10; ++i) {
printf("h_c[%zu] = %f\n", i, h_c[i]);
}
CUDA_CHECK(cudaFree(d_a));
CUDA_CHECK(cudaFree(d_b));
CUDA_CHECK(cudaFree(d_c));
return 0;
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <vector>
#define CUDA_CHECK(call) \
do { \
cudaError_t err__ = (call); \
if (err__ != cudaSuccess) { \
fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \
cudaGetErrorString(err__)); \
exit(EXIT_FAILURE); \
} \
} while (0)
// GPU kernel：逐元素相加
__global__ void vec_add(const float* a, const float* b, float* c, size_t n) {
size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
int main() {
const size_t SIZE = 1 << 20; // 1M elements
size_t bytes = SIZE * sizeof(float);
std::vector<float> h_a(SIZE, 1.0f);
std::vector<float> h_b(SIZE, 2.0f);
std::vector<float> h_c(SIZE, 0.0f);
float *d_a, *d_b, *d_c;
CUDA_CHECK(cudaMalloc((void**)&d_a, bytes));
CUDA_CHECK(cudaMalloc((void**)&d_b, bytes));
CUDA_CHECK(cudaMalloc((void**)&d_c, bytes));
CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyHostToDevice));
// 配置 kernel 启动参数
int threadsPerBlock = 256;
int blocksPerGrid = (SIZE + threadsPerBlock - 1) / threadsPerBlock;
// 启动 kernel
vec_add<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, SIZE);
// 检查 kernel 启动错误
CUDA_CHECK(cudaGetLastError());
// 等待 GPU 执行完成
CUDA_CHECK(cudaDeviceSynchronize());
// 把结果拷回主机
CUDA_CHECK(cudaMemcpy(h_c.data(), d_c, bytes, cudaMemcpyDeviceToHost));
// 简单验证
for (size_t i = 0; i < 10; ++i) {
printf("h_c[%zu] = %f\n", i, h_c[i]);
}
CUDA_CHECK(cudaFree(d_a));
CUDA_CHECK(cudaFree(d_b));
CUDA_CHECK(cudaFree(d_c));
return 0;
}
nvcc add_cuda.cu -o add_cuda
./add_cuda
nsys profile -t cuda,nvtx,osrt -o add_cuda -f true ./add_cuda
nsys stats add_cuda.nsys-rep#include <iostream>
#include <cuda_runtime.h>
int main() {
int device_id = 0; // 想查询的 GPU 编号
cudaDeviceProp prop;
cudaError_t err = cudaGetDeviceProperties(&prop, device_id);
if (err != cudaSuccess) {
std::cerr << "cudaGetDeviceProperties failed: "
<< cudaGetErrorString(err) << std::endl;
return 1;
}
std::cout << "Device name: " << prop.name << std::endl;
std::cout << "Total global memory: " << prop.totalGlobalMem << std::endl;
std::cout << "SM count: " << prop.multiProcessorCount << std::endl;
std::cout << "Max threads per block: " << prop.maxThreadsPerBlock << std::endl;
std::cout << "Max threads per SM: " << prop.maxThreadsPerMultiProcessor << std::endl;
std::cout << "Max threads dim: (" << prop.maxThreadsDim[0] << ", "
<< prop.maxThreadsDim[1] << ", " << prop.maxThreadsDim[2] << ")"
<< std::endl;
std::cout << "Compute capability: "
<< prop.major << "." << prop.minor << std::endl;
return 0;
}