vyre-conform 0.1.0

Conformance suite for vyre backends — proves byte-identical output to CPU reference
Documentation 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40

use vyre::ir::{AtomicOp, BufferDecl, DataType, Expr, Node, Program};
use vyre_conform::{reference::interp, spec::value::Value};

/// ADVERSARIAL: atomic_load hardcodes stride 4 instead of element stride — vector: an atomic
/// on a U64 buffer at index 1 writes to byte offset 4 (inside the first element) instead of
/// offset 8, corrupting adjacent data.
#[test]
fn atomic_on_u64_buffer_uses_element_stride() {
    let program = Program {
        entry_op_id: None,
        buffers: vec![BufferDecl::read_write("u64_buf", 0, DataType::U64)],
        workgroup_size: [1, 1, 1],
        entry: vec![Node::store(
            "u64_buf",
            Expr::u32(0),
            Expr::Atomic {
                op: AtomicOp::Add,
                buffer: "u64_buf".to_string(),
                index: Box::new(Expr::u32(1)),
                expected: None,
                value: Box::new(Expr::u32(1)),
            },
        )],
    };

    // U64 buffer with two elements: [0x0807060504030201, 0x0000000000000000]
    let bytes = vec![
        0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
        0x00,
    ];
    let outputs = interp::run(&program, &[Value::Bytes(bytes)])
        .expect("atomic add must use correct element stride");

    // Expected: first element unchanged, second element += 1.
    let expected = vec![
        0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
        0x00,
    ];
    assert_eq!(outputs, vec![Value::Bytes(expected)]);
}