某显卡公司的面试题目

[guevara]

某显卡公司的面试题目



https://ift.tt/ENdbtMA



Rainmaker's Notebook

前言

卖艺难，卖身更难，今年求职情况可谓惨不忍睹。去年年底我申了 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 或其他很常见的算法题出现在任何一场面试中。

数据结构

1. 写一个数据结构支持 void enqueue(一个函数指针，void *arg, size_t arg_size) 和 void drain(调用所有 enqueue 的函数和 arg), 每次 enqueue 的时候只能发生一次 malloc. 更进一步，实现的代码能否支持在传进 enqueue 的函数里调用 enqueue, 如果要支持需要怎么实现。

2. 有一个发送器 @ 1GHz, 有一个接收器 @ 0.8 GHz. 假设发送器每 100 cycles 里面有 80 个发包 20 个闲置，如果采用一个 FIFO buffer, 这个 buffer 最小要有多大？

3. 一个浮点数数组从前往后加和从后往前加得到的结果是否一样？

4. 给一千万个 FP16 数值进行排序，最快的串行算法是什么，其时间和空间复杂度是什么？

系统

1. Pinned/unpinned host memory access 的区别是什么？

2. 用 compare-and-swap 原子指令实现一个 atomic add.

3. 用原子操作指令实现一个简单的 mutex.

4. sched_yield() 和 sched_resume() 这两个函数的用法。

5. 实现一个无锁队列支持多生产者多消费者模型。进一步，如果每次入队的数据量超过一个 cache line 的大小，应该如何正确实现？

6. 如何设计 GPU-aware 的 MPI_Send()/MPI_Recv() （即这可以直接将 GPU 内存指针传入这两个函数）? 进一步，如果我们需要多次发送，每次都初始化 IPC handle 会很浪费时间，如何解决？ 如果还是多次发送，但是 buffer 会变，如何处理？

7. 给一个异构系统写一个 malloc()/free() 函数，可以在 CPU 或者 GPU 上面分配和释放内存。但是硬件架构可能是 CPU 直连 GPU，也可能是通过 Root Complex 和 Switch 连接。仅可以调用系统的 mmap()/munmap() 和 GPU 上类似的函数。

并行计算

1. 什么是 Amdahl’s law? 如果一个串行程序有 95% 运行时间对应的部分可以并行化，此程序最大的并行加速比是多少？

2. 有 $N$ 个不同的不可拆分的计算任务，第 $i$ 个任务耗时 $i$ 单位，如果有足够多核心，最小运行时间多少？如果可以选核心数，最快但是最不空置核心的并行方法是什么？

3. CPU 和 GPU 硬件架构的主要区别是什么？

4. GPU 编程有什么要注意的影响性能的地方？

5. 设计一个算法在 GPU 上并行解压 (n_repeat, char) 这样的 token 对。每个 stream 有许多这些 token 对，可能有一个或者多个 stream. 例： (3, ab), (5, c) 解压后的输出为 abababccccc. 这个算法有什么性能问题？

6. 写一个 GPU 上的并行矩阵转置函数，给出优化版本和说明优化原因。

7. 写一个 GPU 上的并行排序函数。

硬件

1. 一个 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 有一些新的可用数据？

2. 同上一问的硬件拓扑结构。若 NIC 1 向 GPU 2 的显存进行 RMDA put, 有无数据需要经过 CPU? （注：考点疑似是 IOMMU）

3. 什么是 Base Address Register (BAR), 其作用是什么？

4. 单级 cache CPU, cache size 4MB, 初始化一个 2MB 数组和一个 8MB 数组，哪一个的运行时间更短？

5. 8-way associated 和全相连 cache 的区别是什么，它们各自有什么优缺点？

6. 假设在一个 FP32 峰值性能为 10 TFlops 的 GPU 上做单精度矩阵乘法，并且已知输入输出矩阵所需内存远大于 GPU 显存。GPU 通过 16GB/s 的 PCIE 连接到 CPU, 理论上最少需要多少显存才足够让 GPU 能够跑满浮点数峰值计算性能？

7. 写一个 CPU 上的单线程矩阵转置函数。这个实现会产生多少对主内存的访问？

8. 为什么 L1 缓存即使不考虑价格也不能无限大？

网络通信

1. MPI 中单边 (get, put) 和双边 (send, recv) 操作的主要区别是什么？什么时候选择单边/双边操作？

2. MPI 中的小鹰模式和约会模式的区别是什么？

3. MPI one-sided 和 SHMEM 的主要区别是什么？

4. 用 MPI one-sided 操作和 MPI_Barrier() 实现一个 MPI_Allreduce(), 可以假设进程数是二的幂。分析其时间复杂度。如果不允许使用 MPI_Barrier(), 如何实现？

一般性的编程问题

1. C++ 虚函数是什么，编译器如何处理虚函数的跳转？ （注：考点应该是vtable）

2. 交换 a b 两个整数，不用第三个寄存器，最少要多少次什么操作？进一步，将 a b c 变为 c a b, 不用第四个寄存器，要多少次什么操作？

3. $N$ 个数里找第 $k$ 小的数（$k$ 是一个预先知道的固定的很小的正整数），应该用什么算法，时间复杂度是什么？

4. 排序 $N$ 个取值范围是 $[0, 1000]$ 的整数，时间复杂度是什么？假设有 $P$ 个核心来并行化这一操作，其时间复杂度是什么？

5. 堆栈模拟递归实现和真递归实现哪个快？堆栈模拟的入堆元素比递归的少，少多少？哪些寄存器不需要入堆？

via Rainmaker's Notebook

September 25, 2024 at 07:15PM