ROCm Stack Atlas

Code / 题目

C[row,col] = Σ A[row,k] × B[k,col]

C[0,0] = 1×5 + 2×7 = 19
C[0,1] = 1×6 + 2×8 = 22
C[1,0] = 3×5 + 4×7 = 43
C[1,1] = 3×6 + 4×8 = 50

State / 状态变化

Why / 为什么

A tiny input makes every number inspectable. It teaches execution semantics; it does not demonstrate GPU speed.小输入让每个数都能人工检查。它适合学习执行语义，但不能证明 GPU 更快。

Observe / 如何观察

Prediction: [19, 22, 43, 50]. Do not open the stack yet.先写下预测：[19, 22, 43, 50]。此时无需看技术栈。

Code / 命令

# replace with the architecture reported by rocminfo
hipcc --offload-arch=gfx942 matmul2.hip -o matmul2

__global__ void matmul2(...);  // device code
int main() { ... }                 // host code

State / 状态变化

Why / 纠正常见误解

The CPU will not send each GPU instruction at launch time. It later submits a work descriptor that points to this already-compiled device code.启动时 CPU 不会逐条发送 GPU 指令。它之后提交的是一个工作描述符，指向这里提前编译好的设备代码。

Observe / 如何观察

rocminfo | grep -m1 gfx

Compiler is a build-time side path. Hide it for normal runtime launches.Compiler 是编译期侧路；普通运行时启动中应当隐藏。

Code / 你写的代码

hipSetDevice(0);
// Initialization can happen lazily on this or a later HIP call.
// The exact trigger is runtime-version dependent.

State / 状态变化

Why / 为什么

This setup makes later dispatch cheap. A HIP stream is an ordered runtime abstraction; do not assume it maps one-to-one to a fixed hardware queue.这次设置让后续派发更轻量。HIP stream 是运行时的有序抽象，不应假设它与某条固定硬件队列一一对应。

Observe / 如何观察

rocminfo

Look for the GPU agent and gfx target. First-use overhead can appear in a trace, but exact initialization timing varies.查看 GPU agent 与 gfx 目标。trace 里可能看到首次调用开销，但初始化发生的精确时刻会变化。

Code / 你写的代码

constexpr size_t bytes = 4 * sizeof(float);
float *dA, *dB, *dC;
hipMalloc(reinterpret_cast<void**>(&dA), bytes);
hipMalloc(reinterpret_cast<void**>(&dB), bytes);
hipMalloc(reinterpret_cast<void**>(&dC), bytes);

State / 状态变化

Why / 精确说法

Call this a device allocation or GPU-accessible memory. Do not promise that every allocation must physically reside in VRAM; placement policy and platform can differ.准确说法是设备 allocation 或 GPU 可访问内存。不要断言每次 allocation 都必然物理驻留在 VRAM；放置策略与平台可能不同。

Observe / 如何观察

rocprofv3 --hip-trace -- ./matmul2

The HIP API trace can show the allocation calls and their order.HIP API trace 可以显示 allocation 调用及其先后顺序。

Code / 你写的代码

hipMemcpy(dA, hA.data(), 16, hipMemcpyHostToDevice);
hipMemcpy(dB, hB.data(), 16, hipMemcpyHostToDevice);

// dC exists, but has no valid result yet.

State / 状态变化

Why / 纠正常见误解

The AQL dispatch packet will carry pointers, not the matrix payload. The 16-byte matrices live in device-accessible memory.AQL 派发包携带的是指针，不是矩阵本身。16 字节的矩阵数据位于设备可访问内存。

Observe / 如何观察

rocprofv3 --hip-trace -- ./matmul2

Expect two H2D API calls. Do not infer a particular copy engine without lower-level evidence.应看到两次 H2D API 调用。没有更底层证据时，不要断言具体使用了哪种 copy engine。

Code / 你写的代码

dim3 grid(1, 1);
dim3 block(2, 2);  // four work-items

hipLaunchKernelGGL(matmul2, grid, block, 0, 0,
                   dA, dB, dC);

State / 状态变化

Why / 纠正常见误解

