test/tint/bug/tint/1113.wgsl.expected.ir.msl

SKIP: FAILED

../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: Uniforms = struct @align(16) {
  numTriangles:u32 @offset(0)
  gridSize:u32 @offset(4)
  pad1:u32 @offset(8)
  pad2:u32 @offset(12)
  bbMin:vec3<f32> @offset(16)
  bbMax:vec3<f32> @offset(32)
}

U32s = struct @align(4) {
  values:array<u32> @offset(0)
}

F32s = struct @align(4) {
  values:array<f32> @offset(0)
}

AU32s = struct @align(4) {
  values:array<atomic<u32>> @offset(0)
}

AI32s = struct @align(4) {
  values:array<atomic<i32>> @offset(0)
}

Dbg = struct @align(4) {
  offsetCounter:atomic<u32> @offset(0)
  pad0:u32 @offset(4)
  pad1:u32 @offset(8)
  pad2:u32 @offset(12)
  value0:u32 @offset(16)
  value1:u32 @offset(20)
  value2:u32 @offset(24)
  value3:u32 @offset(28)
  value_f32_0:f32 @offset(32)
  value_f32_1:f32 @offset(36)
  value_f32_2:f32 @offset(40)
  value_f32_3:f32 @offset(44)
}

$B1: {  # root
  %uniforms:ptr<uniform, Uniforms, read> = var @binding_point(0, 0)
  %indices:ptr<storage, U32s, read_write> = var @binding_point(0, 10)
  %positions:ptr<storage, F32s, read_write> = var @binding_point(0, 11)
  %counters:ptr<storage, AU32s, read_write> = var @binding_point(0, 20)
  %LUT:ptr<storage, AI32s, read_write> = var @binding_point(0, 21)
  %dbg:ptr<storage, Dbg, read_write> = var @binding_point(0, 50)
}

