← 返回部落格列表 NVIDIA Coding Questions 高頻題型分布：CUDA × C++ × 系統級題逐個拆解
OA NVIDIA

NVIDIA Coding Questions 高頻題型分布：CUDA × C++ × 系統級題逐個拆解

2026-06-02

NVIDIA 的 coding 面試和大多數網路公司不太一樣：純 LeetCode-style 演算法題只佔 30%-40%，其餘的題目偏向 CUDA 核心語意、C++ 記憶體模型、數值穩定性、以及編譯器/驅動級的系統題。如果你按「FAANG 演算法刷題」路線準備，進了 onsite 很容易在第二三輪 freeze。

這篇文章把 NVIDIA HackerRank OA 與 VO 中真正高頻的 coding 題按五大類歸併，逐類給一道代表題 + 解法 + 職位線側重對照。讀完你應該能根據自己的目標職位（Compiler / DL Frameworks / Driver / Robotics / SWE）做精準刷題。

五大題型分布

題型 佔比 平台 決策權
CUDA 核心語意 25% HackerRank C++ + 白板 DL Frameworks 主輪
C++ 記憶體模型 30% CoderPad C++ Driver / Compiler 主輪
數值計算穩定性 15% HackerRank Python/C++ DL Frameworks 副輪
系統級演算法題 20% HackerRank SWE / Robotics 主輪
字串與編譯器題 10% HackerRank Compiler 副輪

關鍵差異：和站內已有的 NVIDIA HackerRank 文聚焦「四類高頻題型」不同，本篇按職位線決策權重重新切片，方便定向投遞。

題型一：CUDA 核心語意

代表題：Block 內 reduce sum

實作 __device__ float blockReduceSum(float val)，對一個 block 內所有執行緒的 val 求和，使用 warp shuffle 優化。

思路：

  1. 先用 __shfl_down_sync 在 warp 內 reduce
  2. 再用 shared memory 把每個 warp 的結果聚合
  3. 最後再做一次 warp reduce
__inline__ __device__ float warpReduceSum(float val) {
    for (int offset = 16; offset > 0; offset /= 2)
        val += __shfl_down_sync(0xffffffff, val, offset);
    return val;
}

__inline__ __device__ float blockReduceSum(float val) {
    static __shared__ float shared[32];
    int lane = threadIdx.x % 32;
    int wid = threadIdx.x / 32;
    val = warpReduceSum(val);
    if (lane == 0) shared[wid] = val;
    __syncthreads();
    val = (threadIdx.x < blockDim.x / 32) ? shared[lane] : 0;
    if (wid == 0) val = warpReduceSum(val);
    return val;
}

關鍵點：

題型二：C++ 記憶體模型

代表題：實作一個無鎖 SPSC queue

實作 single-producer single-consumer 無鎖佇列，支援 push / pop，禁止使用 mutex。

#include <atomic>
#include <vector>

template <typename T>
class SPSCQueue {
public:
    explicit SPSCQueue(size_t cap)
        : data_(cap), cap_(cap), head_(0), tail_(0) {}

    bool push(const T& v) {
        size_t t = tail_.load(std::memory_order_relaxed);
        size_t next = (t + 1) % cap_;
        if (next == head_.load(std::memory_order_acquire)) return false;
        data_[t] = v;
        tail_.store(next, std::memory_order_release);
        return true;
    }

    bool pop(T& out) {
        size_t h = head_.load(std::memory_order_relaxed);
        if (h == tail_.load(std::memory_order_acquire)) return false;
        out = data_[h];
        head_.store((h + 1) % cap_, std::memory_order_release);
        return true;
    }

private:
    std::vector<T> data_;
    size_t cap_;
    std::atomic<size_t> head_, tail_;
};

關鍵點：

題型三：數值計算穩定性

代表題：Log-Sum-Exp 數值穩定實作

給一個 float 陣列，求 log(sum(exp(x)))，避免 overflow / underflow。

import math

def logsumexp(xs):
    if not xs:
        return float("-inf")
    m = max(xs)
    if m == float("-inf"):
        return m
    s = sum(math.exp(x - m) for x in xs)
    return m + math.log(s)

關鍵點：

題型四：系統級演算法題

代表題：Driver 記憶體配置器（Buddy System）

實作一個 Buddy Allocator，支援 alloc(size) / free(ptr)，最大空間 2^N，配置按 2 的冪向上取整。

