Repository navigation
Expand file tree
/
Copy pathZLUDA_20260914.html
More file actions
898 lines (835 loc) · 114 KB
/
Copy pathZLUDA_20260914.html
File metadata and controls
898 lines (835 loc) · 114 KB
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
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
667
668
669
670
671
672
673
674
675
676
677
678
679
680
681
682
683
684
685
686
687
688
689
690
691
692
693
694
695
696
697
698
699
700
701
702
703
704
705
706
707
708
709
710
711
712
713
714
715
716
717
718
719
720
721
722
723
724
725
726
727
728
729
730
731
732
733
734
735
736
737
738
739
740
741
742
743
744
745
746
747
748
749
750
751
752
753
754
755
756
757
758
759
760
761
762
763
764
765
766
767
768
769
770
771
772
773
774
775
776
777
778
779
780
781
782
783
784
785
786
787
788
789
790
791
792
793
794
795
796
797
798
799
800
801
802
803
804
805
806
807
808
809
810
811
812
813
814
815
816
817
818
819
820
821
822
823
824
825
826
827
828
829
830
831
832
833
834
835
836
837
838
839
840
841
842
843
844
845
846
847
848
849
850
851
852
853
854
855
856
857
858
859
860
861
862
863
864
865
866
867
868
869
870
871
872
873
874
875
876
877
878
879
880
881
882
883
884
885
886
887
888
889
890
891
892
893
894
895
896
897
898
<!DOCTYPE html>
<html lang="en">
<head>
<meta charset="utf-8">
<meta name="viewport" content="width=device-width,initial-scale=1">
<meta name="color-scheme" content="dark">
<title>Code Audit — Interactive</title>
<style>
*,*::before,*::after{box-sizing:border-box;margin:0;padding:0}
:root{color-scheme:dark}
body{font-family:-apple-system,BlinkMacSystemFont,"Segoe UI",Roboto,sans-serif;background:#0f172a;color:#e5e7eb;font-size:14px;line-height:1.5}
.app{max-width:1100px;margin:0 auto;padding:1.5rem 1rem}
.top{display:flex;justify-content:space-between;align-items:flex-start;gap:1rem;margin-bottom:1rem;flex-wrap:wrap}
h1{font-size:1.5rem;font-weight:700;color:#f1f5f9}
.meta{font-size:.8rem;color:#6b7280;margin-top:.2rem}
.top-right{display:flex;align-items:center;gap:.75rem;flex-wrap:wrap}
.project-info{margin-bottom:1rem;padding:.75rem .9rem;background:#1e293b;border:1px solid #334155;border-radius:8px}
.project-info .pi-name{font-size:.95rem;font-weight:700;color:#f1f5f9}
.project-info .pi-desc{font-size:.82rem;color:#d1d5db;margin-top:.25rem}
.project-info .pi-row{display:flex;flex-wrap:wrap;gap:.4rem .9rem;margin-top:.45rem;font-size:.78rem;color:#9ca3af}
.project-info .pi-row .pi-k{color:#6b7280}
.project-info a{color:#60a5fa;text-decoration:none}
.project-info a:hover{text-decoration:underline}
.project-info code{font-family:"JetBrains Mono","Fira Code",Consolas,monospace;font-size:.74rem;color:#93c5fd}
.pi-repo-btn{display:inline-flex;align-items:center;gap:.3rem;padding:.18rem .5rem;background:#0f172a;border:1px solid #3b82f6;border-radius:6px;color:#60a5fa !important;font-weight:600}
.pi-repo-btn:hover{background:#1d283a;text-decoration:none !important}
#search{background:#1e293b;border:1px solid #334155;color:#e5e7eb;padding:.4rem .75rem;border-radius:6px;font-size:.875rem;width:220px;outline:none}
#search:focus{border-color:#3b82f6}
.group-label{font-size:.8rem;color:#9ca3af;display:flex;align-items:center;gap:.4rem}
select{background:#1e293b;border:1px solid #334155;color:#e5e7eb;padding:.35rem .6rem;border-radius:6px;font-size:.8rem;outline:none;cursor:pointer}
select:focus{border-color:#3b82f6}
.filterbar{display:flex;flex-direction:column;gap:.4rem;margin-bottom:.6rem;padding:.6rem .75rem;background:#1e293b;border:1px solid #334155;border-radius:8px}
.filter-row{display:flex;flex-wrap:wrap;gap:.3rem;align-items:center}
.filter-lbl{font-size:.65rem;font-weight:700;text-transform:uppercase;letter-spacing:.08em;color:#4b5563;white-space:nowrap;margin-right:.25rem}
.filter-toggle{font-size:.65rem;font-weight:600;padding:.18rem .5rem;border-radius:4px;border:1px solid #334155;background:#0f172a;color:#9ca3af;cursor:pointer;margin-left:auto;white-space:nowrap;letter-spacing:.04em}
.filter-toggle:hover{background:#1e293b;color:#e5e7eb;border-color:#4b5563}
.badge{font-size:.68rem;font-weight:700;padding:.18rem .5rem;border-radius:4px;border:1px solid transparent;cursor:pointer;letter-spacing:.03em;transition:opacity .12s,filter .12s;user-select:none}
.badge.off{opacity:.28;filter:grayscale(.5)}
.sev-critical{background:#450a0a;color:#fca5a5;border-color:#991b1b}
.sev-high {background:#431407;color:#fdba74;border-color:#9a3412}
.sev-medium {background:#422006;color:#fcd34d;border-color:#854d0e}
.sev-low {background:#0c1a3a;color:#93c5fd;border-color:#1d4ed8}
.sev-negligible{background:#1e293b;color:#6b7280;border-color:#334155}
.type-badge{background:#0c1e33;color:#7dd3fc;border-color:#0369a1}
.diff-very-easy{background:#052e16;color:#86efac;border-color:#16a34a}
.diff-easy{background:#064e3b;color:#6ee7b7;border-color:#059669}
.diff-medium{background:#1e293b;color:#9ca3af;border-color:#4b5563}
.diff-hard{background:#2e1065;color:#c4b5fd;border-color:#7c3aed}
.st-open{background:#1e293b;color:#4b5563;border-color:#334155}
.st-fixed{background:#052e16;color:#4ade80;border-color:#16a34a}
.st-wont{background:#431407;color:#fdba74;border-color:#9a3412}
.st-fp{background:#2e1065;color:#c4b5fd;border-color:#7c3aed}
#progress-wrap{display:flex;align-items:center;gap:.75rem;margin-bottom:.6rem;flex-wrap:wrap}
#progress-bar{flex:1;min-width:120px;height:7px;background:#1e293b;border-radius:4px;overflow:hidden;border:1px solid #334155}
#progress-fill{height:100%;background:#22c55e;border-radius:4px;transition:width .35s ease;width:0}
#progress-text{font-size:.75rem;color:#6b7280;white-space:nowrap}
.action-btn{font-size:.72rem;padding:.25rem .6rem;border-radius:5px;border:1px solid #334155;background:#1e293b;color:#9ca3af;cursor:pointer;white-space:nowrap}
.action-btn:hover{background:#334155;color:#e5e7eb}
.action-btn.active{background:#1d4ed8;color:#fff;border-color:#2563eb}
.action-btn.active:hover{background:#2563eb}
.search-wrap{display:flex;align-items:center;gap:.35rem}
#hidden-bar{margin-bottom:.6rem}
.hidden-inner{display:flex;flex-wrap:wrap;gap:.35rem;align-items:center}
.hidden-lbl{font-size:.72rem;color:#4b5563}
.show-btn{font-size:.7rem;padding:.15rem .45rem;border-radius:4px;background:#1e293b;color:#9ca3af;border:1px solid #334155;cursor:pointer}
.show-btn:hover{background:#334155;color:#e5e7eb}
.show-all-btn{font-size:.7rem;padding:.15rem .45rem;border-radius:4px;background:#2563eb;color:#fff;border:none;cursor:pointer}
.show-all-btn:hover{background:#1d4ed8}
.group-section{margin-bottom:1.5rem}
.group-header{display:flex;align-items:center;gap:.5rem;margin-bottom:.5rem;padding-bottom:.35rem;border-bottom:1px solid #334155}
.group-title{font-size:.95rem;font-weight:700;color:#d1d5db}
.group-title.all-done{color:#4ade80}
.group-count{font-size:.72rem;color:#4b5563;background:#1e293b;padding:.1rem .4rem;border-radius:10px;white-space:nowrap}
.icon-btn{background:transparent;border:none;color:#334155;cursor:pointer;font-size:.85rem;padding:.05rem .2rem;line-height:1}
.icon-btn:hover{color:#6b7280}
.hide-btn{margin-left:auto;font-size:.68rem;padding:.12rem .45rem;border-radius:4px;background:transparent;color:#334155;border:1px solid #1e293b;cursor:pointer}
.hide-btn:hover{background:#1e293b;color:#6b7280}
.card{background:#1e293b;border:1px solid #334155;border-radius:8px;margin-bottom:.6rem;overflow:hidden;transition:opacity .15s}
.card.resolved{opacity:.55;border-left:3px solid #16a34a}
.card.resolved:hover{opacity:.85}
.card-header{display:flex;align-items:center;gap:.5rem;padding:.45rem .85rem;background:#162032;border-bottom:1px solid #334155;flex-wrap:wrap;cursor:pointer;user-select:none}
.card-header:hover{background:#1a2840}
.card-arrow{font-size:.65rem;color:#334155;flex-shrink:0}
.card-id{font-size:.75rem;font-weight:700;color:#6b7280;font-family:monospace;letter-spacing:.04em}
.card-loc{font-size:.72rem;color:#6b7280;font-family:"JetBrains Mono","Fira Code",Consolas,monospace;letter-spacing:.02em;white-space:nowrap;overflow:hidden;text-overflow:ellipsis;min-width:0}
.status-btn{margin-left:auto;font-size:.68rem;font-weight:600;padding:.15rem .5rem;border-radius:4px;border:1px solid transparent;cursor:pointer;white-space:nowrap;transition:filter .1s}
.status-btn:hover{filter:brightness(1.25)}
.copy-btn{font-size:.68rem;font-weight:600;padding:.15rem .5rem;border-radius:4px;border:1px solid #334155;background:#0f172a;color:#9ca3af;cursor:pointer;white-space:nowrap;transition:background .12s,color .12s}
.copy-btn:hover{background:#334155;color:#e5e7eb}
.card-body{padding:.85rem .9rem}
.field-lbl{font-size:.85rem;font-weight:700;text-transform:uppercase;letter-spacing:.1em;color:#7dd3fc;margin:1rem 0 .35rem;padding-bottom:.18rem;border-bottom:1px solid #1e293b}
.field-lbl:first-child{margin-top:0}
.desc,.fix-txt{font-size:.875rem;color:#d1d5db;line-height:1.55}
.desc p,.fix-txt p{margin:.4rem 0}
.desc p:first-child,.fix-txt p:first-child{margin-top:0}
.desc p:last-child,.fix-txt p:last-child{margin-bottom:0}
.loc-block{margin:.3rem 0 .6rem}
.loc-ref{font-size:.8rem;color:#38bdf8;font-family:"JetBrains Mono","Fira Code",Consolas,monospace;font-weight:600;margin-bottom:.22rem}
a.loc-link{color:inherit;text-decoration:none}
a.loc-link:hover{text-decoration:underline}
.loc-lines{color:#4b5563;font-weight:400}
pre{background:#0f172a;color:#cdd6f4;border-radius:5px;padding:.6rem .8rem;overflow-x:auto;white-space:pre;font-size:.8rem;line-height:1.55;font-family:"JetBrains Mono","Fira Code",Consolas,monospace;border:1px solid #1e293b;max-height:320px}
.desc pre.md-code,.fix-txt pre.md-code{margin:.5rem 0;font-size:.82rem;color:#e5e7eb;background:#0b1224;border:1px solid #1e293b;max-height:none}
code.md-inline{font-family:"JetBrains Mono","Fira Code",Consolas,monospace;font-size:.86em;background:#0b1224;color:#fca5a5;border:1px solid #1e293b;border-radius:3px;padding:.04rem .32rem;white-space:pre-wrap;word-break:break-word}
.not-found{font-size:.8rem;color:#334155;font-style:italic}
mark{background:#854d0e;color:#fef9c3;border-radius:2px;padding:0 .1rem}
.no-results{text-align:center;color:#334155;padding:3rem 1rem;font-size:.9rem}
</style>
</head>
<body>
<script type="application/json" id="D">[
{"id":"PANIC_1","type":"panic","type_label":"Panic","severity":"high","difficulty":"medium","description":"In `run_method` in `convert_32bit_to_64bit.rs`, `argument_sizes` is computed for every kernel input argument via `arg.info.v_type.layout().unwrap()`. `Type::layout()` (defined in `ptx_parser/src/ast.rs`) explicitly returns `None` for `Type::Texref` (`Type::Texref => return None`). PTX allows old-style texture-reference kernel parameters declared as `.param .texref` (parsed via `multi_variable` in `ptx_parser/src/lib.rs`, which produces `VariableInfo { v_type: Type::Texref, .. }` for `.param` state space). Since `convert_32bit_to_64bit::run` is invoked unconditionally for every 32-bit module (`ptx/src/pass/mod.rs`, `if ast.address_size == 32 { convert_32bit_to_64bit::run(...) }`), any 32-bit kernel with a `.texref` parameter causes `.unwrap()` to panic on `None`, crashing the whole compilation instead of returning a `TranslateError`.","fix":"Replace the `.unwrap()` with proper error propagation, e.g. `let layout = arg.info.v_type.layout().ok_or_else(error_unreachable)?;` (or `error_todo()` if `.texref` kernel parameters are not yet supported by the 32-bit pass), so malformed/unsupported input produces a `TranslateError` instead of a panic.","locations":[{"ref":"ptx/src/pass/convert_32bit_to_64bit.rs:190-199","code":" 190 | let argument_sizes = method\n 191 | .input_arguments\n 192 | .iter()\n 193 | .map(|arg| {\n 194 | let layout = arg.info.v_type.layout().unwrap();\n 195 | let size = layout.size() as u32;\n 196 | let align = layout.align() as u32;\n 197 | (size, align)\n 198 | })\n 199 | .collect();","lines":10,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx/src/pass/convert_32bit_to_64bit.rs#L190-L199"}]},
{"id":"PANIC_2","type":"panic","type_label":"Panic","severity":"high","difficulty":"easy","description":"`is_addressable` in `insert_implicit_conversions.rs` uses `todo!()` for `ast::StateSpace::SharedCluster`, `ast::StateSpace::SharedCta`, and `ast::StateSpace::ParamFunc`, which unconditionally panics if reached. It is called from `default_implicit_conversion` (used on every operand of every instruction that isn't itself in `Reg` space) whenever `operand_space != ast::StateSpace::Reg` and the instruction is in `Reg` space - e.g. taking the address of a `.shared::cta` or `.shared::cluster` variable (a legitimate, modern PTX ISA feature used with cluster-scoped shared memory on newer GPUs). Note that the sibling function `coerces_to_generic` in the same file treats `SharedCta` and `SharedCluster` as normal, valid address spaces (grouped with `Shared`), showing `is_addressable` is simply missing the equivalent arms rather than these spaces being genuinely unsupported. Unlike other unimplemented-feature paths in this codebase which return a graceful `TranslateError` via `error_todo()`/`error_unreachable()`, this uses a hard `todo!()` panic that will crash the whole compiler on otherwise-valid PTX input.","fix":"Add `ast::StateSpace::SharedCluster | ast::StateSpace::SharedCta => true` (mirroring `coerces_to_generic`) and handle `ast::StateSpace::ParamFunc` explicitly (`true` or `false` depending on intended semantics) instead of falling through to `todo!()`, so `is_addressable` never panics on a valid `ast::StateSpace` value.","locations":[{"ref":"ptx/src/pass/insert_implicit_conversions.rs:184-196","code":" 184 | fn is_addressable(this: ast::StateSpace) -> bool {\n 185 | match this {\n 186 | ast::StateSpace::Const\n 187 | | ast::StateSpace::Generic\n 188 | | ast::StateSpace::Global\n 189 | | ast::StateSpace::ParamEntry\n 190 | | ast::StateSpace::Local\n 191 | | ast::StateSpace::Shared => true,\n 192 | ast::StateSpace::Param | ast::StateSpace::Reg => false,\n 193 | ast::StateSpace::SharedCluster\n 194 | | ast::StateSpace::SharedCta\n 195 | | ast::StateSpace::ParamFunc => todo!(),\n 196 | }","lines":13,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx/src/pass/insert_implicit_conversions.rs#L184-L196"}]},
{"id":"PANIC_3","type":"panic","type_label":"Panic","severity":"high","difficulty":"easy","description":"`zluda_common::drop_checked` transmutes the incoming `T::CudaHandle` directly into a `ManuallyDrop<Box<LiveCheck<T>>>` and immediately calls `LiveCheck::drop_checked` on it, with no null check. This is inconsistent with the sibling function `as_ref` (`zluda_common/src/lib.rs:1088-1094`), which explicitly checks `mem::transmute_copy::<_, usize>(handle) == 0` and returns `None` before dereferencing. Because callers such as `zluda::impl::module::unload` pass the raw `CUmodule` handle straight through (via `from_cuda_nop!`, which performs no null check for bare handle types like `CUmodule`), a caller invoking e.g. `cuModuleUnload(0)` with a NULL handle causes `drop_checked` to dereference a null pointer inside `LiveCheck::drop_checked`'s `self.cookie == T::COOKIE` check, crashing the process.","fix":"Add the same null check used in `as_ref` to `drop_checked` before transmuting the handle, e.g. return `Err(T::LIVENESS_FAIL)` if `mem::transmute_copy::<_, usize>(&handle) == 0`.","locations":[{"ref":"zluda_common/src/lib.rs:1089-1095","code":" 1089 | pub fn as_ref<'a, T: ZludaObject>(handle: &'a T::CudaHandle) -> Option<&'a LiveCheck<T>> {\n 1090 | if unsafe { mem::transmute_copy::<_, usize>(handle) } == 0 {\n 1091 | return None;\n 1092 | }\n 1093 | let option_box = unsafe { mem::transmute::<_, &'a ManuallyDrop<Box<LiveCheck<T>>>>(handle) };\n 1094 | Some(option_box)\n 1095 | }","lines":7,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda_common/src/lib.rs#L1089-L1095"},{"ref":"zluda_common/src/lib.rs:1101-1107","code":" 1101 | pub fn drop_checked<T: ZludaObject>(handle: T::CudaHandle) -> Result<(), T::Error> {\n 1102 | let mut wrapped_object: ManuallyDrop<Box<LiveCheck<T>>> =\n 1103 | unsafe { mem::transmute_copy(&handle) };\n 1104 | let underlying_error = LiveCheck::drop_checked(&mut wrapped_object)?;\n 1105 | unsafe { ManuallyDrop::drop(&mut wrapped_object) };\n 1106 | underlying_error\n 1107 | }","lines":7,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda_common/src/lib.rs#L1101-L1107"}]},
{"id":"PANIC_4","type":"panic","type_label":"Panic","severity":"high","difficulty":"medium","description":"`cuGetExportTable` is generated by the `override_fn_full` macro (`zluda_trace/src/lib.rs:149-159`), which calls `cuGetExportTable_impl` directly with no `catch_unwind` protection, unlike every other traced entry point which goes through `GlobalState2::under_lock` whose `under_lock_impl` closure is wrapped in `std::panic::catch_unwind` (see `zluda_trace/src/lib.rs:1067-1078`). Inside `cuGetExportTable_impl`, the `LateInit::Error => panic!()` arm and the following `.unwrap()` (`zluda_trace/src/lib.rs:691` and `:696`) are reachable in normal operation: if a previous CUDA call already failed to load the real driver library (e.g. `ZLUDA_CUDA_LIB` points at a missing/invalid library, or `GlobalDelayedState::new` otherwise fails), `global_state.delayed_state` becomes `LateInit::Error`, and any subsequent call to `cuGetExportTable` hits the explicit `panic!()`. Because this panic unwinds across an `extern \"system\"` FFI boundary with no `catch_unwind` guard, the process aborts (Rust aborts on unwind across a non-`C-unwind` extern boundary) instead of gracefully returning an error the way every other intercepted function does.","fix":"Wrap the body of `cuGetExportTable_impl` in `std::panic::catch_unwind` (mirroring `GlobalState2::under_lock`), or better, refactor `cuGetExportTable_impl` to reuse `GlobalState2::under_lock` so all entry points share the same panic-safety guarantees.","locations":[{"ref":"zluda_trace/src/lib.rs:149-159","code":" 149 | macro_rules! override_fn_full {\n 150 | ($($abi:literal fn $fn_name: ident ( $($arg_id:ident : $arg_type:ty),* ) -> $ret_type:ty;)*) => {\n 151 | $(\n 152 | #[no_mangle]\n 153 | #[allow(non_snake_case)]\n 154 | pub unsafe extern $abi fn $fn_name ( $( $arg_id : $arg_type),* ) -> $ret_type {\n 155 | paste!{ [<$fn_name _impl >] ( $($arg_id),* ) }\n 156 | }\n 157 | )*\n 158 | }\n 159 | }","lines":11,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda_trace/src/lib.rs#L149-L159"},{"ref":"zluda_trace/src/lib.rs:688-698","code":" 688 | let delayed_state = match global_state.delayed_state {\n 689 | LateInit::Success(ref mut delayed_state) => delayed_state,\n 690 | // There's no libcuda to load, so we might as well panic\n 691 | LateInit::Error => panic!(),\n 692 | LateInit::Unitialized => {\n 693 | global_state.delayed_state = GlobalDelayedState::new(panic_guard.writer, &mut logger);\n 694 | // `global_state.delayed_state` could be LateInit::Error,\n 695 | // we can crash in this case since there's no libcuda\n 696 | global_state.delayed_state.as_mut().unwrap()\n 697 | }\n 698 | };","lines":11,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda_trace/src/lib.rs#L688-L698"},{"ref":"zluda_trace/src/lib.rs:928","code":" 928 | override_fn_full <= [cuGetExportTable],","lines":1,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda_trace/src/lib.rs#L928"}]},
{"id":"PANIC_5","type":"panic","type_label":"Panic","severity":"medium","difficulty":"hard","description":"In `cu_launch_kernel` (`zluda32/src/lib.rs:513-528`), the `.scan(...)` closure that builds `kernel_params` calls `layout.extend(arg_info.to_layout().unwrap()).unwrap()` - two chained `unwrap()`s per kernel argument. `arg_info` comes from `state.function_args`, which was populated by the `zludaGetFunctionArgs` IPC response (`zludaGetFunctionArgsOut.args`, a `Vec<FunctionArgInfo>` deserialized from the 64-bit server, see `cu_module_get_function` at `zluda32/src/lib.rs:703-735`). If `arg_info.to_layout()` fails (e.g. an invalid/zero alignment or a size that isn't a multiple of a power-of-two alignment) the first `unwrap()` panics; if the running `Layout::extend` computation overflows `isize::MAX` while accumulating many/huge argument sizes, the second `unwrap()` panics. Because `FunctionArgInfo` values ultimately originate from parsing the loaded module/kernel's argument metadata (which for a corrupted, adversarial, or simply buggy PTX/cubin module can contain unusual sizes/alignments), a crafted kernel signature can crash the 32-bit `zluda32` client process via this reachable panic instead of surfacing a `CUerror`.","fix":"Replace both `unwrap()` calls with proper error propagation, e.g. `let layout = arg_info.to_layout().map_err(|_| CUerror::INVALID_VALUE)?; let (new_layout, offset) = layout_acc.extend(layout).map_err(|_| CUerror::INVALID_VALUE)?;`, changing the closure to return `Result<_, CUerror>` for every branch.","locations":[{"ref":"zluda32/src/lib.rs:523-538","code":" 523 | let kernel_params = args\n 524 | .iter()\n 525 | .scan(\n 526 | unsafe { Layout::from_size_align_unchecked(0, 1) },\n 527 | |layout, arg_info| {\n 528 | let (new_layout, offset) = layout.extend(arg_info.to_layout().unwrap()).unwrap();\n 529 | total_size = offset + arg_info.size as usize;\n 530 | *layout = new_layout;\n 531 | let start = arg_ptr.wrapping_byte_add(offset as usize).cast::<u8>();\n 532 | Some(Ok(unsafe {\n 533 | std::slice::from_raw_parts(start, arg_info.size as usize)\n 534 | }\n 535 | .to_vec()))\n 536 | },\n 537 | )\n 538 | .collect::<Result<Vec<_>, CUerror>>()?;","lines":16,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda32/src/lib.rs#L523-L538"}]},
{"id":"PANIC_6","type":"panic","type_label":"Panic","severity":"low","difficulty":"easy","description":"In `ptx_parser_macros_impl/src/lib.rs`, `ArgumentField::parse_block` parses the `{ type: ..., space: ..., repr: ..., dst: ..., relaxed_type_check: ... }` block form of an argument declaration used inside `arguments<T>: { ... }` in the `generate_instruction_type!` DSL (consumed from `ptx_parser/src/ast.rs`). The `repr` key is optional at the syntax level (nothing enforces its presence while parsing `all_fields`), but the function ends with `Ok((repr.unwrap(), type_, space, is_dst, relaxed_type_check))`. If a developer writes a field block that omits `repr:` (e.g. only specifies `type:`/`space:`), `repr` stays `None` and `.unwrap()` panics during macro expansion instead of reporting a clean error at the relevant span. The same anti-pattern also occurs in `InstructionVariant::emit_visit_map`, where `Some(Arguments::Decl(_)) => { let map = self.map.as_ref().unwrap(); ... }` panics if a variant using the `arguments: SomeType` (`Decl`) form omits the required `map:` property that `emit_visit_map` needs.","fix":"Replace both `.unwrap()` calls with proper `syn::Error`s: in `parse_block`, if `repr` is `None` after the loop, return `Err(syn::Error::new(content.span(), \"missing required `repr` key in argument block\"))` instead of `Ok((repr.unwrap(), ...))`; in `emit_visit_map`, when `self.map` is `None` for a `Arguments::Decl` variant, return a `syn::Error` from the enclosing fallible parse/emit path (or validate eagerly in `InstructionVariant::parse`) instead of calling `.unwrap()` on `self.map.as_ref()`.","locations":[{"ref":"ptx_parser_macros_impl/src/lib.rs:573-624","code":" 573 | fn parse_block(\n 574 | input: syn::parse::ParseStream,\n 575 | ) -> syn::Result<(Type, Option<Expr>, Option<Expr>, Option<bool>, bool)> {\n 576 | let content;\n 577 | braced!(content in input);\n 578 | let all_fields =\n 579 | Punctuated::<ExprOrPath, Token![,]>::parse_terminated_with(&content, |content| {\n 580 | let lookahead = content.lookahead1();\n 581 | Ok(if lookahead.peek(Token![type]) {\n 582 | content.parse::<Token![type]>()?;\n 583 | content.parse::<Token![:]>()?;\n 584 | ExprOrPath::Type(content.parse::<Expr>()?)\n 585 | } else if lookahead.peek(Ident) {\n 586 | let name_ident = content.parse::<Ident>()?;\n 587 | content.parse::<Token![:]>()?;\n 588 | match &*name_ident.to_string() {\n 589 | \"relaxed_type_check\" => {\n 590 | ExprOrPath::RelaxedTypeCheck(content.parse::<LitBool>()?.value)\n 591 | }\n 592 | \"repr\" => ExprOrPath::Repr(content.parse::<Type>()?),\n 593 | \"space\" => ExprOrPath::Space(content.parse::<Expr>()?),\n 594 | \"dst\" => {\n 595 | let ident = content.parse::<LitBool>()?;\n 596 | ExprOrPath::Dst(ident.value)\n 597 | }\n 598 | name => {\n 599 | return Err(syn::Error::new(\n 600 | name_ident.span(),\n 601 | format!(\"Unexpected key `{}`, expected `repr` or `space\", name),\n 602 | ))\n 603 | }\n 604 | }\n 605 | } else {\n 606 | return Err(lookahead.error());\n 607 | })\n 608 | })?;\n 609 | let mut repr = None;\n 610 | let mut type_ = None;\n 611 | let mut space = None;\n 612 | let mut is_dst = None;\n 613 | let mut relaxed_type_check = false;\n 614 | for exp_or_path in all_fields {\n 615 | match exp_or_path {\n 616 | ExprOrPath::Repr(r) => repr = Some(r),\n 617 | ExprOrPath::Type(t) => type_ = Some(t),\n 618 | ExprOrPath::Space(s) => space = Some(s),\n 619 | ExprOrPath::Dst(x) => is_dst = Some(x),\n 620 | ExprOrPath::RelaxedTypeCheck(relaxed) => relaxed_type_check = relaxed,\n 621 | }\n 622 | }\n 623 | Ok((repr.unwrap(), type_, space, is_dst, relaxed_type_check))\n 624 | }","lines":52,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx_parser_macros_impl/src/lib.rs#L573-L624"},{"ref":"ptx_parser_macros_impl/src/lib.rs:311-327","code":" 311 | fn emit_visit_map(&self, enum_: &Ident, tokens: &mut TokenStream) {\n 312 | let name = &self.name;\n 313 | let data = &self.data.as_ref().map(|_| quote! { data,});\n 314 | let arguments = match self.arguments {\n 315 | None => None,\n 316 | Some(Arguments::Decl(_)) => {\n 317 | let map = self.map.as_ref().unwrap();\n 318 | quote! {\n 319 | #enum_ :: #name { #data arguments } => {\n 320 | #map\n 321 | }\n 322 | }\n 323 | .to_tokens(tokens);\n 324 | return;\n 325 | }\n 326 | Some(Arguments::Def(ref def)) => Some(def),\n 327 | };","lines":17,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx_parser_macros_impl/src/lib.rs#L311-L327"}]},
{"id":"PANIC_7","type":"panic","type_label":"Panic","severity":"low","difficulty":"easy","description":"In `ptx_parser_macros_impl/src/parser.rs`, `Arguments::parse` handles the `[...]` (bracketed) argument syntax used for e.g. `[a{, b}, c]`. When the bracket content does not start with `|` (the `pre_pipe` case), it recursively calls `Self::parse(&bracketed)` into `sub_args`, then unconditionally does `sub_args.0.first_mut().unwrap().pre_bracket = true;` and `sub_args.0.last_mut().unwrap().post_bracket = true;`. If the bracketed content is empty or contains only a stray comma (e.g. a typo like `[]` or `[,]` in the `derive_parser!`/`generate_instruction_type!` DSL input inside `ptx_parser/src/lib.rs` or `ptx_parser/src/ast.rs`), the recursive `Arguments::parse` call returns an empty `Vec` (its parse loop hits the trailing `else { break; }` immediately), and `first_mut()`/`last_mut()` return `None`, so `.unwrap()` panics. This turns a simple typo in the macro's DSL input into an unhelpful Rust panic during macro expansion instead of a clean `syn::Error` pointing at the offending source location.","fix":"Before calling `.unwrap()`, check whether `sub_args.0` is empty and return a `syn::Error` (e.g. `syn::Error::new(bracketed.span(), \"expected at least one argument inside `[...]`\")`) instead of panicking, so malformed DSL input produces a normal macro-expansion diagnostic.","locations":[{"ref":"ptx_parser_macros_impl/src/parser.rs:551-563","code":" 551 | } else if lookahead.peek(token::Bracket) {\n 552 | let bracketed;\n 553 | bracketed!(bracketed in input);\n 554 | if bracketed.peek(Token![|]) {\n 555 | optional = true;\n 556 | bracketed.parse::<Token![|]>()?;\n 557 | pre_pipe = true;\n 558 | ident = bracketed.parse::<HyphenatedIdent>()?;\n 559 | } else {\n 560 | let mut sub_args = Self::parse(&bracketed)?;\n 561 | sub_args.0.first_mut().unwrap().pre_bracket = true;\n 562 | sub_args.0.last_mut().unwrap().post_bracket = true;\n 563 | if peek_brace_token(input, Token![.]) {","lines":13,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx_parser_macros_impl/src/parser.rs#L551-L563"}]},
{"id":"MEM_1","type":"memory_leak","type_label":"Memory Leak","severity":"medium","difficulty":"easy","description":"`host_alloc` calls `hipHostMalloc(pp, bytesize, flags)?` to allocate pinned host memory and then `add_allocation(*pp, bytesize, context)?` to record it in the global allocation table. If `add_allocation` fails (e.g. `global_state().allocations.lock()` returns an error because the mutex was poisoned by an earlier panic elsewhere in the process), `host_alloc` returns `Err` without freeing the memory that `hipHostMalloc` already allocated into `*pp`. This both leaks the pinned host allocation and leaves `*pp` pointing at a live-but-untracked buffer even though the function reports failure, so the caller (having received an error) will never call `cuMemFreeHost` on it. This is inconsistent with the sibling function `alloc_v2`, which wraps the equivalent registration step in a `DropGuard` that calls `hipFree` if `add_allocation` fails.","fix":"Mirror the `DropGuard` pattern used in `alloc_v2`/`fill_with_zero_and_register`: on `add_allocation` failure, call `hipFreeHost(*pp)` before propagating the error, e.g. wrap the call in a guard that frees `*pp` unless `add_allocation` succeeds.","locations":[{"ref":"zluda/src/impl/memory.rs:159-168","code":" 159 | pub(crate) unsafe fn host_alloc(\n 160 | pp: &mut *mut ::core::ffi::c_void,\n 161 | bytesize: usize,\n 162 | flags: ::std::os::raw::c_uint,\n 163 | ) -> CUresult {\n 164 | let context = context::get_current_context()?;\n 165 | hipHostMalloc(pp, bytesize, flags)?;\n 166 | add_allocation(*pp, bytesize, context)?;\n 167 | Ok(())\n 168 | }","lines":10,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda/src/impl/memory.rs#L159-L168"}]},
{"id":"OVF_1","type":"overflow_underflow","type_label":"Overflow Underflow","severity":"high","difficulty":"medium","description":"`Allocator::alloc_range` in `zluda64_server/src/main.rs` computes `size.next_multiple_of(Self::ALLOCATION_UNIT)` where `size: u32` comes directly from the client-controlled `cuMemAlloc_v2In.bytesize` field (forwarded unchecked from the 32-bit client's `cu_mem_alloc_v2`). `u32::next_multiple_of` rounds up via `self + (rhs - r)`, which overflows `u32` for any `size` within `Self::ALLOCATION_UNIT` (256) of `u32::MAX` (e.g. `size = u32::MAX` gives `4294967295 % 256 == 255`, so it computes `4294967295 + 1`, which overflows). In a release build (no overflow checks) this silently wraps to a tiny value, so `units` ends up far smaller than the requested `size`; `alloc_range`/`alloc` then hands back a device pointer backed by a much smaller real allocation than the caller asked for and believes it received. Any subsequent memcpy/memset that uses the original (huge) `bytesize` against that pointer will write/read past the actual backing allocation inside the server's 512 MiB device memory pool, corrupting adjacent allocations. In a debug build the overflow instead panics, crashing the 64-bit server. Either way, a `cuMemAlloc_v2` call with a `bytesize` near `u32::MAX` (a value fully representable and reachable through the 32-bit client's `usize` argument) triggers this.","fix":"Use a checked computation, e.g. `size.checked_add(Self::ALLOCATION_UNIT - 1).and_then(|s| s.checked_div(Self::ALLOCATION_UNIT)).ok_or(CUerror::OUT_OF_MEMORY)?`, and return `CUerror::OUT_OF_MEMORY`/`CUerror::INVALID_VALUE` instead of letting the rounding arithmetic overflow.","locations":[{"ref":"zluda64_server/src/main.rs:211-217","code":" 211 | fn alloc_range(&mut self, size: u32) -> Result<Range<u32>, CUerror> {\n 212 | self.get_or_allocate_device_ptr()?;\n 213 | let units = size.next_multiple_of(Self::ALLOCATION_UNIT) / Self::ALLOCATION_UNIT;\n 214 | self.allocator\n 215 | .allocate_range(units)\n 216 | .map_err(|_| CUerror::OUT_OF_MEMORY)\n 217 | }","lines":7,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda64_server/src/main.rs#L211-L217"},{"ref":"zluda64_server/src/main.rs:1027-1035","code":" 1027 | fn cu_mem_alloc_v2(\n 1028 | state: &mut State,\n 1029 | input: &ArchivedcuMemAlloc_v2In,\n 1030 | ) -> Result<cuMemAlloc_v2Out, CUerror> {\n 1031 | let fake_ptr = state.devmemory.alloc(input.bytesize.to_native())?;\n 1032 | Ok(cuMemAlloc_v2Out {\n 1033 | dptr: u32_le::from_native(fake_ptr),\n 1034 | })\n 1035 | }","lines":9,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda64_server/src/main.rs#L1027-L1035"}]},
{"id":"OVF_2","type":"overflow_underflow","type_label":"Overflow Underflow","severity":"medium","difficulty":"easy","description":"The `s32` function in `ptx_parser/src/lib.rs` parses a signed 32-bit immediate by first parsing the *unsigned magnitude* text with `i32::from_str_radix(text, radix)` and only negating the result afterwards if a `Minus` token was seen. Because `i32::from_str_radix` on the unsigned magnitude only accepts values up to `i32::MAX` (2147483647), the legitimate value `i32::MIN` (-2147483648, magnitude 2147483648) fails to parse: `i32::from_str_radix(\"2147483648\", 10)` returns `Err(PosOverflow)`, so `s32` falls back to returning `0` and records a bogus `PtxError::ParseInt` (`PosOverflow`) instead of the correct value. This was confirmed empirically: parsing the token stream for `-2147483648` with `s32` yields `Ok(0)` plus a `ParseInt`/`PosOverflow` error, instead of `Ok(-2147483648)`. `s32` is used to parse the signed offset in register-offset operands (`preceded(Token::Plus, s32)` inside `ParsedOperand::parse`'s `ident_operands`), so a legitimate PTX memory operand such as `[%rd1+-2147483648]` silently gets mis-parsed into an offset of `0` plus a spurious integer-parse error. The sibling function `int_immediate` avoids this bug by formatting the sign into the string first (`format!(\"-{num}\")`) before calling `i64::from_str_radix`, which correctly handles `i64::MIN`.","fix":"Mirror the approach used in `int_immediate`: build the signed string first and parse it in one call, e.g. `let text = if sign.is_some() { format!(\"-{text}\") } else { text.to_string() }; match i32::from_str_radix(&text, radix) { Ok(x) => Ok(x), Err(err) => Err((0, PtxError::from(err))) }`, so the full signed range (including `i32::MIN`) parses correctly in a single `from_str_radix` call instead of parsing the magnitude as unsigned-range-limited `i32` and negating afterwards.","locations":[{"ref":"ptx_parser/src/lib.rs:316-325","code":" 316 | fn s32<'a, 'input>(stream: &mut PtxParser<'a, 'input>) -> PResult<i32> {\n 317 | take_error((opt(Token::Minus), num).map(|(sign, x)| {\n 318 | let (text, radix, _) = x;\n 319 | match i32::from_str_radix(text, radix) {\n 320 | Ok(x) => Ok(if sign.is_some() { -x } else { x }),\n 321 | Err(err) => Err((0, PtxError::from(err))),\n 322 | }\n 323 | }))\n 324 | .parse_next(stream)\n 325 | }","lines":10,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx_parser/src/lib.rs#L316-L325"}]},
{"id":"OVF_3","type":"overflow_underflow","type_label":"Overflow Underflow","severity":"medium","difficulty":"easy","description":"`cu_device_get_name` in `zluda64_server/src/main.rs` does `let mut name = vec![0u8; input.len.to_native() as usize];` where `input.len` is the `cuDeviceGetNameIn.len` field, which on the client side (`zluda32/src/lib.rs:376-396`, `cu_device_get_name`) is taken directly from the caller-supplied `len: ::core::ffi::c_int` argument of `cuDeviceGetName` and forwarded over IPC unmodified. `c_int` is signed. If a caller (or a corrupted/malformed IPC message) supplies a negative `len`, `input.len.to_native() as usize` wraps around to a huge value close to `usize::MAX` on the 64-bit server, so `vec![0u8; ...]` attempts a near-unbounded allocation, aborting the 64-bit server process (and therefore breaking the whole 32-bit-to-64-bit proxy for the running application) instead of returning `CUerror::INVALID_VALUE`.","fix":"Validate `input.len.to_native()` is non-negative (and reasonably bounded) before using it as an allocation size, e.g. `let len = input.len.to_native(); if len < 0 { return Err(CUerror::INVALID_VALUE); } let mut name = vec![0u8; len as usize];`.","locations":[{"ref":"zluda64_server/src/main.rs:1127-1146","code":" 1127 | fn cu_device_get_name(\n 1128 | state: &mut State,\n 1129 | input: &ArchivedcuDeviceGetNameIn,\n 1130 | ) -> Result<cuDeviceGetNameOut, CUerror> {\n 1131 | if input.dev.to_native() != state.device {\n 1132 | return Err(CUerror::INVALID_DEVICE);\n 1133 | }\n 1134 | let mut name = vec![0u8; input.len.to_native() as usize];\n 1135 | unsafe {\n 1136 | cuDeviceGetName(\n 1137 | name.as_mut_ptr().cast(),\n 1138 | input.len.to_native(),\n 1139 | state.device,\n 1140 | )\n 1141 | }?;\n 1142 | if let Some(pos) = name.iter().copied().position(|c| c == 0) {\n 1143 | name.truncate(pos);\n 1144 | }\n 1145 | Ok(cuDeviceGetNameOut { name })\n 1146 | }","lines":20,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda64_server/src/main.rs#L1127-L1146"}]},
{"id":"OVF_4","type":"overflow_underflow","type_label":"Overflow Underflow","severity":"medium","difficulty":"medium","description":"`get_elf_size` in `zluda_trace/src/trace.rs` reads the ELF header (`*start.cast::<elf32::header::Header>()` or `*start.cast::<elf64::header::Header>()`, 52/64 bytes) and then walks program/section headers purely based on offsets found inside that header, without ever validating that the underlying `raw_image` buffer is actually that large. `raw_image` originates from application-controlled arguments to `cuModuleLoadData`/`cuModuleLoadDataEx`/`cuModuleLoadFatBinary`/`cuLibraryLoadData`, which `record_new_library` (`zluda_trace/src/trace.rs:70-117`) passes straight into `zluda_common::CodeLibraryRef::try_load`. `CodeLibraryRef::try_load` only checks the first 4 bytes for the ELF magic (`zluda_common/src/lib.rs:1173-1187`) before classifying the buffer as `CodeModuleRef::Elf`; it performs no length validation. If the application supplies a module image that starts with the ELF magic bytes but is truncated (fewer than 52/64 bytes, or with corrupted/out-of-range program/section header offsets), `get_elf_size` reads past the end of the allocation, an out-of-bounds heap read that can crash the tracing process or leak adjacent memory into the dumped trace.","fix":"Thread the known/declared image size (when available from a wrapping fatbin/cubin header) into `get_elf_size` and bound every offset/read against it before dereferencing, or at minimum validate that program/section header offsets plus their sizes are not obviously insane (e.g. sanity-check against a conservative maximum) before doing the pointer arithmetic in `get_max_end_for`.","locations":[{"ref":"zluda_trace/src/trace.rs:168-192","code":" 168 | unsafe fn get_elf_size(start: *const c_void) -> Option<usize> {\n 169 | let start = start.cast::<u8>();\n 170 | let ei_class = start.add(mem::size_of_val(elf::header::ELFMAG));\n 171 | let (header, header_size, is_64bit): (elf::header::Header, _, _) = match *ei_class {\n 172 | elf::header::ELFCLASS32 => (\n 173 | (*start.cast::<elf32::header::Header>()).into(),\n 174 | mem::size_of::<elf32::header::Header>() as u64,\n 175 | false,\n 176 | ),\n 177 | elf::header::ELFCLASS64 => (\n 178 | (*start.cast::<elf64::header::Header>()).into(),\n 179 | mem::size_of::<elf64::header::Header>() as u64,\n 180 | true,\n 181 | ),\n 182 | _ => return None,\n 183 | };\n 184 | let mut max_end = header_size;\n 185 | max_end = max_end.max(get_max_end_for::<elf::program_header::ProgramHeader>(\n 186 | start, header, is_64bit,\n 187 | )?);\n 188 | max_end = max_end.max(get_max_end_for::<elf::section_header::SectionHeader>(\n 189 | start, header, is_64bit,\n 190 | )?);\n 191 | Some(max_end as usize)\n 192 | }","lines":25,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda_trace/src/trace.rs#L168-L192"}]},
{"id":"LOGIC_1","type":"logic_error","type_label":"Logic Error","severity":"critical","difficulty":"medium","description":"`Allocator::translate` and `Allocator::translate_range` in `zluda64_server/src/main.rs` compute `base_ptr.wrapping_byte_add(offset)` from a client-supplied `offset` (a fake `CUdeviceptr` value coming straight off the IPC channel from the 32-bit `zluda32` process) with **no bounds check whatsoever** against any real allocation - not even a check that `offset < Allocator::ALLOCATOR_SIZE`. This is inconsistent with `Allocator::get_range`, which does look up `allocation_ends` (a `BTreeMap`) to validate the offset falls within a live allocation; that safe pattern exists in the same file but is only used by `cu_mem_get_address_range_v2`. Every other consumer of `translate`/`translate_range` - `cu_memcpy_dto_d_v2` (dstDevice/srcDevice), `cu_memcpy_dto_d_async_v2`, `cu_memset_d8_v2` (dstDevice), `cu_memcpy_dtoh_async_v2` (src_device), `cu_memcpy_hto_d_async_v2` (dst_device), and `cu_tex_ref_set_address_v2` (dptr) - forwards the raw client-supplied offset (and, for the memcpy/memset opcodes, an equally client-supplied `byte_count`/`n`) straight to the real HIP call (`cuMemcpyDtoD_v2`, `cuMemsetD8_v2`, `cuMemcpyDtoHAsync_v2`, `cuMemcpyHtoDAsync_v2`) without ever checking that `[offset, offset + size)` stays inside a real, currently-allocated range. A 32-bit client that sends a `cuMemcpyDtoD_v2`/`cuMemsetD8_v2`/`cuMemcpyDtoHAsync_v2` message with a bogus or stale `dptr`/`srcDevice`/`dstDevice` offset (e.g. after the corresponding `cuMemFree_v2` freed it, or simply an out-of-range/garbage offset from a corrupted or malicious IPC message) causes the 64-bit server to read or write far outside the actual HIP allocation. For `cu_memcpy_dtoh_async_v2` this is an out-of-bounds *read* whose contents get copied back into the reply and sent to the 32-bit client - an information-disclosure primitive - while for the write-side opcodes it is out-of-bounds memory corruption in the server process.","fix":"Give `Allocator::translate`/`translate_range` the same validation `get_range` already performs: look up `allocation_ends` for the allocation containing `offset`, and reject the call (return `CUerror::INVALID_VALUE`) if the allocation is missing or if `offset + size` exceeds the allocation's end. Thread the requested copy/memset size into `translate` (or add a `translate_checked(offset, size)` helper) so `cu_memcpy_dto_d_v2`, `cu_memcpy_dto_d_async_v2`, `cu_memset_d8_v2`, `cu_memcpy_dtoh_async_v2`, `cu_memcpy_hto_d_async_v2`, and `cu_tex_ref_set_address_v2` all validate the full byte range before calling into HIP.","locations":[{"ref":"zluda64_server/src/main.rs:235-241","code":" 235 | fn translate(&self, offset: u32) -> Result<*mut c_void, CUerror> {\n 236 | if offset == 0 {\n 237 | return Ok(ptr::null_mut());\n 238 | }\n 239 | let base_ptr = self.get_device_ptr()?;\n 240 | Ok(base_ptr.wrapping_byte_add(offset as usize))\n 241 | }","lines":7,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda64_server/src/main.rs#L235-L241"},{"ref":"zluda64_server/src/main.rs:729-738","code":" 729 | fn cu_memcpy_dto_d_v2(\n 730 | state: &mut State,\n 731 | input: &ArchivedcuMemcpyDtoD_v2In,\n 732 | ) -> Result<cuMemcpyDtoD_v2Out, CUerror> {\n 733 | let dst_device = CUdeviceptr_v2(state.devmemory.translate(input.dstDevice.to_native())?);\n 734 | let src_device = CUdeviceptr_v2(state.devmemory.translate(input.srcDevice.to_native())?);\n 735 | let byte_count = input.ByteCount.to_native();\n 736 | unsafe { cuMemcpyDtoD_v2(dst_device, src_device, byte_count as usize) }?;\n 737 | Ok(cuMemcpyDtoD_v2Out {})\n 738 | }","lines":10,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda64_server/src/main.rs#L729-L738"},{"ref":"zluda64_server/src/main.rs:740-750","code":" 740 | fn cu_memcpy_dto_d_async_v2(\n 741 | state: &mut State,\n 742 | input: &ArchivedcuMemcpyDtoDAsync_v2In,\n 743 | ) -> Result<cuMemcpyDtoDAsync_v2Out, CUerror> {\n 744 | let dst_device = CUdeviceptr_v2(state.devmemory.translate(input.dstDevice.to_native())?);\n 745 | let src_device = CUdeviceptr_v2(state.devmemory.translate(input.srcDevice.to_native())?);\n 746 | let byte_count = input.ByteCount.to_native();\n 747 | let cu_stream = CUstream(state.handles.get(input.hStream.to_native())?);\n 748 | unsafe { cuMemcpyDtoDAsync_v2(dst_device, src_device, byte_count as usize, cu_stream) }?;\n 749 | Ok(cuMemcpyDtoDAsync_v2Out {})\n 750 | }","lines":11,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda64_server/src/main.rs#L740-L750"},{"ref":"zluda64_server/src/main.rs:752-761","code":" 752 | fn cu_memset_d8_v2(\n 753 | state: &mut State,\n 754 | input: &ArchivedcuMemsetD8_v2In,\n 755 | ) -> Result<cuMemsetD8_v2Out, CUerror> {\n 756 | let dptr = CUdeviceptr_v2(state.devmemory.translate(input.dstDevice.to_native())?);\n 757 | let uc = input.uc;\n 758 | let n = input.N.to_native();\n 759 | unsafe { cuMemsetD8_v2(dptr, uc, n as usize) }?;\n 760 | Ok(cuMemsetD8_v2Out {})\n 761 | }","lines":10,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda64_server/src/main.rs#L752-L761"},{"ref":"zluda64_server/src/main.rs:857-875","code":" 857 | fn cu_tex_ref_set_address_v2(\n 858 | state: &mut State,\n 859 | input: &ArchivedcuTexRefSetAddress_v2In,\n 860 | ) -> Result<cuTexRefSetAddress_v2Out, CUerror> {\n 861 | let mut byte_offset = 0;\n 862 | let dptr = input.dptr.to_native();\n 863 | let dptr = CUdeviceptr_v2(state.devmemory.translate(dptr)?);\n 864 | unsafe {\n 865 | cuTexRefSetAddress_v2(\n 866 | &mut byte_offset,\n 867 | state.handles.get(input.hTexRef.to_native())?,\n 868 | dptr,\n 869 | input.bytes.to_native() as usize,\n 870 | )\n 871 | }?;\n 872 | Ok(cuTexRefSetAddress_v2Out {\n 873 | ByteOffset: (byte_offset as u32).into(),\n 874 | })\n 875 | }","lines":19,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda64_server/src/main.rs#L857-L875"},{"ref":"zluda64_server/src/main.rs:891-909","code":" 891 | fn cu_memcpy_dtoh_async_v2(\n 892 | state: &mut State,\n 893 | input: &ArchivedcuMemcpyDtoHAsync_v2In,\n 894 | ) -> Result<cuMemcpyDtoHAsync_v2Out, CUerror> {\n 895 | let devptr = state.devmemory.translate(input.src_device.to_native())?;\n 896 | let stream = input.stream.to_native();\n 897 | let stream = CUstream(state.handles.get(stream)?);\n 898 | let mut dst_host = vec![0u8; input.byte_count.to_native() as usize];\n 899 | unsafe {\n 900 | cuMemcpyDtoHAsync_v2(\n 901 | dst_host.as_mut_ptr().cast(),\n 902 | CUdeviceptr_v2(devptr),\n 903 | input.byte_count.to_native() as usize,\n 904 | stream,\n 905 | )\n 906 | }?;\n 907 | unsafe { cuStreamSynchronize(stream) }?;\n 908 | Ok(cuMemcpyDtoHAsync_v2Out { dst_host })\n 909 | }","lines":19,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda64_server/src/main.rs#L891-L909"},{"ref":"zluda64_server/src/main.rs:1010-1025","code":" 1010 | fn cu_memcpy_hto_d_async_v2(\n 1011 | state: &mut State,\n 1012 | input: &ArchivedcuMemcpyHtoDAsync_v2In,\n 1013 | ) -> Result<cuMemcpyHtoDAsync_v2Out, CUerror> {\n 1014 | let devptr = state.devmemory.translate(input.dst_device.to_native())?;\n 1015 | let stream = CUstream(state.handles.get(input.stream.to_native())?);\n 1016 | unsafe {\n 1017 | cuMemcpyHtoDAsync_v2(\n 1018 | CUdeviceptr_v2(devptr),\n 1019 | input.src_host.as_ptr().cast(),\n 1020 | input.src_host.len(),\n 1021 | stream,\n 1022 | )\n 1023 | }?;\n 1024 | Ok(cuMemcpyHtoDAsync_v2Out {})\n 1025 | }","lines":16,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda64_server/src/main.rs#L1010-L1025"}]},
{"id":"LOGIC_2","type":"logic_error","type_label":"Logic Error","severity":"critical","difficulty":"hard","description":"`GlobalState2::under_lock` (via `under_lock_impl`) acquires `GLOBAL_STATE2` (a `ReentrantMutex<RefCell<GlobalState2>>`) once at the top of the function and holds it for the *entire* duration of the call, including the call to `inner_call(pre_ptr)` which invokes the real, potentially long-blocking CUDA driver function (e.g. `cuStreamSynchronize`, `cuEventSynchronize`, `cuCtxSynchronize`, which are routed through this same path via the `extern_redirect` macro). Since virtually every intercepted CUDA entry point in `zluda_trace` goes through `GlobalState2::under_lock`, this means all CUDA calls across *all* threads in the traced process are fully serialized on one global lock for as long as the underlying driver call takes. Because `ReentrantMutex` only allows re-entry from the *same* thread, this can deadlock a real multi-threaded application: if thread A calls a blocking function like `cuStreamSynchronize` (holding `GLOBAL_STATE2` for the whole wait), and the work it is waiting for must be submitted by thread B via another traced function such as `cuLaunchKernel`, thread B will block forever trying to acquire `GLOBAL_STATE2`, so the awaited work is never submitted and thread A never wakes up. This directly contradicts the function's own design comment stating 'Thread-safety: CUDA functions can be called from multiple threads'.","fix":"Do not hold the `GLOBAL_STATE2` lock across the call to `inner_call`. Copy/extract everything needed from `GlobalDelayedState`/`FnCallLog` before releasing the lock, drop the `ReentrantMutexGuard` before invoking `inner_call(pre_ptr)`, perform the actual (possibly blocking) CUDA call without holding the global lock, and then re-acquire the lock only for the short post-call logging phase.","locations":[{"ref":"zluda_trace/src/lib.rs:947-948","code":" 947 | static GLOBAL_STATE2: LazyLock<ReentrantMutex<RefCell<GlobalState2>>> =\n 948 | LazyLock::new(|| ReentrantMutex::new(RefCell::new(GlobalState2::new())));","lines":2,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda_trace/src/lib.rs#L947-L948"},{"ref":"zluda_trace/src/lib.rs:1007-1066","code":" 1007 | ) -> InnerResult {\n 1008 | let global_state = GLOBAL_STATE2.lock();\n 1009 | let global_state_ref_cell = &*global_state;\n 1010 | let (pre_state, pre_ptr) = {\n 1011 | let mut global_state_ref_mut = global_state_ref_cell.borrow_mut();\n 1012 | let global_state = &mut *global_state_ref_mut;\n 1013 | let panic_guard = OuterCallGuard {\n 1014 | writer: &mut global_state.log_writer,\n 1015 | log_root: &global_state.log_stack,\n 1016 | };\n 1017 | let mut logger = RefMut::map(global_state.log_stack.borrow_mut(), |log_stack| {\n 1018 | log_stack.enter()\n 1019 | });\n 1020 | logger.name = name.clone();\n 1021 | let delayed_state = match global_state.delayed_state {\n 1022 | LateInit::Success(ref mut delayed_state) => delayed_state,\n 1023 | // There's no libcuda to load, so we might as well panic\n 1024 | LateInit::Error => panic!(),\n 1025 | LateInit::Unitialized => {\n 1026 | global_state.delayed_state =\n 1027 | GlobalDelayedState::new(panic_guard.writer, &mut logger);\n 1028 | // `global_state.delayed_state` could be LateInit::Error,\n 1029 | // we can crash in this case since there's no libcuda\n 1030 | global_state.delayed_state.as_mut().unwrap()\n 1031 | }\n 1032 | };\n 1033 | let fn_ptr = pre_call(delayed_state, &mut logger);\n 1034 | match fn_ptr {\n 1035 | Some(fn_ptr) => {\n 1036 | mem::forget(panic_guard);\n 1037 | fn_ptr\n 1038 | }\n 1039 | None => {\n 1040 | logger.log(ErrorEntry::FunctionNotFound(name));\n 1041 | return internal_error;\n 1042 | }\n 1043 | }\n 1044 | };\n 1045 | let panic_guard = InnerCallGuard(global_state_ref_cell);\n 1046 | let inner_result = inner_call(pre_ptr);\n 1047 | let global_state = &mut *global_state_ref_cell.borrow_mut();\n 1048 | mem::forget(panic_guard);\n 1049 | let _drop_guard = OuterCallGuard {\n 1050 | writer: &mut global_state.log_writer,\n 1051 | log_root: &global_state.log_stack,\n 1052 | };\n 1053 | let mut logger = RefMut::map(global_state.log_stack.borrow_mut(), |log_stack| {\n 1054 | log_stack.resume()\n 1055 | });\n 1056 | logger.args = args.map(|args| args());\n 1057 | logger.output = Some(format_status(inner_result));\n 1058 | post_call(\n 1059 | global_state.delayed_state.as_mut().unwrap(),\n 1060 | &mut logger,\n 1061 | pre_state,\n 1062 | pre_ptr,\n 1063 | inner_result,\n 1064 | );\n 1065 | inner_result\n 1066 | }","lines":60,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda_trace/src/lib.rs#L1007-L1066"}]},
{"id":"LOGIC_3","type":"logic_error","type_label":"Logic Error","severity":"high","difficulty":"medium","description":"In `InsertMemSSAVisitor::visit_mov` in `insert_explicit_load_store.rs`, whenever a `mov` reads from a variable that was remapped from `.param`/kernel-argument space to `.local` (`RemapAction::LDStSpaceChange`), the pass synthesizes an implicit `Cvta` (`ExplicitToGeneric`) instruction with `size: 64` hardcoded unconditionally. This pass (`insert_explicit_load_store::run`) is invoked unconditionally in `ptx/src/pass/mod.rs` for both 32-bit and 64-bit modules (it runs before the `if ast.address_size == 32 { convert_32bit_to_64bit::run(...) }` branch, and nothing downstream ever rewrites the `size` field of previously-inserted `Cvta` instructions). For a 32-bit PTX module (`ast.address_size == 32`), pointers/addresses are meant to be 32-bit, matching the parser's own `cvta.to.space.u32`/`cvta.space.u32` encoding (`ptx_parser/src/lib.rs` around line 2764-2801, where `size` comes from the actual `.u32`/`.u64` modifier on the source instruction). By hardcoding `size: 64` here, `emit_cvta` in `ptx/src/pass/llvm/emit.rs` builds a 64-bit `LLVMBuildPtrToInt` result for what should be a 32-bit register (whose LLVM type was registered based on the 32-bit PTX type), producing a width-mismatched value that either fails LLVM module verification (`debug_assertions`) or silently miscompiles addressing for `.local`/`.param`-turned-`.local` variables in 32-bit kernels.","fix":"Thread the module's address size (like `insert_implicit_conversions::run` already receives `is_32bit: bool`) into `insert_explicit_load_store::run` / `InsertMemSSAVisitor`, and set `size: if is_32bit { 32 } else { 64 }` instead of the hardcoded `size: 64` in the synthesized `ast::CvtaDetails`.","locations":[{"ref":"ptx/src/pass/insert_explicit_load_store.rs:302-332","code":" 302 | fn visit_mov(\n 303 | &mut self,\n 304 | data: ptx_parser::MovDetails,\n 305 | mut arguments: ptx_parser::MovArgs<SpirvWord>,\n 306 | ) -> ast::Instruction<SpirvWord> {\n 307 | if let Some(remap) = self.variables.get(&arguments.src) {\n 308 | match remap {\n 309 | RemapAction::PreLdPostSt { .. } => {}\n 310 | RemapAction::LDStSpaceChange {\n 311 | name, new_space, ..\n 312 | } => {\n 313 | let generic_var = self\n 314 | .resolver\n 315 | .register_unnamed(Some((data.typ.clone(), ast::StateSpace::Reg)));\n 316 | self.pre.push(ast::Instruction::Cvta {\n 317 | data: ast::CvtaDetails {\n 318 | state_space: *new_space,\n 319 | direction: ast::CvtaDirection::ExplicitToGeneric,\n 320 | size: 64,\n 321 | },\n 322 | arguments: ast::CvtaArgs {\n 323 | dst: generic_var,\n 324 | src: *name,\n 325 | },\n 326 | });\n 327 | arguments.src = generic_var;\n 328 | }\n 329 | }\n 330 | }\n 331 | ast::Instruction::Mov { data, arguments }\n 332 | }","lines":31,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx/src/pass/insert_explicit_load_store.rs#L302-L332"}]},
{"id":"LOGIC_4","type":"logic_error","type_label":"Logic Error","severity":"high","difficulty":"medium","description":"`resolve_function_pointers::run` (via `run_directive`) walks `directives` in a single forward pass, inserting each non-kernel `method.name` into `functions` only after it is reached, and immediately calling `run_method` on that same directive using the `functions` set as it stands at that point. `run_statement` recognizes a function-pointer `Mov` only if `functions.contains(&src_reg)` is already true. Because the set is populated incrementally in file order, any `mov.u64 %rd, some_func;` that takes the address of a `.func` whose `Directive2::Method` directive appears *later* in `directives` than the directive currently being processed will not be recognized as a function pointer (the name is not yet in `functions`), so the statement is left as a plain `Statement::Instruction(ast::Instruction::Mov ...)` instead of being rewritten to `UnconditionalStatement::FunctionPointer(FunctionPointerDetails { .. })`. Downstream passes (e.g. `remove_unreachable_basic_blocks` and LLVM emission in `ptx/src/pass/llvm/emit.rs`, which special-cases `Statement::FunctionPointer`) never see this case as a function pointer, so a forward reference to a function pointer silently miscompiles instead of being resolved or reported as an error.","fix":"Do two passes over `directives`: first collect every non-kernel `method.name` into `functions` (via `Directive2::Method` matching without recursing into bodies), then do a second pass calling `run_method`/`run_statement` with the now-complete `functions` set. For example, change `run` to first `directives.iter().for_each(|d| if let Directive2::Method(m) = d { if !m.is_kernel() { functions.insert(m.name); } })` before mapping `run_directive` over the directives (and drop the `functions.insert` call currently inside `run_directive`).","locations":[{"ref":"ptx/src/pass/resolve_function_pointers.rs:5-28","code":" 5 | pub(crate) fn run<'input>(\n 6 | directives: Vec<UnconditionalDirective>,\n 7 | ) -> Result<Vec<UnconditionalDirective>, TranslateError> {\n 8 | let mut functions = FxHashSet::default();\n 9 | directives\n 10 | .into_iter()\n 11 | .map(|directive| run_directive(&mut functions, directive))\n 12 | .collect::<Result<Vec<_>, _>>()\n 13 | }\n 14 | \n 15 | fn run_directive<'input>(\n 16 | functions: &mut FxHashSet<SpirvWord>,\n 17 | directive: UnconditionalDirective,\n 18 | ) -> Result<UnconditionalDirective, TranslateError> {\n 19 | Ok(match directive {\n 20 | var @ Directive2::Variable(..) => var,\n 21 | Directive2::Method(method) => {\n 22 | if !method.is_kernel() {\n 23 | functions.insert(method.name);\n 24 | }\n 25 | Directive2::Method(run_method(functions, method)?)\n 26 | }\n 27 | })\n 28 | }","lines":24,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx/src/pass/resolve_function_pointers.rs#L5-L28"}]},
{"id":"LOGIC_5","type":"logic_error","type_label":"Logic Error","severity":"high","difficulty":"medium","description":"`context::set_current` (which implements both `cuCtxSetCurrent` and, via `push_current`, `cuCtxPushCurrent`) always *pushes* a new `(CUcontext, hipDevice_t)` entry onto the per-thread `STACK` when `raw_ctx` is non-null, instead of replacing the top of the stack as the real CUDA driver does for `cuCtxSetCurrent`. Per NVIDIA's documented behavior, `cuCtxSetCurrent` replaces the top of the context stack; it does not grow it. Because ZLUDA routes `cuCtxSetCurrent` straight to `set_current` (same function used for `cuCtxPushCurrent`), every call to `cuCtxSetCurrent(ctx)` adds a new frame to the thread-local `STACK` (a `Vec<(CUcontext, hipDevice_t)>`). A caller that legitimately calls `cuCtxSetCurrent` repeatedly (a very common pattern for binding the current context before an operation) causes `STACK` to grow without bound - an unbounded per-thread memory leak - and desynchronizes `cuCtxPopCurrent`/`cuCtxGetCurrent` from actual CUDA semantics: a later `cuCtxPopCurrent` reveals a stale context that a real CUDA driver would never have retained, since `SetCurrent` should have simply overwritten it.","fix":"Split `set_current` into two distinct behaviors: keep the current push semantics for `push_current`/`cuCtxPushCurrent`, and give `cuCtxSetCurrent` its own implementation that replaces the top-of-stack entry (`stack.last_mut()` update, or pop-then-push a single frame) instead of unconditionally calling `stack.push(...)`, matching the documented `cuCtxSetCurrent` contract of replacing rather than stacking the current context.","locations":[{"ref":"zluda/src/impl/context.rs:191-222","code":" 191 | pub(crate) fn set_current(raw_ctx: CUcontext) -> CUresult {\n 192 | let new_device = if raw_ctx.0 == ptr::null_mut() {\n 193 | STACK.with(|stack| {\n 194 | let mut stack = stack.borrow_mut();\n 195 | if let Some((_, old_device)) = stack.pop() {\n 196 | if let Some((_, new_device)) = stack.last() {\n 197 | if old_device != *new_device {\n 198 | return Some(*new_device);\n 199 | }\n 200 | }\n 201 | }\n 202 | None\n 203 | })\n 204 | } else {\n 205 | let ctx: &Context = FromCuda::<_, CUerror>::from_cuda(&raw_ctx)?;\n 206 | let device = ctx.device;\n 207 | STACK.with(move |stack| {\n 208 | let mut stack = stack.borrow_mut();\n 209 | let last_device = stack.last().map(|(_, dev)| *dev);\n 210 | stack.push((raw_ctx, device));\n 211 | match last_device {\n 212 | None => Some(device),\n 213 | Some(last_device) if last_device != device => Some(device),\n 214 | _ => None,\n 215 | }\n 216 | })\n 217 | };\n 218 | if let Some(dev) = new_device {\n 219 | unsafe { hipSetDevice(dev)? };\n 220 | }\n 221 | Ok(())\n 222 | }","lines":32,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda/src/impl/context.rs#L191-L222"}]},
{"id":"LOGIC_6","type":"logic_error","type_label":"Logic Error","severity":"medium","difficulty":"easy","description":"`impl std::fmt::Display for LinkingDirective` in `ptx_parser/src/ast.rs` uses an `if ... else if ... else if ...` chain to decide which linking keyword to print (`.extern`, `.visible`, `.weak`), even though `LinkingDirective` is a `bitflags` type and the parser (`linking_directive` in `ptx_parser/src/lib.rs`, which folds tokens with `x | y`) can legitimately produce values with more than one bit set, e.g. `.weak .visible .func foo(...)` yields `LinkingDirective::WEAK | LinkingDirective::VISIBLE`. Because the branches are `else if`, only the first matching flag (`EXTERN`, then `VISIBLE`, then `WEAK`) is ever pushed into `directives`, so a value with both `VISIBLE` and `WEAK` set prints only `.visible` and silently drops `.weak`. This function is exercised by `ptx/src/pass/test/mod.rs` (`format!(\"{} \", function.linkage)`), so any dumped/round-tripped PTX for a weak-and-visible function loses the `.weak` modifier.","fix":"Replace the `else if` chain with independent `if` checks so every set flag is emitted, e.g.:\n```rust\nif self.contains(LinkingDirective::EXTERN) {\n directives.push(\".extern\");\n}\nif self.contains(LinkingDirective::VISIBLE) {\n directives.push(\".visible\");\n}\nif self.contains(LinkingDirective::WEAK) {\n directives.push(\".weak\");\n}\n```","locations":[{"ref":"ptx_parser/src/ast.rs:1971-1983","code":" 1971 | impl std::fmt::Display for LinkingDirective {\n 1972 | fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {\n 1973 | let mut directives = vec![];\n 1974 | if self.contains(LinkingDirective::EXTERN) {\n 1975 | directives.push(\".extern\");\n 1976 | } else if self.contains(LinkingDirective::VISIBLE) {\n 1977 | directives.push(\".visible\");\n 1978 | } else if self.contains(LinkingDirective::WEAK) {\n 1979 | directives.push(\".weak\");\n 1980 | }\n 1981 | write!(f, \"{}\", directives.join(\" \"))\n 1982 | }\n 1983 | }","lines":13,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx_parser/src/ast.rs#L1971-L1983"}]},
{"id":"LOGIC_7","type":"logic_error","type_label":"Logic Error","severity":"medium","difficulty":"hard","description":"`ContextState::reset` is invoked from inside `Context::with_state_mut`'s closure, i.e. while the `Context`'s `state: Mutex<ContextState>` is already locked by the current thread. Inside `reset`, for every stored context-local-storage entry it directly invokes the user-supplied destructor callback `_cb(data.handle, ...)` (registered via the dark-API `context_local_storage_put`) while still holding that same lock. Because `std::sync::Mutex` is not reentrant, if the callback calls back into the CUDA driver on the same context - for example via `cuCtxSetFlags`, `cuCtxGetContextLocalStorage`/`cuCtxSetContextLocalStorage` (`context_local_storage_get`/`_put`/`_delete` in `driver.rs`), or `cuDevicePrimaryCtxGetState` on the same primary context - the thread deadlocks trying to re-lock `self.state`. The code even has a `TODO` acknowledging this exact risk but does not prevent it.","fix":"Collect the `(handle, key, value, reset_cb)` tuples into a temporary `Vec` first, drop the `state` lock (return from the locked closure with the pending callbacks), and invoke the callbacks afterwards, outside of `with_state_mut`, so a callback that re-enters the driver on the same context cannot deadlock on `self.state`.","locations":[{"ref":"zluda/src/impl/context.rs:41-61","code":" 41 | pub(crate) fn reset(&mut self) -> CUresult {\n 42 | for (key, data) in self.storage.iter_mut() {\n 43 | if let Some(_cb) = data.reset_cb {\n 44 | // TODO: check that these callbacks do not call into the CUDA driver\n 45 | // since this could result in a recursive mutex lock.\n 46 | _cb(data.handle, *key as *mut c_void, data.value as *mut c_void);\n 47 | }\n 48 | }\n 49 | self.ref_count = 0;\n 50 | self.flags = 0;\n 51 | // drop all modules and return first error if any\n 52 | let result = self.modules.drain().fold(Ok(()), |res: CUresult, hmod| {\n 53 | match (res, zluda_common::drop_checked::<module::Module>(hmod)) {\n 54 | (Err(e), _) => Err(e),\n 55 | (_, Err(e)) => Err(e),\n 56 | _ => Ok(()),\n 57 | }\n 58 | });\n 59 | self.storage.clear();\n 60 | result\n 61 | }","lines":21,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda/src/impl/context.rs#L41-L61"}]},
{"id":"LOGIC_8","type":"logic_error","type_label":"Logic Error","severity":"low","difficulty":"easy","description":"`FUNC(nanosleep_u32)` (the device-side implementation of PTX `nanosleep.u32`) computes `uint64_t cycles_per_ns = ATTR(CLOCK_RATE) / 1000000;` before multiplying by `nanoseconds`. Since `ATTR(CLOCK_RATE)` is the GPU clock rate in kHz, any device clocked below 1 GHz (`CLOCK_RATE < 1000000`) makes `cycles_per_ns` truncate to `0`, so `cycles = nanoseconds * cycles_per_ns` is always `0` regardless of the requested sleep duration, and `nanosleep.u32` becomes a near no-op (only the fixed `+= 63` / `>> 6` rounding contributes a fixed few cycles) instead of sleeping for the requested time. Even for clock rates above 1 GHz that aren't exact multiples of 1000000 kHz, dividing before multiplying systematically under-counts cycles versus computing `nanoseconds * ATTR(CLOCK_RATE) / 1000000`.","fix":"Compute cycles with the multiplication before the division to preserve precision, e.g. `uint64_t cycles = (uint64_t)nanoseconds * ATTR(CLOCK_RATE) / 1000000;`, avoiding the intermediate truncated `cycles_per_ns`.","locations":[{"ref":"ptx/lib/zluda_ptx_impl.cpp:343-364","code":" 343 | void FUNC(nanosleep_u32)(uint32_t nanoseconds)\n 344 | {\n 345 | // clock_rate is in kHz\n 346 | uint64_t cycles_per_ns = ATTR(CLOCK_RATE) / 1000000;\n 347 | uint64_t cycles = nanoseconds * cycles_per_ns;\n 348 | // Avoid small sleep values resulting in s_sleep 0\n 349 | cycles += 63;\n 350 | // s_sleep N sleeps for 64 * N cycles\n 351 | uint64_t sleep_amount = cycles / 64;\n 352 | \n 353 | // The argument to s_sleep must be a constant\n 354 | for (size_t i = 0; i < sleep_amount >> 4; i++)\n 355 | __builtin_amdgcn_s_sleep(16);\n 356 | if (sleep_amount & 8U)\n 357 | __builtin_amdgcn_s_sleep(8);\n 358 | if (sleep_amount & 4U)\n 359 | __builtin_amdgcn_s_sleep(4);\n 360 | if (sleep_amount & 2U)\n 361 | __builtin_amdgcn_s_sleep(2);\n 362 | if (sleep_amount & 1U)\n 363 | __builtin_amdgcn_s_sleep(1);\n 364 | }","lines":22,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx/lib/zluda_ptx_impl.cpp#L343-L364"}]},
{"id":"LOGIC_9","type":"logic_error","type_label":"Logic Error","severity":"low","difficulty":"medium","description":"`ref_set_address_v2` in `zluda/src/impl/tex.rs` takes an output parameter `byte_offset: *mut usize` that callers rely on being written (matching CUDA's `cuTexRefSetAddress`, which always writes `*ByteOffset`). When `dev_ptr` is null, the function takes an early-return shortcut calling `hipUnbindTexture(texref)` instead of `hipTexRefSetAddress`, and never writes to `*byte_offset`. Any caller that passes an uninitialized `usize` for `byte_offset` (as is typical, since the driver is expected to fill it in) and then uses the returned offset for subsequent address computation will silently use a stale/uninitialized value instead of the `0` (or otherwise correct) offset a real driver would report for this case.","fix":"In the `dev_ptr.0.is_null()` branch, write `*byte_offset = 0` (or the appropriate value) before calling `hipUnbindTexture(texref)`, so the out parameter is always initialized the same way it is on the normal `hipTexRefSetAddress` path.","locations":[{"ref":"zluda/src/impl/tex.rs:98-108","code":" 98 | pub(crate) unsafe fn ref_set_address_v2(\n 99 | byte_offset: *mut usize,\n 100 | texref: *mut textureReference,\n 101 | dev_ptr: hipDeviceptr_t,\n 102 | bytes: usize,\n 103 | ) -> hipError_t {\n 104 | if dev_ptr.0.is_null() {\n 105 | return hipUnbindTexture(texref);\n 106 | }\n 107 | hipTexRefSetAddress(byte_offset, texref, dev_ptr, bytes)\n 108 | }","lines":11,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda/src/impl/tex.rs#L98-L108"}]},
{"id":"CPY_1","type":"copy_paste","type_label":"Copy Paste","severity":"high","difficulty":"easy","description":"In `occupancy_max_active_blocks_per_multiprocessor_with_flags`, after calling `hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` to compute `num_blocks` (the number of resident *blocks* per multiprocessor), the code clamps it with `*num_blocks = (*num_blocks).max(1).min(max_block_size);` where `max_block_size` was obtained from `hipFuncGetAttribute(..., HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, ...)`, i.e. the maximum *threads per block*. This mixes two unrelated units: clamping a block count against a thread-per-block limit is meaningless and appears copy-pasted from the neighboring `occupancy_max_potential_block_size` function (lines 671-694), where clamping `block_size` (a thread count) against `max_block_size` is correct. Additionally, the `.max(1)` forces the reported occupancy to be at least 1 block even when HIP legitimately determines that 0 blocks can be resident (e.g. the kernel's shared memory or register usage is too large for the given `block_size`), silently misreporting an infeasible launch configuration as feasible to any caller of `cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` doing occupancy-based tuning.","fix":"Remove the incorrect `.max(1).min(max_block_size)` clamp on `num_blocks` (or replace it with a real bound, if any, expressed in blocks-per-multiprocessor, not threads-per-block). Return the value from `hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` as-is, including the legitimate `0` case.","locations":[{"ref":"zluda/src/impl/driver.rs:647-669","code":" 647 | pub(crate) unsafe fn occupancy_max_active_blocks_per_multiprocessor_with_flags(\n 648 | num_blocks: &mut ::core::ffi::c_int,\n 649 | func: &function::Function,\n 650 | block_size: ::core::ffi::c_int,\n 651 | dynamic_smem_size: usize,\n 652 | flags: ::core::ffi::c_uint,\n 653 | ) -> hipError_t {\n 654 | hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(\n 655 | num_blocks,\n 656 | func.base,\n 657 | block_size,\n 658 | dynamic_smem_size,\n 659 | flags,\n 660 | )?;\n 661 | let mut max_block_size = 0;\n 662 | hipFuncGetAttribute(\n 663 | &mut max_block_size,\n 664 | hipFunction_attribute::HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK,\n 665 | func.base,\n 666 | )?;\n 667 | *num_blocks = (*num_blocks).max(1).min(max_block_size);\n 668 | Ok(())\n 669 | }","lines":23,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda/src/impl/driver.rs#L647-L669"}]},
{"id":"CPY_2","type":"copy_paste","type_label":"Copy Paste","severity":"high","difficulty":"easy","description":"In `apply_detours` (`zluda_redirect/src/lib.rs`), the block that calls `DetourAttach(std::ptr::from_mut(&mut get_cubin_size).cast(), get_ptx_size as _)` (lines 421-427) is immediately followed by an identical, duplicated block doing the exact same `DetourAttach` call again (lines 428-434). This is clearly a copy-paste error: the intent was almost certainly to attach a third pair (likely `get_cubin`/`get_ptx` was meant to be attached once and something else a second time, or the duplicate should have been removed). Calling `DetourAttach` twice on the same trampoline pointer without an intervening `DetourDetach` fails in Microsoft Detours (returns a non-`NO_ERROR` code), so the second, duplicated call causes `apply_detours` to return `false`. In `detour_nvrtc`, `apply_detours` returning `false` triggers `DetourTransactionAbort()`, which rolls back *all* detours attached so far in the transaction - including the legitimate `get_cubin` -> `get_ptx` attach done earlier at line 418. The net effect is that `nvrtcGetCUBIN`/`nvrtcGetCUBINSize`/`nvrtcCompileProgram` redirection silently never actually applies for any loaded `nvrtc*.dll`, breaking the core mechanism this file exists for (making apps that call `nvrtcGetCUBIN` receive PTX data instead), with no user-visible error.","fix":"Remove the duplicated `DetourAttach` block at lines 428-434, or replace it with the actually-intended additional attach (e.g. attaching `get_ptx`/`get_ptx_size` themselves, or another entry point) if a third redirection was meant to be installed there.","locations":[{"ref":"zluda_redirect/src/lib.rs:421-434","code":" 421 | if DetourAttach(\n 422 | std::ptr::from_mut(&mut get_cubin_size).cast(),\n 423 | get_ptx_size as _,\n 424 | ) != NO_ERROR as i32\n 425 | {\n 426 | return false;\n 427 | }\n 428 | if DetourAttach(\n 429 | std::ptr::from_mut(&mut get_cubin_size).cast(),\n 430 | get_ptx_size as _,\n 431 | ) != NO_ERROR as i32\n 432 | {\n 433 | return false;\n 434 | }","lines":14,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/zluda_redirect/src/lib.rs#L421-L434"}]},
{"id":"ERR_1","type":"incorrect_error_handling","type_label":"Incorrect Error Handling","severity":"medium","difficulty":"hard","description":"`take_till_end_of_statement` in `ptx_parser/src/lib.rs` is the statement-level error-recovery scanner used by `statement`'s `with_recovery` to find the end of a malformed/unrecognized statement. It tracks brace `depth` correctly for `Token::LBrace`/`Token::RBrace` (only erroring on an `RBrace` when `depth == 0`), but the `Token::Semicolon | Token::Colon` arm terminates the scan unconditionally, ignoring the current `depth`. So if a malformed statement contains a `{ ... }` group with a semicolon inside it (e.g. a stray/garbled `unknown_op { a; b };`), recovery stops at the *inner* semicolon instead of the statement-ending one, leaving the rest (including the now-unmatched closing `}`) in the token stream. This was confirmed empirically with the input `unknown_op { a; b };` inside a function body: `parse_module_checked` returns two errors, `UnrecognizedStatement(\"unknown_op { a;\")` (truncated at the inner semicolon, as predicted) and a second, uninformative `Parser(ContextError { context: [], cause: None })`, i.e. the leftover unmatched `}` breaks the surrounding grammar and produces an opaque top-level parse failure in addition to the mis-scoped recovery message, even though only one statement in the source was malformed.","fix":"Guard the `Token::Semicolon | Token::Colon` arm with the same `depth == 0` check used for `Token::RBrace`, e.g. only treat a semicolon/colon as the end of the malformed statement when `depth == 0`; otherwise fall through to the default `_ => {}` arm so the scan continues past semicolons/colons that are nested inside an unbalanced `{ ... }` group.","locations":[{"ref":"ptx_parser/src/lib.rs:882-919","code":" 882 | fn take_till_end_of_statement<\n 883 | 'a,\n 884 | I: Stream<Token = (Token<'a>, std::ops::Range<usize>)>,\n 885 | E: ParserError<I>,\n 886 | >() -> impl Parser<I, <I as Stream>::Slice, E> {\n 887 | trace(\"take_till_end_of_statement\", move |stream: &mut I| {\n 888 | let mut depth = 0;\n 889 | \n 890 | let mut iterator = stream.iter_offsets().peekable();\n 891 | while let Some((current_offset, (token, _))) = iterator.next() {\n 892 | match token {\n 893 | Token::LBrace => {\n 894 | depth += 1;\n 895 | }\n 896 | Token::RBrace => {\n 897 | if depth == 0 {\n 898 | return Err(ErrMode::from_error_kind(\n 899 | stream,\n 900 | winnow::error::ErrorKind::Token,\n 901 | ));\n 902 | }\n 903 | depth -= 1;\n 904 | }\n 905 | Token::Semicolon | Token::Colon => {\n 906 | let offset = if let Some((next_offset, _)) = iterator.peek() {\n 907 | *next_offset\n 908 | } else {\n 909 | current_offset\n 910 | };\n 911 | return Ok(stream.next_slice(offset));\n 912 | }\n 913 | _ => {}\n 914 | }\n 915 | }\n 916 | \n 917 | Err(ParserError::from_error_kind(stream, ErrorKind::Eof))\n 918 | })\n 919 | }","lines":38,"link":"https://github.com/vosen/ZLUDA/blob/9c8b43f242985150f86a7f485218b7b82c3e96ca/ptx_parser/src/lib.rs#L882-L919"}]}
]</script>
<script type="application/json" id="PI">{"name":"ZLUDA","git_url":"https://github.com/vosen/ZLUDA.git","repo_url":"https://github.com/vosen/ZLUDA","commit":"9c8b43f242985150f86a7f485218b7b82c3e96ca","commit_short":"9c8b43f242","commit_url":"https://github.com/vosen/ZLUDA/commit/9c8b43f242985150f86a7f485218b7b82c3e96ca","branch":"master","generated":"2026-09-14","description":"CUDA on non-NVIDIA GPUs"}</script>
<div class="app">
<div class="top">
<div>
<h1>Code Audit Report — ZLUDA</h1>
<div class="meta">Generated 2026-09-14 — <span id="stats"></span></div>
</div>
<div class="top-right">
<div class="search-wrap">
<input type="search" id="search" placeholder="Search findings…" autocomplete="off">
<button id="search-scope" class="action-btn" type="button" title="Toggle: search everywhere vs. only in file/folder paths">📁 Paths only</button>
</div>
<label class="group-label">Group by
<select id="groupby">
<option value="type">Type</option>
<option value="severity">Severity</option>
<option value="file">File</option>
<option value="folder">Folder</option>
<option value="none">None (flat)</option>
</select>
</label>
</div>
</div>
<div class="project-info"><div class="pi-name">ZLUDA</div><div class="pi-desc">CUDA on non-NVIDIA GPUs</div><div class="pi-row"><a class="pi-repo-btn" href="https://github.com/vosen/ZLUDA" target="_blank" rel="noopener">🔗 https://github.com/vosen/ZLUDA</a><span><span class="pi-k">branch</span> <code>master</code></span><span><span class="pi-k">commit</span> <a href="https://github.com/vosen/ZLUDA/commit/9c8b43f242985150f86a7f485218b7b82c3e96ca" target="_blank" rel="noopener"><code>9c8b43f242</code></a></span></div></div>
<div class="filterbar" id="filterbar"></div>
<div id="progress-wrap">
<div id="progress-bar"><div id="progress-fill"></div></div>
<span id="progress-text"></span>
<button class="action-btn" data-action="collapse-all">⊟ Collapse all</button>
<button class="action-btn" data-action="expand-all">⊞ Expand all</button>
<button class="action-btn" data-action="download-md">⬇ MD</button>
<button class="action-btn" data-action="download-md-short">⬇ MD Short</button>
<button class="action-btn" data-action="download-toml">⬇ TOML</button>
</div>
<div id="hidden-bar"></div>
<div id="content"></div>
</div>
<script>
(function(){
'use strict';
// Data
const FINDINGS = JSON.parse(document.getElementById('D').textContent);
const PROJECT_INFO = JSON.parse(document.getElementById('PI').textContent);
const SEVS = ['critical','high','medium','low','negligible'];
const STATUSES = ['open','fixed','wont-fix','false-positive'];
const ALL_TYPES = [...new Set(FINDINGS.map(f => f.type))].sort(cmpTypes);
// Build lookup map for O(1) access
const FINDING_BY_ID = new Map(FINDINGS.map(f => [f.id, f]));
// Only severities actually present in the report — the filter bar hides the rest.
const PRESENT_SEVS = SEVS.filter(s => FINDINGS.some(f => (f.severity || 'medium') === s));
// Whether any finding is flagged easy — the "Quick wins" filter row is hidden otherwise.
const DIFFICULTY_ORDER = ['very_easy','easy','medium','hard'];
const DIFFICULTY_LBL = {very_easy:'⚡ Very Easy', easy:'Easy', medium:'Medium', hard:'Hard'};
const DIFFICULTY_CLS = {very_easy:'diff-very-easy', easy:'diff-easy', medium:'diff-medium', hard:'diff-hard'};
const PRESENT_DIFFS = DIFFICULTY_ORDER.filter(d => FINDINGS.some(f => f.difficulty === d));
const HAS_DIFFICULTY = PRESENT_DIFFS.length > 0;
const STORE_KEY = 'audit-status:' + location.href;
const STATUS_CLS = {open:'st-open', fixed:'st-fixed', 'wont-fix':'st-wont', 'false-positive':'st-fp'};
const STATUS_LBL = {open:'● open', fixed:'✓ fixed', 'wont-fix':"— won't fix", 'false-positive':'✗ false pos.'};
const STATUS_NEXT = {open:'fixed', fixed:'wont-fix', 'wont-fix':'false-positive', 'false-positive':'open'};
// State
const S = {
search: '', searchScope: 'all', groupBy: 'type',
hiddenSev: new Set(), hiddenType: new Set(),
hiddenStatus: new Set(), hiddenDiff: new Set(), hiddenGroup: new Set(),
collapsed: new Set(),
};
const statuses = {};
// Persistence
function loadStatuses() {
try { Object.assign(statuses, JSON.parse(localStorage.getItem(STORE_KEY) || '{}')); } catch {}
}
function saveStatuses() {
try { localStorage.setItem(STORE_KEY, JSON.stringify(statuses)); } catch {}
}
function statusOf(id) { return statuses[id] || 'open'; }
// URL hash — debounced so rapid filter changes don't thrash history
let _hashTimer = 0;
function syncHash() {
clearTimeout(_hashTimer);
_hashTimer = setTimeout(() => {
const p = new URLSearchParams();
if (S.groupBy !== 'type') p.set('g', S.groupBy);
if (S.hiddenSev.size) p.set('hs', [...S.hiddenSev].join(','));
if (S.hiddenType.size) p.set('ht', [...S.hiddenType].join(','));
if (S.hiddenStatus.size) p.set('hst', [...S.hiddenStatus].join(','));
if (S.hiddenDiff.size) p.set('hd', [...S.hiddenDiff].join(','));
if (S.search) p.set('q', S.search);
if (S.searchScope !== 'all') p.set('qs', S.searchScope);
const str = p.toString();
history.replaceState(null, '', str ? '#' + str : location.pathname + location.search);
}, 300);
}
function loadHash() {
const p = new URLSearchParams(location.hash.slice(1));
if (p.has('g')) { S.groupBy = p.get('g'); document.getElementById('groupby').value = S.groupBy; }
if (p.has('q')) { S.search = p.get('q'); document.getElementById('search').value = S.search; }
if (p.has('qs')) { S.searchScope = p.get('qs'); }
document.getElementById('search-scope').classList.toggle('active', S.searchScope === 'locations');
if (p.has('hs')) p.get('hs') .split(',').filter(Boolean).forEach(v => S.hiddenSev.add(v));
if (p.has('ht')) p.get('ht') .split(',').filter(Boolean).forEach(v => S.hiddenType.add(v));
if (p.has('hst')) p.get('hst').split(',').filter(Boolean).forEach(v => S.hiddenStatus.add(v));
if (p.has('hd')) p.get('hd') .split(',').filter(Boolean).forEach(v => S.hiddenDiff.add(v));
}
// Helpers
function esc(s) {
return String(s).replace(/&/g,'&').replace(/</g,'<').replace(/>/g,'>').replace(/"/g,'"');
}
function highlight(text, q) {
if (!q) return esc(text);
const re = new RegExp('(' + q.replace(/[.*+?^${}()|[\]\\]/g,'\\$&') + ')', 'gi');
return String(text).split(re).map((p, i) => i % 2 ? '<mark>' + esc(p) + '</mark>' : esc(p)).join('');
}
// Render text with ```fenced``` code blocks and `inline` backticks,
// also paragraph-splitting on blank lines. Output is HTML-safe.
function renderMd(text, q) {
const src = String(text == null ? '' : text);
const FENCE = /```[ \t]*([A-Za-z0-9_+-]*)[ \t]*\r?\n([\s\S]*?)```/g;
const parts = [];
let last = 0, m;
while ((m = FENCE.exec(src)) !== null) {
parts.push({kind: 'text', value: src.slice(last, m.index)});
parts.push({kind: 'code', value: m[2].replace(/\r?\n$/, '')});
last = m.index + m[0].length;
}
parts.push({kind: 'text', value: src.slice(last)});
function renderInline(s) {
const re = /`([^`\n]+)`/g;
let out = '', i = 0, mm;
while ((mm = re.exec(s)) !== null) {
out += highlight(s.slice(i, mm.index), q);
out += '<code class="md-inline">' + highlight(mm[1], q) + '</code>';
i = mm.index + mm[0].length;
}
out += highlight(s.slice(i), q);
return out;
}
function renderTextChunk(s) {
const paragraphs = s.split(/\r?\n[ \t]*\r?\n/);
const out = [];
for (const p of paragraphs) {
const trimmed = p.replace(/^\s+|\s+$/g, '');
if (!trimmed) continue;
out.push('<p>' + renderInline(trimmed).replace(/\r?\n/g, '<br>') + '</p>');
}
return out.join('');
}
let html = '';
for (const p of parts) {
if (p.kind === 'text') html += renderTextChunk(p.value);
else html += '<pre class="md-code">' + highlight(p.value, q) + '</pre>';
}
return html;
}
function sevCls(s) { return SEVS.includes(s) ? 'sev-'+s : 'sev-negligible'; }
function sevLbl(s) { return s ? s[0].toUpperCase()+s.slice(1) : 'Medium'; }
function typeLabel(t) { return t.replace(/_/g,' ').replace(/\b\w/g, c => c.toUpperCase()); }
function isCustom(t) { return String(t).toLowerCase() === 'custom'; }
function typeCls(t) { return 'type-badge' + (isCustom(t) ? ' custom-type' : ''); }
// 'custom' findings are the user's own externally-sourced ones; pin them first
// everywhere types are listed so they never blend into the audit's own categories.
function cmpTypes(a, b) {
if (isCustom(a) !== isCustom(b)) return isCustom(a) ? -1 : 1;
return String(a).localeCompare(String(b));
}
// First sentence of text (up to and including the first period).
function firstSentence(text) {
const m = String(text).match(/\.(\s|$)/);
return m ? text.slice(0, text.indexOf(m[0]) + 1) : text;
}
// Render a "**Label:** value" line. GitHub/CommonMark only recognizes a
// ``` fence when it starts at the beginning of a line, so if the value
// itself opens with a fence, put it on its own line instead of right
// after the inline label - otherwise the fence marker gets glued onto
// "**Label:** ```" and never renders as a code block.
function mdField(label, value) {
value = value || '';
return /^```/.test(value.trim())
? '**' + label + ':**\n\n' + value
: '**' + label + ':** ' + value;
}
// Render the "**Locations:**" block. When a location has a commit-pinned
// permalink (GitHub/GitLab), just paste the bare URL on its own line -
// GitHub/GitLab auto-expand a line-range blob link into an inline code
// preview, so re-pasting the snippet as a fenced block would duplicate it.
// Forges without a detected permalink shape (or missing commit/repo info)
// fall back to the fenced ref+snippet so the location is still readable.
function mdLocations(locations) {
if (!locations.length) return '';
let out = '**Locations:**\n\n';
for (const l of locations) {
if (l.link) out += l.link + '\n\n';
else if (l.code != null) out += '```\n' + l.ref + '\n' + l.code + '\n```\n\n';
else out += '- `' + l.ref + '` _(file not found)_\n\n';
}
return out;
}
// Build full MD report for in-browser download.
function buildMdFull() {
const SORD = ['critical','high','medium','low','negligible'];
const today = new Date().toISOString().slice(0, 10);
let out = '# Code Audit Report\n\n_Generated ' + today + '_\n\n';
const typeMap = new Map();
for (const f of FINDINGS) {
if (!typeMap.has(f.type)) typeMap.set(f.type, []);
typeMap.get(f.type).push(f);
}
const sortedTypes = [...typeMap.keys()].sort(cmpTypes);
out += '## Summary\n\n| Type | Severity | Count |\n|------|----------|-------|\n';
for (const t of sortedTypes) {
const items = typeMap.get(t);
const topSev = items.map(f => f.severity || 'medium').reduce((best, s) => {
const bi = SORD.indexOf(best), si = SORD.indexOf(s);
return (si >= 0 && (bi < 0 || si < bi)) ? s : best;
}, 'negligible');
out += '| ' + typeLabel(t) + ' | ' + topSev + ' | ' + items.length + ' |\n';
}
for (const t of sortedTypes) {
out += '\n## ' + typeLabel(t) + '\n\n';
for (const f of typeMap.get(t)) {
const sevStr = f.severity ? ' `' + f.severity.toUpperCase() + '`' : '';
out += '### ' + f.id + sevStr + '\n\n';
out += mdField('Description', f.description) + '\n\n';
out += mdLocations(f.locations);
out += mdField('Fix', f.fix) + '\n\n---\n\n';
}
}
return out;
}
// Build MD short checklist for in-browser download.
function buildMdShort() {
const lines = [];
const ids = [];
const sorted = [...FINDINGS].sort((a, b) => a.type.localeCompare(b.type));
for (const f of sorted) {
ids.push(f.id);
const desc = firstSentence(f.description || '');
const locs = f.locations.map(l => l.ref);
const locStr = locs.length ? ' - [' + locs.join(', ') + ']' : '';
const sevStr = f.severity ? ' [' + f.severity.toUpperCase() + ']' : '';
lines.push('- [ ] - ' + f.id + sevStr + ' - ' + desc + locStr);
}
lines.push('', '## IDs only', '');
ids.forEach(id => lines.push('- [ ] - ' + id));
return lines.join('\n') + '\n';
}
function downloadText(text, filename, mime) {
const a = document.createElement('a');
a.href = URL.createObjectURL(new Blob([text], {type: mime || 'text/plain'}));
a.download = filename;
a.click();
setTimeout(() => URL.revokeObjectURL(a.href), 1000);
}
function tomlStr(s) {
return '"' + String(s == null ? '' : s)
.replace(/\\/g, '\\\\')
.replace(/"/g, '\\"')
.replace(/\n/g, '\\n')
.replace(/\r/g, '\\r')
.replace(/\t/g, '\\t') + '"';
}
// Build a [project] header block from embedded PROJECT_INFO, so the exported
// findings.toml is self-contained (doesn't depend on a sibling project-info.toml
// that a second, separately-triggered download could silently fail to save).
// '' when there is no project info at all.
function buildProjectToml() {
const ORDER = ['name','description','git_url','repo_url','commit','commit_short','commit_url','branch','generated'];
const pi = PROJECT_INFO || {};
if (!Object.keys(pi).length) return '';
const lines = ['[project]'];
for (const k of ORDER) {
if (k in pi) lines.push(k + ' = ' + tomlStr(pi[k]));
}
return lines.join('\n') + '\n\n';
}
// Build TOML matching the original findings.toml format (re-parseable by toml_to_report.py).
function buildFindingsToml() {
function pad(key) { return (key + ' ').slice(0, 12); }
const chunks = [];
for (const f of FINDINGS) {
const lns = ['[[finding]]'];
// Pin the id so re-rendering this export - or filtering it down later - can't
// renumber findings that were already referenced by id elsewhere.
if (f.id && f.id !== '?') lns.push(pad('id') + '= ' + tomlStr(f.id));
lns.push(pad('type') + '= ' + tomlStr(f.type));
lns.push(pad('severity') + '= ' + tomlStr(f.severity || 'medium'));
if (f.difficulty) lns.push(pad('difficulty') + '= ' + tomlStr(f.difficulty));
lns.push(pad('description') + '= ' + tomlStr(f.description || ''));
const refs = f.locations.map(l => tomlStr(l.ref));
lns.push(pad('locations') + '= [' + refs.join(', ') + ']');
lns.push(pad('fix') + '= ' + tomlStr(f.fix || ''));
chunks.push(lns.join('\n'));
}
return buildProjectToml() + chunks.join('\n\n') + '\n';
}
// Build the GitHub-ready Markdown for one finding (mirrors findings.md).
function findingToMd(f) {
let out = '### ' + f.id + (f.severity ? ' `' + f.severity.toUpperCase() + '`' : '') + '\n\n';
out += mdField('Description', f.description) + '\n\n';
out += mdLocations(f.locations);
out += mdField('Fix', f.fix) + '\n';
return out;
}
function flashBtn(btn, msg) {
btn.textContent = msg;
setTimeout(() => { btn.textContent = '⧉ Copy'; }, 1200);
}
function copyText(text, btn) {
const ok = () => flashBtn(btn, '✓ Copied');
const no = () => flashBtn(btn, '✗ Failed');
const fallback = () => {
try {
const ta = document.createElement('textarea');
ta.value = text; ta.style.position = 'fixed'; ta.style.top = '-9999px';
document.body.appendChild(ta); ta.select();
const r = document.execCommand('copy');
document.body.removeChild(ta);
r ? ok() : no();
} catch { no(); }
};
if (navigator.clipboard && navigator.clipboard.writeText)
navigator.clipboard.writeText(text).then(ok, fallback);
else fallback();
}
// Filter logic
function getKey(f) {
if (S.groupBy === 'severity') return f.severity || 'medium';
if (S.groupBy === 'file') return f.locations.length ? f.locations[0].ref.replace(/:.*/,'') : '(no location)';
if (S.groupBy === 'folder') {
if (!f.locations.length) return '(no location)';
const path = f.locations[0].ref.replace(/:.*/,'');
const idx = path.lastIndexOf('/');
return idx === -1 ? '(root)' : path.slice(0, idx);
}
if (S.groupBy === 'none') return '';
return f.type_label;
}
// Pre-compute search haystacks per finding (cleared on full re-render).
// Two variants: everything (default), and file/folder paths only - for when a
// search term (e.g. a directory name like 'dlls/abc') also happens to appear
// inside description/fix prose and you only want path matches.
const _hay = new Map();
const _hayLoc = new Map();
function hayOf(f) {
if (!_hay.has(f.id)) {
_hay.set(f.id, [f.id, f.type_label, f.severity, f.description, f.fix,
f.difficulty ? (DIFFICULTY_LBL[f.difficulty] || f.difficulty) : '',
...f.locations.map(l => l.ref)].join(' ').toLowerCase());
}
return _hay.get(f.id);
}
function hayLocOf(f) {
if (!_hayLoc.has(f.id)) {
_hayLoc.set(f.id, f.locations.map(l => l.ref).join(' ').toLowerCase());
}
return _hayLoc.get(f.id);
}
function matches(f) {
if (S.hiddenDiff.size > 0 && f.difficulty && S.hiddenDiff.has(f.difficulty)) return false;
if (S.hiddenSev.has(f.severity || 'medium')) return false;
if (S.hiddenType.has(f.type)) return false;
const st = statusOf(f.id);
if (S.hiddenStatus.has(st)) return false;
if (S.search) {
const q = S.search.toLowerCase();
if (S.searchScope === 'locations') {
if (!hayLocOf(f).includes(q)) return false;
} else if (!hayOf(f).includes(q) && !st.includes(q)) {
return false;
}
}
return true;
}
function sortKeys(keys) {
if (S.groupBy === 'severity') {
return [...keys].sort((a, b) => {
const ai = SEVS.indexOf(a), bi = SEVS.indexOf(b);
return (ai < 0 ? 99 : ai) - (bi < 0 ? 99 : bi);
});
}
if (S.groupBy === 'type') return [...keys].sort(cmpTypes);
return [...keys].sort();
}
// Build body HTML for an expanded card (called lazily on first expand).
function renderCardBody(f, q) {
let locs = '';
for (const l of f.locations) {
const lineInfo = l.lines != null
? ' <span class="loc-lines">(' + l.lines + ' line' + (l.lines === 1 ? '' : 's') + ')</span>'
: '';
const refHtml = l.link
? '<a class="loc-link" href="' + esc(l.link) + '" target="_blank" rel="noopener" title="Open at this commit">' + highlight(l.ref, q) + ' ↗</a>'
: highlight(l.ref, q);
locs += '<div class="loc-block"><div class="loc-ref">' + refHtml + lineInfo + '</div>';
locs += l.code != null ? '<pre>' + esc(l.code) + '</pre>' : '<div class="not-found">(file not found)</div>';
locs += '</div>';
}
return (
'<div class="field-lbl">Description</div>' +
'<div class="desc">' + renderMd(f.description, q) + '</div>' +
(locs ? '<div class="field-lbl">Locations</div>' + locs : '') +
'<div class="field-lbl">Fix</div>' +
'<div class="fix-txt">' + renderMd(f.fix, q) + '</div>'
);
}
// Render card shell only — header + empty or populated body.
// When collapsed the body is empty (display:none) — avoids renderMd for 1800+ cards.
function renderCard(f, q) {
const sc = sevCls(f.severity || 'medium');
const sl = sevLbl(f.severity || 'medium');
const st = statusOf(f.id);
const sc2 = STATUS_CLS[st] || 'st-open';
const sl2 = STATUS_LBL[st] || '● open';
const isCollapsed = S.collapsed.has(f.id);
const isResolved = st !== 'open';
const locSummary = f.locations.length
? f.locations[0].ref + (f.locations.length > 1 ? ' \u2026' : '')
: '';
const bodyInner = isCollapsed ? '' : renderCardBody(f, q);
return (
'<div class="card' + (isResolved ? ' resolved' : '') + '" data-card-id="' + esc(f.id) + '">' +
'<div class="card-header" data-action="toggle-card" data-id="' + esc(f.id) + '">' +
'<span class="card-arrow">' + (isCollapsed ? '\u25b6' : '\u25bc') + '</span>' +
'<span class="card-id">' + esc(f.id) + '</span>' +
'<span class="badge ' + sc + '" style="cursor:default">' + esc(sl) + '</span>' +
'<span class="badge ' + typeCls(f.type) + '" style="cursor:default">' + esc(f.type_label) + '</span>' +
(f.difficulty ? '<span class="badge ' + (DIFFICULTY_CLS[f.difficulty]||'') + '" style="cursor:default" title="Difficulty: ' + esc(DIFFICULTY_LBL[f.difficulty]||f.difficulty) + '">' + esc(DIFFICULTY_LBL[f.difficulty]||f.difficulty) + '</span>' : '') +
(locSummary ? '<span class="card-loc" title="' + esc(locSummary) + '">' + esc(locSummary) + '</span>' : '') +
'<button class="status-btn ' + sc2 + '" data-action="cycle-status" data-id="' + esc(f.id) + '">' + sl2 + '</button>' +
'<button class="copy-btn" data-action="copy-md" data-id="' + esc(f.id) + '" title="Copy as Markdown for GitHub">\u29c9 Copy</button>' +
'</div>' +
'<div class="card-body" data-body-id="' + esc(f.id) + '"' + (isCollapsed ? ' style="display:none"' : '') + '>' +
bodyInner +
'</div></div>'
);
}
function initFilterBar() {
const fb = document.getElementById('filterbar');
let h = '<div class="filter-row"><span class="filter-lbl">Severity:</span>';
for (const s of PRESENT_SEVS)
h += '<button class="badge ' + sevCls(s) + '" data-sev="' + s + '">' + sevLbl(s) + '</button>';
h += '<button class="filter-toggle" data-toggle-all="sev"></button></div>';
h += '<div class="filter-row"><span class="filter-lbl">Type:</span>';
for (const t of ALL_TYPES)
h += '<button class="badge ' + typeCls(t) + '" data-type="' + esc(t) + '">' + esc(typeLabel(t)) + '</button>';
h += '<button class="filter-toggle" data-toggle-all="type"></button></div>';
h += '<div class="filter-row"><span class="filter-lbl">Status:</span>';
for (const st of STATUSES)
h += '<button class="badge ' + (STATUS_CLS[st]||'st-open') + '" data-status="' + st + '">' + (STATUS_LBL[st]||st) + '</button>';
h += '<button class="filter-toggle" data-toggle-all="status"></button></div>';
if (HAS_DIFFICULTY) {
h += '<div class="filter-row"><span class="filter-lbl">Difficulty:</span>';
for (const d of PRESENT_DIFFS)
h += '<button class="badge ' + (DIFFICULTY_CLS[d]||'') + '" data-diff="' + d + '">' + esc(DIFFICULTY_LBL[d]||d) + '</button>';
h += '<button class="filter-toggle" data-toggle-all="diff"></button></div>';
}
fb.innerHTML = h;
}
function updateToggles() {
for (const b of document.querySelectorAll('[data-sev]')) b.classList.toggle('off', S.hiddenSev.has(b.dataset.sev));
for (const b of document.querySelectorAll('[data-type]')) b.classList.toggle('off', S.hiddenType.has(b.dataset.type));
for (const b of document.querySelectorAll('[data-status]')) b.classList.toggle('off', S.hiddenStatus.has(b.dataset.status));
for (const b of document.querySelectorAll('[data-diff]')) b.classList.toggle('off', S.hiddenDiff.has(b.dataset.diff));
for (const b of document.querySelectorAll('[data-toggle-all]')) {
const k = b.dataset.toggleAll;
const set = k === 'sev' ? S.hiddenSev : k === 'type' ? S.hiddenType : k === 'status' ? S.hiddenStatus : k === 'diff' ? S.hiddenDiff : null;
if (set) b.textContent = set.size === 0 ? 'Hide all' : 'Show all';
}
}
function updateProgress() {
const total = FINDINGS.length;
const done = FINDINGS.filter(f => statusOf(f.id) !== 'open').length;
const pct = total ? Math.round(done / total * 100) : 0;
document.getElementById('progress-fill').style.width = pct + '%';
document.getElementById('progress-text').textContent = done + ' / ' + total + ' resolved (' + pct + '%)';
}
// Surgical toggle: flip one card without touching the rest of the DOM.
function toggleCard(id) {
const wasCollapsed = S.collapsed.has(id);
if (wasCollapsed) S.collapsed.delete(id); else S.collapsed.add(id);
const body = document.querySelector('[data-body-id="' + id + '"]');
const arrow = document.querySelector('[data-card-id="' + id + '"] .card-arrow');
if (!body) { render(); return; }
if (wasCollapsed) {
if (!body.firstChild) {
const f = FINDING_BY_ID.get(id);
if (f) body.innerHTML = renderCardBody(f, S.search);
}
body.style.display = '';
if (arrow) arrow.textContent = '\u25bc';
} else {
body.style.display = 'none';
if (arrow) arrow.textContent = '\u25b6';
}
}
// Surgical status update: update one card's button + resolved class.
function cycleStatus(id) {
statuses[id] = STATUS_NEXT[statusOf(id)] || 'open';
saveStatuses();
const st = statusOf(id);
const btn = document.querySelector('.status-btn[data-id="' + id + '"]');
if (!btn) { render(); return; }
btn.className = 'status-btn ' + (STATUS_CLS[st] || 'st-open');
btn.textContent = STATUS_LBL[st] || '\u25cf open';
const card = document.querySelector('[data-card-id="' + id + '"]');
if (card) card.classList.toggle('resolved', st !== 'open');
updateProgress();
}
// Surgical collapse/expand for a group: toggle all bodies in the group without full re-render.
function collapseGroup(key, collapse) {
FINDINGS.forEach(f => {
if (getKey(f) !== key) return;
if (collapse) S.collapsed.add(f.id); else S.collapsed.delete(f.id);
const body = document.querySelector('[data-body-id="' + f.id + '"]');
const arrow = document.querySelector('[data-card-id="' + f.id + '"] .card-arrow');
if (!body) return;
if (collapse) {
body.style.display = 'none';
if (arrow) arrow.textContent = '\u25b6';
} else {
if (!body.firstChild) {
const ff = FINDING_BY_ID.get(f.id);
if (ff) body.innerHTML = renderCardBody(ff, S.search);
}
body.style.display = '';
if (arrow) arrow.textContent = '\u25bc';
}
});
}
// Full re-render: used for filter/group-by/search changes.
// Batched via requestAnimationFrame.
let _rafPending = false;
function scheduleRender() {
if (_rafPending) return;
_rafPending = true;
requestAnimationFrame(() => { _rafPending = false; render(); });
}
function render() {
_hay.clear();
_hayLoc.clear();
updateToggles();
updateProgress();
const q = S.search;
const visible = FINDINGS.filter(matches);
document.getElementById('stats').textContent = visible.length + ' / ' + FINDINGS.length + ' findings';
// Build groups
const map = new Map();
for (const f of visible) {
const k = getKey(f);
if (!map.has(k)) map.set(k, []);
map.get(k).push(f);
}
// Hidden-groups bar
const hbar = document.getElementById('hidden-bar');
if (S.hiddenGroup.size > 0) {
let hh = '<div class="hidden-inner"><span class="hidden-lbl">Hidden:</span>';
for (const g of [...S.hiddenGroup].sort())
hh += '<button class="show-btn" data-action="show-group" data-key="' + esc(g) + '">' + esc(g || '(all findings)') + '</button>';
hh += '<button class="show-all-btn" data-action="show-all">Show all</button></div>';
hbar.innerHTML = hh;
} else {
hbar.innerHTML = '';
}
// Content
const content = document.getElementById('content');
if (!visible.length) {
content.innerHTML = '<div class="no-results">No findings match the current filters.</div>';
syncHash();
return;
}
let hh = '';
for (const k of sortKeys([...map.keys()])) {
if (S.hiddenGroup.has(k)) continue;
const items = map.get(k);
const title = k || 'All Findings';
const done = items.filter(f => statusOf(f.id) !== 'open').length;
const allDone = done === items.length;
hh += '<div class="group-section">' +
'<div class="group-header">' +
'<span class="group-title' + (allDone ? ' all-done' : '') + '">' + esc(title) + '</span>' +
'<span class="group-count">' + done + '/' + items.length + '</span>' +
'<button class="icon-btn" data-action="collapse-group" data-key="' + esc(k) + '" title="Collapse all">\u229f</button>' +
'<button class="icon-btn" data-action="expand-group" data-key="' + esc(k) + '" title="Expand all">\u229e</button>' +
'<button class="hide-btn" data-action="hide-group" data-key="' + esc(k) + '">Hide</button>' +
'</div>';
for (const f of items) hh += renderCard(f, q);
hh += '</div>';
}
content.innerHTML = hh;
syncHash();
}
// Long-press solo: hold a filter badge 450 ms to isolate that value (hide all others in the group).
let _soloTimer = 0;
let _soloFired = false;
function _applySolo(b) {
_soloFired = true;
const sev = b.dataset.sev, typ = b.dataset.type, sta = b.dataset.status, dif = b.dataset.diff;
if (sev) { S.hiddenSev.clear(); PRESENT_SEVS.filter(v => v !== sev).forEach(v => S.hiddenSev.add(v)); }
else if (typ) { S.hiddenType.clear(); ALL_TYPES.filter(v => v !== typ).forEach(v => S.hiddenType.add(v)); }
else if (sta) { S.hiddenStatus.clear(); STATUSES.filter(v => v !== sta).forEach(v => S.hiddenStatus.add(v)); }
else if (dif) { S.hiddenDiff.clear(); PRESENT_DIFFS.filter(v => v !== dif).forEach(v => S.hiddenDiff.add(v)); }
scheduleRender();
}
const _fb = document.getElementById('filterbar');
_fb.addEventListener('pointerdown', e => {
const b = e.target.closest('[data-sev],[data-type],[data-status],[data-diff]');
if (!b) return;
_soloFired = false;
clearTimeout(_soloTimer);
_soloTimer = setTimeout(() => _applySolo(b), 450);
}, {passive: true});
_fb.addEventListener('pointerup', () => clearTimeout(_soloTimer), {passive: true});
_fb.addEventListener('pointercancel', () => clearTimeout(_soloTimer), {passive: true});
_fb.addEventListener('pointermove', e => {
if (!_soloTimer) return;
if (!e.target.closest('[data-sev],[data-type],[data-status],[data-diff]')) clearTimeout(_soloTimer);
}, {passive: true});
// Event delegation
_fb.addEventListener('click', e => {
if (_soloFired) { _soloFired = false; return; }
const b = e.target.closest('[data-sev],[data-type],[data-status],[data-toggle-all],[data-diff]');
if (!b) return;
if (b.dataset.toggleAll) {
const k = b.dataset.toggleAll;
const set = k === 'sev' ? S.hiddenSev : k === 'type' ? S.hiddenType : k === 'status' ? S.hiddenStatus : k === 'diff' ? S.hiddenDiff : null;
const full = k === 'sev' ? PRESENT_SEVS : k === 'type' ? ALL_TYPES : k === 'status' ? STATUSES : k === 'diff' ? PRESENT_DIFFS : [];
if (!set) return;
if (set.size === 0) full.forEach(v => set.add(v));
else set.clear();
}
else if (b.dataset.sev) { const s = S.hiddenSev; s.has(b.dataset.sev) ? s.delete(b.dataset.sev) : s.add(b.dataset.sev); }
else if (b.dataset.type) { const s = S.hiddenType; s.has(b.dataset.type) ? s.delete(b.dataset.type) : s.add(b.dataset.type); }
else if (b.dataset.status) { const s = S.hiddenStatus; s.has(b.dataset.status) ? s.delete(b.dataset.status) : s.add(b.dataset.status); }
else if (b.dataset.diff) { const s = S.hiddenDiff; s.has(b.dataset.diff) ? s.delete(b.dataset.diff) : s.add(b.dataset.diff); }
scheduleRender();
});
document.addEventListener('click', e => {
const b = e.target.closest('[data-action]');
if (!b) return;
const { action, key, id } = b.dataset;
if (action === 'hide-group') { S.hiddenGroup.add(key); scheduleRender(); }
else if (action === 'show-group') { S.hiddenGroup.delete(key); scheduleRender(); }
else if (action === 'show-all') { S.hiddenGroup.clear(); scheduleRender(); }
else if (action === 'collapse-all') { FINDINGS.forEach(f => S.collapsed.add(f.id)); scheduleRender(); }
else if (action === 'expand-all') { S.collapsed.clear(); scheduleRender(); }
else if (action === 'copy-md') {
e.stopPropagation();
const f = FINDING_BY_ID.get(id);
if (f) copyText(findingToMd(f), b);
}
else if (action === 'download-md') { downloadText(buildMdFull(), 'findings.md'); }
else if (action === 'download-md-short') { downloadText(buildMdShort(), 'findings-short.md'); }
else if (action === 'download-toml') {
// Single download (was two) - a second auto-triggered download is commonly
// blocked by the browser unless the user explicitly allows it, which silently
// drops project-info.toml. Embedding [project] in findings.toml itself avoids
// that failure mode entirely.
downloadText(buildFindingsToml(), 'findings.toml', 'application/toml');
}
else if (action === 'cycle-status') {
e.stopPropagation();
cycleStatus(id);
}
else if (action === 'toggle-card') {
toggleCard(id);
}
else if (action === 'collapse-group') {
collapseGroup(key, true);
}
else if (action === 'expand-group') {
collapseGroup(key, false);
}
});
// Debounced search — 180 ms after last keystroke
let _searchTimer = 0;
document.getElementById('search').addEventListener('input', e => {
S.search = e.target.value;
clearTimeout(_searchTimer);
_searchTimer = setTimeout(scheduleRender, 180);
});
document.getElementById('groupby').addEventListener('change', e => {
S.groupBy = e.target.value;
S.hiddenGroup.clear();
scheduleRender();
});
document.getElementById('search-scope').addEventListener('click', e => {
S.searchScope = S.searchScope === 'locations' ? 'all' : 'locations';
e.currentTarget.classList.toggle('active', S.searchScope === 'locations');
scheduleRender();
});
// Init
loadStatuses();
initFilterBar();
loadHash();
FINDINGS.forEach(f => S.collapsed.add(f.id));
render();
})();
</script>
</body></html>