jayinlab

AMD PM4 개요 — Type-3 패킷 구조와 Dispatch 패밀리 opcode

“OpenCL enqueue → Vulkan dispatch → 그다음은?”
그다음이 PM4다. 드라이버가 하드웨어에 전달하는 패킷 기반 command stream.
세부 비트필드 암기보다 명령 스트림 mental model을 먼저 만든다.

PM4를 한 줄로

PM4는 GPU에 전달되는 패킷 기반 command stream 포맷이다.

CPU/드라이버 → PM4 패킷 작성 → Ring Buffer에 기록
                                    ↓
                           GPU CP(Command Processor)가 패킷 해석 → 실행

앱이 직접 PM4를 쓰지 않는다. Vulkan API 호출이 드라이버에 의해 최종적으로 PM4 패킷 시퀀스로 변환된다.

Type-3 패킷 구조

PM4에는 여러 타입이 있지만, Type-3이 실질적인 동작 명령을 담는다.

[ Header: IT_TYPE=3 | Opcode | Count ]
[ Payload Word 0                     ]
[ Payload Word 1                     ]
[ ...                                ]

Dispatch 패밀리 주요 opcode

OpenCL → Vulkan → PM4 연결 사다리

1. clEnqueueNDRangeKernel (OpenCL)
        ↓
2. vkCmdDispatch (Vulkan)
        ↓
3. 드라이버가 command stream 준비
        ↓
4. IT_SET_SH_REG (레지스터 설정 패킷들)
   IT_DISPATCH_DIRECT (dispatch 트리거)
        ↓
5. GPU CP가 패킷 해석 → Compute Engine에 작업 전달
        ↓
6. Wavefront 스케줄링 → Shader 실행

핵심 mental model 3줄

1. PM4 = 패킷 스트림 — 개별 명령이 아니라 연속된 패킷의 흐름
2. Type-3 = 실제 동작 명령 핵심 클래스 — 상태 설정 + dispatch 트리거
3. Dispatch는 상태 설정 시퀀스 뒤에 트리거됨 — dispatch 패킷 하나만 보면 안 된다

자주 하는 오해

오해: “dispatch 패킷만 보면 된다”
→ dispatch 앞에 shader 레지스터 설정, 리소스 바인딩 설정 등 많은 준비 패킷이 붙는다.
“dispatch 비용"은 앞단 상태 전환 비용과 함께 봐야 한다.

오해: “PM4는 GPU마다 같다”
→ AMD 내에서도 세대별 차이가 있다. 지금은 공통 mental model을 만들고, 이후 세대별 디테일로 들어간다.

이해 확인 질문

Q1. PM4는 API 계층에서 어디쯤에 위치하나?

Q2. 왜 Type-3를 먼저 잡아야 하나?

Q3. dispatch 패킷만 보면 충분하지 않은 이유는?

Q4. PM4 학습에서 먼저 외울 것은 비트필드인가, 시퀀스 구조인가?

Q5. PM4 mental model 3줄을 적어봐.

관련 글

• PM4 제출 흐름 — vkQueueSubmit → PM4 → ring → GPU 7단계 animation
• PM4 Indirect Buffer — IT_INDIRECT_BUFFER 패킷과 IB 계층 구조
• clFinish 내부 — IT_EVENT_WRITE와 fence 연결

관련 용어

[[pm4-packet]], [[ring-buffer]], [[wavefront]], [[command-buffer]]