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

softgpu_core/
aql.rs