%toVoxelPos = func(%position:vec3<f32>):vec3<f32> {
  $B2: {
    %9:ptr<uniform, vec3<f32>, read> = access %uniforms, 4u
    %10:f32 = load_vector_element %9, 0u
    %11:ptr<uniform, vec3<f32>, read> = access %uniforms, 4u
    %12:f32 = load_vector_element %11, 1u
    %13:ptr<uniform, vec3<f32>, read> = access %uniforms, 4u
    %14:f32 = load_vector_element %13, 2u
    %15:vec3<f32> = construct %10, %12, %14
    %bbMin:ptr<function, vec3<f32>, read_write> = var, %15
    %17:ptr<uniform, vec3<f32>, read> = access %uniforms, 5u
    %18:f32 = load_vector_element %17, 0u
    %19:ptr<uniform, vec3<f32>, read> = access %uniforms, 5u
    %20:f32 = load_vector_element %19, 1u
    %21:ptr<uniform, vec3<f32>, read> = access %uniforms, 5u
    %22:f32 = load_vector_element %21, 2u
    %23:vec3<f32> = construct %18, %20, %22
    %bbMax:ptr<function, vec3<f32>, read_write> = var, %23
    %25:vec3<f32> = load %bbMax
    %26:vec3<f32> = load %bbMin
    %27:vec3<f32> = sub %25, %26
    %bbSize:ptr<function, vec3<f32>, read_write> = var, %27
    %29:f32 = load_vector_element %bbSize, 0u
    %30:f32 = load_vector_element %bbSize, 1u
    %31:f32 = max %29, %30
    %32:f32 = load_vector_element %bbSize, 2u
    %33:f32 = max %31, %32
    %cubeSize:ptr<function, f32, read_write> = var, %33
    %35:ptr<uniform, u32, read> = access %uniforms, 1u
    %36:u32 = load %35
    %37:f32 = convert %36
    %gridSize:ptr<function, f32, read_write> = var, %37
    %39:f32 = load %gridSize
    %40:f32 = access %position, 0u
    %41:ptr<uniform, vec3<f32>, read> = access %uniforms, 4u
    %42:f32 = load_vector_element %41, 0u
    %43:f32 = sub %40, %42
    %44:f32 = mul %39, %43
    %45:f32 = load %cubeSize
    %46:f32 = div %44, %45
    %gx:ptr<function, f32, read_write> = var, %46
    %48:f32 = load %gridSize
    %49:f32 = access %position, 1u
    %50:ptr<uniform, vec3<f32>, read> = access %uniforms, 4u
    %51:f32 = load_vector_element %50, 1u
    %52:f32 = sub %49, %51
    %53:f32 = mul %48, %52
    %54:f32 = load %cubeSize
    %55:f32 = div %53, %54
    %gy:ptr<function, f32, read_write> = var, %55
    %57:f32 = load %gridSize
    %58:f32 = access %position, 2u
    %59:ptr<uniform, vec3<f32>, read> = access %uniforms, 4u
    %60:f32 = load_vector_element %59, 2u
    %61:f32 = sub %58, %60
    %62:f32 = mul %57, %61
    %63:f32 = load %cubeSize
    %64:f32 = div %62, %63
    %gz:ptr<function, f32, read_write> = var, %64
    %66:f32 = load %gx
    %67:f32 = load %gy
    %68:f32 = load %gz
    %69:vec3<f32> = construct %66, %67, %68
    ret %69
  }
}
%toIndex1D = func(%gridSize_1:u32, %voxelPos:vec3<f32>):u32 {  # %gridSize_1: 'gridSize'
  $B3: {
    %73:vec3<u32> = call %tint_v3f32_to_v3u32, %voxelPos
    %icoord:ptr<function, vec3<u32>, read_write> = var, %73
    %76:u32 = load_vector_element %icoord, 0u
    %77:u32 = load_vector_element %icoord, 1u
    %78:u32 = mul %gridSize_1, %77
    %79:u32 = add %76, %78
    %80:u32 = mul %gridSize_1, %gridSize_1
    %81:u32 = load_vector_element %icoord, 2u
    %82:u32 = mul %80, %81
    %83:u32 = add %79, %82
    ret %83
  }
}
%toIndex3D = func(%gridSize_2:u32, %index:u32):vec3<u32> {  # %gridSize_2: 'gridSize'
  $B4: {
    %87:u32 = mul %gridSize_2, %gridSize_2
    %88:u32 = call %tint_div_u32, %index, %87
    %z:ptr<function, u32, read_write> = var, %88
    %91:u32 = mul %gridSize_2, %gridSize_2
    %92:u32 = load %z
    %93:u32 = mul %91, %92
    %94:u32 = sub %index, %93
    %95:u32 = call %tint_div_u32, %94, %gridSize_2
    %y:ptr<function, u32, read_write> = var, %95
    %97:u32 = call %tint_mod_u32, %index, %gridSize_2
    %x:ptr<function, u32, read_write> = var, %97
    %100:u32 = load %x
    %101:u32 = load %y
    %102:u32 = load %z
    %103:vec3<u32> = construct %100, %101, %102
    ret %103
  }
}
%loadPosition = func(%vertexIndex:u32):vec3<f32> {
  $B5: {
    %106:u32 = mul 3u, %vertexIndex
    %107:u32 = add %106, 0u
    %108:ptr<storage, f32, read_write> = access %positions, 0u, %107
    %109:f32 = load %108
    %110:u32 = mul 3u, %vertexIndex
    %111:u32 = add %110, 1u
    %112:ptr<storage, f32, read_write> = access %positions, 0u, %111
    %113:f32 = load %112
    %114:u32 = mul 3u, %vertexIndex
    %115:u32 = add %114, 2u
    %116:ptr<storage, f32, read_write> = access %positions, 0u, %115
    %117:f32 = load %116
    %118:vec3<f32> = construct %109, %113, %117
    %position_1:ptr<function, vec3<f32>, read_write> = var, %118  # %position_1: 'position'
    %120:vec3<f32> = load %position_1
    ret %120
  }
}
%doIgnore = func():void {
  $B6: {
    %122:ptr<uniform, u32, read> = access %uniforms, 0u
    %123:u32 = load %122
    %g42:ptr<function, u32, read_write> = var, %123
    %125:ptr<storage, u32, read_write> = access %dbg, 5u
    %126:u32 = load %125
    %kj6:ptr<function, u32, read_write> = var, %126
    %128:ptr<storage, atomic<u32>, read_write> = access %counters, 0u, 0i
    %129:u32 = atomicLoad %128
    %b53:ptr<function, u32, read_write> = var, %129
    %131:ptr<storage, u32, read_write> = access %indices, 0u, 0i
    %132:u32 = load %131
    %rwg:ptr<function, u32, read_write> = var, %132
    %134:ptr<storage, f32, read_write> = access %positions, 0u, 0i
    %135:f32 = load %134
    %rb5:ptr<function, f32, read_write> = var, %135
    %137:ptr<storage, atomic<i32>, read_write> = access %LUT, 0u, 0i
    %138:i32 = atomicLoad %137
    %g55:ptr<function, i32, read_write> = var, %138
    ret
  }
}
%main_count = @compute @workgroup_size(128, 1, 1) func(%GlobalInvocationID:vec3<u32> [@global_invocation_id]):void {
  $B7: {
    %142:u32 = access %GlobalInvocationID, 0u
    %triangleIndex:ptr<function, u32, read_write> = var, %142
    %144:u32 = load %triangleIndex
    %145:ptr<uniform, u32, read> = access %uniforms, 0u
    %146:u32 = load %145
    %147:bool = gte %144, %146
    if %147 [t: $B8] {  # if_1
      $B8: {  # true
        ret
      }
    }
    %148:void = call %doIgnore
    %149:u32 = load %triangleIndex
    %150:u32 = mul 3u, %149
    %151:u32 = add %150, 0u
    %152:ptr<storage, u32, read_write> = access %indices, 0u, %151
    %153:u32 = load %152
    %i0:ptr<function, u32, read_write> = var, %153
    %155:u32 = load %triangleIndex
    %156:u32 = mul 3u, %155
    %157:u32 = add %156, 1u
    %158:ptr<storage, u32, read_write> = access %indices, 0u, %157
    %159:u32 = load %158
    %i1:ptr<function, u32, read_write> = var, %159
    %161:u32 = load %triangleIndex
    %162:u32 = mul 3u, %161
    %163:u32 = add %162, 2u
    %164:ptr<storage, u32, read_write> = access %indices, 0u, %163
    %165:u32 = load %164
    %i2:ptr<function, u32, read_write> = var, %165
    %167:u32 = load %i0
    %168:vec3<f32> = call %loadPosition, %167
    %p0:ptr<function, vec3<f32>, read_write> = var, %168
    %170:u32 = load %i1
    %171:vec3<f32> = call %loadPosition, %170
    %p1:ptr<function, vec3<f32>, read_write> = var, %171
    %173:u32 = load %i2
    %174:vec3<f32> = call %loadPosition, %173
    %p2:ptr<function, vec3<f32>, read_write> = var, %174
    %176:vec3<f32> = load %p0
    %177:vec3<f32> = load %p1
    %178:vec3<f32> = add %176, %177
    %179:vec3<f32> = load %p2
    %180:vec3<f32> = add %178, %179
    %181:vec3<f32> = div %180, 3.0f
    %center:ptr<function, vec3<f32>, read_write> = var, %181
    %183:vec3<f32> = load %center
    %184:vec3<f32> = call %toVoxelPos, %183
    %voxelPos_1:ptr<function, vec3<f32>, read_write> = var, %184  # %voxelPos_1: 'voxelPos'
    %186:ptr<uniform, u32, read> = access %uniforms, 1u
    %187:u32 = load %186
    %188:vec3<f32> = load %voxelPos_1
    %189:u32 = call %toIndex1D, %187, %188
    %voxelIndex:ptr<function, u32, read_write> = var, %189
    %191:u32 = load %voxelIndex
    %192:ptr<storage, atomic<u32>, read_write> = access %counters, 0u, %191
    %193:u32 = atomicAdd %192, 1u
    %acefg:ptr<function, u32, read_write> = var, %193
    %195:u32 = load %triangleIndex
    %196:bool = eq %195, 0u
    if %196 [t: $B9] {  # if_2
      $B9: {  # true
        %197:ptr<storage, u32, read_write> = access %dbg, 4u
        %198:ptr<uniform, u32, read> = access %uniforms, 1u
        %199:u32 = load %198
        store %197, %199
        %200:ptr<storage, f32, read_write> = access %dbg, 8u
        %201:f32 = load_vector_element %center, 0u
        store %200, %201
        %202:ptr<storage, f32, read_write> = access %dbg, 9u
        %203:f32 = load_vector_element %center, 1u
        store %202, %203
        %204:ptr<storage, f32, read_write> = access %dbg, 10u
        %205:f32 = load_vector_element %center, 2u
        store %204, %205
        exit_if  # if_2
      }
    }
    ret
  }
}
%main_create_lut = @compute @workgroup_size(128, 1, 1) func(%GlobalInvocationID_1:vec3<u32> [@global_invocation_id]):void {  # %GlobalInvocationID_1: 'GlobalInvocationID'
  $B10: {
    %208:u32 = access %GlobalInvocationID_1, 0u
    %voxelIndex_1:ptr<function, u32, read_write> = var, %208  # %voxelIndex_1: 'voxelIndex'
    %210:void = call %doIgnore
    %211:ptr<uniform, u32, read> = access %uniforms, 1u
    %212:u32 = load %211
    %213:ptr<uniform, u32, read> = access %uniforms, 1u
    %214:u32 = load %213
    %215:u32 = mul %212, %214
    %216:ptr<uniform, u32, read> = access %uniforms, 1u
    %217:u32 = load %216
    %218:u32 = mul %215, %217
    %maxVoxels:ptr<function, u32, read_write> = var, %218
    %220:u32 = load %voxelIndex_1
    %221:u32 = load %maxVoxels
    %222:bool = gte %220, %221
    if %222 [t: $B11] {  # if_3
      $B11: {  # true
        ret
      }
    }
    %223:u32 = load %voxelIndex_1
    %224:ptr<storage, atomic<u32>, read_write> = access %counters, 0u, %223
    %225:u32 = atomicLoad %224
    %numTriangles:ptr<function, u32, read_write> = var, %225
    %offset:ptr<function, i32, read_write> = var, -1i
    %228:u32 = load %numTriangles
    %229:bool = gt %228, 0u
    if %229 [t: $B12] {  # if_4
      $B12: {  # true
        %230:ptr<storage, atomic<u32>, read_write> = access %dbg, 0u
        %231:u32 = load %numTriangles
        %232:u32 = atomicAdd %230, %231
        %233:i32 = convert %232
        store %offset, %233
        exit_if  # if_4
      }
    }
    %234:u32 = load %voxelIndex_1
    %235:ptr<storage, atomic<i32>, read_write> = access %LUT, 0u, %234
    %236:i32 = load %offset
    %237:void = atomicStore %235, %236
    ret
  }
}
%main_sort_triangles = @compute @workgroup_size(128, 1, 1) func(%GlobalInvocationID_2:vec3<u32> [@global_invocation_id]):void {  # %GlobalInvocationID_2: 'GlobalInvocationID'
  $B13: {
    %240:u32 = access %GlobalInvocationID_2, 0u
    %triangleIndex_1:ptr<function, u32, read_write> = var, %240  # %triangleIndex_1: 'triangleIndex'
    %242:void = call %doIgnore
    %243:u32 = load %triangleIndex_1
    %244:ptr<uniform, u32, read> = access %uniforms, 0u
    %245:u32 = load %244
    %246:bool = gte %243, %245
    if %246 [t: $B14] {  # if_5
      $B14: {  # true
        ret
      }
    }
    %247:u32 = load %triangleIndex_1
    %248:u32 = mul 3u, %247
    %249:u32 = add %248, 0u
    %250:ptr<storage, u32, read_write> = access %indices, 0u, %249
    %251:u32 = load %250
    %i0_1:ptr<function, u32, read_write> = var, %251  # %i0_1: 'i0'
    %253:u32 = load %triangleIndex_1
    %254:u32 = mul 3u, %253
    %255:u32 = add %254, 1u
    %256:ptr<storage, u32, read_write> = access %indices, 0u, %255
    %257:u32 = load %256
    %i1_1:ptr<function, u32, read_write> = var, %257  # %i1_1: 'i1'
    %259:u32 = load %triangleIndex_1
    %260:u32 = mul 3u, %259
    %261:u32 = add %260, 2u
    %262:ptr<storage, u32, read_write> = access %indices, 0u, %261
    %263:u32 = load %262
    %i2_1:ptr<function, u32, read_write> = var, %263  # %i2_1: 'i2'
    %265:u32 = load %i0_1
    %266:vec3<f32> = call %loadPosition, %265
    %p0_1:ptr<function, vec3<f32>, read_write> = var, %266  # %p0_1: 'p0'
    %268:u32 = load %i1_1
    %269:vec3<f32> = call %loadPosition, %268
    %p1_1:ptr<function, vec3<f32>, read_write> = var, %269  # %p1_1: 'p1'
    %271:u32 = load %i2_1
    %272:vec3<f32> = call %loadPosition, %271
    %p2_1:ptr<function, vec3<f32>, read_write> = var, %272  # %p2_1: 'p2'
    %274:vec3<f32> = load %p0_1
    %275:vec3<f32> = load %p1_1
    %276:vec3<f32> = add %274, %275
    %277:vec3<f32> = load %p2_1
    %278:vec3<f32> = add %276, %277
    %279:vec3<f32> = div %278, 3.0f
    %center_1:ptr<function, vec3<f32>, read_write> = var, %279  # %center_1: 'center'
    %281:vec3<f32> = load %center_1
    %282:vec3<f32> = call %toVoxelPos, %281
    %voxelPos_2:ptr<function, vec3<f32>, read_write> = var, %282  # %voxelPos_2: 'voxelPos'
    %284:ptr<uniform, u32, read> = access %uniforms, 1u
    %285:u32 = load %284
    %286:vec3<f32> = load %voxelPos_2
    %287:u32 = call %toIndex1D, %285, %286
    %voxelIndex_2:ptr<function, u32, read_write> = var, %287  # %voxelIndex_2: 'voxelIndex'
    %289:u32 = load %voxelIndex_2
    %290:ptr<storage, atomic<i32>, read_write> = access %LUT, 0u, %289
    %291:i32 = atomicAdd %290, 1i
    %triangleOffset:ptr<function, i32, read_write> = var, %291
    ret
  }
}
%tint_div_u32 = func(%lhs:u32, %rhs:u32):u32 {
  $B15: {
    %295:bool = eq %rhs, 0u
    %296:u32 = select %rhs, 1u, %295
    %297:u32 = div %lhs, %296
    ret %297
  }
}
%tint_mod_u32 = func(%lhs_1:u32, %rhs_1:u32):u32 {  # %lhs_1: 'lhs', %rhs_1: 'rhs'
  $B16: {
    %300:bool = eq %rhs_1, 0u
    %301:u32 = select %rhs_1, 1u, %300
    %302:u32 = let %301
    %303:u32 = div %lhs_1, %302
    %304:u32 = mul %303, %302
    %305:u32 = sub %lhs_1, %304
    ret %305
  }
}
%tint_v3f32_to_v3u32 = func(%value:vec3<f32>):vec3<u32> {
  $B17: {
    %307:vec3<u32> = convert %value
    %308:vec3<bool> = gte %value, vec3<f32>(0.0f)
    %309:vec3<u32> = select vec3<u32>(0u), %307, %308
    %310:vec3<bool> = lte %value, vec3<f32>(4294967040.0f)
    %311:vec3<u32> = select vec3<u32>(4294967295u), %309, %310
    ret %311
  }
}

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.  *
********************************************************************