某显卡公司的面试题目
Published:
前言
卖艺难，卖身更难，今年求职情况可谓惨不忍睹。去年年底我申了 24 所北美大学的 CS 教职，喜提全拒德（严格意义上来说，还有 20 所至今没有给我发拒信）。去年和今年我都去系里的 job talk 蹭了不少饭，眼看着应聘者的引用数从 $[200, 2000]$ 的区间暴涨到 $[500, 5000]$, 而且所有来做 job talk 的人做的东西都和 AI 沾边。二月份赶紧跑路找工业界打螺丝的岗位，先后投了如下公司：
- 硬件公司：AMD, Intel, NVIDIA, HPE
- 互联网公司：Amazon, Google, Microsoft, Apple, ByteDance
- 工业公司：ANSYS, DE Shaw Research
只有三个公司给了我面试机会，都是因为有强力内推才拿到的面试。最后我面完了某显卡公司的两个职位，一个面了九场，一个面了八场外加两个线下四十八小时代码优化测试。下面是这两个职位的面试中的一些题目，仅供后来诸君参考。由于所有题目均为面试以后回忆记录，可能部分题目描述不够充分。另外，这两个职位的面试都很有针对性，没有任何 LeetCode 或其他很常见的算法题出现在任何一场面试中。
数据结构
写一个数据结构支持
void enqueue(一个函数指针，void *arg, size_t arg_size)和void drain(调用所有 enqueue 的函数和 arg), 每次 enqueue 的时候只能发生一次 malloc. 更进一步，实现的代码能否支持在传进 enqueue 的函数里调用 enqueue, 如果要支持需要怎么实现。有一个发送器 @ 1GHz, 有一个接收器 @ 0.8 GHz. 假设发送器每 100 cycles 里面有 80 个发包 20 个闲置，如果采用一个 FIFO buffer, 这个 buffer 最小要有多大？
一个浮点数数组从前往后加和从后往前加得到的结果是否一样？
给一千万个 FP16 数值进行排序，最快的串行算法是什么，其时间和空间复杂度是什么？
系统
Pinned/unpinned host memory access 的区别是什么？
用 compare-and-swap 原子指令实现一个 atomic add.
用原子操作指令实现一个简单的 mutex.
sched_yield()和sched_resume()这两个函数的用法。实现一个无锁队列支持多生产者多消费者模型。进一步，如果每次入队的数据量超过一个 cache line 的大小，应该如何正确实现？
如何设计 GPU-aware 的
MPI_Send()/MPI_Recv()（即这可以直接将 GPU 内存指针传入这两个函数）? 进一步，如果我们需要多次发送，每次都初始化 IPC handle 会很浪费时间，如何解决？ 如果还是多次发送，但是 buffer 会变，如何处理？给一个异构系统写一个
malloc()/free()函数，可以在 CPU 或者 GPU 上面分配和释放内存。但是硬件架构可能是 CPU 直连 GPU，也可能是通过 Root Complex 和 Switch 连接。仅可以调用系统的mmap()/munmap()和 GPU 上类似的函数。
并行计算
什么是 Amdahl’s law? 如果一个串行程序有 95% 运行时间对应的部分可以并行化，此程序最大的并行加速比是多少？
有 $N$ 个不同的不可拆分的计算任务，第 $i$ 个任务耗时 $i$ 单位，如果有足够多核心，最小运行时间多少？如果可以选核心数，最快但是最不空置核心的并行方法是什么？
CPU 和 GPU 硬件架构的主要区别是什么？
GPU 编程有什么要注意的影响性能的地方？
设计一个算法在 GPU 上并行解压 (n_repeat, char) 这样的 token 对。每个 stream 有许多这些 token 对，可能有一个或者多个 stream. 例： (3, ab), (5, c) 解压后的输出为 abababccccc. 这个算法有什么性能问题？
写一个 GPU 上的并行矩阵转置函数，给出优化版本和说明优化原因。
写一个 GPU 上的并行排序函数。
硬件
一个 CPU 连着 PCIE switch 1 和 PCIE switch 2. 每个 PCIE switch 分别连着一个 NIC 和一个 GPU. 两个 NIC 之间有网线，可以 RDMA. 若 PCIE switch 1 上的 NIC 向 CPU 报告其完成了一个将 GPU 1 显存中的数据进行 RDMA put 到 GPU 2 显存上的操作，CPU 是否立即可以告知 GPU 2 有一些新的可用数据？
同上一问的硬件拓扑结构。若 NIC 1 向 GPU 2 的显存进行 RMDA put, 有无数据需要经过 CPU? （注：考点疑似是 IOMMU）
什么是 Base Address Register (BAR), 其作用是什么？
单级 cache CPU, cache size 4MB, 初始化一个 2MB 数组和一个 8MB 数组，哪一个的运行时间更短？
8-way associated 和全相连 cache 的区别是什么，它们各自有什么优缺点？
假设在一个 FP32 峰值性能为 10 TFlops 的 GPU 上做单精度矩阵乘法，并且已知输入输出矩阵所需内存远大于 GPU 显存。GPU 通过 16GB/s 的 PCIE 连接到 CPU, 理论上最少需要多少显存才足够让 GPU 能够跑满浮点数峰值计算性能？
写一个 CPU 上的单线程矩阵转置函数。这个实现会产生多少对主内存的访问？
为什么 L1 缓存即使不考虑价格也不能无限大？
网络通信
MPI 中单边 (get, put) 和双边 (send, recv) 操作的主要区别是什么？什么时候选择单边/双边操作？
MPI 中的小鹰模式和约会模式的区别是什么？
MPI one-sided 和 SHMEM 的主要区别是什么？
用 MPI one-sided 操作和
MPI_Barrier()实现一个MPI_Allreduce(), 可以假设进程数是二的幂。分析其时间复杂度。如果不允许使用MPI_Barrier(), 如何实现？
一般性的编程问题
C++ 虚函数是什么，编译器如何处理虚函数的跳转？ （注：考点应该是vtable）
交换 a b 两个整数，不用第三个寄存器，最少要多少次什么操作？进一步，将 a b c 变为 c a b, 不用第四个寄存器，要多少次什么操作？
$N$ 个数里找第 $k$ 小的数（$k$ 是一个预先知道的固定的很小的正整数），应该用什么算法，时间复杂度是什么？
排序 $N$ 个取值范围是 $[0, 1000]$ 的整数，时间复杂度是什么？假设有 $P$ 个核心来并行化这一操作，其时间复杂度是什么？
堆栈模拟递归实现和真递归实现哪个快？堆栈模拟的入堆元素比递归的少，少多少？哪些寄存器不需要入堆？