从CPU到GPU：手把手教你用CUDA在Jetson Nano上加速矩阵乘法（附完整代码）

从CPU到GPU：手把手教你用CUDA在Jetson Nano上加速矩阵乘法（附完整代码）

当你在Jetson Nano这样的嵌入式设备上运行复杂的矩阵运算时，是否曾被CPU的缓慢计算速度困扰？想象一下，一个1000×1000的矩阵乘法在CPU上可能需要数秒才能完成，而同样的计算在GPU上可能只需要几十毫秒。这就是并行计算的魔力——而CUDA正是打开这扇大门的钥匙。

Jetson Nano作为一款强大的边缘计算设备，其内置的128核Maxwell架构GPU提供了惊人的并行计算能力。本文将带你从零开始，不仅理解CUDA并行计算的核心思想，更会手把手教你如何在资源受限的嵌入式环境中实现高效的矩阵乘法加速。不同于普通的教程，我们会特别关注Jetson Nano上的实际限制和优化技巧，比如内存管理、线程配置和编译器特性等实际问题。

1. 环境准备与基础概念

在开始编写代码之前，我们需要确保Jetson Nano上的开发环境已经正确配置。首先检查你的Jetpack版本和CUDA工具包：

cat /etc/nv_tegra_release # 查看L4T版本 nvcc --version # 查看CUDA编译器版本

Jetson Nano的内存架构有其特殊性。它采用共享内存设计，CPU和GPU物理上共用同一块内存。这种统一内存架构(Unified Memory)带来了便利，但也需要特别注意内存分配策略：

• 页锁定内存(Pinned Memory)：使用cudaMallocHost分配，可提高传输速度
• 设备内存(Device Memory)：使用常规cudaMalloc分配
• 内存拷贝开销：在嵌入式设备上尤为明显，需要尽量减少Host-Device传输

理解CUDA的线程组织模型是并行编程的基础。CUDA使用网格(Grid)、线程块(Block)和线程(Thread)的三层结构：

2. CPU串行实现：基准与对比

我们先实现一个标准的CPU矩阵乘法作为性能基准。这个三重循环版本虽然简单，但能清晰展示计算复杂度：

void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) { for (int i = 0; i < m; ++i) { // 遍历A的行 for (int j = 0; j < k; ++j) { // 遍历B的列 int tmp = 0; for (int h = 0; h < n; ++h) { // 内积计算 tmp += h_a[i * n + h] * h_b[h * k + j]; } h_result[i * k + j] = tmp; } } }

这个实现的时间复杂度是O(m×n×k)，对于1000×1000的矩阵，需要10亿次乘加运算。在Jetson Nano的CPU上，这可能需要几秒钟的时间。

性能测试数据对比：

注意：实际性能会受内存布局、缓存利用等多种因素影响。在嵌入式设备上，合理的内存管理往往比单纯算法优化更重要。

3. CUDA并行实现详解

现在让我们转向GPU并行实现。CUDA核函数是运行在GPU上的并行函数，以下是一个基本的矩阵乘法核函数：

__global__ void gpu_matrix_mult(int *a, int *b, int *c, int m, int n, int k) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; if(col < k && row < m) { int sum = 0; for(int i = 0; i < n; i++) { sum += a[row * n + i] * b[i * k + col]; } c[row * k + col] = sum; } }

这个核函数的调用方式很关键，需要合理配置线程网格：

dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); dim3 dimGrid((k + BLOCK_SIZE - 1) / BLOCK_SIZE, (m + BLOCK_SIZE - 1) / BLOCK_SIZE); gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k);

在Jetson Nano上，由于硬件限制，BLOCK_SIZE的选择尤为重要。经过测试，16×16的块大小通常能获得最佳性能：

• 太小：无法充分利用GPU的多处理器
• 太大：可能导致寄存器溢出和共享内存不足

Jetson Nano特有的优化技巧：

1. 使用cudaMallocHost分配页锁定主机内存，提高传输速度
2. 减少不必要的主机-设备数据传输
3. 适当增加每个Block的线程数(但不超过硬件限制)
4. 利用共享内存(Shared Memory)减少全局内存访问

4. 完整实现与常见问题

以下是完整的矩阵乘法实现，包含了错误处理和资源清理：

