| SKIP: FAILED |
| |
| ../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: Tables = struct @align(4) { |
| edges:array<u32, 256> @offset(0) |
| tris:array<i32, 4096> @offset(1024) |
| } |
| |
| IsosurfaceVolume = struct @align(16) { |
| min:vec3<f32> @offset(0) |
| max:vec3<f32> @offset(16) |
| step:vec3<f32> @offset(32) |
| size:vec3<u32> @offset(48) |
| threshold:f32 @offset(60) |
| values:array<f32> @offset(64) |
| } |
| |
| PositionBuffer = struct @align(4) { |
| values:array<f32> @offset(0) |
| } |
| |
| NormalBuffer = struct @align(4) { |
| values:array<f32> @offset(0) |
| } |
| |
| IndexBuffer = struct @align(4) { |
| tris:array<u32> @offset(0) |
| } |
| |
| DrawIndirectArgs = struct @align(4) { |
| vc:u32 @offset(0) |
| vertexCount:atomic<u32> @offset(4) |
| firstVertex:u32 @offset(8) |
| firstInstance:u32 @offset(12) |
| indexCount:atomic<u32> @offset(16) |
| indexedInstanceCount:u32 @offset(20) |
| indexedFirstIndex:u32 @offset(24) |
| indexedBaseVertex:u32 @offset(28) |
| indexedFirstInstance:u32 @offset(32) |
| } |
| |
| $B1: { # root |
| %tables:ptr<storage, Tables, read> = var @binding_point(0, 0) |
| %volume:ptr<storage, IsosurfaceVolume, read_write> = var @binding_point(0, 1) |
| %positionsOut:ptr<storage, PositionBuffer, read_write> = var @binding_point(0, 2) |
| %normalsOut:ptr<storage, NormalBuffer, read_write> = var @binding_point(0, 3) |
| %indicesOut:ptr<storage, IndexBuffer, read_write> = var @binding_point(0, 4) |
| %drawOut:ptr<storage, DrawIndirectArgs, read_write> = var @binding_point(0, 5) |
| %positions:ptr<private, array<vec3<f32>, 12>, read_write> = var |
| %normals:ptr<private, array<vec3<f32>, 12>, read_write> = var |
| %indices:ptr<private, array<u32, 12>, read_write> = var |
| %cubeVerts:ptr<private, u32, read_write> = var, 0u |
| } |
| |
| %valueAt = func(%index:vec3<u32>):f32 { |
| $B2: { |
| %13:ptr<storage, vec3<u32>, read_write> = access %volume, 3u |
| %14:vec3<u32> = load %13 |
| %15:vec3<bool> = gte %index, %14 |
| %16:bool = any %15 |
| if %16 [t: $B3] { # if_1 |
| $B3: { # true |
| ret 0.0f |
| } |
| } |
| %17:u32 = access %index, 0u |
| %18:u32 = access %index, 1u |
| %19:ptr<storage, vec3<u32>, read_write> = access %volume, 3u |
| %20:u32 = load_vector_element %19, 0u |
| %21:u32 = mul %18, %20 |
| %22:u32 = add %17, %21 |
| %23:u32 = access %index, 2u |
| %24:ptr<storage, vec3<u32>, read_write> = access %volume, 3u |
| %25:u32 = load_vector_element %24, 0u |
| %26:u32 = mul %23, %25 |
| %27:ptr<storage, vec3<u32>, read_write> = access %volume, 3u |
| %28:u32 = load_vector_element %27, 1u |
| %29:u32 = mul %26, %28 |
| %30:u32 = add %22, %29 |
| %valueIndex:u32 = let %30 |
| %32:ptr<storage, f32, read_write> = access %volume, 5u, %valueIndex |
| %33:f32 = load %32 |
| ret %33 |
| } |
| } |
| %positionAt = func(%index_1:vec3<u32>):vec3<f32> { # %index_1: 'index' |
| $B4: { |
| %36:ptr<storage, vec3<f32>, read_write> = access %volume, 0u |
| %37:vec3<f32> = load %36 |
| %38:vec3<f32> = let %37 |
| %39:ptr<storage, vec3<f32>, read_write> = access %volume, 2u |
| %40:vec3<f32> = load %39 |
| %41:vec3<f32> = let %40 |
| %42:vec3<u32> = swizzle %index_1, xyz |
| %43:vec3<f32> = convert %42 |
| %44:vec3<f32> = mul %41, %43 |
| %45:vec3<f32> = add %38, %44 |
| ret %45 |
| } |
| } |
| %normalAt = func(%index_2:vec3<u32>):vec3<f32> { # %index_2: 'index' |
| $B5: { |
| %48:vec3<u32> = sub %index_2, vec3<u32>(1u, 0u, 0u) |
| %49:f32 = call %valueAt, %48 |
| %50:f32 = let %49 |
| %51:vec3<u32> = add %index_2, vec3<u32>(1u, 0u, 0u) |
| %52:f32 = call %valueAt, %51 |
| %53:f32 = sub %50, %52 |
| %54:f32 = let %53 |
| %55:vec3<u32> = sub %index_2, vec3<u32>(0u, 1u, 0u) |
| %56:f32 = call %valueAt, %55 |
| %57:f32 = let %56 |
| %58:vec3<u32> = add %index_2, vec3<u32>(0u, 1u, 0u) |
| %59:f32 = call %valueAt, %58 |
| %60:f32 = sub %57, %59 |
| %61:f32 = let %60 |
| %62:vec3<u32> = sub %index_2, vec3<u32>(0u, 0u, 1u) |
| %63:f32 = call %valueAt, %62 |
| %64:f32 = let %63 |
| %65:vec3<u32> = add %index_2, vec3<u32>(0u, 0u, 1u) |
| %66:f32 = call %valueAt, %65 |
| %67:f32 = sub %64, %66 |
| %68:vec3<f32> = construct %54, %61, %67 |
| ret %68 |
| } |
| } |
| %interpX = func(%index_3:u32, %i:vec3<u32>, %va:f32, %vb:f32):void { # %index_3: 'index' |
| $B6: { |
| %74:ptr<storage, f32, read_write> = access %volume, 4u |
| %75:f32 = load %74 |
| %76:f32 = sub %75, %va |
| %77:f32 = sub %vb, %va |
| %78:f32 = div %76, %77 |
| %mu:f32 = let %78 |
| %80:u32 = load %cubeVerts |
| %81:ptr<private, vec3<f32>, read_write> = access %positions, %80 |
| %82:ptr<private, vec3<f32>, read_write> = let %81 |
| %83:vec3<f32> = call %positionAt, %i |
| %84:vec3<f32> = let %83 |
| %85:ptr<storage, vec3<f32>, read_write> = access %volume, 2u |
| %86:f32 = load_vector_element %85, 0u |
| %87:f32 = mul %86, %mu |
| %88:vec3<f32> = construct %87, 0.0f, 0.0f |
| %89:vec3<f32> = add %84, %88 |
| store %82, %89 |
| %90:vec3<f32> = call %normalAt, %i |
| %na:vec3<f32> = let %90 |
| %92:vec3<u32> = add %i, vec3<u32>(1u, 0u, 0u) |
| %93:vec3<f32> = call %normalAt, %92 |
| %nb:vec3<f32> = let %93 |
| %95:u32 = load %cubeVerts |
| %96:ptr<private, vec3<f32>, read_write> = access %normals, %95 |
| %97:ptr<private, vec3<f32>, read_write> = let %96 |
| %98:vec3<f32> = construct %mu, %mu, %mu |
| %99:vec3<f32> = mix %na, %nb, %98 |
| store %97, %99 |
| %100:ptr<private, u32, read_write> = access %indices, %index_3 |
| %101:u32 = load %cubeVerts |
| store %100, %101 |
| %102:u32 = load %cubeVerts |
| %103:u32 = add %102, 1u |
| store %cubeVerts, %103 |
| ret |
| } |
| } |
| %interpY = func(%index_4:u32, %i_1:vec3<u32>, %va_1:f32, %vb_1:f32):void { # %index_4: 'index', %i_1: 'i', %va_1: 'va', %vb_1: 'vb' |
| $B7: { |
| %109:ptr<storage, f32, read_write> = access %volume, 4u |
| %110:f32 = load %109 |
| %111:f32 = sub %110, %va_1 |
| %112:f32 = sub %vb_1, %va_1 |
| %113:f32 = div %111, %112 |
| %mu_1:f32 = let %113 # %mu_1: 'mu' |
| %115:u32 = load %cubeVerts |
| %116:ptr<private, vec3<f32>, read_write> = access %positions, %115 |
| %117:ptr<private, vec3<f32>, read_write> = let %116 |
| %118:vec3<f32> = call %positionAt, %i_1 |
| %119:vec3<f32> = let %118 |
| %120:ptr<storage, vec3<f32>, read_write> = access %volume, 2u |
| %121:f32 = load_vector_element %120, 1u |
| %122:f32 = mul %121, %mu_1 |
| %123:vec3<f32> = construct 0.0f, %122, 0.0f |
| %124:vec3<f32> = add %119, %123 |
| store %117, %124 |
| %125:vec3<f32> = call %normalAt, %i_1 |
| %na_1:vec3<f32> = let %125 # %na_1: 'na' |
| %127:vec3<u32> = add %i_1, vec3<u32>(0u, 1u, 0u) |
| %128:vec3<f32> = call %normalAt, %127 |
| %nb_1:vec3<f32> = let %128 # %nb_1: 'nb' |
| %130:u32 = load %cubeVerts |
| %131:ptr<private, vec3<f32>, read_write> = access %normals, %130 |
| %132:ptr<private, vec3<f32>, read_write> = let %131 |
| %133:vec3<f32> = construct %mu_1, %mu_1, %mu_1 |
| %134:vec3<f32> = mix %na_1, %nb_1, %133 |
| store %132, %134 |
| %135:ptr<private, u32, read_write> = access %indices, %index_4 |
| %136:u32 = load %cubeVerts |
| store %135, %136 |
| %137:u32 = load %cubeVerts |
| %138:u32 = add %137, 1u |
| store %cubeVerts, %138 |
| ret |
| } |
| } |
| %interpZ = func(%index_5:u32, %i_2:vec3<u32>, %va_2:f32, %vb_2:f32):void { # %index_5: 'index', %i_2: 'i', %va_2: 'va', %vb_2: 'vb' |
| $B8: { |
| %144:ptr<storage, f32, read_write> = access %volume, 4u |
| %145:f32 = load %144 |
| %146:f32 = sub %145, %va_2 |
| %147:f32 = sub %vb_2, %va_2 |
| %148:f32 = div %146, %147 |
| %mu_2:f32 = let %148 # %mu_2: 'mu' |
| %150:u32 = load %cubeVerts |
| %151:ptr<private, vec3<f32>, read_write> = access %positions, %150 |
| %152:ptr<private, vec3<f32>, read_write> = let %151 |
| %153:vec3<f32> = call %positionAt, %i_2 |
| %154:vec3<f32> = let %153 |
| %155:ptr<storage, vec3<f32>, read_write> = access %volume, 2u |
| %156:f32 = load_vector_element %155, 2u |
| %157:f32 = mul %156, %mu_2 |
| %158:vec3<f32> = construct 0.0f, 0.0f, %157 |
| %159:vec3<f32> = add %154, %158 |
| store %152, %159 |
| %160:vec3<f32> = call %normalAt, %i_2 |
| %na_2:vec3<f32> = let %160 # %na_2: 'na' |
| %162:vec3<u32> = add %i_2, vec3<u32>(0u, 0u, 1u) |
| %163:vec3<f32> = call %normalAt, %162 |
| %nb_2:vec3<f32> = let %163 # %nb_2: 'nb' |
| %165:u32 = load %cubeVerts |
| %166:ptr<private, vec3<f32>, read_write> = access %normals, %165 |
| %167:ptr<private, vec3<f32>, read_write> = let %166 |
| %168:vec3<f32> = construct %mu_2, %mu_2, %mu_2 |
| %169:vec3<f32> = mix %na_2, %nb_2, %168 |
| store %167, %169 |
| %170:ptr<private, u32, read_write> = access %indices, %index_5 |
| %171:u32 = load %cubeVerts |
| store %170, %171 |
| %172:u32 = load %cubeVerts |
| %173:u32 = add %172, 1u |
| store %cubeVerts, %173 |
| ret |
| } |
| } |
| %computeMain = @compute @workgroup_size(4, 4, 4) func(%global_id:vec3<u32> [@global_invocation_id]):void { |
| $B9: { |
| %i0:vec3<u32> = let %global_id |
| %177:vec3<u32> = add %global_id, vec3<u32>(1u, 0u, 0u) |
| %i1:vec3<u32> = let %177 |
| %179:vec3<u32> = add %global_id, vec3<u32>(1u, 1u, 0u) |
| %i2:vec3<u32> = let %179 |
| %181:vec3<u32> = add %global_id, vec3<u32>(0u, 1u, 0u) |
| %i3:vec3<u32> = let %181 |
| %183:vec3<u32> = add %global_id, vec3<u32>(0u, 0u, 1u) |
| %i4:vec3<u32> = let %183 |
| %185:vec3<u32> = add %global_id, vec3<u32>(1u, 0u, 1u) |
| %i5:vec3<u32> = let %185 |
| %187:vec3<u32> = add %global_id, vec3<u32>(1u) |
| %i6:vec3<u32> = let %187 |
| %189:vec3<u32> = add %global_id, vec3<u32>(0u, 1u, 1u) |
| %i7:vec3<u32> = let %189 |
| %191:f32 = call %valueAt, %i0 |
| %v0:f32 = let %191 |
| %193:f32 = call %valueAt, %i1 |
| %v1:f32 = let %193 |
| %195:f32 = call %valueAt, %i2 |
| %v2:f32 = let %195 |
| %197:f32 = call %valueAt, %i3 |
| %v3:f32 = let %197 |
| %199:f32 = call %valueAt, %i4 |
| %v4:f32 = let %199 |
| %201:f32 = call %valueAt, %i5 |
| %v5:f32 = let %201 |
| %203:f32 = call %valueAt, %i6 |
| %v6:f32 = let %203 |
| %205:f32 = call %valueAt, %i7 |
| %v7:f32 = let %205 |
| %cubeIndex:ptr<function, u32, read_write> = var, 0u |
| %208:ptr<storage, f32, read_write> = access %volume, 4u |
| %209:f32 = load %208 |
| %210:bool = lt %v0, %209 |
| if %210 [t: $B10] { # if_2 |
| $B10: { # true |
| %211:u32 = load %cubeIndex |
| %212:u32 = or %211, 1u |
| store %cubeIndex, %212 |
| exit_if # if_2 |
| } |
| } |
| %213:ptr<storage, f32, read_write> = access %volume, 4u |
| %214:f32 = load %213 |
| %215:bool = lt %v1, %214 |
| if %215 [t: $B11] { # if_3 |
| $B11: { # true |
| %216:u32 = load %cubeIndex |
| %217:u32 = or %216, 2u |
| store %cubeIndex, %217 |
| exit_if # if_3 |
| } |
| } |
| %218:ptr<storage, f32, read_write> = access %volume, 4u |
| %219:f32 = load %218 |
| %220:bool = lt %v2, %219 |
| if %220 [t: $B12] { # if_4 |
| $B12: { # true |
| %221:u32 = load %cubeIndex |
| %222:u32 = or %221, 4u |
| store %cubeIndex, %222 |
| exit_if # if_4 |
| } |
| } |
| %223:ptr<storage, f32, read_write> = access %volume, 4u |
| %224:f32 = load %223 |
| %225:bool = lt %v3, %224 |
| if %225 [t: $B13] { # if_5 |
| $B13: { # true |
| %226:u32 = load %cubeIndex |
| %227:u32 = or %226, 8u |
| store %cubeIndex, %227 |
| exit_if # if_5 |
| } |
| } |
| %228:ptr<storage, f32, read_write> = access %volume, 4u |
| %229:f32 = load %228 |
| %230:bool = lt %v4, %229 |
| if %230 [t: $B14] { # if_6 |
| $B14: { # true |
| %231:u32 = load %cubeIndex |
| %232:u32 = or %231, 16u |
| store %cubeIndex, %232 |
| exit_if # if_6 |
| } |
| } |
| %233:ptr<storage, f32, read_write> = access %volume, 4u |
| %234:f32 = load %233 |
| %235:bool = lt %v5, %234 |
| if %235 [t: $B15] { # if_7 |
| $B15: { # true |
| %236:u32 = load %cubeIndex |
| %237:u32 = or %236, 32u |
| store %cubeIndex, %237 |
| exit_if # if_7 |
| } |
| } |
| %238:ptr<storage, f32, read_write> = access %volume, 4u |
| %239:f32 = load %238 |
| %240:bool = lt %v6, %239 |
| if %240 [t: $B16] { # if_8 |
| $B16: { # true |
| %241:u32 = load %cubeIndex |
| %242:u32 = or %241, 64u |
| store %cubeIndex, %242 |
| exit_if # if_8 |
| } |
| } |
| %243:ptr<storage, f32, read_write> = access %volume, 4u |
| %244:f32 = load %243 |
| %245:bool = lt %v7, %244 |
| if %245 [t: $B17] { # if_9 |
| $B17: { # true |
| %246:u32 = load %cubeIndex |
| %247:u32 = or %246, 128u |
| store %cubeIndex, %247 |
| exit_if # if_9 |
| } |
| } |
| %248:u32 = load %cubeIndex |
| %249:ptr<storage, u32, read> = access %tables, 0u, %248 |
| %250:u32 = load %249 |
| %edges:u32 = let %250 |
| %252:u32 = and %edges, 1u |
| %253:bool = neq %252, 0u |
| if %253 [t: $B18] { # if_10 |
| $B18: { # true |
| %254:void = call %interpX, 0u, %i0, %v0, %v1 |
| exit_if # if_10 |
| } |
| } |
| %255:u32 = and %edges, 2u |
| %256:bool = neq %255, 0u |
| if %256 [t: $B19] { # if_11 |
| $B19: { # true |
| %257:void = call %interpY, 1u, %i1, %v1, %v2 |
| exit_if # if_11 |
| } |
| } |
| %258:u32 = and %edges, 4u |
| %259:bool = neq %258, 0u |
| if %259 [t: $B20] { # if_12 |
| $B20: { # true |
| %260:void = call %interpX, 2u, %i3, %v3, %v2 |
| exit_if # if_12 |
| } |
| } |
| %261:u32 = and %edges, 8u |
| %262:bool = neq %261, 0u |
| if %262 [t: $B21] { # if_13 |
| $B21: { # true |
| %263:void = call %interpY, 3u, %i0, %v0, %v3 |
| exit_if # if_13 |
| } |
| } |
| %264:u32 = and %edges, 16u |
| %265:bool = neq %264, 0u |
| if %265 [t: $B22] { # if_14 |
| $B22: { # true |
| %266:void = call %interpX, 4u, %i4, %v4, %v5 |
| exit_if # if_14 |
| } |
| } |
| %267:u32 = and %edges, 32u |
| %268:bool = neq %267, 0u |
| if %268 [t: $B23] { # if_15 |
| $B23: { # true |
| %269:void = call %interpY, 5u, %i5, %v5, %v6 |
| exit_if # if_15 |
| } |
| } |
| %270:u32 = and %edges, 64u |
| %271:bool = neq %270, 0u |
| if %271 [t: $B24] { # if_16 |
| $B24: { # true |
| %272:void = call %interpX, 6u, %i7, %v7, %v6 |
| exit_if # if_16 |
| } |
| } |
| %273:u32 = and %edges, 128u |
| %274:bool = neq %273, 0u |
| if %274 [t: $B25] { # if_17 |
| $B25: { # true |
| %275:void = call %interpY, 7u, %i4, %v4, %v7 |
| exit_if # if_17 |
| } |
| } |
| %276:u32 = and %edges, 256u |
| %277:bool = neq %276, 0u |
| if %277 [t: $B26] { # if_18 |
| $B26: { # true |
| %278:void = call %interpZ, 8u, %i0, %v0, %v4 |
| exit_if # if_18 |
| } |
| } |
| %279:u32 = and %edges, 512u |
| %280:bool = neq %279, 0u |
| if %280 [t: $B27] { # if_19 |
| $B27: { # true |
| %281:void = call %interpZ, 9u, %i1, %v1, %v5 |
| exit_if # if_19 |
| } |
| } |
| %282:u32 = and %edges, 1024u |
| %283:bool = neq %282, 0u |
| if %283 [t: $B28] { # if_20 |
| $B28: { # true |
| %284:void = call %interpZ, 10u, %i2, %v2, %v6 |
| exit_if # if_20 |
| } |
| } |
| %285:u32 = and %edges, 2048u |
| %286:bool = neq %285, 0u |
| if %286 [t: $B29] { # if_21 |
| $B29: { # true |
| %287:void = call %interpZ, 11u, %i3, %v3, %v7 |
| exit_if # if_21 |
| } |
| } |
| %288:u32 = load %cubeIndex |
| %289:u32 = and 4u, 31u |
| %290:u32 = shl %288, %289 |
| %291:u32 = add %290, 1u |
| %triTableOffset:u32 = let %291 |
| %293:u32 = sub %triTableOffset, 1u |
| %294:ptr<storage, i32, read> = access %tables, 1u, %293 |
| %295:i32 = load %294 |
| %296:u32 = convert %295 |
| %indexCount:u32 = let %296 |
| %298:ptr<storage, atomic<u32>, read_write> = access %drawOut, 1u |
| %299:u32 = load %cubeVerts |
| %300:u32 = atomicAdd %298, %299 |
| %firstVertex:ptr<function, u32, read_write> = var, %300 |
| %302:u32 = access %global_id, 0u |
| %303:u32 = access %global_id, 1u |
| %304:ptr<storage, vec3<u32>, read_write> = access %volume, 3u |
| %305:u32 = load_vector_element %304, 0u |
| %306:u32 = mul %303, %305 |
| %307:u32 = add %302, %306 |
| %308:u32 = access %global_id, 2u |
| %309:ptr<storage, vec3<u32>, read_write> = access %volume, 3u |
| %310:u32 = load_vector_element %309, 0u |
| %311:u32 = mul %308, %310 |
| %312:ptr<storage, vec3<u32>, read_write> = access %volume, 3u |
| %313:u32 = load_vector_element %312, 1u |
| %314:u32 = mul %311, %313 |
| %315:u32 = add %307, %314 |
| %bufferOffset:u32 = let %315 |
| %317:u32 = mul %bufferOffset, 15u |
| %firstIndex:u32 = let %317 |
| loop [i: $B30, b: $B31, c: $B32] { # loop_1 |
| $B30: { # initializer |
| %i_3:ptr<function, u32, read_write> = var, 0u # %i_3: 'i' |
| next_iteration # -> $B31 |
| } |
| $B31: { # body |
| %320:u32 = load %i_3 |
| %321:u32 = load %cubeVerts |
| %322:bool = lt %320, %321 |
| if %322 [t: $B33, f: $B34] { # if_22 |
| $B33: { # true |
| exit_if # if_22 |
| } |
| $B34: { # false |
| exit_loop # loop_1 |
| } |
| } |
| %323:u32 = load %firstVertex |
| %324:u32 = mul %323, 3u |
| %325:u32 = load %i_3 |
| %326:u32 = mul %325, 3u |
| %327:u32 = add %324, %326 |
| %328:ptr<storage, f32, read_write> = access %positionsOut, 0u, %327 |
| %329:u32 = load %i_3 |
| %330:ptr<private, vec3<f32>, read_write> = access %positions, %329 |
| %331:f32 = load_vector_element %330, 0u |
| store %328, %331 |
| %332:u32 = load %firstVertex |
| %333:u32 = mul %332, 3u |
| %334:u32 = load %i_3 |
| %335:u32 = mul %334, 3u |
| %336:u32 = add %333, %335 |
| %337:u32 = add %336, 1u |
| %338:ptr<storage, f32, read_write> = access %positionsOut, 0u, %337 |
| %339:u32 = load %i_3 |
| %340:ptr<private, vec3<f32>, read_write> = access %positions, %339 |
| %341:f32 = load_vector_element %340, 1u |
| store %338, %341 |
| %342:u32 = load %firstVertex |
| %343:u32 = mul %342, 3u |
| %344:u32 = load %i_3 |
| %345:u32 = mul %344, 3u |
| %346:u32 = add %343, %345 |
| %347:u32 = add %346, 2u |
| %348:ptr<storage, f32, read_write> = access %positionsOut, 0u, %347 |
| %349:u32 = load %i_3 |
| %350:ptr<private, vec3<f32>, read_write> = access %positions, %349 |
| %351:f32 = load_vector_element %350, 2u |
| store %348, %351 |
| %352:u32 = load %firstVertex |
| %353:u32 = mul %352, 3u |
| %354:u32 = load %i_3 |
| %355:u32 = mul %354, 3u |
| %356:u32 = add %353, %355 |
| %357:ptr<storage, f32, read_write> = access %normalsOut, 0u, %356 |
| %358:u32 = load %i_3 |
| %359:ptr<private, vec3<f32>, read_write> = access %normals, %358 |
| %360:f32 = load_vector_element %359, 0u |
| store %357, %360 |
| %361:u32 = load %firstVertex |
| %362:u32 = mul %361, 3u |
| %363:u32 = load %i_3 |
| %364:u32 = mul %363, 3u |
| %365:u32 = add %362, %364 |
| %366:u32 = add %365, 1u |
| %367:ptr<storage, f32, read_write> = access %normalsOut, 0u, %366 |
| %368:u32 = load %i_3 |
| %369:ptr<private, vec3<f32>, read_write> = access %normals, %368 |
| %370:f32 = load_vector_element %369, 1u |
| store %367, %370 |
| %371:u32 = load %firstVertex |
| %372:u32 = mul %371, 3u |
| %373:u32 = load %i_3 |
| %374:u32 = mul %373, 3u |
| %375:u32 = add %372, %374 |
| %376:u32 = add %375, 2u |
| %377:ptr<storage, f32, read_write> = access %normalsOut, 0u, %376 |
| %378:u32 = load %i_3 |
| %379:ptr<private, vec3<f32>, read_write> = access %normals, %378 |
| %380:f32 = load_vector_element %379, 2u |
| store %377, %380 |
| continue # -> $B32 |
| } |
| $B32: { # continuing |
| %381:u32 = load %i_3 |
| %382:u32 = add %381, 1u |
| store %i_3, %382 |
| next_iteration # -> $B31 |
| } |
| } |
| loop [i: $B35, b: $B36, c: $B37] { # loop_2 |
| $B35: { # initializer |
| %i_4:ptr<function, u32, read_write> = var, 0u # %i_4: 'i' |
| next_iteration # -> $B36 |
| } |
| $B36: { # body |
| %384:u32 = load %i_4 |
| %385:bool = lt %384, %indexCount |
| if %385 [t: $B38, f: $B39] { # if_23 |
| $B38: { # true |
| exit_if # if_23 |
| } |
| $B39: { # false |
| exit_loop # loop_2 |
| } |
| } |
| %386:u32 = load %i_4 |
| %387:u32 = add %triTableOffset, %386 |
| %388:ptr<storage, i32, read> = access %tables, 1u, %387 |
| %389:i32 = load %388 |
| %index_6:i32 = let %389 # %index_6: 'index' |
| %391:u32 = load %i_4 |
| %392:u32 = add %firstIndex, %391 |
| %393:ptr<storage, u32, read_write> = access %indicesOut, 0u, %392 |
| %394:u32 = load %firstVertex |
| %395:ptr<private, u32, read_write> = access %indices, %index_6 |
| %396:u32 = load %395 |
| %397:u32 = add %394, %396 |
| store %393, %397 |
| continue # -> $B37 |
| } |
| $B37: { # continuing |
| %398:u32 = load %i_4 |
| %399:u32 = add %398, 1u |
| store %i_4, %399 |
| next_iteration # -> $B36 |
| } |
| } |
| loop [i: $B40, b: $B41, c: $B42] { # loop_3 |
| $B40: { # initializer |
| %i_5:ptr<function, u32, read_write> = var, %indexCount # %i_5: 'i' |
| next_iteration # -> $B41 |
| } |
| $B41: { # body |
| %401:u32 = load %i_5 |
| %402:bool = lt %401, 15u |
| if %402 [t: $B43, f: $B44] { # if_24 |
| $B43: { # true |
| exit_if # if_24 |
| } |
| $B44: { # false |
| exit_loop # loop_3 |
| } |
| } |
| %403:u32 = load %i_5 |
| %404:u32 = add %firstIndex, %403 |
| %405:ptr<storage, u32, read_write> = access %indicesOut, 0u, %404 |
| %406:u32 = load %firstVertex |
| store %405, %406 |
| continue # -> $B42 |
| } |
| $B42: { # continuing |
| %407:u32 = load %i_5 |
| %408:u32 = add %407, 1u |
| store %i_5, %408 |
| next_iteration # -> $B41 |
| } |
| } |
| ret |
| } |
| } |
| |
| unhandled variable address space |
| ******************************************************************** |
| * The tint shader compiler has encountered an unexpected error. * |
| * * |
| * Please help us fix this issue by submitting a bug report at * |
| * crbug.com/tint with the source program that triggered the bug. * |
| ******************************************************************** |