This API call is not four CPU-to-GPU instruction streams. It specifies a kernel, grid, work-group and arguments; the runtime lowers that description next.这个 API 调用不是四条由 CPU 发往 GPU 的指令流。它描述 kernel、grid、工作组和参数；下一步由运行时把描述降到底层。

Observe / 如何观察

A launch can return before C is ready. Timing only the host call is not kernel timing.launch 返回时 C 可能尚未算好。只测主机 API 调用时间，不等于 kernel 执行时间。

Internal / 概念化内部动作

// conceptual — not application code
packet.kernel_object = matmul2_code;
packet.grid_size      = {2, 2, 1};
packet.workgroup_size = {2, 2, 1};
packet.kernarg        = {dA, dB, dC};
packet.completion     = signal;

State / 状态变化

Why / 精确说法

For this simple direct kernel, one dispatch is a useful mental model. Do not generalize that every high-level API always becomes exactly one packet; a tuned library call can enqueue several operations.对这个简单的直接 kernel，一次派发是合适的心智模型。但不要推广成“每个高层 API 永远恰好一个 packet”；调优库调用可能排入多个操作。

Observe / 如何观察

A profiler can relate a HIP API call to a dispatch. Ordinary application traces usually do not expose the raw packet bytes.profiler 可以把 HIP API 调用与 dispatch 对应起来；普通应用 trace 通常看不到原始 packet 字节。

Internal / 概念化内部动作

// conceptual — the runtime owns this
index = reserve_queue_index();
fill_packet_body(ring[index]);
publish_header_release(ring[index]);
mapped_doorbell = index;

State / 状态变化

Why / 关键边界

After queue setup, steady-state direct dispatch can avoid entering KFD for every launch. But KFD can still reappear for queue lifecycle, GPUVM faults or migration, preemption, reset and recovery.队列建好后，稳态直接派发通常无需每次进入 KFD。但队列生命周期、GPUVM 缺页或迁移、抢占、重置与恢复仍可能再次进入 KFD。

Observe / 如何观察

Mapped-MMIO doorbell writes are part of the mental model; a normal HIP trace may show the dispatch without exposing this individual store.映射 MMIO doorbell 写入属于底层心智模型；普通 HIP trace 可能显示 dispatch，却不一定暴露这次具体 store。

Device code / GPU 执行

__global__ void matmul2(const float* A,
                        const float* B, float* C) {
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  float sum = 0;
  for (int k = 0; k < 2; ++k)
    sum += A[row*2+k] * B[k*2+col];
  C[row*2+col] = sum;
}

State / 状态变化

Why / 性能边界

This is the only frame where the matrix arithmetic happens. A 2 × 2 launch uses four work-items and leaves most wave lanes idle; launch and copy overhead dominate, so it is a teaching case only.只有这一帧真正进行矩阵运算。2 × 2 只有四个工作项，多数 wave lane 都空闲，启动与拷贝开销会占主导，因此它只是教学案例。

Observe / 如何观察

Use a kernel trace for dispatch duration; use ROCm Compute Profiler on a realistic larger matrix for waves, occupancy and memory counters.用 kernel trace 看 dispatch 时长；要分析 wave、occupancy 与内存计数器，应换成真实的大矩阵再使用 ROCm Compute Profiler。

Code / 你写的代码

hipDeviceSynchronize();
hipMemcpy(hC.data(), dC, 16, hipMemcpyDeviceToHost);

assert(hC[0] == 19 && hC[1] == 22);
assert(hC[2] == 43 && hC[3] == 50);
hipFree(dA); hipFree(dB); hipFree(dC);

State / 状态变化

Why / 最后一处误解

A completion signal says the work is done; it does not automatically copy C to CPU memory. Synchronization makes correctness and timing boundaries explicit.完成信号只说明工作做完了，不会自动把 C 拷到 CPU 内存。同步让正确性边界和计时边界变得明确。

Observe / 最终证据

PASS · C = [[19, 22], [43, 50]]

The loop is closed: prediction, GPU result and CPU verification agree.闭环完成：预测、GPU 结果与 CPU 校验一致。