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
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
//! Address Space Validator - validates correct address space usage
use crate::bugs::Severity;
use crate::parser::types::{Modifier, Opcode};
use crate::parser::{Instruction, Operand, PtxModule, SourceLocation, Statement};
use std::collections::HashSet;
/// Bug: Generic addressing of shared memory
#[derive(Debug, Clone)]
pub struct GenericSharedBug {
/// Source location
pub location: SourceLocation,
/// Instruction that triggered the bug
pub instruction: Instruction,
/// Severity
pub severity: Severity,
/// Fix suggestion
pub fix: String,
}
/// Address Space Validator
pub struct AddressSpaceValidator {
/// Registers holding cvta.shared results (generic shared addresses)
shared_base_regs: HashSet<String>,
}
impl AddressSpaceValidator {
/// Create a new address space validator
pub fn new() -> Self {
Self {
shared_base_regs: HashSet::new(),
}
}
/// Detect generic addressing of shared memory (F021)
///
/// WRONG: cvta.shared.u64 %rd, smem; ld.u32 [%rd]
/// RIGHT: ld.shared.u32 [smem_offset]
pub fn detect_generic_shared_access(&mut self, module: &PtxModule) -> Vec<GenericSharedBug> {
let mut bugs = Vec::new();
for kernel in &module.kernels {
self.shared_base_regs.clear();
for stmt in &kernel.body {
if let Statement::Instruction(instr) = stmt {
// Track cvta.shared destinations
if instr.opcode == Opcode::Cvta && self.has_shared_modifier(instr) {
if let Some(Operand::Register(dest)) = instr.operands.first() {
self.shared_base_regs.insert(dest.clone());
}
}
// Detect generic ld/st using tracked registers
if (instr.opcode == Opcode::Ld || instr.opcode == Opcode::St)
&& !self.has_space_modifier(instr)
{
// Check if address operand uses a generic shared register
let addr_operand = if instr.opcode == Opcode::Ld {
instr.operands.get(1)
} else {
instr.operands.first()
};
if let Some(operand) = addr_operand {
if self.uses_generic_shared_reg(operand) {
bugs.push(GenericSharedBug {
location: instr.location.clone(),
instruction: instr.clone(),
severity: Severity::Critical,
fix: "Use ld.shared with 32-bit offset instead".into(),
});
}
}
}
}
}
}
bugs
}
/// Detect shared memory using 64-bit addresses (F022)
pub fn detect_shared_mem_u64(&self, module: &PtxModule) -> Vec<GenericSharedBug> {
let mut bugs = Vec::new();
for kernel in &module.kernels {
for stmt in &kernel.body {
if let Statement::Instruction(instr) = stmt {
// Check for ld.shared.u64 or st.shared.u64 with 64-bit address
if self.has_shared_modifier(instr) && self.has_u64_modifier(instr) {
bugs.push(GenericSharedBug {
location: instr.location.clone(),
instruction: instr.clone(),
severity: Severity::High,
fix: "Use 32-bit offset for shared memory addressing".into(),
});
}
}
}
}
bugs
}
/// Detect cvta.shared inside loops (F083)
pub fn detect_loop_cvta_shared(&self, module: &PtxModule) -> Vec<GenericSharedBug> {
let mut bugs = Vec::new();
for kernel in &module.kernels {
let mut in_loop = false;
let mut loop_start_labels = HashSet::new();
for stmt in &kernel.body {
match stmt {
Statement::Label(label) => {
// Simple heuristic: label containing "loop" starts a loop
if label.to_lowercase().contains("loop") {
in_loop = true;
loop_start_labels.insert(label.clone());
}
}
Statement::Instruction(instr) => {
// Check for backward branch (loop end)
if instr.opcode == Opcode::Bra {
for operand in &instr.operands {
if let Operand::Label(target) = operand {
if loop_start_labels.contains(target) {
in_loop = false;
}
}
}
}
// Detect cvta.shared inside loop
if in_loop
&& instr.opcode == Opcode::Cvta
&& self.has_shared_modifier(instr)
{
bugs.push(GenericSharedBug {
location: instr.location.clone(),
instruction: instr.clone(),
severity: Severity::High,
fix: "Move cvta.shared outside loop to reduce register pressure"
.into(),
});
}
}
_ => {}
}
}
}
bugs
}
fn has_shared_modifier(&self, instr: &Instruction) -> bool {
instr
.modifiers
.iter()
.any(|m| matches!(m, Modifier::Shared))
}
fn has_space_modifier(&self, instr: &Instruction) -> bool {
instr
.modifiers
.iter()
.any(|m| m.as_address_space().is_some())
}
fn has_u64_modifier(&self, instr: &Instruction) -> bool {
instr
.modifiers
.iter()
.any(|m| matches!(m, Modifier::U64 | Modifier::B64))
}
fn uses_generic_shared_reg(&self, operand: &Operand) -> bool {
match operand {
Operand::Register(reg) => self.shared_base_regs.contains(reg),
Operand::Memory(addr) => {
// Check if the memory address contains a generic shared register
self.shared_base_regs.iter().any(|reg| addr.contains(reg))
}
_ => false,
}
}
}
impl Default for AddressSpaceValidator {
fn default() -> Self {
Self::new()
}
}
#[cfg(test)]
mod tests {
use super::*;
use crate::parser::Parser;
// F021: No cvta.shared followed by generic ld/st
#[test]
fn f021_no_generic_shared_access() {
let ptx = r#"
.version 8.0
.target sm_70
.address_size 64
.entry test()
{
.reg .u32 %r<10>;
ld.shared.u32 %r0, [%r1];
ret;
}
"#;
let mut parser = Parser::new(ptx).expect("parser creation should succeed");
let module = parser.parse().expect("parsing should succeed");
let mut validator = AddressSpaceValidator::new();
let bugs = validator.detect_generic_shared_access(&module);
assert!(
bugs.is_empty(),
"F021: Should have no generic shared access bugs"
);
}
// F023: Direct .shared addressing preferred
#[test]
fn f023_direct_shared_addressing() {
let ptx = r#"
.version 8.0
.target sm_70
.address_size 64
.entry test()
{
.reg .u32 %r<10>;
ld.shared.u32 %r0, [%r1];
st.shared.u32 [%r2], %r0;
ret;
}
"#;
let mut parser = Parser::new(ptx).expect("parser creation should succeed");
let module = parser.parse().expect("parsing should succeed");
let mut validator = AddressSpaceValidator::new();
let bugs = validator.detect_generic_shared_access(&module);
assert!(
bugs.is_empty(),
"F023: Direct shared addressing should not trigger bugs"
);
}
}