CC BY 4.0 （除特别声明或转载文章外）
如果这篇博客帮助到你，可以请我喝一杯咖啡~
引子
大四保研之后除了打 ASC 之外并没有别的事情做，有亿点想在研究生之前找一段实习。正好上学期快接近期末的时候，和字节 AILab 的几位同学（悦航 & elfin）加了微信聊了一下，感觉他们需要做的东西和我研究生想做的方向比较符合，于是答应在 ASC 决赛名单出了之后考虑一下后面的时间安排再投一投简历。
其实我本来是打算三月份左右再开始找一些暑期实习的，因此还没有准备好简历，临时在 GitHub 上拉了一个看起来不错的模板 billryan/resume。做简历的时候有点想吐槽自己好像没有什么能拿出手的论文或者项目经历，大部分都是课内作业和比赛项目…唯一有一点 star 的还是前端项目，也有丶离谱。
19 号上午，ASC 赛会公布了决赛名单，我们还在上面，这意味着我在四月份到五月中的一段时间要同时准备决赛和毕业论文答辩。当天晚上，我把自己的简历发了过去。字节的流程非常快，当晚就有 HR 同学加我的微信。第二天上午，我在回家的高铁上和 HR 约了一面的时间在四天后的下午三点。
一月二十四日，一面
我是在线上面试，点开邮件里的链接后会跳到牛客的线上面试平台。除了和面试官视频的窗口外，还有一个代码块编辑器和一键运行的按钮，显然面试会有线上 coding 环节。自我介绍完之后，果然马上就要编程，不过并不是考察算法，而是要写 CUDA。
输入一个 $M\times N$ 的矩阵，输出其各行向量规约后得到的向量 $Y_i = \sum_{j=0}^{N-1}A_{i,j}$
拿到这个问题有一点点出乎我的意料，最初和悦航加好友就是因为一年之前刚学 CUDA 时候写的两篇区间规约的博客（用 Shuffle 加速 CUDA 上的 Reduce 操作、再谈 CUDA 上的 Reduce 操作），但我确实有很长时间没有做过类似的操作了。所以在开始码代码的时候，回想之前是怎么做的同时，我也同时问了几个问题：
- 数据范围大约是多大？（大约 $2048\times 2048$ 这个数量级）
- 被归约的数据类型是什么？问到一半，我想了想，顺手写了
template。 - 洗牌指令是 CUDA 6.5 引入的，但是现在的显卡基本都支持了，可不可以用？（可以）
- 可以通过多级规约/多次调用 cublas 的 asum 过程实现？（希望只启动一次 kernel ，减少多次启动的开销）
一边问一边写的差不多了。不过还是有些头大，有些 API（__syncthreads）现场忘了，也不好去查文档。好在这里面试官说只要意思到了就可以了。
template <typename T, int WARPSIZE, int BLOCKSIZE>
__global__ void reduce(T *A, int lda, int n, T *y)
{
A += lda * blockIdx.x;
__shared__ T smem[BLOCKSIZE];
{
T val = 0;
for (int id = threadIdx.x; id < n; id += BLOCKSIZE)
val += A[id];
smem[threadIdx.x] = val;
}
#pragma unroll
for (int offset = BLOCKSIZE >> 1; offset > (WARPSIZE >> 1); offset >>= 1)
{
__syncthreads();
if (threadIdx.x < offset)
smem[threadIdx.x] += smem[threadIdx.x ^ offset];
}
if (threadIdx.x < WARPSIZE)
{
T val = smem[threadIdx.x];
#pragma unroll
for (int offset = WARPSIZE >> 1; offset > 0; offset >>= 1)
val += __shfl_xor_sync(0xffffff, val, offset, WARPSIZE);
if (threadIdx.x == 0)
y[blockIdx.x] = val;
}
}
码完之后，面试官又继续问了这个核函数启动时候的 blockDim 和 gridDim 怎么设置（前者 (1024, 0, 0)，后者 (m, 0, 0)）；答上来之后接着要我介绍了一下 grid，block，warp 等一些常见概念的意思（任务分发的最小单位，线程通信的最小单位，线程调度的最小单位）。
然后面试官又继续问了一下，是否有更好的方案？比如这里用了很多的 shared memory，shared memory 过大会降低 SM 占用率。我算了一下，如果这里规约的类型是 8 字节的 double，即使 BLOCKSIZE 设置成 1024，每个 block 也只用了 8KB shared memory，按照 V100 96KB per SM 的容量来算，每个 SM 做多可以启动 12 个 block，而实际上每个 SM 最多只有 2048 个活跃线程，因此此处是完全不会影响 SM 占用率的。
不过我好像会错意了，这里面试官的意思是说，使用 shared memory 的数量可以再少一些，可以每个 warp 之间用 shuffle 指令规约后，使用 shared memory 分享结果，再交给第一个 warp 规约，这样可以减少需要 shared memory 的数量减少 WARPSIZE 倍。我觉得见仁见智，因为这样做实际上增大了线程的计算量（我写的代码里，下标大的线程在前几次规约后就可以空转了）；同时这也要求 BLOCKSIZE == WARPSIZE * WARPSIZE ，并不具有很好的可移植性（在 ROCm 中 WARPSIZE 通常是 64）。
事后想一想，区间 reduce 这个案例确实是很适合面试现场问的一个问题，不像矩阵向量乘、矩阵乘法这样有太多细节容易写错，同时也有足够多优化的空间。
再之后时间还有很多空余，就又来了一道传统的算法题。
输入十进数 M < 1e8 , 输出最小的一个由 0、1 组成的十进数 N，满足 N > 0, N % M = 0， 如不存在输出 -1. 例子： Input : 2 output : 10 Input : 3 output : 111
一开始我在什么时候不存在这个解的情况纠结了一会儿。不过，面试官一直在引导我，先不考虑这些情况，直接去做呢？直接做的话那就是敲一个 BFS 了。太久没写 cpp，我甚至一开始忘记了 string 的下标是从前往后还是从后往前。此外，线上 coding 还是有点不习惯，因为全角半角括号没看出来导致一开始编译一直不过。
跌跌撞撞总算是敲出了一个能跑的版本。之后，面试官问我，这个算法的空间复杂度上还能不能优化？我好像又把问题想复杂了（我回答「这里搜索状态是用 string 表示的一个大整数，在搜索队列拉长的时候状态也会变得很大。可以用一个整数表示这是当前搜索到的第几个大整数，然后从这个整数的二进制位来还原大整数，从而减少空间上的占用」）。面试官继续提示我，这里每次搜索结果只和我的余数 sum 有关。于是我明白了他的意思，对这个余数判重就好了。不过确实，受数据范围 M < 1e8 的影响，我一开始并没有往这个方向上想（判重需要上百 M 的空间，一些算法竞赛里的内存限制是 64M ~ 128M）。如果数据范围是 M < 1e6 的话可能会好一些叭。
#include <bits/stdc++.h>
using namespace std;
int main()
{
int m;
cin >> m;
vector<char> flag(m, 0);
for (deque<string> q(1, "1"); !q.empty(); q.pop_front())
{
string cur = q.front();
int sum = 0;
for (int i = 0; i < cur.size(); ++i)
{
sum = (sum * 10 + (cur[i] - '0')) % m;
}
if (sum == 0)
{
cout << cur;
return 0;
}
if (!flag[sum])
{
flag[sum] = 1;
q.push_back(cur + "0");
q.push_back(cur + "1");
}
}
cout << -1;
}
再之后的面试基本上就以聊天为主，依稀记得有聊到，除了字节还想去哪里吗？我就老实说了，研究生还想去绿厂（NVIDIA）康康，毕竟要在 GPU 上做研究的话还是他家比较在行；不过在那之前我想先从「客户」的视角看看工业界是怎么用 GPU 的；接着又聊了聊字节这边的一些情况。面试结束后也有和面试官加微信，总体来说感觉人挺好 d。
一月二十四日，二面
二面上来又是 coding ！
$n, x, y, z$ 都是非负整数，求针对任意输入 $n$ ，满足 $n = 5x + 2y + z$ 的解 $(x, y, z)$ 的个数。 例如 $n = 5$ 的时候有如下四种情况：
n x y z 5 0 0 5 5 0 1 3 5 0 2 1 5 1 0 0
是我比较讨厌的推公式题（一般丢给队友做），不过真在面试碰上了也只能硬着头皮来了。实际上我自己做这类题一般的习惯是先写个暴力打表，然后找规律或者直接查 OEIS…感觉上这个题会有一个 $O(1)$ 的公式，不过不是很直观。只好先写一个 $O(\frac{n}{5})$ 枚举 $x$ 的代码。现场打了前 100 的表出来，我也并没有一眼看出来；组合数学的角度也没有好思路。
这时候面试官提醒了一下，叠加的时候每次加的表达式都是一个只关于变量 x 的函数。大概明白意思了，把表达式里的公共项提出来即可，然后可以使用数列求和公式。不过 5*x/2 这个式子里有整除，并不是等差数列，因此要按照 x 分奇偶讨论。写完奇数部分的代码已经很繁琐，让人不想再看了，偶数部分刚要开始写，面试官说可以了。
#include <iostream>
using namespace std;
int main()
{
int n;
cin >> n;
//for(int n = 1; n < 100; ++n)
//{
long long ans = 0;
/*
for(int x=0; 5*x<=n; x+=2)
{
ans += n /2 - 5*x/2 + 1;
}
*/
int t = n / (5 * x) / 2 + 1;
ans = (n / 2) * t + t + (0 + 5 * t) * t / 2;
for (int x = 1; 5 * x <= n; x += 2)
{
//ans += (n - 5*x) /2 + 1;
ans += (n - 1) / 2 - 5 * x / 2 + 1;
}
ans += ...;
cout << ans << ' ';
//}
}
再之后的面试感觉也是以聊为主。比如开始写 CUDA 有多少时间（从大三上的高性能计算课开始算的话是一年出头）；介绍一些自己写 CUDA 的经验（我记得我说了「CUDA 写起来结构感很强，因此如果要写出高性能的实现的话代码里一般会出现很多嵌套的 for 循环，包括 NVIDIA 自己给出来的 CUTLASS 库也是这样的；另外也建议多看 NVIDIA 的白皮书，练一些官方推荐的优化案例，比如矩阵向量乘、区间规约、矩阵乘法」）。又聊了些别的话题，记得有要我介绍自己做过的一个比较大的项目经历，就说了之前打 ASC 时候做 QuEST 优化的一些经历，感觉还行。然后话题又往超算队聊了一聊。
二面将近结束的时候，面试官说，他之前已经加我微信和我聊过了（我楞了一下，问是悦航吗？是 elfin），可能是电脑外放太差我并没有从声音上听出来吧 hh 在之后又聊了聊入职之后想做什么，我说还是想做一些比较底层的东西，于是我们又提了一下之前聊到的 nvrtc。
总之两轮技术面就这样过了，确实和传闻中的一样，到哪都要直接写代码。不过我觉得题目难度并不高（也许是因为我只是实习面试？），大概是 CF div2 B~C 题的难度，正常去做的话二十分钟应该可以做出来，不需要再特地去刷题（感觉最好还是提前练那么一两道找手感）。或者说，面试官在做题的时候其实也不会就在那里等着，会给一些提示和引导；刚好我做题时有这样那样啰里吧嗦的习惯，一眼没看出好的思路的时候也没有感觉特别尴尬。
一月二十六日，三面
第二天，之前的 HR 同学就和我说过了技术面，之后还有 HR 面。于是约了在之后一天去面试。由于不需要线上编程，这次邮件点开之后是飞书的视频面试界面。
嗯，三面感觉还是唠嗑。我记得聊的话题有，介绍一下简历上打的比赛（还是介绍了超算的一些比赛）；你在赛队里起到了什么作用（队长，内培/赛题分工/和学院老师联系/GPU 方向优化）；你感觉工作量最大的事情是什么（介绍了招新相关的一些事情）。总之，我觉得参加过超算比赛还是有挺多内容可以说的，倒是挺好的。然后 HR 和我说，前两面面试官反馈都挺好的，应该可以拿到 offer。于是也聊了聊入职之后的安排（三月底到五月中需要一个半月准备 ASC 决赛和毕业论文答辩；这部分可以请一个长假）。最后 HR 也有问我有什么想问的吗？我回答，之前和 elfin 聊的已经比较多了，对这个实习在做什么还有能学到什么已经有一个大致的估计了。
总之，年轻人的第一次实习面试感觉来的太快，毫无准备。好在过了一两天，确定了拿到 offer，辣就完结撒花叭~
你们说我一个在广东上学的安徽人，怎么就到北京来实习了呢