AMD AI Chip Software Engineer Interview: CUDA, Operator Development, and Chip Architecture

Background

I have a bachelor's in Computer Science and a master's focused on High-Performance Computing. After graduation, I spent 2 years as a software engineer at an AI chip startup, primarily working on CUDA operator development and model adaptation, also participating in inference engine optimization. Cambricon has always been a benchmark company in China's AI chip industry, so when I saw they were hiring AI chip software engineers, I applied immediately.

Honestly, I was pretty anxious after applying — Cambricon's technical requirements are high, and AI chip software is a competitive field. About a week and a half later, HR contacted me to schedule three technical rounds. I spent a week systematically reviewing CUDA programming, operator development, and chip architecture, and organized my past operator optimization projects into a presentation for easy reference during interviews.

Interview Process Review

Round 1: C/C++ + CUDA Fundamentals (~1 hour)

Round 1 was with an engineer in their early 30s, likely a technical lead. After a self-introduction, we dove into fundamentals.

C/C++ Section:

The first question surprised me — "Explain the implementation principle of C++ virtual function tables, and how they change under multiple inheritance." I'd reviewed this before, so I walked through vptr pointing to vtable, vtable layout under single inheritance, and multiple vptrs pointing to different vtables under multiple inheritance. The interviewer followed up on vtable changes under virtual inheritance — I didn't answer this completely, only mentioning that virtual inheritance introduces vbptr. The interviewer supplemented the specific structure of the virtual base class table.

Then memory management — "What's the difference between new/delete and malloc/free in C++? How do you use placement new?" I covered constructor/destructor calls, type safety, and memory allocation failure handling. For placement new, I gave an example of constructing objects on pre-allocated memory, commonly used in memory pools.

Smart pointers came up — "How is shared_ptr's reference count implemented? Is it thread-safe? What problem does weak_ptr solve?" I explained that shared_ptr maintains reference counts through a control block, the reference count itself is thread-safe, but the pointed-to object isn't. weak_ptr solves circular reference problems without incrementing the reference count.

CUDA Section:

The first CUDA question was about thread hierarchy — "What's CUDA's thread hierarchy? How do grid, block, and thread relate?" I started from SM hardware architecture, moved to the software-level grid-block-thread mapping, and explained the concept of warps. The interviewer followed up with "What is warp divergence? How does it affect performance?" I explained that if threads within the same warp execute different branch paths, they execute serially, reducing efficiency.

Then memory hierarchy — "What memory types exist in CUDA? What are their respective access speeds and use cases?" I listed global memory, shared memory, constant memory, texture memory, and registers, covering access latency and use cases for each. The interviewer specifically asked about shared memory bank conflicts — I explained that shared memory is divided into 32 banks, and if multiple threads in the same warp access different addresses in the same bank, conflicts occur, which can be avoided through padding.

Streams and events — "What's the purpose of CUDA streams? How do multiple streams parallelize?" I explained that streams are execution queues for GPU operations, and operations in different streams can execute in parallel. The interviewer followed up with "How do you synchronize between streams?" I mentioned cudaStreamSynchronize and cudaEventSynchronize.

The last question had me hand-write a simple CUDA kernel — vector addition. This was basic, and I finished quickly. The interviewer asked me to optimize it for memory coalescing, so I added __restrict__ qualifiers and explained aligned access.

After Round 1, the interviewer said "fundamentals are decent, but CUDA optimization experience needs strengthening," and told me to wait for Round 2.

Round 2: Operator Development + Performance Optimization (~1.5 hours)

Round 2 was with a more senior technical expert. The questions were noticeably deeper and more closely tied to actual work.

It opened with a practical question — "Tell me about the most challenging operator optimization project you've done." I chose a FlashAttention operator optimization, starting from the problems with the naive implementation — high memory usage, low memory access efficiency — then explaining how to optimize through tiling, online softmax, and recomputation strategies. The interviewer probed each aspect — "How did you determine the tile size? How do you ensure numerical stability with online softmax? How do you balance the trade-off between recomputation and storing intermediate results?"

Then performance analysis — "What tools do you use for CUDA performance analysis? How do you identify performance bottlenecks?" I discussed using Nsight Compute and Nsight Systems, and how to analyze metrics like occupancy, memory throughput, and compute throughput. The interviewer followed up with "If you find a kernel is memory-bound, how would you optimize it?" I covered memory coalescing, shared memory usage, and data prefetching.

Operator fusion came up — "What's the principle behind operator fusion? What are common fusion patterns?" I explained that operator fusion improves performance by reducing intermediate result writes and reads to global memory. Common patterns include element-wise fusion, reduce fusion, and conv+bn+relu fusion. The interviewer asked "What should you watch out for when fusing?" I mentioned data dependencies, computation precision, and debugging difficulties.

Tensor Core was also covered — "What's the programming model for Tensor Core? How do you use the WMMA API?" I introduced Tensor Core's matrix multiply-accumulate operations and APIs like wmma::load_matrix_sync and wmma::mma_sync. The interviewer followed up with "What data layout requirements does Tensor Core have?" I explained row-major and column-major alignment requirements and how to convert layouts using the ldmatrix instruction.

