dan sinclair | f1f381a | 2023-11-22 09:44:15 +0000 | [diff] [blame] | 1 | SKIP: 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 | |
| 8 | IsosurfaceVolume = 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 | |
| 17 | PositionBuffer = struct @align(4) { |
| 18 | values:array<f32> @offset(0) |
| 19 | } |
| 20 | |
| 21 | NormalBuffer = struct @align(4) { |
| 22 | values:array<f32> @offset(0) |
| 23 | } |
| 24 | |
| 25 | IndexBuffer = struct @align(4) { |
| 26 | tris:array<u32> @offset(0) |
| 27 | } |
| 28 | |
| 29 | DrawIndirectArgs = 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 | |
| 610 | unhandled 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 | ******************************************************************** |