思路：

#include <vector>
#include <set>
#include <map>

class BuddyAllocator {
public:
    BuddyAllocator(int n): N(n) {
        free_lists.resize(n + 1);
        free_lists[n].insert(0);
    }
    int alloc(int size) {
        int k = 0;
        while ((1 << k) < size) k++;
        int j = k;
        while (j <= N && free_lists[j].empty()) j++;
        if (j > N) return -1;
        int addr = *free_lists[j].begin();
        free_lists[j].erase(free_lists[j].begin());
        while (j > k) {
            j--;
            free_lists[j].insert(addr + (1 << j));
        }
        used[addr] = k;
        return addr;
    }
    void free(int addr) {
        int k = used[addr];
        used.erase(addr);
        while (k < N) {
            int buddy = addr ^ (1 << k);
            auto it = free_lists[k].find(buddy);
            if (it == free_lists[k].end()) break;
            free_lists[k].erase(it);
            addr = std::min(addr, buddy);
            k++;
        }
        free_lists[k].insert(addr);
    }
private:
    int N;
    std::vector<std::set<int>> free_lists;
    std::map<int, int> used;
};

關鍵點：

題型五：字串與編譯器題

代表題：Token 流的簡單 expression evaluator

輸入 ["3", "+", "4", "*", "(", "2", "+", "1", ")"]，求值。

def evaluate(tokens):
    def precedence(op):
        return {"+": 1, "-": 1, "*": 2, "/": 2}.get(op, 0)
    def apply(a, b, op):
        return {"+": a+b, "-": a-b, "*": a*b, "/": a//b}[op]

    nums, ops = [], []
    for t in tokens:
        if t == "(":
            ops.append(t)
        elif t == ")":
            while ops and ops[-1] != "(":
                b, a = nums.pop(), nums.pop()
                nums.append(apply(a, b, ops.pop()))
            ops.pop()
        elif t in "+-*/":
            while ops and precedence(ops[-1]) >= precedence(t):
                b, a = nums.pop(), nums.pop()
                nums.append(apply(a, b, ops.pop()))
            ops.append(t)
        else:
            nums.append(int(t))
    while ops:
        b, a = nums.pop(), nums.pop()
        nums.append(apply(a, b, ops.pop()))
    return nums[0]

關鍵點：

職位線 × 題型權重

職位 CUDA C++ 記憶體 數值 系統 編譯器
DL Frameworks ★★★ ★★ ★★★ ★ ★
Driver ★ ★★★ ★ ★★★ ★★
Compiler ★ ★★ ★★ ★★ ★★★
Robotics ★ ★★ ★★ ★★★ ★
SWE General ★ ★ ★ ★★★ ★

備考策略

階段 內容
第 1 週 鎖定目標職位，按權重表挑出主刷題型
第 2 週 把 LeetCode "NVIDIA" tag 與 CUDA samples 各刷 30 道
第 3 週 系統級題型 + 職位特定 follow-up
模考 Karat-style C++ mock × 2

FAQ

Q1：NVIDIA OA 是 HackerRank 還是 Karat？ HackerRank 是 OA 主平台，Karat 是 phone screen 階段。兩者題風差異很大。

Q2：CUDA 不熟能投 NVIDIA SWE 嗎？ 能。SWE General 職位 CUDA 題很少，主要看 C++ 與系統題。但 DL Frameworks / Driver 必須會 CUDA。

Q3：Coding 面試是 C++ 強制嗎？ DL Frameworks / Driver 強制；Compiler 偏好 C++/Rust；SWE 與 Robotics 接受 Python，但 C++ 加分。

Q4：白板 CUDA 題怎麼準備？ 先把 reduce / scan / matmul / convolution 四個經典核心手寫一遍。面試時面試官會逐行追問 sync 與 race condition。

Q5：從台灣投 NVIDIA 美國職位，OA 難度會變嗎？ 題目相同，但時區與英語溝通 是隱性門檻。建議先做 1-2 次英文 mock。

正在準備 NVIDIA Coding 面試？

如果你想做 CUDA 核心題逐行 review、目標職位精準定位，或希望面試日有真人 OA代面 / VO代面 全程陪跑，可以聊聊看完整的 OA輔助 / VO輔助 方案。

聯絡方式

需要面試真題與客製備戰計畫？立刻聯絡微信 Coding0201，獲取真題。

Email: [email protected] Telegram: @OAVOProxy