Round 2 also included an open-ended system design question — "If you were to design a general-purpose operator library, how would you architect it?" I covered operator registration, auto-tuning, multi-backend support, and computation graph optimization. The interviewer was particularly interested in auto-tuning, asking "How do you define the search space? What search strategy do you use?"

At the end of Round 2, the interviewer said "good project experience, but system design skills need improvement" — a fair assessment.

Round 3: Chip Architecture + Project Deep Dive (~1.5 hours)

Round 3 was with the department's technical lead, focusing more on architectural understanding and big-picture thinking.

First, they asked me to detail Cambricon MLU's architecture characteristics. Honestly, I wasn't well-prepared for this part — I only knew that MLU uses a GPGPU-like architecture with dedicated AI acceleration units. Seeing my limited architectural knowledge, the interviewer pivoted — "What AI accelerator architectures are you familiar with? How do they differ from GPUs?" I compared TPU's systolic array, Huawei Da Vinci's Cube unit, and Cambricon's MLUcore. The interviewer seemed satisfied with this answer.

Then model deployment — "What steps does a model go through from training to deployment on an AI chip?" I covered model export, operator adaptation, precision calibration, performance tuning, and end-to-end verification. The interviewer followed up with "How do you do precision calibration? How do you evaluate INT8 quantization accuracy loss?" I detailed PTQ and QAT methods and the differences between per-channel and per-tensor quantization.

Compilers came up — "How much do you know about deep learning compilers like TVM and TensorRT?" I discussed TVM's Relay IR, operator scheduling, and auto-tuning mechanism, as well as TensorRT's layer fusion, precision calibration, and automatic kernel selection. The interviewer asked "If TVM doesn't have an implementation for a particular operator, how do you add one?" I explained the custom operator registration process.

Finally, several open-ended questions — "What do you think is the biggest challenge in AI chip software stacks?" and "How do you see MoE models affecting chip architecture?" I shared my views, and the interviewer kept probing for details, testing depth of thought.

About a week after Round 3, HR notified me that I passed. The overall process was fairly smooth.

Key Questions Summary

C/C++:

1. Virtual function table implementation and changes under multiple inheritance

2. Differences between new/delete and malloc/free

3. Usage of placement new

4. Smart pointer implementation and thread safety

5. What problem weak_ptr solves

CUDA:

6. Thread hierarchy: grid, block, thread, warp

7. Warp divergence and its performance impact

8. Memory types and access speed comparison

9. Shared memory bank conflicts

10. CUDA streams and events

11. Memory coalescing optimization

Operator Development:

12. FlashAttention operator optimization approach

13. CUDA performance analysis tools and methods

14. Memory-bound kernel optimization strategies

15. Operator fusion principles and common patterns

16. Tensor Core programming model and WMMA API

17. General-purpose operator library architecture design

Chip Architecture & Deployment:

18. AI accelerator architecture comparison

19. Model deployment pipeline

20. INT8 quantization and precision calibration

21. TVM/TensorRT compiler principles

Key Takeaways

1. CUDA fundamentals must be solid. Cambricon's AI chip software position has high CUDA requirements — it's not enough to just write kernels; you need to understand the underlying hardware architecture and optimization principles. I recommend thoroughly studying the CUDA C++ Programming Guide and CUDA Best Practices Guide.

2. Have hands-on operator optimization experience. Interviewers will dig deep into your operator optimization projects, and every detail could be probed. I recommend completing at least 2-3 in-depth operator optimizations.

3. Understand the target chip's architecture. While interviewers won't expect you to know Cambricon's MLU inside out, you should at least understand its basic architectural characteristics and how it differs from GPUs. This demonstrates your interest in the position.

4. System design skills matter. Rounds 2 and 3 both involve system design questions, testing whether you can think from a holistic perspective. I recommend studying open-source project architectures.

5. Stay current with cutting-edge technology. The AI chip field evolves rapidly, and interviewers will ask about frontier topics. Keep up with top conference papers and industry developments.

FAQ

Q: Does Cambricon require specific educational credentials?

A: Technical rounds don't directly ask about education, but a master's degree is essentially the threshold. During my interview, I felt the interviewers valued project experience and practical ability more.

Q: Can you pass without AI chip experience?

A: If you have GPU programming experience, it's manageable, but you need to demonstrate your understanding of AI chips. I recommend studying Cambricon's public documentation and CNToolkit usage beforehand.

Q: Do they ask you to write code on the spot?

A: Yes. Round 1 had me hand-write a CUDA kernel; Round 2 required writing operator optimization pseudocode. Practice CUDA whiteboard coding in advance.

Q: What's the work intensity like?

A: From what I understand, Cambricon's work intensity is above average in the chip industry. Overtime is common but not extreme.

Q: How's the compensation?

A: Cambricon's compensation is above average in the AI chip industry. With stock options, the overall package is quite good.