ARTICLE · INTELLIGENCE

战地情报 · 详情页

来自尧图项目组的一线实战观察与深度解析

CS149并行计算入门:从多线程到CUDA的完整避坑指南

CS149并行计算入门:从多线程到CUDA的完整避坑指南 1. 写在前面这门课到底值不值得啃先说结论Stanford CS149Parallel Computing是我见过最适合入门并行计算的一门课没有之一。它不像很多学校那种照本宣科的并行课上来就念线程库API而是真的从硬件架构出发把为什么要有并行并行到底解决了什么问题讲得明明白白。我当时是在准备秋招的时候发现这门课的。算法岗、后端岗面试里动不动就聊多线程、并发优化、性能调优光靠八股文根本撑不住深挖面试官一问你实际写过哪些并行程序就露馅。CS149恰好补齐了这个短板它有五个编程作业从最基础的多线程矩阵乘法到SIMD向量化、OpenMP并行循环、CUDA GPU编程最后还有一个没有标准答案的性能优化大作业做完之后你对程序为什么慢怎么让它快的理解完全不一样。这门课适合谁我认真说有C/C基础起码能写指针、结构体知道makefile大概是怎么回事的新手对GPU编程、CPU多核优化感兴趣但没有系统性学过的人准备面试但项目里缺性能和并发亮点的人哪怕只是好奇我电脑8个核为什么程序只用一个核的人有一个心理预期要先摆正这门课的视频是英文的作业要求也是英文slides里有大量密集的硬件和系统知识。但别被吓住它比MIT 6.824那种硬核系统课友好太多了作业的README写得非常详细你缺的只是几个关键的系统性知识点这篇文章会把我踩过的坑和梳理出来的重点全部分享给你。2. 环境配置是我踩过最大的坑直接决定你前20小时的心情CS149的编程作业用的是C官方建议你有一台能跑Linux或macOS的机器。我先说结论如果你用的是Windows请务必装WSLWindows Subsystem for Linux双系统那是不怕折腾的人干的事虚拟机性能损失看运气。当时我天真地以为Windows下的Visual Studio能解决一切结果第一个作业直接教我做人。作业里的makefile是按照Linux环境写的用了很多POSIX线程相关的库pthread.hWindows原生根本不支持哪怕你用MinGW强行编过了后面的OpenMP和CUDA作业也是全线崩盘。2.1 我这套配置方案macOS/Linux/Windows三个平台实测LinuxUbuntu 20.04以上环境配置如果你本身就是Ubuntu用户恭喜你基本什么都不用装。官方要求的环境就是gcc/g、make、cmake。用一条命令检查gcc --version g --version make --version cmake --version缺哪个就装哪个sudo apt update sudo apt install build-essential cmakebuild-essential这个包会把gcc、g、make、libc-dev这些都一起带进来不用一个个装。后面如果作业需要OpenMPgcc自带libgomp不需要额外配置。macOSIntel芯片和M1/M2都试过环境配置macOS自带Clang编译器但注意CS149检查作业用的主要是GCC行为个别地方Clang和GCC在OpenMP支持上不太一样。最好还是装上gccbrew install gccM系列芯片的macOS上作业3的CUDA部分跑不了因为CUDA不支持macOS。我当时是在这个作业上换了一台Linux台式机或者用学校/公司的服务器跑的。如果你只有Mac建议提前规划好CUDA作业怎么解决不要等到deadline前一天才发现跑不了。Windows环境配置新手推荐方案强烈建议装WSL2这是我在Windows上折腾了很久之后的最优解。在PowerShell管理员模式里执行wsl --install然后装Ubuntu 22.04 LTS直接在Microsoft Store里搜就行。装好之后进WSL按上面Ubuntu的步骤装编译环境。VSCode装一个WSL插件就能直接在Windows里写代码、在WSL里编译运行体验非常顺滑。2.2 VSCode连接WSL的配置细节实测最舒服的方案很多新手在这个环节会卡住VSCode里写了代码但不知道文件在Windows和WSL之间怎么共享。用WSL插件之后VSCode左侧的远程资源管理器里能看到你的WSL发行版点连接就会自动在WSL环境里打开一个新的VSCode窗口底栏会显示WSL: Ubuntu。然后在WSL里打开你的CS149作业目录mkdir -p ~/cs149 cd ~/cs149 # 把你从GitHub拉的作业代码放进来C/C扩展在WSL里要重新装一遍因为WSL是一个独立的远程环境VSCode的扩展分为本地扩展和WSL扩展两种。你在WSL窗口里点扩展面板它会提示你推荐安装在WSL里的扩展把C/C和CMake Tools装上就行。关于编译调试我最推荐的组合是CMake Tools C/C扩展。终端里make编译一遍VSCode的CMake工具能帮你识别CMakeLists.txtF5就能启动调试打断点、看变量都比较顺手。2.3 CUDA环境配置只有做到作业4才需要作业4开始涉及CUDA编程。如果你用的是Linux需要有NVIDIA显卡或者能连到带GPU的服务器。装CUDA工具链sudo apt install nvidia-cuda-toolkit # 检查是否装好 nvcc --version注意作业评分主要在Linux上nvcc版本不用追最新10.x以上都够用。上课用的其实是斯坦福自己的Linux服务器otter cluster如果你能搞到带NVIDIA GPU的Linux机器就行不需要完全复刻。如果你没有GPU服务器云GPU是按小时计费的做作业的时候租几小时完全够用。我当时用了一台云GPU服务器编译和运行都极其顺畅。3. 课程内容拆解五大编程作业循序渐进CS149有6个章节视频加5个编程作业1个期末考试。视频大约20个左右每个20-50分钟不等。真正让这门课封神的是它的编程作业设计从最简单的pthreads到最复杂的CUDA性能优化每一步都能感受到自己的能力在肉眼可见地提升。3.1 作业0和作业1热身与多线程基础作业0不用写代码就是搭环境、拉代码、跑一下测试确认编译工具链没问题。重点看它的makefile结构后面所有作业都是在这个模板上叠加的。作业1是实现一个并行程序给定一个输入图统计每个节点的连通分量数量。这个作业的精髓在于你需要自己写pthreads的线程管理自己设计线程间的任务分配方式。最容易踩的坑有这几个一是线程安全。当多个线程同时累加一个全局计数器时必须用互斥锁或原子操作保护。我当时因为没加锁跑了半天结果都是错的而且错误是随机出现的排查起来非常痛苦。用GCC的-fsanitizethread选项编译一遍能直接告诉你是哪一行的数据竞争g -fsanitizethread -g -O1 program.cpp -o program -lpthread二是线程数量。你可能会觉得线程越多越快实际不是这样。线程创建和销毁本身有开销当任务量不够大时多线程反而比单线程慢。这个作业有个关键测试指标——加速比speedup单线程运行时间除以多线程运行时间理想值是线程数实际能达到50%就算不错了。我在8核机器上跑四线程加速比能到3.2左右再往上就上不去了瓶颈在锁竞争和内存带宽。3.2 作业2SIMD向量化感受指令级的并行作业2要求在1万维的浮点数组上做向量运算用SIMD指令优化。这里的核心是理解数据级并行DLP和一条指令处理多个数据Single Instruction Multiple Data的区别。这个作业有两个从零到一的跨度第一你要把C代码编译成AVX2指令集g -mavx2 -O3 -stdc11 simd.cpp -o simd-mavx2是让编译器允许生成AVX2指令-O3是最高优化级别。当时我对SIMD的理解是从一个矩阵乘法的例子开始的假设你要把8个float相加普通做法是循环8次一条条加SIMD的做法是用一条指令同时把8个float塞进256位的寄存器一条加法指令搞定8个数。第二你必须处理内存对齐。AVX2的load指令要求内存地址按32字节对齐否则程序直接崩溃。分配数组时用float* data (float*)aligned_alloc(32, size * sizeof(float));而不是普通的malloc或new。这个坑我踩了整整一个晚上最后在讲义第4节的alignment requirements段落里找到的答案。3.3 作业3OpenMP并行循环体验自动并行化作业3是把几种不同的并行算法例如图像处理算法、就是给你的数组做模糊、锐化等操作用OpenMP指令并行化。OpenMP的好处是它在编译器层面帮你管理线程你只需要在要并行的循环前面加一行#pragma omp parallel for for (int i 0; i size; i) { // 对数组做处理 }最核心的是理解reduction归约的机制。如果你在一个并行循环里对同一个变量做累加必须声明成reduction#pragma omp parallel for reduction(:sum) for (int i 0; i size; i) { sum data[i]; }编译器会把每个线程的sum复制一份线程各自累加最后自动汇总。如果不加reduction会出现严重的竞态条件。这个作业的收获在于你不需要手动管线程但要理解schedule调度策略。OpenMP有static静态划分、dynamic动态划分和guided指导性划分三种调度方式每种方式在不同数据分布下的性能差距能达到3倍。建议把三种都试一遍记录运行时间用数据说话。3.4 作业4CUDA GPU编程打开新世界的大门作业4是写一个简单的GPU内核理解GPU为何能跑几千个线程同时工作。CUDA编程的核心是三个层次线程thread、块block、网格grid。一个典型的内核长这样__global__ void vector_add(float* a, float* b, float* out, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) { out[i] a[i] b[i]; } } // 调用 int threads_per_block 256; int blocks (n threads_per_block - 1) / threads_per_block; vector_addblocks, threads_per_block(a, b, out, n);blockIdx.x、blockDim.x、threadIdx.x是内置变量分别代表当前块在网格中的索引、每块的线程数、当前线程在块中的索引。这个公式i blockIdx.x * blockDim.x threadIdx.x是所有CUDA程序的基石每个线程通过它算出自己负责处理哪个数据。最大的坑是内存拷贝。GPU不能直接访问CPU内存主机内存你需要用cudaMemcpy把数据拷到GPU显存算完再拷回来。我第一次写的时候忘了拷数据结果GPU端读到的全是乱码。还有就是block size的选择。block太小比如32会导致线程调度开销大block太大比如1024会超过硬件限制。256或512是大多数场景下的经验值但实际性能要看profile结果CS149作业里就有要求你用nvprof性能分析工具看代码的热点。3.5 作业5性能优化大作业没有标准答案的挑战这是压轴的大作业给定一个算法函数你自己想办法让它跑得更快。我当时处理的是一个图遍历算法初始版本跑1000轮要40毫秒我优化到6.8毫秒提速将近6倍。主要的优化手段调整数据结构把链表改成连续数组减少缓存不命中用__restrict__关键字告知编译器指针不重叠让编译器更好地向量化手动展开循环减少循环开销和分支预测失败用OpenMP把最外层循环并行化尝试prefetch预取数据到缓存这个作业最值钱的地方在于它没有标准答案你必须自己想办法用性能分析工具找瓶颈然后去优化。这个过程正是工作里高性能系统开发的日常。做完这个作业之后你再看别人的代码这里可以更快的感觉会自然浮现。4. 学习过程中的典型问题和排查思路4.1 编译报错混乱怎么定位问题新手最怕的就是几百行编译错误刷满整个终端。我的建议是先看第一个错误不要看后面的。编译器报错往往是一堆级联错误第一个才是根因。常见的第一行错误fatal error: pthread.h: No such file or directory缺少pthread库编译命令里加-lpthreadundefined reference to pthread_*链接阶段缺少pthread库同样加-lpthreadfatal error: omp.h: No such file or directory没用-fopenmp编译选项error: ‘thread’ does not name a type忘记加#include thread编译命令用对是很重要的比如作业2的makefile里实际上会有完整的参数组合你只要看懂了makefile编译一般不会出问题。4.2 程序运行结果正确但速度不达标这是CS149作业最常见的情况。性能不够往往不是代码逻辑错了而是你忽略了编译器的优化潜力。务必确认编译选项里有-O2或-O3没有优化等级的编译性能可以直接差3-5倍。另外检查你的计时方式。作业要求测的是纯计算耗时不是整个程序从启动到结束的时间。测试时记得对同一份输入跑多次取平均值首次运行时可能涉及冷缓存和内存页分配会拉低成绩。4.3 多线程结果不稳定如果你发现结果有时候对、有时候不对八成是数据竞争。先不要急着看代码用Sanitizer工具检测# ThreadSanitizer检测数据竞争和线程错误使用 g -fsanitizethread -g -O1 your_program.cpp -o your_program -lpthread ./your_program这个工具会精确到哪一行代码访问了哪个共享变量极大概率帮你锁定bug。另一个常见场景你程序本身没错但打印输出的顺序是乱的这并不是bug因为线程调度顺序本来就不确定。作业只检查最终结果的正确性不需要输出一致。4.4 作业4 CUDA相关的问题CUDA程序最常见的问题之一是内存越界。线程数n不是block size的整数倍时if (i n)这个保护条件一定要加否则GPU会访问无效地址轻则报错重则整个驱动崩溃。调试建议用cuda-memcheck或新版compute-sanitizercompute-sanitizer ./your_program它能指出是哪个线程、访问了什么非法地址。CUDA程序的printf输出在GPU端调用cudaDeviceSynchronize()且内核正常返回后输出才会刷新到控制台。4.5 一个不容忽视的问题多线程程序的时间测量我在作业1里踩的坑是用std::chrono::steady_clock计时的时候测的是墙钟时间如果你在循环里打印调试信息时间会暴涨。作业要求测的是纯并行循环的执行时间里面的打印输出全部删掉最好把计时部分单独隔离。另外后台如果跑着Chrome、微信这类软件性能测试会受到干扰。跑性能测试之前关几样占内存的程序次数跑满、取中位数这样数据才可靠。我当时写了一个脚本自动跑10次取中位数可以直接复制去用#!/bin/bash for i in {1..10}; do ./your_program 21 | grep time done5. 我的一些个人建议和避坑心得视频要配合slides看不要只看视频。cs149的slides做得极其精致每个复杂的硬件机制都有图示只看视频反而容易错过细节。我当时把每个Lecture的视频在1.25倍速看一遍然后花两倍的时间精读对应slides效果最好。视频里Kayvon老师讲课的节奏比较慢很多内容也不需要看第二遍但slides里的代码示例和架构图值得反复研读。作业一定独立完成。这门课是有官方Release和评测系统的但你如果只是把答案扒下来跑通收获的大概只有会抄这一项技能。CS149的作业都是阶梯式的最初版本能跑通但你要自己优化到性能达标线这个过程练到的是分析和排查的能力是面试能讲出细节的经历。第一周比较痛苦是因为环境没搭好、术语不熟但扛过前两个作业之后你会开始期待每个新作业——因为它们像一个一个的关卡每次通关你都能感觉到自己在变强。顺便说一句CS149的所有视频、slides和作业代码都在公开的课程网站上不需要任何账号或校内权限就能获取。课程全程用的是免费/开源工具没有额外花费。唯一可能需要花钱的场景是CUDA作业需要GPU服务器但云GPU按小时收费很多人几个小时内就能完成。最后再分享一个小技巧这个课程的名字在这一季可能叫CS149但在斯坦福的课程列表里它还有一个更长的名字叫Parallel Computing搜的时候如果看到全称不用怀疑就是这门课。祝学习顺利有问题可以随时交流。
RELATED READING

延伸阅读

更多一线实战笔记与深度复盘,助您持续精进