softgpu_core/
aql.rs

1//! AQL packet decode/validate for SoftGPU Phase 4 (no kernel execution).
2//!
3//! Parses HSA kernel-dispatch (and minimal barrier) packets from owned bytes.
4//! Layout provenance: pinned ROCR `hsa.h` (`hsa_kernel_dispatch_packet_t`
5//! under `HSA_LARGE_MODEL`, 64 bytes).
6
7use crate::queue::{AQL_PACKET_BYTES, PACKET_TYPE_INVALID, PACKET_TYPE_KERNEL_DISPATCH};
8
9/// `HSA_PACKET_TYPE_BARRIER_AND`.
10pub const PACKET_TYPE_BARRIER_AND: u16 = 3;
11/// `HSA_PACKET_TYPE_AGENT_DISPATCH`.
12pub const PACKET_TYPE_AGENT_DISPATCH: u16 = 4;
13/// `HSA_PACKET_TYPE_BARRIER_OR`.
14pub const PACKET_TYPE_BARRIER_OR: u16 = 5;
15
16/// SoftGPU experimental no-execution completion contract label.
17pub const DIAGNOSTIC_COMPLETE_NO_EXECUTION: &str = "diagnostic_complete_no_execution";
18pub const DIAGNOSTIC_REJECTED: &str = "diagnostic_rejected";
19
20#[derive(Debug, Clone, Copy, PartialEq, Eq)]
21pub enum PacketType {
22    Invalid,
23    KernelDispatch,
24    BarrierAnd,
25    AgentDispatch,
26    BarrierOr,
27    VendorOrUnknown(u16),
28}
29
30impl PacketType {
31    pub fn from_header_type(ty: u16) -> Self {
32        match ty & 0xff {
33            1 => Self::Invalid,
34            2 => Self::KernelDispatch,
35            3 => Self::BarrierAnd,
36            4 => Self::AgentDispatch,
37            5 => Self::BarrierOr,
38            other => Self::VendorOrUnknown(other),
39        }
40    }
41
42    pub fn as_u16(self) -> u16 {
43        match self {
44            Self::Invalid => PACKET_TYPE_INVALID,
45            Self::KernelDispatch => PACKET_TYPE_KERNEL_DISPATCH,
46            Self::BarrierAnd => PACKET_TYPE_BARRIER_AND,
47            Self::AgentDispatch => PACKET_TYPE_AGENT_DISPATCH,
48            Self::BarrierOr => PACKET_TYPE_BARRIER_OR,
49            Self::VendorOrUnknown(v) => v,
50        }
51    }
52}
53
54/// How SoftGPU classifies a kernarg pointer without claiming contents.
55#[derive(Debug, Clone, PartialEq, Eq)]
56pub enum KernargClass {
57    Null,
58    /// SoftGPU-tracked allocation (alloc_id when known).
59    SoftGpu {
60        alloc_id: Option<u64>,
61        addr: u64,
62    },
63    /// Non-null pointer SoftGPU does not own — opaque for replay.
64    ForeignOpaque {
65        addr: u64,
66    },
67}
68
69/// Normalized dispatch descriptor (vendor-neutral SoftGPU view).
70#[derive(Debug, Clone, PartialEq, Eq)]
71pub struct DispatchDescriptor {
72    pub packet_index: u64,
73    pub packet_type: PacketType,
74    pub header: u16,
75    pub setup: u16,
76    pub dimensions: u16,
77    pub workgroup_size: [u16; 3],
78    pub grid_size: [u32; 3],
79    pub private_segment_size: u32,
80    pub group_segment_size: u32,
81    pub kernel_object: u64,
82    pub kernarg: KernargClass,
83    pub completion_signal: u64,
84    /// Owned AQL bytes for replay without living host pointers.
85    pub captured_bytes: [u8; AQL_PACKET_BYTES],
86}
87
88#[derive(Debug, Clone, PartialEq, Eq)]
89pub enum AqlParseError {
90    WrongLength { got: usize },
91    StillInvalid,
92    UnsupportedType { packet_type: u16 },
93    ZeroWorkgroup { dim: usize },
94    GridSmallerThanWorkgroup { dim: usize },
95    BadDimensions { dimensions: u16 },
96    DimConstraint { detail: String },
97}
98
99impl std::fmt::Display for AqlParseError {
100    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
101        write!(f, "{self:?}")
102    }
103}
104
105fn read_u16(bytes: &[u8], off: usize) -> u16 {
106    u16::from_le_bytes([bytes[off], bytes[off + 1]])
107}
108
109fn read_u32(bytes: &[u8], off: usize) -> u32 {
110    u32::from_le_bytes([bytes[off], bytes[off + 1], bytes[off + 2], bytes[off + 3]])
111}
112
113fn read_u64(bytes: &[u8], off: usize) -> u64 {
114    u64::from_le_bytes([
115        bytes[off],
116        bytes[off + 1],
117        bytes[off + 2],
118        bytes[off + 3],
119        bytes[off + 4],
120        bytes[off + 5],
121        bytes[off + 6],
122        bytes[off + 7],
123    ])
124}
125
126/// Parse a Phase 4-supported packet (kernel-dispatch, or minimal barrier).
127pub fn parse_supported_packet(
128    bytes: &[u8],
129    packet_index: u64,
130    classify_kernarg: impl FnOnce(u64) -> KernargClass,
131) -> Result<DispatchDescriptor, AqlParseError> {
132    if bytes.len() != AQL_PACKET_BYTES {
133        return Err(AqlParseError::WrongLength { got: bytes.len() });
134    }
135    let header = read_u16(bytes, 0);
136    let ty = header & 0xff;
137    match PacketType::from_header_type(ty) {
138        PacketType::Invalid => Err(AqlParseError::StillInvalid),
139        PacketType::KernelDispatch => parse_kernel_dispatch(bytes, packet_index, classify_kernarg),
140        PacketType::BarrierAnd | PacketType::BarrierOr => {
141            parse_barrier_minimal(bytes, packet_index)
142        }
143        other => Err(AqlParseError::UnsupportedType {
144            packet_type: other.as_u16(),
145        }),
146    }
147}
148
149/// Minimal barrier accept: header type + completion signal only (deps unchecked).
150pub fn parse_barrier_minimal(
151    bytes: &[u8],
152    packet_index: u64,
153) -> Result<DispatchDescriptor, AqlParseError> {
154    if bytes.len() != AQL_PACKET_BYTES {
155        return Err(AqlParseError::WrongLength { got: bytes.len() });
156    }
157    let header = read_u16(bytes, 0);
158    let ty = header & 0xff;
159    let packet_type = PacketType::from_header_type(ty);
160    match packet_type {
161        PacketType::BarrierAnd | PacketType::BarrierOr => {}
162        PacketType::Invalid => return Err(AqlParseError::StillInvalid),
163        other => {
164            return Err(AqlParseError::UnsupportedType {
165                packet_type: other.as_u16(),
166            });
167        }
168    }
169    let mut captured = [0u8; AQL_PACKET_BYTES];
170    captured.copy_from_slice(bytes);
171    Ok(DispatchDescriptor {
172        packet_index,
173        packet_type,
174        header,
175        setup: 0,
176        dimensions: 0,
177        workgroup_size: [0, 0, 0],
178        grid_size: [0, 0, 0],
179        private_segment_size: 0,
180        group_segment_size: 0,
181        kernel_object: 0,
182        kernarg: KernargClass::Null,
183        completion_signal: read_u64(bytes, 56),
184        captured_bytes: captured,
185    })
186}
187
188/// Parse and validate a kernel-dispatch packet from exactly 64 bytes.
189pub fn parse_kernel_dispatch(
190    bytes: &[u8],
191    packet_index: u64,
192    classify_kernarg: impl FnOnce(u64) -> KernargClass,
193) -> Result<DispatchDescriptor, AqlParseError> {
194    if bytes.len() != AQL_PACKET_BYTES {
195        return Err(AqlParseError::WrongLength { got: bytes.len() });
196    }
197    let header = read_u16(bytes, 0);
198    let ty = header & 0xff;
199    let packet_type = PacketType::from_header_type(ty);
200    match packet_type {
201        PacketType::Invalid => return Err(AqlParseError::StillInvalid),
202        PacketType::KernelDispatch => {}
203        other => {
204            return Err(AqlParseError::UnsupportedType {
205                packet_type: other.as_u16(),
206            });
207        }
208    }
209
210    let setup = read_u16(bytes, 2);
211    // HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS width 2 at bit 0.
212    let dimensions = setup & 0b11;
213    if !(1..=3).contains(&dimensions) {
214        return Err(AqlParseError::BadDimensions { dimensions });
215    }
216
217    let workgroup_size = [read_u16(bytes, 4), read_u16(bytes, 6), read_u16(bytes, 8)];
218    let grid_size = [
219        read_u32(bytes, 12),
220        read_u32(bytes, 16),
221        read_u32(bytes, 20),
222    ];
223    let private_segment_size = read_u32(bytes, 24);
224    let group_segment_size = read_u32(bytes, 28);
225    let kernel_object = read_u64(bytes, 32);
226    let kernarg_addr = read_u64(bytes, 40);
227    let completion_signal = read_u64(bytes, 56);
228
229    for d in 0..3 {
230        if workgroup_size[d] == 0 {
231            return Err(AqlParseError::ZeroWorkgroup { dim: d });
232        }
233        if grid_size[d] < u32::from(workgroup_size[d]) {
234            return Err(AqlParseError::GridSmallerThanWorkgroup { dim: d });
235        }
236    }
237    if dimensions == 1 && (workgroup_size[1] != 1 || workgroup_size[2] != 1) {
238        return Err(AqlParseError::DimConstraint {
239            detail: "dims=1 requires workgroup y=z=1".into(),
240        });
241    }
242    if dimensions == 1 && (grid_size[1] != 1 || grid_size[2] != 1) {
243        return Err(AqlParseError::DimConstraint {
244            detail: "dims=1 requires grid y=z=1".into(),
245        });
246    }
247    if dimensions == 2 && (workgroup_size[2] != 1 || grid_size[2] != 1) {
248        return Err(AqlParseError::DimConstraint {
249            detail: "dims=2 requires workgroup/grid z=1".into(),
250        });
251    }
252
253    let mut captured = [0u8; AQL_PACKET_BYTES];
254    captured.copy_from_slice(bytes);
255
256    Ok(DispatchDescriptor {
257        packet_index,
258        packet_type,
259        header,
260        setup,
261        dimensions,
262        workgroup_size,
263        grid_size,
264        private_segment_size,
265        group_segment_size,
266        kernel_object,
267        kernarg: classify_kernarg(kernarg_addr),
268        completion_signal,
269        captured_bytes: captured,
270    })
271}
272
273/// Replay parser on captured bytes (no live host pointers required).
274pub fn replay_dispatch(
275    bytes: &[u8],
276    packet_index: u64,
277) -> Result<DispatchDescriptor, AqlParseError> {
278    parse_supported_packet(bytes, packet_index, |addr| {
279        if addr == 0 {
280            KernargClass::Null
281        } else {
282            // Replay cannot re-resolve SoftGPU ownership; keep opaque.
283            KernargClass::ForeignOpaque { addr }
284        }
285    })
286}
287
288/// Build a lawful minimal 1D kernel-dispatch packet (SoftGPU test fixture).
289pub fn golden_kernel_dispatch_1d(
290    workgroup_x: u16,
291    grid_x: u32,
292    kernel_object: u64,
293    kernarg: u64,
294    completion_signal: u64,
295) -> [u8; AQL_PACKET_BYTES] {
296    let mut b = [0u8; AQL_PACKET_BYTES];
297    // header: type = KERNEL_DISPATCH
298    b[0] = PACKET_TYPE_KERNEL_DISPATCH as u8;
299    // setup: dimensions = 1
300    b[2] = 1;
301    b[4..6].copy_from_slice(&workgroup_x.to_le_bytes());
302    b[6..8].copy_from_slice(&1u16.to_le_bytes());
303    b[8..10].copy_from_slice(&1u16.to_le_bytes());
304    b[12..16].copy_from_slice(&grid_x.to_le_bytes());
305    b[16..20].copy_from_slice(&1u32.to_le_bytes());
306    b[20..24].copy_from_slice(&1u32.to_le_bytes());
307    b[32..40].copy_from_slice(&kernel_object.to_le_bytes());
308    b[40..48].copy_from_slice(&kernarg.to_le_bytes());
309    b[56..64].copy_from_slice(&completion_signal.to_le_bytes());
310    b
311}
312
313#[cfg(test)]
314mod tests {
315    use super::*;
316
317    #[test]
318    fn golden_parses_and_replays() {
319        let bytes = golden_kernel_dispatch_1d(64, 256, 0xABCDu64, 0, 0x1111);
320        let d = parse_kernel_dispatch(&bytes, 0, |_| KernargClass::Null).unwrap();
321        assert_eq!(d.dimensions, 1);
322        assert_eq!(d.workgroup_size, [64, 1, 1]);
323        assert_eq!(d.grid_size, [256, 1, 1]);
324        assert_eq!(d.kernel_object, 0xABCD);
325        assert_eq!(d.completion_signal, 0x1111);
326        let replayed = replay_dispatch(&d.captured_bytes, 0).unwrap();
327        assert_eq!(replayed.grid_size, d.grid_size);
328        assert_eq!(replayed.kernel_object, d.kernel_object);
329    }
330
331    #[test]
332    fn rejects_zero_workgroup() {
333        let mut bytes = golden_kernel_dispatch_1d(0, 256, 1, 0, 0);
334        // force wg_x = 0 already
335        let err = parse_kernel_dispatch(&bytes, 0, |_| KernargClass::Null).unwrap_err();
336        assert!(matches!(err, AqlParseError::ZeroWorkgroup { dim: 0 }));
337        bytes = golden_kernel_dispatch_1d(64, 32, 1, 0, 0); // grid < wg
338        let err = parse_kernel_dispatch(&bytes, 0, |_| KernargClass::Null).unwrap_err();
339        assert!(matches!(
340            err,
341            AqlParseError::GridSmallerThanWorkgroup { dim: 0 }
342        ));
343    }
344
345    #[test]
346    fn rejects_invalid_and_unsupported() {
347        let mut bytes = [0u8; AQL_PACKET_BYTES];
348        bytes[0] = PACKET_TYPE_INVALID as u8;
349        assert!(matches!(
350            parse_kernel_dispatch(&bytes, 0, |_| KernargClass::Null),
351            Err(AqlParseError::StillInvalid)
352        ));
353        bytes[0] = PACKET_TYPE_AGENT_DISPATCH as u8;
354        assert!(matches!(
355            parse_kernel_dispatch(&bytes, 0, |_| KernargClass::Null),
356            Err(AqlParseError::UnsupportedType { .. })
357        ));
358    }
359}
softgpu_core/aql.rs

softgpu_core/
aql.rs
