Interview Experience | GPU Kernels / Triton | π₯π₯ | AMD
2026 (Sep) Sr. Software Engineer - AI Triton Kernels @ AMD β San Jose, CA β OA + 3 interviews | π Neutral | Hard | Offer
An OA and three interviews. The posting is about writing Triton / Gluon kernels (matmul, attention, MoE), so there was no ordinary system design round. This is the loop for my req. I've heard other teams run it differently, so take it as one data point.
Round 1 β OA
I only remember one question. Given an original shape and a target shape that has at most one -1, work out the -1 dimension, or report an error if it can't be worked out.
Some testcases failed. One of them had a 0 in the dimensions. The statement didn't say what -1 should mean next to a 0; I returned an error, and I don't know if that was the intended answer.
After the OA the recruiter asked for my availability and then went quiet for a long time.
Round 2 β Technical Interview 1
The interviewer joined late and said up front that time was cut, so there would be only one question.
He pasted a Triton kernel into a shared doc. Row-wise softmax, and the output didn't match PyTorch. Find out why.
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)
I suspected fp16 accumulation first. He said the input was fp32, so that was a detour. What got me to the real problem was his question about what the lanes beyond n_cols read when BLOCK is bigger than n_cols. They read 0.0, and they still go into the max and the sum. I changed other to -float("inf").
Follow-up: what if one row doesn't fit in a single BLOCK. Loop over chunks, keep a running max and a running sum, and when the max changes multiply the old sum by exp(old max β new max). I wrote it that way and he didn't say whether it was right. Whether the output pass has to read x a second time: I said yes, he asked if there was a way around it, and I didn't get to an answer.
For choosing BLOCK I said autotune over BLOCK and num_warps. He brought up waves_per_eu. I haven't used it and said so.
Round 3 β Technical Interview 2
Mostly Q&A.
First one was memory ordering. One thread writes data and then a flag. Another thread sees the flag and reads data. Why can it still get the old value? He mixed barrier, fence and cache coherence while explaining the question and I lost the thread. I asked him to say it again, still didn't follow, and he moved on.
Second was LDS bank conflicts. Accesses to the same bank get serialized, that part I had. Then how many banks the LDS has on AMD. No answer.
There were one or two more about occupancy and register usage. I don't remember how they were phrased.
Round 4 β Hiring Manager
Mostly conversation. He asked why Triton, and I ended up telling a benchmark story. My numbers didn't match a coworker's, I spent a while on it, and it turned out the two sides had built with different compile flags. Nothing to do with the kernel. He asked how the baseline had been picked; I got halfway through and he moved to the next thing.
He mentioned in passing that the role doesn't do visa sponsorship. The kernels on the team are mostly matmul, attention and MoE, and which one I'd start on wasn't said. I asked how many days a week hybrid means. He said the recruiter would go through that with me.