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
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161

//! Catalog entry for `detect_command_injection`.

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_command_injection"
archetype = "rule-bytes-to-bool"
display_name = "Detect Command Injection"
summary = "Returns true for shell metacharacters paired with command names."
category = "C"

[intrinsic]
wgsl = "security_detection_detect_command_injection"

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

laws = []
equivalence_classes = ["semicolon_command", "subshell_command", "plain_argument", "t47_cap"]
workgroup_size = [64, 1, 1]
tags = ["security-detection", "command-injection", "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_semicolon_id"
input = "name=alice; id"
expected = true

[[case]]
name = "negative_ampersand_text"
input = "company=Research & Development"
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 hex_value(byte: u32) -> u32 {
    if (byte >= 48u && byte <= 57u) {
        return byte - 48u;
    }
    let lowered = lower_ascii(byte);
    if (lowered >= 97u && lowered <= 102u) {
        return lowered - 87u;
    }
    return 255u;
}

fn normalized_at(index: u32) -> u32 {
    if (index + 2u < params.input_len && input[index] == 37u) {
        let hi = hex_value(input[index + 1u]);
        let lo = hex_value(input[index + 2u]);
        if (hi < 16u && lo < 16u) {
            return lower_ascii(hi * 16u + lo);
        }
    }
    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, len: u32) -> bool {
    if (index + len > params.input_len) {
        return false;
    }
    let needle = array<u32, 8>(n0, n1, n2, n3, n4, n5, n6, n7);
    for (var offset = 0u; offset < len; offset = offset + 1u) {
        if (normalized_at(index + offset) != needle[offset]) {
            return false;
        }
    }
    return true;
}

fn contains(lane: u32, n0: u32, n1: u32, n2: u32, n3: u32, n4: u32, n5: u32, n6: u32, n7: 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, len)) {
            return true;
        }
    }
    return false;
}

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

@compute @workgroup_size(64)
fn security_detection_detect_command_injection(@builtin(local_invocation_id) lid: vec3<u32>) {
    let lane = lid.x;
    var flags = 0u;
    if (contains(lane, 59u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 1u) || contains(lane, 38u, 38u, 0u, 0u, 0u, 0u, 0u, 0u, 2u) ||
        contains(lane, 124u, 124u, 0u, 0u, 0u, 0u, 0u, 0u, 2u) || contains(lane, 96u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 1u) ||
        contains(lane, 36u, 40u, 0u, 0u, 0u, 0u, 0u, 0u, 2u) || contains(lane, 124u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 1u)) {
        flags = flags | 1u;
    }
    if (contains(lane, 32u, 99u, 97u, 116u, 32u, 0u, 0u, 0u, 5u) || contains(lane, 32u, 105u, 100u, 0u, 0u, 0u, 0u, 0u, 3u) ||
        contains(lane, 32u, 119u, 104u, 111u, 97u, 109u, 105u, 0u, 7u) || contains(lane, 32u, 99u, 117u, 114u, 108u, 32u, 0u, 0u, 6u) ||
        contains(lane, 32u, 119u, 103u, 101u, 116u, 32u, 0u, 0u, 6u) || contains(lane, 32u, 110u, 99u, 32u, 0u, 0u, 0u, 0u, 4u) ||
        contains(lane, 32u, 47u, 98u, 105u, 110u, 47u, 115u, 104u, 8u) || contains(lane, 112u, 111u, 119u, 101u, 114u, 115u, 104u, 101u, 8u)) {
        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 shell metacharacters paired with command names.
///
/// # Errors
///
/// Returns `Fix: ...` when input exceeds 64 MiB.
pub fn detect_command_injection(input: &[u8]) -> Result<bool, DetectionError> {
    attacks::detect_command_injection(input)
}

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

    /// 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;
        }
    }
}