aprender-explain 0.70.1

PTX/SIMD/wgpu visualization and tracing CLI for Trueno
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
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71

use super::super::analyzer::*;
use super::super::types::*;

#[test]
fn test_shared_mem_u64_detection() {
    let ptx = r"
.visible .entry test() {
    .reg .u64 %rd<5>;
    .reg .f32 %f<2>;
    .shared .b8 smem[4096];
    st.shared.f32 [%rd0], %f0;
    ret;
}
";
    let result = PtxBugAnalyzer::new().analyze(ptx);
    assert!(result.has_bug(&PtxBugClass::SharedMemU64Addressing));
}

#[test]
fn test_shared_mem_u32_valid() {
    let ptx = r"
.visible .entry test() {
    .reg .u32 %r<5>;
    .reg .f32 %f<2>;
    .shared .b8 smem[4096];
    st.shared.f32 [%r0], %f0;
    ret;
}
";
    let result = PtxBugAnalyzer::new().analyze(ptx);
    assert!(!result.has_bug(&PtxBugClass::SharedMemU64Addressing));
}

#[test]
fn test_missing_barrier_sync_strict() {
    let ptx = r"
.visible .entry test() {
    .shared .b8 smem[1024];
    st.shared.f32 [%r0], %f0;
    ld.shared.f32 %f1, [%r1];
    ret;
}
";
    // Non-strict mode: no warning
    let normal_result = PtxBugAnalyzer::new().analyze(ptx);
    assert!(!normal_result.has_bug(&PtxBugClass::MissingBarrierSync));

    // Strict mode: warning
    let strict_result = PtxBugAnalyzer::strict().analyze(ptx);
    assert!(strict_result.has_bug(&PtxBugClass::MissingBarrierSync));
}

#[test]
fn test_barrier_present_valid() {
    let ptx = r"
.visible .entry test() {
    .shared .b8 smem[1024];
    st.shared.f32 [%r0], %f0;
    bar.sync 0;
    ld.shared.f32 %f1, [%r1];
    ret;
}
";
    let result = PtxBugAnalyzer::strict().analyze(ptx);
    // Should not have the broad "no bar.sync" warning
    let missing_barrier_bugs: Vec<_> = result.bugs_of_class(&PtxBugClass::MissingBarrierSync);
    // The specific st/ld pattern should not trigger since bar.sync is present
    assert!(missing_barrier_bugs
        .iter()
        .all(|b| !b.message.contains("ld.shared follows st.shared")));
}