blob: 335443171fd87fd66263e13c146bac558e771c84 [file] [log] [blame]
dan sinclairf1f381a2023-11-22 09:44:15 +00001SKIP: FAILED
2
3<dawn>/src/tint/lang/msl/writer/printer/printer.cc:355 internal compiler error: Tables = struct @align(4) {
4 edges:array<u32, 256> @offset(0)
5 tris:array<i32, 4096> @offset(1024)
6}
7
8IsosurfaceVolume = struct @align(16) {
9 min:vec3<f32> @offset(0)
10 max:vec3<f32> @offset(16)
11 step:vec3<f32> @offset(32)
12 size:vec3<u32> @offset(48)
13 threshold:f32 @offset(60)
14 values:array<f32> @offset(64)
15}
16
17PositionBuffer = struct @align(4) {
18 values:array<f32> @offset(0)
19}
20
21NormalBuffer = struct @align(4) {
22 values:array<f32> @offset(0)
23}
24
25IndexBuffer = struct @align(4) {
26 tris:array<u32> @offset(0)
27}
28
29DrawIndirectArgs = struct @align(4) {
30 vc:u32 @offset(0)
31 vertexCount:atomic<u32> @offset(4)
32 firstVertex:u32 @offset(8)
33 firstInstance:u32 @offset(12)
34 indexCount:atomic<u32> @offset(16)
35 indexedInstanceCount:u32 @offset(20)
36 indexedFirstIndex:u32 @offset(24)
37 indexedBaseVertex:u32 @offset(28)
38 indexedFirstInstance:u32 @offset(32)
39}
40
41%b1 = block { # root
42 %tables:ptr<storage, Tables, read> = var @binding_point(0, 0)
43 %volume:ptr<storage, IsosurfaceVolume, read_write> = var @binding_point(0, 1)
44 %positionsOut:ptr<storage, PositionBuffer, read_write> = var @binding_point(0, 2)
45 %normalsOut:ptr<storage, NormalBuffer, read_write> = var @binding_point(0, 3)
46 %indicesOut:ptr<storage, IndexBuffer, read_write> = var @binding_point(0, 4)
47 %drawOut:ptr<storage, DrawIndirectArgs, read_write> = var @binding_point(0, 5)
48 %positions:ptr<private, array<vec3<f32>, 12>, read_write> = var
49 %normals:ptr<private, array<vec3<f32>, 12>, read_write> = var
50 %indices:ptr<private, array<u32, 12>, read_write> = var
51 %cubeVerts:ptr<private, u32, read_write> = var, 0u
52}
53
54%valueAt = func(%index:vec3<u32>):f32 -> %b2 {
55 %b2 = block {
56 %13:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
57 %14:vec3<u32> = load %13
58 %15:vec3<bool> = gte %index, %14
59 %16:bool = any %15
60 if %16 [t: %b3] { # if_1
61 %b3 = block { # true
62 ret 0.0f
63 }
64 }
65 %17:u32 = access %index, 0u
66 %18:u32 = access %index, 1u
67 %19:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
68 %20:u32 = load_vector_element %19, 0u
69 %21:u32 = mul %18, %20
70 %22:u32 = add %17, %21
71 %23:u32 = access %index, 2u
72 %24:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
73 %25:u32 = load_vector_element %24, 0u
74 %26:u32 = mul %23, %25
75 %27:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
76 %28:u32 = load_vector_element %27, 1u
77 %29:u32 = mul %26, %28
78 %valueIndex:u32 = add %22, %29
79 %31:ptr<storage, f32, read_write> = access %volume, 5u, %valueIndex
80 %32:f32 = load %31
81 ret %32
82 }
83}
84%positionAt = func(%index_1:vec3<u32>):vec3<f32> -> %b4 { # %index_1: 'index'
85 %b4 = block {
86 %35:ptr<storage, vec3<f32>, read_write> = access %volume, 0u
87 %36:vec3<f32> = load %35
88 %37:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
89 %38:vec3<f32> = load %37
90 %39:vec3<u32> = swizzle %index_1, xyz
91 %40:vec3<f32> = convert %39
92 %41:vec3<f32> = mul %38, %40
93 %42:vec3<f32> = add %36, %41
94 ret %42
95 }
96}
97%normalAt = func(%index_2:vec3<u32>):vec3<f32> -> %b5 { # %index_2: 'index'
98 %b5 = block {
99 %45:vec3<u32> = sub %index_2, vec3<u32>(1u, 0u, 0u)
100 %46:f32 = call %valueAt, %45
101 %47:vec3<u32> = add %index_2, vec3<u32>(1u, 0u, 0u)
102 %48:f32 = call %valueAt, %47
103 %49:f32 = sub %46, %48
104 %50:vec3<u32> = sub %index_2, vec3<u32>(0u, 1u, 0u)
105 %51:f32 = call %valueAt, %50
106 %52:vec3<u32> = add %index_2, vec3<u32>(0u, 1u, 0u)
107 %53:f32 = call %valueAt, %52
108 %54:f32 = sub %51, %53
109 %55:vec3<u32> = sub %index_2, vec3<u32>(0u, 0u, 1u)
110 %56:f32 = call %valueAt, %55
111 %57:vec3<u32> = add %index_2, vec3<u32>(0u, 0u, 1u)
112 %58:f32 = call %valueAt, %57
113 %59:f32 = sub %56, %58
114 %60:vec3<f32> = construct %49, %54, %59
115 ret %60
116 }
117}
118%interpX = func(%index_3:u32, %i:vec3<u32>, %va:f32, %vb:f32):void -> %b6 { # %index_3: 'index'
119 %b6 = block {
120 %66:ptr<storage, f32, read_write> = access %volume, 4u
121 %67:f32 = load %66
122 %68:f32 = sub %67, %va
123 %69:f32 = sub %vb, %va
124 %mu:f32 = div %68, %69
125 %71:u32 = load %cubeVerts
126 %72:ptr<private, vec3<f32>, read_write> = access %positions, %71
127 %73:vec3<f32> = call %positionAt, %i
128 %74:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
129 %75:f32 = load_vector_element %74, 0u
130 %76:f32 = mul %75, %mu
131 %77:vec3<f32> = construct %76, 0.0f, 0.0f
132 %78:vec3<f32> = add %73, %77
133 store %72, %78
134 %na:vec3<f32> = call %normalAt, %i
135 %80:vec3<u32> = add %i, vec3<u32>(1u, 0u, 0u)
136 %nb:vec3<f32> = call %normalAt, %80
137 %82:u32 = load %cubeVerts
138 %83:ptr<private, vec3<f32>, read_write> = access %normals, %82
139 %84:vec3<f32> = construct %mu, %mu, %mu
140 %85:vec3<f32> = mix %na, %nb, %84
141 store %83, %85
142 %86:ptr<private, u32, read_write> = access %indices, %index_3
143 %87:u32 = load %cubeVerts
144 store %86, %87
145 %88:u32 = load %cubeVerts
146 %89:u32 = add %88, 1u
147 store %cubeVerts, %89
148 ret
149 }
150}
151%interpY = func(%index_4:u32, %i_1:vec3<u32>, %va_1:f32, %vb_1:f32):void -> %b7 { # %index_4: 'index', %i_1: 'i', %va_1: 'va', %vb_1: 'vb'
152 %b7 = block {
153 %95:ptr<storage, f32, read_write> = access %volume, 4u
154 %96:f32 = load %95
155 %97:f32 = sub %96, %va_1
156 %98:f32 = sub %vb_1, %va_1
157 %mu_1:f32 = div %97, %98 # %mu_1: 'mu'
158 %100:u32 = load %cubeVerts
159 %101:ptr<private, vec3<f32>, read_write> = access %positions, %100
160 %102:vec3<f32> = call %positionAt, %i_1
161 %103:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
162 %104:f32 = load_vector_element %103, 1u
163 %105:f32 = mul %104, %mu_1
164 %106:vec3<f32> = construct 0.0f, %105, 0.0f
165 %107:vec3<f32> = add %102, %106
166 store %101, %107
167 %na_1:vec3<f32> = call %normalAt, %i_1 # %na_1: 'na'
168 %109:vec3<u32> = add %i_1, vec3<u32>(0u, 1u, 0u)
169 %nb_1:vec3<f32> = call %normalAt, %109 # %nb_1: 'nb'
170 %111:u32 = load %cubeVerts
171 %112:ptr<private, vec3<f32>, read_write> = access %normals, %111
172 %113:vec3<f32> = construct %mu_1, %mu_1, %mu_1
173 %114:vec3<f32> = mix %na_1, %nb_1, %113
174 store %112, %114
175 %115:ptr<private, u32, read_write> = access %indices, %index_4
176 %116:u32 = load %cubeVerts
177 store %115, %116
178 %117:u32 = load %cubeVerts
179 %118:u32 = add %117, 1u
180 store %cubeVerts, %118
181 ret
182 }
183}
184%interpZ = func(%index_5:u32, %i_2:vec3<u32>, %va_2:f32, %vb_2:f32):void -> %b8 { # %index_5: 'index', %i_2: 'i', %va_2: 'va', %vb_2: 'vb'
185 %b8 = block {
186 %124:ptr<storage, f32, read_write> = access %volume, 4u
187 %125:f32 = load %124
188 %126:f32 = sub %125, %va_2
189 %127:f32 = sub %vb_2, %va_2
190 %mu_2:f32 = div %126, %127 # %mu_2: 'mu'
191 %129:u32 = load %cubeVerts
192 %130:ptr<private, vec3<f32>, read_write> = access %positions, %129
193 %131:vec3<f32> = call %positionAt, %i_2
194 %132:ptr<storage, vec3<f32>, read_write> = access %volume, 2u
195 %133:f32 = load_vector_element %132, 2u
196 %134:f32 = mul %133, %mu_2
197 %135:vec3<f32> = construct 0.0f, 0.0f, %134
198 %136:vec3<f32> = add %131, %135
199 store %130, %136
200 %na_2:vec3<f32> = call %normalAt, %i_2 # %na_2: 'na'
201 %138:vec3<u32> = add %i_2, vec3<u32>(0u, 0u, 1u)
202 %nb_2:vec3<f32> = call %normalAt, %138 # %nb_2: 'nb'
203 %140:u32 = load %cubeVerts
204 %141:ptr<private, vec3<f32>, read_write> = access %normals, %140
205 %142:vec3<f32> = construct %mu_2, %mu_2, %mu_2
206 %143:vec3<f32> = mix %na_2, %nb_2, %142
207 store %141, %143
208 %144:ptr<private, u32, read_write> = access %indices, %index_5
209 %145:u32 = load %cubeVerts
210 store %144, %145
211 %146:u32 = load %cubeVerts
212 %147:u32 = add %146, 1u
213 store %cubeVerts, %147
214 ret
215 }
216}
217%computeMain = @compute @workgroup_size(4, 4, 4) func(%global_id:vec3<u32> [@global_invocation_id]):void -> %b9 {
218 %b9 = block {
219 %i0:vec3<u32> = let %global_id
220 %i1:vec3<u32> = add %global_id, vec3<u32>(1u, 0u, 0u)
221 %i2:vec3<u32> = add %global_id, vec3<u32>(1u, 1u, 0u)
222 %i3:vec3<u32> = add %global_id, vec3<u32>(0u, 1u, 0u)
223 %i4:vec3<u32> = add %global_id, vec3<u32>(0u, 0u, 1u)
224 %i5:vec3<u32> = add %global_id, vec3<u32>(1u, 0u, 1u)
225 %i6:vec3<u32> = add %global_id, vec3<u32>(1u)
226 %i7:vec3<u32> = add %global_id, vec3<u32>(0u, 1u, 1u)
227 %v0:f32 = call %valueAt, %i0
228 %v1:f32 = call %valueAt, %i1
229 %v2:f32 = call %valueAt, %i2
230 %v3:f32 = call %valueAt, %i3
231 %v4:f32 = call %valueAt, %i4
232 %v5:f32 = call %valueAt, %i5
233 %v6:f32 = call %valueAt, %i6
234 %v7:f32 = call %valueAt, %i7
235 %cubeIndex:ptr<function, u32, read_write> = var, 0u
236 %167:ptr<storage, f32, read_write> = access %volume, 4u
237 %168:f32 = load %167
238 %169:bool = lt %v0, %168
239 if %169 [t: %b10] { # if_2
240 %b10 = block { # true
241 %170:u32 = load %cubeIndex
242 %171:u32 = or %170, 1u
243 store %cubeIndex, %171
244 exit_if # if_2
245 }
246 }
247 %172:ptr<storage, f32, read_write> = access %volume, 4u
248 %173:f32 = load %172
249 %174:bool = lt %v1, %173
250 if %174 [t: %b11] { # if_3
251 %b11 = block { # true
252 %175:u32 = load %cubeIndex
253 %176:u32 = or %175, 2u
254 store %cubeIndex, %176
255 exit_if # if_3
256 }
257 }
258 %177:ptr<storage, f32, read_write> = access %volume, 4u
259 %178:f32 = load %177
260 %179:bool = lt %v2, %178
261 if %179 [t: %b12] { # if_4
262 %b12 = block { # true
263 %180:u32 = load %cubeIndex
264 %181:u32 = or %180, 4u
265 store %cubeIndex, %181
266 exit_if # if_4
267 }
268 }
269 %182:ptr<storage, f32, read_write> = access %volume, 4u
270 %183:f32 = load %182
271 %184:bool = lt %v3, %183
272 if %184 [t: %b13] { # if_5
273 %b13 = block { # true
274 %185:u32 = load %cubeIndex
275 %186:u32 = or %185, 8u
276 store %cubeIndex, %186
277 exit_if # if_5
278 }
279 }
280 %187:ptr<storage, f32, read_write> = access %volume, 4u
281 %188:f32 = load %187
282 %189:bool = lt %v4, %188
283 if %189 [t: %b14] { # if_6
284 %b14 = block { # true
285 %190:u32 = load %cubeIndex
286 %191:u32 = or %190, 16u
287 store %cubeIndex, %191
288 exit_if # if_6
289 }
290 }
291 %192:ptr<storage, f32, read_write> = access %volume, 4u
292 %193:f32 = load %192
293 %194:bool = lt %v5, %193
294 if %194 [t: %b15] { # if_7
295 %b15 = block { # true
296 %195:u32 = load %cubeIndex
297 %196:u32 = or %195, 32u
298 store %cubeIndex, %196
299 exit_if # if_7
300 }
301 }
302 %197:ptr<storage, f32, read_write> = access %volume, 4u
303 %198:f32 = load %197
304 %199:bool = lt %v6, %198
305 if %199 [t: %b16] { # if_8
306 %b16 = block { # true
307 %200:u32 = load %cubeIndex
308 %201:u32 = or %200, 64u
309 store %cubeIndex, %201
310 exit_if # if_8
311 }
312 }
313 %202:ptr<storage, f32, read_write> = access %volume, 4u
314 %203:f32 = load %202
315 %204:bool = lt %v7, %203
316 if %204 [t: %b17] { # if_9
317 %b17 = block { # true
318 %205:u32 = load %cubeIndex
319 %206:u32 = or %205, 128u
320 store %cubeIndex, %206
321 exit_if # if_9
322 }
323 }
324 %207:u32 = load %cubeIndex
325 %208:ptr<storage, u32, read> = access %tables, 0u, %207
326 %edges:u32 = load %208
327 %210:u32 = and %edges, 1u
328 %211:bool = neq %210, 0u
329 if %211 [t: %b18] { # if_10
330 %b18 = block { # true
331 %212:void = call %interpX, 0u, %i0, %v0, %v1
332 exit_if # if_10
333 }
334 }
335 %213:u32 = and %edges, 2u
336 %214:bool = neq %213, 0u
337 if %214 [t: %b19] { # if_11
338 %b19 = block { # true
339 %215:void = call %interpY, 1u, %i1, %v1, %v2
340 exit_if # if_11
341 }
342 }
343 %216:u32 = and %edges, 4u
344 %217:bool = neq %216, 0u
345 if %217 [t: %b20] { # if_12
346 %b20 = block { # true
347 %218:void = call %interpX, 2u, %i3, %v3, %v2
348 exit_if # if_12
349 }
350 }
351 %219:u32 = and %edges, 8u
352 %220:bool = neq %219, 0u
353 if %220 [t: %b21] { # if_13
354 %b21 = block { # true
355 %221:void = call %interpY, 3u, %i0, %v0, %v3
356 exit_if # if_13
357 }
358 }
359 %222:u32 = and %edges, 16u
360 %223:bool = neq %222, 0u
361 if %223 [t: %b22] { # if_14
362 %b22 = block { # true
363 %224:void = call %interpX, 4u, %i4, %v4, %v5
364 exit_if # if_14
365 }
366 }
367 %225:u32 = and %edges, 32u
368 %226:bool = neq %225, 0u
369 if %226 [t: %b23] { # if_15
370 %b23 = block { # true
371 %227:void = call %interpY, 5u, %i5, %v5, %v6
372 exit_if # if_15
373 }
374 }
375 %228:u32 = and %edges, 64u
376 %229:bool = neq %228, 0u
377 if %229 [t: %b24] { # if_16
378 %b24 = block { # true
379 %230:void = call %interpX, 6u, %i7, %v7, %v6
380 exit_if # if_16
381 }
382 }
383 %231:u32 = and %edges, 128u
384 %232:bool = neq %231, 0u
385 if %232 [t: %b25] { # if_17
386 %b25 = block { # true
387 %233:void = call %interpY, 7u, %i4, %v4, %v7
388 exit_if # if_17
389 }
390 }
391 %234:u32 = and %edges, 256u
392 %235:bool = neq %234, 0u
393 if %235 [t: %b26] { # if_18
394 %b26 = block { # true
395 %236:void = call %interpZ, 8u, %i0, %v0, %v4
396 exit_if # if_18
397 }
398 }
399 %237:u32 = and %edges, 512u
400 %238:bool = neq %237, 0u
401 if %238 [t: %b27] { # if_19
402 %b27 = block { # true
403 %239:void = call %interpZ, 9u, %i1, %v1, %v5
404 exit_if # if_19
405 }
406 }
407 %240:u32 = and %edges, 1024u
408 %241:bool = neq %240, 0u
409 if %241 [t: %b28] { # if_20
410 %b28 = block { # true
411 %242:void = call %interpZ, 10u, %i2, %v2, %v6
412 exit_if # if_20
413 }
414 }
415 %243:u32 = and %edges, 2048u
416 %244:bool = neq %243, 0u
417 if %244 [t: %b29] { # if_21
418 %b29 = block { # true
419 %245:void = call %interpZ, 11u, %i3, %v3, %v7
420 exit_if # if_21
421 }
422 }
423 %246:u32 = load %cubeIndex
424 %247:u32 = shiftl %246, 4u
425 %triTableOffset:u32 = add %247, 1u
426 %249:u32 = sub %triTableOffset, 1u
427 %250:ptr<storage, i32, read> = access %tables, 1u, %249
428 %251:i32 = load %250
429 %indexCount:u32 = convert %251
430 %253:ptr<storage, atomic<u32>, read_write> = access %drawOut, 1u
431 %254:u32 = load %cubeVerts
432 %255:u32 = atomicAdd %253, %254
433 %firstVertex:ptr<function, u32, read_write> = var, %255
434 %257:u32 = access %global_id, 0u
435 %258:u32 = access %global_id, 1u
436 %259:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
437 %260:u32 = load_vector_element %259, 0u
438 %261:u32 = mul %258, %260
439 %262:u32 = add %257, %261
440 %263:u32 = access %global_id, 2u
441 %264:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
442 %265:u32 = load_vector_element %264, 0u
443 %266:u32 = mul %263, %265
444 %267:ptr<storage, vec3<u32>, read_write> = access %volume, 3u
445 %268:u32 = load_vector_element %267, 1u
446 %269:u32 = mul %266, %268
447 %bufferOffset:u32 = add %262, %269
448 %firstIndex:u32 = mul %bufferOffset, 15u
449 loop [i: %b30, b: %b31, c: %b32] { # loop_1
450 %b30 = block { # initializer
451 %i_3:ptr<function, u32, read_write> = var, 0u # %i_3: 'i'
452 next_iteration %b31
453 }
454 %b31 = block { # body
455 %273:u32 = load %i_3
456 %274:u32 = load %cubeVerts
457 %275:bool = lt %273, %274
458 if %275 [t: %b33, f: %b34] { # if_22
459 %b33 = block { # true
460 exit_if # if_22
461 }
462 %b34 = block { # false
463 exit_loop # loop_1
464 }
465 }
466 %276:u32 = load %firstVertex
467 %277:u32 = mul %276, 3u
468 %278:u32 = load %i_3
469 %279:u32 = mul %278, 3u
470 %280:u32 = add %277, %279
471 %281:ptr<storage, f32, read_write> = access %positionsOut, 0u, %280
472 %282:u32 = load %i_3
473 %283:ptr<private, vec3<f32>, read_write> = access %positions, %282
474 %284:f32 = load_vector_element %283, 0u
475 store %281, %284
476 %285:u32 = load %firstVertex
477 %286:u32 = mul %285, 3u
478 %287:u32 = load %i_3
479 %288:u32 = mul %287, 3u
480 %289:u32 = add %286, %288
481 %290:u32 = add %289, 1u
482 %291:ptr<storage, f32, read_write> = access %positionsOut, 0u, %290
483 %292:u32 = load %i_3
484 %293:ptr<private, vec3<f32>, read_write> = access %positions, %292
485 %294:f32 = load_vector_element %293, 1u
486 store %291, %294
487 %295:u32 = load %firstVertex
488 %296:u32 = mul %295, 3u
489 %297:u32 = load %i_3
490 %298:u32 = mul %297, 3u
491 %299:u32 = add %296, %298
492 %300:u32 = add %299, 2u
493 %301:ptr<storage, f32, read_write> = access %positionsOut, 0u, %300
494 %302:u32 = load %i_3
495 %303:ptr<private, vec3<f32>, read_write> = access %positions, %302
496 %304:f32 = load_vector_element %303, 2u
497 store %301, %304
498 %305:u32 = load %firstVertex
499 %306:u32 = mul %305, 3u
500 %307:u32 = load %i_3
501 %308:u32 = mul %307, 3u
502 %309:u32 = add %306, %308
503 %310:ptr<storage, f32, read_write> = access %normalsOut, 0u, %309
504 %311:u32 = load %i_3
505 %312:ptr<private, vec3<f32>, read_write> = access %normals, %311
506 %313:f32 = load_vector_element %312, 0u
507 store %310, %313
508 %314:u32 = load %firstVertex
509 %315:u32 = mul %314, 3u
510 %316:u32 = load %i_3
511 %317:u32 = mul %316, 3u
512 %318:u32 = add %315, %317
513 %319:u32 = add %318, 1u
514 %320:ptr<storage, f32, read_write> = access %normalsOut, 0u, %319
515 %321:u32 = load %i_3
516 %322:ptr<private, vec3<f32>, read_write> = access %normals, %321
517 %323:f32 = load_vector_element %322, 1u
518 store %320, %323
519 %324:u32 = load %firstVertex
520 %325:u32 = mul %324, 3u
521 %326:u32 = load %i_3
522 %327:u32 = mul %326, 3u
523 %328:u32 = add %325, %327
524 %329:u32 = add %328, 2u
525 %330:ptr<storage, f32, read_write> = access %normalsOut, 0u, %329
526 %331:u32 = load %i_3
527 %332:ptr<private, vec3<f32>, read_write> = access %normals, %331
528 %333:f32 = load_vector_element %332, 2u
529 store %330, %333
530 continue %b32
531 }
532 %b32 = block { # continuing
533 %334:u32 = load %i_3
534 %335:u32 = add %334, 1u
535 store %i_3, %335
536 next_iteration %b31
537 }
538 }
539 loop [i: %b35, b: %b36, c: %b37] { # loop_2
540 %b35 = block { # initializer
541 %i_4:ptr<function, u32, read_write> = var, 0u # %i_4: 'i'
542 next_iteration %b36
543 }
544 %b36 = block { # body
545 %337:u32 = load %i_4
546 %338:bool = lt %337, %indexCount
547 if %338 [t: %b38, f: %b39] { # if_23
548 %b38 = block { # true
549 exit_if # if_23
550 }
551 %b39 = block { # false
552 exit_loop # loop_2
553 }
554 }
555 %339:u32 = load %i_4
556 %340:u32 = add %triTableOffset, %339
557 %341:ptr<storage, i32, read> = access %tables, 1u, %340
558 %index_6:i32 = load %341 # %index_6: 'index'
559 %343:u32 = load %i_4
560 %344:u32 = add %firstIndex, %343
561 %345:ptr<storage, u32, read_write> = access %indicesOut, 0u, %344
562 %346:u32 = load %firstVertex
563 %347:ptr<private, u32, read_write> = access %indices, %index_6
564 %348:u32 = load %347
565 %349:u32 = add %346, %348
566 store %345, %349
567 continue %b37
568 }
569 %b37 = block { # continuing
570 %350:u32 = load %i_4
571 %351:u32 = add %350, 1u
572 store %i_4, %351
573 next_iteration %b36
574 }
575 }
576 loop [i: %b40, b: %b41, c: %b42] { # loop_3
577 %b40 = block { # initializer
578 %i_5:ptr<function, u32, read_write> = var, %indexCount # %i_5: 'i'
579 next_iteration %b41
580 }
581 %b41 = block { # body
582 %353:u32 = load %i_5
583 %354:bool = lt %353, 15u
584 if %354 [t: %b43, f: %b44] { # if_24
585 %b43 = block { # true
586 exit_if # if_24
587 }
588 %b44 = block { # false
589 exit_loop # loop_3
590 }
591 }
592 %355:u32 = load %i_5
593 %356:u32 = add %firstIndex, %355
594 %357:ptr<storage, u32, read_write> = access %indicesOut, 0u, %356
595 %358:u32 = load %firstVertex
596 store %357, %358
597 continue %b42
598 }
599 %b42 = block { # continuing
600 %359:u32 = load %i_5
601 %360:u32 = add %359, 1u
602 store %i_5, %360
603 next_iteration %b41
604 }
605 }
606 ret
607 }
608}
609
610unhandled variable address space
611********************************************************************
612* The tint shader compiler has encountered an unexpected error. *
613* *
614* Please help us fix this issue by submitting a bug report at *
615* crbug.com/tint with the source program that triggered the bug. *
616********************************************************************