#include <stdio.h> #include <stdlib.h> #include <math.h> #include <cuda_runtime.h> #define BLOCK_SIZE 16 #define CHECK(call) { \ const cudaError_t error = call; \ if (error != cudaSuccess) { \ printf("Error: %s:%d, ", __FILE__, __LINE__); \ printf("code:%d, reason: %s\n", error, cudaGetErrorString(error)); \ exit(1); \ } \ } __global__ void gpu_matrix_mult(int *a, int *b, int *c, int m, int n, int k) { // 核函数实现同上 } int main() { int m = 1024, n = 1024, k = 1024; int *h_a, *h_b, *h_c, *h_c_ref; int *d_a, *d_b, *d_c; // 分配页锁定主机内存 CHECK(cudaMallocHost((void**)&h_a, m*n*sizeof(int))); CHECK(cudaMallocHost((void**)&h_b, n*k*sizeof(int))); CHECK(cudaMallocHost((void**)&h_c, m*k*sizeof(int))); CHECK(cudaMallocHost((void**)&h_c_ref, m*k*sizeof(int))); // 初始化数据 for(int i = 0; i < m*n; i++) h_a[i] = rand()%100; for(int i = 0; i < n*k; i++) h_b[i] = rand()%100; // 分配设备内存 CHECK(cudaMalloc((void**)&d_a, m*n*sizeof(int))); CHECK(cudaMalloc((void**)&d_b, n*k*sizeof(int))); CHECK(cudaMalloc((void**)&d_c, m*k*sizeof(int))); // 数据传输 CHECK(cudaMemcpy(d_a, h_a, m*n*sizeof(int), cudaMemcpyHostToDevice)); CHECK(cudaMemcpy(d_b, h_b, n*k*sizeof(int), cudaMemcpyHostToDevice)); // 配置并启动核函数 dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); dim3 dimGrid((k + dimBlock.x - 1)/dimBlock.x, (m + dimBlock.y - 1)/dimBlock.y); cudaEvent_t start, stop; CHECK(cudaEventCreate(&start)); CHECK(cudaEventCreate(&stop)); CHECK(cudaEventRecord(start)); gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); CHECK(cudaEventRecord(stop)); CHECK(cudaEventSynchronize(stop)); float milliseconds = 0; CHECK(cudaEventElapsedTime(&milliseconds, start, stop)); printf("GPU Time: %.3f ms\n", milliseconds); // 回传结果 CHECK(cudaMemcpy(h_c, d_c, m*k*sizeof(int), cudaMemcpyDeviceToHost)); // 验证 cpu_matrix_mult(h_a, h_b, h_c_ref, m, n, k); int errors = 0; for(int i = 0; i < m*k; i++) { if(abs(h_c[i] - h_c_ref[i]) > 1e-5) errors++; } printf("Errors: %d\n", errors); // 释放资源 CHECK(cudaFree(d_a)); CHECK(cudaFree(d_b)); CHECK(cudaFree(d_c)); CHECK(cudaFreeHost(h_a)); CHECK(cudaFreeHost(h_b)); CHECK(cudaFreeHost(h_c)); CHECK(cudaFreeHost(h_c_ref)); return 0; }

Jetson Nano上的常见问题与解决方案：

1. nvcc编译器限制：goto语句后不能有变量声明，这是nvcc的特殊限制
2. 内存不足：大矩阵运算时容易耗尽内存，可考虑分块计算
3. 性能瓶颈：使用nvprof工具分析性能瓶颈
4. 线程配置：过大的Block会导致寄存器溢出，建议从16×16开始尝试

5. 进阶优化技巧

当基本实现运行稳定后，我们可以考虑以下优化策略进一步提升性能：

共享内存优化：通过将数据块加载到共享内存，可以显著减少全局内存访问：

__global__ void gpu_matrix_mult_shared(int *a, int *b, int *c, int m, int n, int k) { __shared__ int s_a[BLOCK_SIZE][BLOCK_SIZE]; __shared__ int s_b[BLOCK_SIZE][BLOCK_SIZE]; int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; int sum = 0; for(int tile = 0; tile < (n + BLOCK_SIZE - 1)/BLOCK_SIZE; tile++) { if(row < m && (tile*BLOCK_SIZE + threadIdx.x) < n) { s_a[threadIdx.y][threadIdx.x] = a[row*n + tile*BLOCK_SIZE + threadIdx.x]; } if(col < k && (tile*BLOCK_SIZE + threadIdx.y) < n) { s_b[threadIdx.y][threadIdx.x] = b[(tile*BLOCK_SIZE + threadIdx.y)*k + col]; } __syncthreads(); for(int i = 0; i < BLOCK_SIZE; i++) { sum += s_a[threadIdx.y][i] * s_b[i][threadIdx.x]; } __syncthreads(); } if(row < m && col < k) { c[row*k + col] = sum; } }

CUDA流(Stream)的应用：对于更大的矩阵，可以使用多个CUDA流实现计算与传输的重叠：

cudaStream_t stream1, stream2; cudaStreamCreate(&stream1); cudaStreamCreate(&stream2); // 分块处理矩阵 for(int i = 0; i < m; i += BLOCK_SIZE) { cudaMemcpyAsync(d_a + i*n, h_a + i*n, BLOCK_SIZE*n*sizeof(int), cudaMemcpyHostToDevice, stream1); // 类似处理其他块... gpu_matrix_mult<<<dimGrid, dimBlock, 0, stream1>>>(...); }

Jetson Nano特有的电源管理：别忘了设备的热设计功率(TDP)限制，适当调整运行模式：

sudo nvpmodel -m 0 # 最大性能模式(10W) sudo nvpmodel -m 1 # 节能模式(5W)

在Jetson Nano上调试CUDA程序时，内存错误往往是最棘手的问题。使用cuda-memcheck工具可以帮助定位内存访问问题：

cuda-memcheck ./matrix_mult

经过这些优化后，在1024×1024矩阵乘法上，我们通常能看到2-3倍的性能提升。但要注意，在嵌入式环境中，优化往往需要在性能和功耗之间找到平衡点。