AI Infra 校招 (1)
Q: CPU与GPU的区别?
CPU设计着重处理单个线程的复杂计算和控制流程,拥有强大的分支预测、乱序执行和大容量缓存。GPU设计用于高密度并行计算,更多晶体管投入数据处理而非缓存和流量控制,ALU数量远多于CPU。CPU适合延迟敏感的串行任务,GPU适合吞吐量导向的大规模并行任务。
Q: CUDA编程中的SM和SP是什么?
- SP(Streaming Processor):流处理器,最基本的计算单元,执行单个线程的指令。
- SM(Streaming Multiprocessor):流多处理器,由多个SP加上共享内存、寄存器文件、Warp调度器等资源组成。SM是GPU调度和执行的基本单元,一个SM可同时管理多个线程块。
Q: CUDA编程的内存模型有哪些层次?
- 寄存器:每个线程私有,速度最快,容量最小。
- 共享内存:Block内线程共享,低延迟(几个周期),每SM几十KB。
- 全局内存(Global Memory):所有线程可访问,容量大(GB级),延迟高(数百周期)。
- 此外还有常量内存(只读,有缓存)、纹理内存(优化空间局部性访问)和本地内存(寄存器溢出)。
Q: CUDA编程的软件模型(线程层次)是怎样的?
- Thread:最小执行单元,执行kernel中的代码。
- Block(线程块):多个线程组成,Block内线程可通过共享内存通信和同步。
- Grid(线程格):多个Block组成,对应一次kernel launch的所有线程。
硬件执行时,Block被分配到SM上,线程以Warp(32线程)为单位调度执行。
Q: CUDA Stream的概念是什么?
Stream是主机发出的在设备中执行的CUDA操作序列(包括kernel执行和主机-设备数据传输)。同一stream中的操作按序执行,不同stream中的操作可以并发执行。利用多stream可以实现计算与数据传输的重叠(overlap),提高GPU利用率。
Q: 使用共享内存时需要注意什么?
- 线程同步:在利用共享内存进行线程间协作前,必须调用
__syncthreads()确保共享内存数据对Block内所有线程准备就绪。 - 避免Bank Conflict:共享内存分为32个bank,同一Warp内多个线程访问同一bank的不同地址时产生冲突,导致串行化。解决方法包括padding和调整访问模式。
Q: 对一个CUDA kernel进行优化可以从哪些角度入手?
- 访存优化:合并全局内存访问(coalesced access)、使用共享内存减少全局内存访问、避免bank conflict。
- 计算优化:减少分支divergence、使用快速数学指令、利用Tensor Core。
- 并行度优化:调整block size提高SM占用率(occupancy)、减少寄存器使用。
- 延迟隐藏:增加活跃warp数、计算与访存重叠。
- 减少同步:减少__syncthreads调用、使用warp-level原语。
Q: GPU L1/L2缓存的特点?
- L1 Cache:每个SM私有,与共享内存共享同一物理存储(可配置分配比例),延迟低。Ampere及以后架构中L1统一管理。
- L2 Cache:所有SM共享,容量更大(数MB),用于缓存全局内存访问,降低DRAM访问频率。L2命中延迟约200+周期,远低于DRAM访问。
Q: 同步Stream和异步Stream的区别?
- 默认Stream(同步):stream 0,操作会与其他stream同步,具有隐式同步行为。
- 非默认Stream(异步):显式创建的stream,操作异步执行,不与其他stream同步(除非显式同步)。可以实现kernel执行与内存传输的overlap。
Q: 手撕:CUDA实现矩阵乘法?
(编程题)
Q: 手撕:CUDA实现Softmax规约?
(编程题)
Q: 手撕:CUDA实现NCHW转NHWC?
(编程题)
Q: 手撕:CUDA统计数组中每个元素出现的频率(元素范围0-256)?
(编程题)
Q: 手撕:将数组中奇数位置的数放在左边、偶数位置的数放在右边(原地操作)?
(编程题)