1. 项目概述为什么要在Jetson Nano上折腾并行计算如果你手头有一块NVIDIA Jetson Nano 2GB的开发板可能已经用它跑过一些图像识别、目标检测的Demo感觉这个“小盒子”还挺能干。但有没有那么一瞬间你会好奇它那颗128核的Maxwell架构GPU到底有多大能耐它和我们在台式机上用的那些“大块头”显卡在并行计算上有什么本质区别这就是我们这次要深入探讨的核心。简单来说这个项目就是一次对Jetson Nano 2GB并行计算能力的“摸底考试”。我们不满足于仅仅调用现成的深度学习框架如TensorFlow、PyTorch去跑模型而是要深入到更底层亲手写一些CUDA C/C代码看看如何直接指挥这128个CUDA核心协同工作解决一些经典的并行计算问题。这就像你不光会开车还想知道发动机的缸内直喷技术是怎么提升效率的。对于嵌入式AI开发者、边缘计算爱好者或者任何想理解GPU如何加速计算的人来说这个过程能帮你建立起对硬件性能最直观的认知让你在后续优化模型、设计算法时心里更有底。Jetson Nano虽然算力有限官方标称472 GFLOPS但其完整的CUDA生态和真实的GPU架构是学习并行编程绝佳的“实验田”。在这里踩过的坑、获得的性能提升百分比其原理和经验可以无缝迁移到更强大的Jetson AGX Orin甚至数据中心GPU上。接下来我们就从环境准备开始一步步解锁它的并行计算性能。2. 环境准备与核心工具链解析工欲善其事必先利其器。在Jetson Nano上搞CUDA编程环境和工具的选择直接决定了开发体验和效率。2.1 JetPack SDK一切的基础Jetson Nano出厂或刷机后其操作系统、CUDA驱动、cuDNN等核心组件都打包在NVIDIA的JetPack SDK中。对于Jetson Nano 2GB最主流且稳定的版本是JetPack 4.6对应L4T 32.6.1。这个版本包含了Ubuntu 18.04、CUDA 10.2、cuDNN 8.0等。虽然CUDA 10.2不是最新版但对于Nano的Maxwell架构和大多数边缘应用来说完全够用且生态兼容性最好。注意切勿在Jetson Nano上尝试安装其他版本的CUDA Toolkit例如从NVIDIA官网下载的x86版本。Jetson系列的CUDA驱动和运行时是深度集成在L4TLinux for Tegra系统中的强行安装其他版本几乎必然导致系统启动失败。所有开发都应基于JetPack提供的环境。验证环境是否就绪打开终端依次执行以下命令# 查看系统版本和L4T版本 cat /etc/nv_tegra_release # 输出类似# R32 (release), REVISION: 6.1, GCID: 27863751, BOARD: t210ref, ... 这确认了是JetPack 4.6。 # 查看CUDA编译器版本 nvcc --version # 应显示Cuda compilation tools, release 10.2, V10.2.89 # 查看GPU状态 sudo apt-get install -y python3-pip pip3 install jetson-stats sudo jtop运行jtop后你可以看到一个漂亮的资源监控界面确认GPU、内存、CPU等信息正常显示。这是后续性能测试时观察负载的利器。2.2 开发工具选型VSCode SSH远程开发在Jetson Nano上直接编码体验并不好。我强烈推荐采用“远程开发”模式将Nano连接到局域网在你的主力电脑Windows/Mac/Linux上使用Visual Studio Code通过SSH远程连接到Nano进行编码、编译和调试。为什么这么选性能与体验主力机的CPU、内存和编辑器性能远胜于Nano编码流畅度天壤之别。生态完善VSCode的C/C扩展、CUDA语法高亮、远程开发插件非常成熟。文件管理方便可以直接在本地和远程之间拖拽文件。具体设置步骤在Nano上确保开启SSH服务sudo systemctl enable ssh sudo systemctl start ssh。在主力机安装VSCode并安装官方扩展“Remote - SSH”。配置SSH连接至Nano的IP地址。连接成功后在VSCode的扩展商店安装“C/C”和“CUDA”相关扩展。现在你可以在主力机上享受智能补全和语法高亮而代码实际在Nano上编译运行。2.3 基础性能基准nvprof与Nsight Systems在开始写自己的CUDA代码前我们需要两个“测量工具”来建立性能基准。nvprof(Legacy): 随CUDA Toolkit安装的命令行性能分析器。虽然NVIDIA正在推动其替代品但在Jetson NanoCUDA 10.2上它仍然是即用且功能强大的工具。它可以给出核函数执行时间、内存拷贝耗时、占用率等关键指标。# 对一个可执行文件进行性能分析 nvprof ./my_cuda_programNsight Systems: 更现代、更强大的系统级性能分析工具。它提供时间线视图能让你看到CPU和GPU活动的重叠情况直观发现是计算瓶颈还是内存瓶颈。对于Jetson Nano你需要从NVIDIA官网下载对应aarch64ARM架构的版本。分析时需要带--force-override参数因为Nano的GPU驱动是/dev/nvhost-ctrl而非标准的/dev/nvidiactl。# 使用Nsight Systems进行命令行分析 nsys profile --force-override true -o my_report ./my_cuda_program生成.qdrep报告文件后可以传输到主力机用Nsight Systems GUI打开进行可视化分析。有了这些工具我们就可以像医生拿着听诊器和CT机一样去“诊断”我们CUDA程序的性能了。3. 并行计算初体验从向量加法开始理论学习再多不如动手写一行代码。我们从一个最经典的并行计算例子——向量加法开始直观感受CUDA编程模型和性能差异。3.1 CPU串行实现与性能基线首先我们建立一个CPU版本的向量加法作为性能基线。这能让我们清楚地看到并行化带来的收益。// vec_add_cpu.cpp #include iostream #include chrono #include vector #include cstdlib void vec_add_cpu(float* A, float* B, float* C, int N) { for (int i 0; i N; i) { C[i] A[i] B[i]; } } int main() { const int N 1 24; // 约1600万个元素 std::vectorfloat A(N, 1.0f); // 初始化为1.0 std::vectorfloat B(N, 2.0f); // 初始化为2.0 std::vectorfloat C(N, 0.0f); auto start std::chrono::high_resolution_clock::now(); vec_add_cpu(A.data(), B.data(), C.data(), N); auto end std::chrono::high_resolution_clock::now(); std::chrono::durationdouble elapsed end - start; std::cout CPU 向量加法耗时: elapsed.count() * 1000 ms std::endl; // 简单验证结果检查前5个元素 for (int i 0; i 5; i) { std::cout C[ i ] C[i] std::endl; // 应输出3.0 } return 0; }用g编译并运行记录下耗时。在Jetson Nano的四核Cortex-A57 CPU上这个操作可能需要几十到上百毫秒。这就是我们的起点。3.2 CUDA并行实现核函数与线程组织现在我们编写CUDA版本的向量加法。核心思想是让GPU上的成千上万个线程每个只负责计算一个或几个元素的加法。// vec_add.cu #include iostream #include cuda_runtime.h // CUDA核函数 (Kernel Function) // 每个线程执行一次这个函数 __global__ void vec_add_kernel(float* A, float* B, float* C, int N) { // 计算当前线程的全局索引 int idx blockIdx.x * blockDim.x threadIdx.x; // 确保索引不越界 if (idx N) { C[idx] A[idx] B[idx]; } } int main() { const int N 1 24; // 1600万 size_t size N * sizeof(float); // 1. 在主机(CPU)上分配并初始化内存 float *h_A new float[N]; float *h_B new float[N]; float *h_C new float[N]; for (int i 0; i N; i) { h_A[i] 1.0f; h_B[i] 2.0f; } // 2. 在设备(GPU)上分配内存 float *d_A, *d_B, *d_C; cudaMalloc((void**)d_A, size); cudaMalloc((void**)d_B, size); cudaMalloc((void**)d_C, size); // 3. 将数据从主机拷贝到设备 (H2D) cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); // 4. 设置线程网格(Grid)和线程块(Block)大小并启动核函数 int threadsPerBlock 256; // 每个Block有256个线程 int blocksPerGrid (N threadsPerBlock - 1) / threadsPerBlock; // 计算需要的Block数量 // 创建CUDA事件用于精确计时 cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start); vec_add_kernelblocksPerGrid, threadsPerBlock(d_A, d_B, d_C, N); cudaEventRecord(stop); cudaEventSynchronize(stop); // 等待核函数执行完成 float milliseconds 0; cudaEventElapsedTime(milliseconds, start, stop); std::cout CUDA 向量加法核函数执行耗时: milliseconds ms std::endl; // 5. 将结果从设备拷贝回主机 (D2H) cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); // 6. 验证结果 for (int i 0; i 5; i) { std::cout h_C[ i ] h_C[i] std::endl; } // 7. 清理设备内存 cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); delete[] h_A; delete[] h_B; delete[] h_C; return 0; }使用nvcc编译nvcc -o vec_add vec_add.cu。运行后你会看到核函数执行时间可能只有几毫秒相比CPU版本的几十毫秒有数十倍的加速比3.3 性能分析与关键洞察用nvprof深入分析一下nvprof ./vec_add你会看到类似下面的输出其中包含几个关键指标XXX Profiling result: Type Time(%) Time Calls Avg Min Max Name GPU activities: 58.21% 2.3440ms 1 2.3440ms 2.3440ms 2.3440ms vec_add_kernel(float*, float*, float*, int) 41.79% 1.6835ms 2 841.77us 820.67us 862.85us [CUDA memcpy HtoD] 0.00% 1.1840us 1 1.1840us 1.1840us 1.1840us [CUDA memcpy DtoH]关键洞察计算 vs. 内存拷贝可以看到实际计算核函数只占了总GPU活动时间的58%而将数据从主机内存拷贝到设备内存H2D却占了42%。这是一个非常重要的信号对于小规模或简单的计算内存拷贝的开销可能完全抵消甚至超过并行计算带来的收益。这就是所谓的“内存墙”问题在边缘设备上的体现。线程配置我们选择了threadsPerBlock 256。为什么是256这是一个经验值。在Jetson Nano的Maxwell架构上每个流多处理器SM的最大线程数是2048而每个线程块Block的线程数最好是32一个Warp的大小的倍数。256是一个能较好平衡占用率和寄存器压力的值。你可以尝试改为128或512用nvprof对比性能理解线程配置的影响。隐藏的内存拷贝我们只计算了核函数的执行时间。如果加上H2D和D2H的内存拷贝总时间整个CUDA程序的耗时可能和CPU版本相差无几甚至更慢。这引出了下一个重要话题如何优化内存访问。4. 深入性能优化内存层次结构与实战策略知道了瓶颈所在我们就要着手优化。在GPU编程中优化内存访问是提升性能最有效的手段之一。4.1 理解Jetson Nano的内存架构Jetson Nano 2GB有一个关键特点CPU和GPU共享同一块物理内存Unified Memory。这与我们常见的台式机独显有独立的显存不同。这种统一内存架构UMA有好有坏好处在某些情况下可以避免显式的内存拷贝通过cudaMallocManaged分配托管内存简化编程。坏处CPU和GPU会竞争同一内存带宽。如果两者同时高强度访问内存性能会相互拖累。对于追求极致性能的场景我们仍然需要显式地管理内存并利用好GPU内部的缓存层次。4.2 优化策略一使用页锁定内存Pinned Memory默认的new或malloc分配的主机内存是“可分页”的。当GPU通过DMA直接内存访问拷贝数据时如果源内存正在被操作系统交换到磁盘或者物理地址不连续会导致额外的延迟和拷贝。 使用**页锁定内存Pinned Memory**可以解决这个问题。它保证物理内存常驻且地址连续从而允许DMA以最高速度进行传输。// 替换原来的 new 分配 float *h_A, *h_B, *h_C; cudaMallocHost((void**)h_A, size); // 分配页锁定主机内存 cudaMallocHost((void**)h_B, size); cudaMallocHost((void**)h_C, size); // ... 初始化 h_A, h_B // 执行内存拷贝和核函数 // ... // 清理时使用 cudaFreeHost(h_A); cudaFreeHost(h_B); cudaFreeHost(h_C);重新编译运行并用nvprof观察[CUDA memcpy HtoD]的时间。你会发现耗时显著降低可能从800多微秒降到400微秒左右。这就是优化内存拷贝的立竿见影的效果。4.3 优化策略二核函数中的合并内存访问GPU的全局内存带宽很高但延迟也很大。为了高效利用带宽GPU喜欢“合并访问”Coalesced Access。即一个线程束Warp32个线程中的所有线程在一次内存事务中访问连续对齐的全局内存地址。在我们的vec_add_kernel中每个线程访问A[idx],B[idx],C[idx]。由于idx是连续的threadIdx.x连续且我们以float4字节为单位访问这天然就是完美的合并访问。这是向量加法性能高的原因之一。反面例子跨步访问如果核函数这样写__global__ void bad_access_kernel(float* A, float* B, float* C, int N, int stride) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx * stride N) { C[idx * stride] A[idx * stride] B[idx * stride]; // 非连续访问 } }当stride较大时一个Warp的32个线程访问的内存地址相隔很远无法合并会导致多次内存事务性能急剧下降。在Jetson Nano这种内存带宽有限的设备上这种影响会被放大。4.4 优化策略三利用共享内存Shared Memory共享内存是GPU上每个线程块Block内部的高速、低延迟的存储空间可以被该Block内的所有线程共享。对于存在数据复用例如矩阵乘法、卷积的算法将数据从全局内存先加载到共享内存可以极大减少对全局内存的访问次数。我们以一个更复杂的例子——矩阵转置来说明。朴素的方法是每个线程读取全局内存中的一个元素然后写到转置后的位置。这会导致对全局内存的写入是非合并的因为写入的地址不连续。// 优化前的朴素转置写入非合并 __global__ void transpose_naive(float* in, float* out, int width, int height) { int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; if (x width y height) { out[y * width x] in[x * height y]; // 对out的写入是跨行的非合并 } }使用共享内存优化// 使用共享内存优化转置读写都合并 __global__ void transpose_shared(float* in, float* out, int width, int height) { // 声明一个二维共享内存数组 __shared__ float tile[BLOCK_DIM][BLOCK_DIM1]; // 1 是为了避免bank conflict int x blockIdx.x * BLOCK_DIM threadIdx.x; int y blockIdx.y * BLOCK_DIM threadIdx.y; // 将输入矩阵的一个块加载到共享内存合并读 if (x width y height) { tile[threadIdx.y][threadIdx.x] in[y * width x]; } __syncthreads(); // 等待块内所有线程完成加载 // 计算转置后的输出坐标 int out_x blockIdx.y * BLOCK_DIM threadIdx.x; int out_y blockIdx.x * BLOCK_DIM threadIdx.y; // 从共享内存写入全局内存合并写 if (out_x height out_y width) { out[out_y * height out_x] tile[threadIdx.x][threadIdx.y]; } }在这个优化版本中线程块协作将全局内存中连续的一块数据BLOCK_DIM x BLOCK_DIM以合并访问的方式读入共享内存tile。经过__syncthreads()同步确保所有数据到位。然后每个线程从共享内存中读取转置后的数据注意tile的下标交换了再以合并访问的方式写入全局内存。通过nvprof对比两个核函数的性能你会发现优化后的版本在Jetson Nano上能有数倍的提升。共享内存是解决非合并访问和实现线程间通信的利器但需要仔细设计数据加载和存储模式并注意线程同步。5. 综合性能测试与CPU及不同线程配置的对比理论说再多不如数据有说服力。我们设计一个简单的测试来量化不同配置下的性能差异。5.1 测试设计不同数据规模与线程块大小我们修改向量加法的程序让它能测试不同数据规模N和不同线程块大小threadsPerBlock下的性能。// benchmark.cu #include iostream #include vector #include cuda_runtime.h #include chrono // ... 核函数定义 ... void run_benchmark(int N, int threadsPerBlock) { // ... 分配内存初始化数据 ... // 预热避免首次启动开销 vec_add_kernelblocksPerGrid, threadsPerBlock(d_A, d_B, d_C, N); cudaDeviceSynchronize(); // 多次运行取平均 int iterations 100; cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start); for (int i 0; i iterations; i) { vec_add_kernelblocksPerGrid, threadsPerBlock(d_A, d_B, d_C, N); } cudaEventRecord(stop); cudaEventSynchronize(stop); float total_ms; cudaEventElapsedTime(total_ms, start, stop); float avg_kernel_ms total_ms / iterations; // ... 计算带宽和吞吐量 ... // 理论带宽Jetson Nano 2GB 内存带宽约 25.6 GB/s // 实际带宽 (数据量 * 2次读 1次写) / 时间 size_t data_transferred (size_t)N * sizeof(float) * 3; // A读B读C写 float bandwidth_GBs (data_transferred / (avg_kernel_ms / 1000.0f)) / 1e9; float percentage_of_peak (bandwidth_GBs / 25.6f) * 100.0f; std::cout N N , BlockSize threadsPerBlock , Time avg_kernel_ms ms , BW bandwidth_GBs GB/s ( percentage_of_peak % of peak) std::endl; // ... 清理 ... }我们测试从1K到16M的不同数据规模以及32, 64, 128, 256, 512等不同的线程块大小。5.2 测试结果分析与解读在Jetson Nano 2GB上运行后我们可能会得到类似下表的数据数值为模拟需实测数据规模 (N)线程块大小核函数平均耗时 (ms)实测带宽 (GB/s)峰值带宽占比1K2560.0050.381.5%64K2560.0326.023.4%1M2560.429.035.2%16M2566.511.846.1%16M1286.811.344.1%16M5127.110.842.2%从数据中我们可以解读出以下几点小数据规模性能极差当N1K时带宽利用率只有1.5%。这是因为核函数启动、线程调度等固定开销占据了主导。启示在Jetson Nano上如果计算任务非常小使用CUDA可能得不偿失CPU反而更快。存在“甜蜜点”随着数据规模增大带宽利用率迅速提升在16M时达到接近50%的峰值带宽。对于向量加法这种内存密集型Memory-Bound操作能达到峰值带宽的40-50%已经是相当不错的表现因为还有指令开销、地址计算等。线程块大小的影响对于这个简单的内存拷贝型核函数256是一个较好的选择。128和512的性能略有下降。原因线程块大小会影响GPU上SM的占用率Occupancy和指令流水线的效率。需要通过实际测试来找到特定算法在特定硬件上的最优值。与CPU对比对于16M的向量加法我们假设CPU串行版本耗时约80ms而CUDA核函数仅6.5ms加速比超过12倍。但请记住这仅仅是核函数执行时间。如果算上H2D和D2H的内存拷贝总时间可能额外增加2-3ms整体加速比会下降到8-10倍。这再次强调了减少主机与设备间数据传输的重要性。6. 实战避坑指南与高级技巧纸上得来终觉浅绝知此事要躬行。在实际操作中你会遇到各种预料之外的问题。这里分享几个我踩过的坑和对应的解决方案。6.1 常见编译与运行时错误错误error: identifier __syncthreads is undefined原因在.cu文件中核函数__global__和设备函数__device__必须用nvcc编译。如果你不小心用g编译了包含CUDA语法的文件就会报这个错。解决确保使用nvcc编译器或者在你的CMakeLists.txt中正确设置CUDA语言支持。错误CUDA error: no kernel image is available for execution on the device原因这是Jetson开发者最常遇到的错误之一。它意味着你编译的CUDA代码特别是PTX和Cubin与当前GPU的架构不兼容。Jetson Nano是Maxwell架构sm_53。解决在编译时显式指定正确的架构和代码生成选项。nvcc -archsm_53 -o my_program my_program.cu或者在CMake中设置set(CUDA_ARCHITECTURES 53)错误CUDA error: out of memory原因Jetson Nano 2GB的共享内存很小。除了GPU显存系统内存也被大量占用。解决使用sudo jtop监控内存使用情况。检查代码中是否有内存泄漏每次cudaMalloc后都要有对应的cudaFree。减少单次处理的数据批量大小Batch Size。考虑使用cudaMallocManaged分配托管内存让系统自动管理但要注意性能可能略有损失。6.2 性能调优进阶流式处理与异步执行为了进一步隐藏内存拷贝开销和核函数执行开销可以使用**CUDA流Stream**实现异步并发执行。基本原理默认情况下所有CUDA操作内存拷贝、核函数都在一个默认流Stream 0中顺序执行。我们可以创建多个流将不同的操作放到不同的流中。只要资源不冲突这些操作就可以在GPU上并发执行。一个典型模式计算与通信重叠假设我们需要处理一个很大的数据可以将其分成若干块。我们创建两个流流1拷贝第2块数据H2D同时执行第1块数据的核函数计算。流2拷贝第1块结果回主机D2H同时执行第2块数据的核函数计算。 如此流水线化可以显著提升整体吞吐量。cudaStream_t stream1, stream2; cudaStreamCreate(stream1); cudaStreamCreate(stream2); // 为每块数据分配独立的内存主机和设备 // ... // 流水线执行 for (int i 0; i num_chunks; i) { int stream_id i % 2; cudaStream_t cur_stream (stream_id 0) ? stream1 : stream2; // 异步内存拷贝 H2D cudaMemcpyAsync(d_A[stream_id], h_A[stream_id], chunk_size, cudaMemcpyHostToDevice, cur_stream); cudaMemcpyAsync(d_B[stream_id], h_B[stream_id], chunk_size, cudaMemcpyHostToDevice, cur_stream); // 在同一个流中启动核函数依赖前面的拷贝完成 vec_add_kernelblocks, threads, 0, cur_stream(d_A[stream_id], d_B[stream_id], d_C[stream_id], chunk_size); // 异步将结果拷贝回主机 D2H cudaMemcpyAsync(h_C[stream_id], d_C[stream_id], chunk_size, cudaMemcpyDeviceToHost, cur_stream); } // 等待所有流完成 cudaDeviceSynchronize();在Jetson Nano上由于GPU计算单元相对较少流式处理带来的收益可能不如大型GPU明显但对于IO密集型的复杂流水线仍然能有效提升资源利用率。6.3 利用Tensor Core别想了但可以了解其原理在更高级的Jetson设备如Jetson AGX Orin含有Tensor Core上你可以通过特定的库如cuBLASLt或编程模型WMMA API来调用Tensor Core进行混合精度矩阵计算获得巨大的性能提升。但在Jetson NanoMaxwell架构上没有Tensor Core。这是一个重要的硬件认知。所以在Nano上做矩阵乘法我们只能优化到使用共享内存和寄存器而无法利用Tensor Core的硬件加速。这决定了Nano的性能天花板。了解这一点能帮助你在算法选型和性能预期上做出更合理的判断。7. 总结与项目延伸思考经过这一系列的实验、分析和优化我们对Jetson Nano 2GB的并行计算性能应该有了一个立体而深刻的认识。它不是一个性能怪兽而是一个麻雀虽小五脏俱全的CUDA教学和实践平台。核心收获并行优势显著但非万能对于计算密集、数据规模大的任务CUDA能带来数量级的加速。但对于小任务或IO密集型任务启动开销和内存拷贝可能使其优势尽失。内存访问是命门在Jetson Nano上优化内存访问模式合并访问、使用共享内存、使用页锁定内存比单纯增加线程数更能提升性能。工具链至关重要nvprof和Nsight Systems是洞察性能瓶颈的“眼睛”熟练使用它们是从“会写CUDA”到“写好CUDA”的关键。硬件特性决定天花板理解Nano的Maxwell架构、128个CUDA核心、共享内存大小、没有Tensor Core等特性是进行有效性能调优和算法设计的前提。项目延伸掌握了这些基础后你可以尝试更具挑战性的项目来巩固和深化实现一个简单的矩阵乘法尝试不同的分块Tiling策略并使用共享内存进行优化对比不同块大小对性能的影响。实现图像卷积或池化操作这是一个经典的Stencil计算模式对共享内存的使用和边界处理有更高要求。集成到实际应用将你优化的CUDA核函数封装成PyTorch或TensorFlow的自定义算子在真实的深度学习推理流水线中调用它体验从底层优化到上层应用的全链路。最后我个人最深的体会是在资源受限的边缘设备上进行并行计算优化就像在螺蛳壳里做道场。每一个字节的内存、每一个时钟周期都需要精打细算。这种“抠门”的优化经验恰恰是理解高性能计算精髓的捷径。当你未来面对更强大的硬件时这段在Jetson Nano上“斤斤计较”的经历会让你对性能的理解更加透彻。