面经 | GPU Kernels / Triton | 🔥🔥 | AMD
2026(9月) Sr. Software Engineer - AI Triton Kernels @ AMD | San Jose, CA · OA + 三轮面试 | 😐 Neutral | Hard | Offer
OA 加三场面试。岗位写的是 Triton / Gluon kernel，matmul、attention、MoE 那一类，所以没有常规的 system design。这是我这个 req 的流程，听说别的组不一样，当个数据点看就行。
Round 1 — OA
只记得一题。给原始 shape 和目标 shape，目标里最多有一个 -1，把 -1 那一维推出来，推不出来就报错。
有几个 testcase 没过，其中一个维度里带 0。-1 碰上 0 该算什么，题面没写，我当时返回了报错，不知道是不是想要的结果。
OA 之后 recruiter 问了 availability，接着很久没动静。
Round 2 — Technical Interview 1
面试官晚进来了一会儿，开口先说时间被压了，只做一题。
他在共享文档里贴了一段 Triton kernel，按行做 softmax，输出和 PyTorch 对不上，让我找原因。
x = tl.load(x_ptr + row * n_cols + offs, mask=mask, other=0.0)
x = x - tl.max(x, axis=0)
num = tl.exp(x)
den = tl.sum(num, axis=0)
tl.store(out_ptr + row * n_cols + offs, num / den, mask=mask)
先怀疑的是 fp16 累加，他说输入就是 fp32，这条路白走了一阵。是他问了一句 BLOCK 比 n_cols 大的时候，多出来那些 lane 读到的是什么，才看到 other=0.0：它们读进来是 0，照样参与 max 和 sum。改成 other=-float("inf")。
follow-up：一行放不进一个 BLOCK 怎么办。分块循环，维护 running max 和 running sum，max 变了就把之前的 sum 乘 exp(旧 max − 新 max)。我是这么写的，他没说对不对。输出那一遍要不要把 x 再读一次，我说要，他问有没有办法不读，这个没答上来。
BLOCK 怎么选，我说 autotune 扫 BLOCK 和 num_warps。他提了 waves_per_eu，没用过，我直接说了没用过。
Round 3 — Technical Interview 2
基本是问答。
第一题 memory ordering：一个线程先写 data 再写 flag，另一个线程看到 flag 之后读 data，为什么还可能读到旧值。他讲题的时候 barrier、fence、cache coherence 几个词混着用，我没跟上，让他再说一遍，还是没懂，他就跳到下一题了。
第二题 LDS bank conflict。同一个 bank 的访问会串行，这个答了。接着问 AMD 上 LDS 一共有多少个 bank，答不上来。
还有一两个 occupancy 和寄存器用量的问题，具体怎么问的不记得了。
Round 4 — Hiring Manager
基本是聊天。问我为什么做 Triton，我讲了个 benchmark 的事：我的数字和同事的对不上，查了一阵，最后是两边编译选项不一样，跟 kernel 没关系。他问 baseline 当时是怎么选的，我讲到一半他转去下一个话题了。
他顺口提了一句这个岗位不做签证 sponsor。组里的 kernel 主要是 matmul、attention、MoE，我具体先做哪个没说。我问 hybrid 一周去几天，他说 recruiter 会跟我对。