vyre 0.4.0

GPU compute intermediate representation with a standard operation library
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
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141

//! Catalog entry for `detect_lfi`.

use crate::ops::security_detection::detector_support::{attacks, DetectionError};

/// Embedded operation spec formerly stored in metadata/spec.toml.
pub const SPEC_TOML: &str = r#"schema_version = 1
id = "security_detection.detect_lfi"
archetype = "rule-bytes-to-bool"
display_name = "Detect LFI"
summary = "Returns true for traversal into local sensitive files."
category = "C"

[intrinsic]
wgsl = "security_detection_detect_lfi"

[signature]
inputs = ["Bytes"]
output = "Bool"

laws = []
equivalence_classes = ["etc_passwd", "windows_file", "plain_path", "t47_cap"]
workgroup_size = [64, 1, 1]
tags = ["security-detection", "lfi", "owasp", "t47"]
fixtures_dir = "fixtures/"
"#;

/// Embedded reference vectors formerly stored in fixtures/reference-vectors.toml.
pub const REFERENCE_VECTORS_TOML: &str = r#"[[case]]
name = "positive_etc_passwd"
input = "file=../../../../etc/passwd"
expected = true

[[case]]
name = "negative_relative_asset"
input = "file=../assets/logo.png"
expected = false
"#;

/// WGSL lowering source for this detector.
pub mod lowering {
    /// Return the detector-specific WGSL source.
    #[must_use]
    pub const fn source() -> &'static str {
        r#"struct Params {
    input_len: u32,
    _pad0: u32,
    _pad1: u32,
    _pad2: u32,
}

@group(0) @binding(0) var<storage, read> input: array<u32>;
@group(0) @binding(1) var<storage, read_write> output: array<u32>;
@group(0) @binding(2) var<uniform> params: Params;

fn lower_ascii(byte: u32) -> u32 {
    if (byte >= 65u && byte <= 90u) {
        return byte + 32u;
    }
    return byte;
}

fn byte_at(index: u32) -> u32 {
    return lower_ascii(input[index]);
}

fn matches_at(index: u32, n0: u32, n1: u32, n2: u32, n3: u32, n4: u32, n5: u32, n6: u32, n7: u32, n8: u32, n9: u32, n10: u32, len: u32) -> bool {
    if (index + len > params.input_len) {
        return false;
    }
    let needle = array<u32, 11>(n0, n1, n2, n3, n4, n5, n6, n7, n8, n9, n10);
    for (var offset = 0u; offset < len; offset = offset + 1u) {
        if (byte_at(index + offset) != needle[offset]) {
            return false;
        }
    }
    return true;
}

fn contains_pattern(lane: u32, n0: u32, n1: u32, n2: u32, n3: u32, n4: u32, n5: u32, n6: u32, n7: u32, n8: u32, n9: u32, n10: u32, len: u32) -> bool {
    for (var index = lane; index < params.input_len; index = index + 64u) {
        if (matches_at(index, n0, n1, n2, n3, n4, n5, n6, n7, n8, n9, n10, len)) {
            return true;
        }
    }
    return false;
}

var<workgroup> lane_flags: array<u32, 64>;

@compute @workgroup_size(64)
fn security_detection_detect_lfi(@builtin(local_invocation_id) lid: vec3<u32>) {
    let lane = lid.x;
    var flags = 0u;
    if (contains_pattern(lane, 46u, 46u, 47u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 3u) ||
        contains_pattern(lane, 46u, 46u, 92u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 3u)) {
        flags = flags | 1u;
    }
    if (contains_pattern(lane, 47u, 101u, 116u, 99u, 47u, 112u, 97u, 115u, 115u, 119u, 100u, 11u) ||
        contains_pattern(lane, 98u, 111u, 111u, 116u, 46u, 105u, 110u, 105u, 0u, 0u, 0u, 8u) ||
        contains_pattern(lane, 119u, 105u, 110u, 46u, 105u, 110u, 105u, 0u, 0u, 0u, 0u, 7u)) {
        flags = flags | 2u;
    }
    lane_flags[lane] = flags;
    workgroupBarrier();
    if (lane == 0u) {
        var combined = 0u;
        for (var i = 0u; i < 64u; i = i + 1u) {
            combined = combined | lane_flags[i];
        }
        output[0] = select(0u, 1u, (combined & 3u) == 3u);
    }
}
"#
    }
}

/// Return true for traversal into local sensitive files.
///
/// # Errors
///
/// Returns `Fix: ...` when input exceeds 64 MiB.
pub fn detect_lfi(input: &[u8]) -> Result<bool, DetectionError> {
    attacks::detect_lfi(input)
}

/// Compatibility surface for the previous generated implementation module.
pub mod implementation {
    pub use super::detect_lfi;
    /// Compatibility module for callers that used the generated kernel path.
    pub mod kernel {
        pub use super::super::detect_lfi;
    }

    /// Compatibility module for callers that used the generated lowering path.
    pub mod lowering {
        /// Compatibility module for callers that used `implementation::lowering::wgsl`.
        pub mod wgsl {
            pub use super::super::super::lowering::source;
        }
    }
}