Skip to content

Commit d94e0a1

Browse files
jlebarjlebar
committed
Minimize m048 reproducer scaffolding
1 parent 6c84f8e commit d94e0a1

3 files changed

Lines changed: 4 additions & 135 deletions

File tree

‎ptx/known-miscompiles/m048-intmax-popc-sub-mask-fold/NOTES.md‎

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -10,9 +10,10 @@ The original saved fuzzer program was:
1010
```
1111

1212
The automated reducer shrank the original 2,644-line program to 180 lines.
13-
Manual cleanup showed that a straight-line scalar version and a simple
14-
structured branch version both compile correctly; the remaining trigger still
15-
needs structured-branch context around the value chain.
13+
Manual cleanup removed unreferenced structured labels, bringing the checked-in
14+
PTX to 114 lines. A straight-line scalar version and a simple structured
15+
branch version both compile correctly; the remaining trigger still needs
16+
structured-branch context around the value chain.
1617

1718
## Scalar Trace
1819

‎ptx/known-miscompiles/m048-intmax-popc-sub-mask-fold/reduced.ptx‎

Lines changed: 0 additions & 66 deletions
Original file line numberDiff line numberDiff line change
@@ -82,96 +82,30 @@ structured_if_1_then:
8282
sub.u32 %r4, 3046743225, %r1;
8383
and.b32 %r23, %r4, %r36;
8484
structured_if_1_else:
85-
structured_if_1_done:
8685
setp.lt.u32 %p2, %r42, 1019145259;
87-
structured_if_2_then:
88-
structured_if_2_else:
89-
structured_if_2_done:
9086
setp.lt.u32 %p3, %r35, %r27;
91-
structured_if_3_then:
9287
setp.ne.u32 %p4, %r43, %r43;
93-
structured_if_4_then:
94-
structured_if_4_else:
95-
structured_if_4_done:
9688
setp.ne.u32 %p5, 2683960277, 4278255360;
97-
structured_if_5_then:
98-
structured_if_5_else:
99-
structured_if_5_done:
100-
structured_if_3_else:
101-
structured_if_3_done:
10289
setp.ge.u32 %p6, %r13, %r11;
103-
structured_if_6_then:
10490
setp.le.u32 %p7, %r41, %r42;
105-
structured_if_7_then:
106-
structured_if_7_else:
107-
structured_if_7_done:
108-
structured_if_6_else:
109-
structured_if_6_done:
11091
setp.le.u32 %p8, %r14, %r25;
111-
structured_if_8_then:
112-
structured_if_8_else:
113-
structured_if_8_done:
11492
structured_if_0_else:
11593
setp.lt.u32 %p9, 32768, %r17;
116-
structured_if_9_then:
11794
setp.le.u32 %p10, %r38, 266618723;
118-
structured_if_10_then:
119-
structured_if_10_else:
120-
structured_if_10_done:
12195
setp.eq.u32 %p11, %r19, %r47;
122-
structured_if_11_then:
123-
structured_if_11_else:
124-
structured_if_11_done:
12596
setp.eq.u32 %p12, %r5, 2;
126-
structured_if_12_then:
127-
structured_if_12_else:
128-
structured_if_12_done:
12997
setp.lt.u32 %p13, %r19, 894635186;
130-
structured_if_13_then:
131-
structured_if_13_else:
132-
structured_if_13_done:
133-
structured_if_9_else:
134-
structured_if_9_done:
13598
setp.eq.u32 %p14, %r46, %r23;
136-
structured_if_14_then:
13799
setp.ne.u32 %p15, %r2, 512;
138-
structured_if_15_then:
139-
structured_if_15_else:
140-
structured_if_15_done:
141-
structured_if_14_else:
142100
setp.gt.u32 %p16, %r4, 3712077568;
143-
structured_if_16_then:
144-
structured_if_16_else:
145-
structured_if_16_done:
146101
setp.lt.u32 %p17, %r46, %r0;
147-
structured_if_17_then:
148102
setp.lt.u32 %p18, %r31, %r34;
149-
structured_if_18_then:
150-
structured_if_18_else:
151103
setp.lt.u32 %p19, %r10, %r5;
152-
structured_if_19_then:
153-
structured_if_19_else:
154-
structured_if_19_done:
155-
structured_if_18_done:
156-
structured_if_17_else:
157-
structured_if_17_done:
158-
structured_if_14_done:
159104
setp.lt.u32 %p20, %r28, 3155773991;
160-
structured_if_20_then:
161-
structured_if_20_else:
162-
structured_if_20_done:
163105
setp.gt.u32 %p21, %r33, %r21;
164-
structured_if_21_then:
165-
structured_if_21_else:
166-
structured_if_21_done:
167-
structured_if_0_done:
168106
setp.le.u32 %p22, %r8, %r39;
169-
structured_if_22_then:
170107
add.u32 %r0, %r23, %r26;
171-
structured_if_22_else:
172-
structured_if_22_done:
173108

174-
exit:
175109
cvta.to.global.u64 %rd4, %rd1;
176110
mul.wide.u32 %rd5, %r48, 16;
177111
add.s64 %rd4, %rd4, %rd5;

‎ptx/known-miscompiles/m048-intmax-popc-sub-mask-fold/repro_nvcc_inline_ptx.cu‎

Lines changed: 0 additions & 66 deletions
Original file line numberDiff line numberDiff line change
@@ -115,95 +115,29 @@ __global__ void repro_kernel(const uint32_t* in, uint32_t* out, uint32_t n) {
115115
"sub.u32 r4, 3046743225, r1;\n\t"
116116
"and.b32 r23, r4, r36;\n\t"
117117
"structured_if_1_else:\n\t"
118-
"structured_if_1_done:\n\t"
119118
"setp.lt.u32 p2, r42, 1019145259;\n\t"
120-
"structured_if_2_then:\n\t"
121-
"structured_if_2_else:\n\t"
122-
"structured_if_2_done:\n\t"
123119
"setp.lt.u32 p3, r35, r27;\n\t"
124-
"structured_if_3_then:\n\t"
125120
"setp.ne.u32 p4, r43, r43;\n\t"
126-
"structured_if_4_then:\n\t"
127-
"structured_if_4_else:\n\t"
128-
"structured_if_4_done:\n\t"
129121
"setp.ne.u32 p5, 2683960277, 4278255360;\n\t"
130-
"structured_if_5_then:\n\t"
131-
"structured_if_5_else:\n\t"
132-
"structured_if_5_done:\n\t"
133-
"structured_if_3_else:\n\t"
134-
"structured_if_3_done:\n\t"
135122
"setp.ge.u32 p6, r13, r11;\n\t"
136-
"structured_if_6_then:\n\t"
137123
"setp.le.u32 p7, r41, r42;\n\t"
138-
"structured_if_7_then:\n\t"
139-
"structured_if_7_else:\n\t"
140-
"structured_if_7_done:\n\t"
141-
"structured_if_6_else:\n\t"
142-
"structured_if_6_done:\n\t"
143124
"setp.le.u32 p8, r14, r25;\n\t"
144-
"structured_if_8_then:\n\t"
145-
"structured_if_8_else:\n\t"
146-
"structured_if_8_done:\n\t"
147125
"structured_if_0_else:\n\t"
148126
"setp.lt.u32 p9, 32768, r17;\n\t"
149-
"structured_if_9_then:\n\t"
150127
"setp.le.u32 p10, r38, 266618723;\n\t"
151-
"structured_if_10_then:\n\t"
152-
"structured_if_10_else:\n\t"
153-
"structured_if_10_done:\n\t"
154128
"setp.eq.u32 p11, r19, r47;\n\t"
155-
"structured_if_11_then:\n\t"
156-
"structured_if_11_else:\n\t"
157-
"structured_if_11_done:\n\t"
158129
"setp.eq.u32 p12, r5, 2;\n\t"
159-
"structured_if_12_then:\n\t"
160-
"structured_if_12_else:\n\t"
161-
"structured_if_12_done:\n\t"
162130
"setp.lt.u32 p13, r19, 894635186;\n\t"
163-
"structured_if_13_then:\n\t"
164-
"structured_if_13_else:\n\t"
165-
"structured_if_13_done:\n\t"
166-
"structured_if_9_else:\n\t"
167-
"structured_if_9_done:\n\t"
168131
"setp.eq.u32 p14, r46, r23;\n\t"
169-
"structured_if_14_then:\n\t"
170132
"setp.ne.u32 p15, r2, 512;\n\t"
171-
"structured_if_15_then:\n\t"
172-
"structured_if_15_else:\n\t"
173-
"structured_if_15_done:\n\t"
174-
"structured_if_14_else:\n\t"
175133
"setp.gt.u32 p16, r4, 3712077568;\n\t"
176-
"structured_if_16_then:\n\t"
177-
"structured_if_16_else:\n\t"
178-
"structured_if_16_done:\n\t"
179134
"setp.lt.u32 p17, r46, r0;\n\t"
180-
"structured_if_17_then:\n\t"
181135
"setp.lt.u32 p18, r31, r34;\n\t"
182-
"structured_if_18_then:\n\t"
183-
"structured_if_18_else:\n\t"
184136
"setp.lt.u32 p19, r10, r5;\n\t"
185-
"structured_if_19_then:\n\t"
186-
"structured_if_19_else:\n\t"
187-
"structured_if_19_done:\n\t"
188-
"structured_if_18_done:\n\t"
189-
"structured_if_17_else:\n\t"
190-
"structured_if_17_done:\n\t"
191-
"structured_if_14_done:\n\t"
192137
"setp.lt.u32 p20, r28, 3155773991;\n\t"
193-
"structured_if_20_then:\n\t"
194-
"structured_if_20_else:\n\t"
195-
"structured_if_20_done:\n\t"
196138
"setp.gt.u32 p21, r33, r21;\n\t"
197-
"structured_if_21_then:\n\t"
198-
"structured_if_21_else:\n\t"
199-
"structured_if_21_done:\n\t"
200-
"structured_if_0_done:\n\t"
201139
"setp.le.u32 p22, r8, r39;\n\t"
202-
"structured_if_22_then:\n\t"
203140
"add.u32 r0, r23, r26;\n\t"
204-
"structured_if_22_else:\n\t"
205-
"structured_if_22_done:\n\t"
206-
"exit:\n\t"
207141
"cvta.to.global.u64 rd4, rd1;\n\t"
208142
"mul.wide.u32 rd5, r48, 16;\n\t"
209143
"add.s64 rd4, rd4, rd5;\n\t"

0 commit comments

Comments
 (0)