test/tint/benchmark/cluster-lights.wgsl.expected.ir.msl

SKIP: FAILED

../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: Camera = struct @align(16) {
  projection:mat4x4<f32> @offset(0)
  inverseProjection:mat4x4<f32> @offset(64)
  view:mat4x4<f32> @offset(128)
  position:vec3<f32> @offset(192)
  time:f32 @offset(204)
  outputSize:vec2<f32> @offset(208)
  zNear:f32 @offset(216)
  zFar:f32 @offset(220)
}

ClusterBounds = struct @align(16) {
  minAABB:vec3<f32> @offset(0)
  maxAABB:vec3<f32> @offset(16)
}

Clusters = struct @align(16) {
  bounds:array<ClusterBounds, 27648> @offset(0)
}

ClusterLights = struct @align(4) {
  offset:u32 @offset(0)
  count:u32 @offset(4)
}

ClusterLightGroup = struct @align(4) {
  offset:atomic<u32> @offset(0)
  lights:array<ClusterLights, 27648> @offset(4)
  indices:array<u32, 1769472> @offset(221188)
}

Light = struct @align(16) {
  position:vec3<f32> @offset(0)
  range:f32 @offset(12)
  color:vec3<f32> @offset(16)
  intensity:f32 @offset(28)
}

GlobalLights = struct @align(16) {
  ambient:vec3<f32> @offset(0)
  dirColor:vec3<f32> @offset(16)
  dirIntensity:f32 @offset(28)
  dirDirection:vec3<f32> @offset(32)
  lightCount:u32 @offset(44)
  lights:array<Light> @offset(48)
}

$B1: {  # root
  %camera:ptr<uniform, Camera, read> = var @binding_point(0, 0)
  %clusters:ptr<storage, Clusters, read> = var @binding_point(0, 1)
  %clusterLights:ptr<storage, ClusterLightGroup, read_write> = var @binding_point(0, 2)
  %globalLights:ptr<storage, GlobalLights, read> = var @binding_point(0, 3)
}

%linearDepth = func(%depthSample:f32):f32 {
  $B2: {
    %7:ptr<uniform, f32, read> = access %camera, 7u
    %8:f32 = load %7
    %9:ptr<uniform, f32, read> = access %camera, 6u
    %10:f32 = load %9
    %11:f32 = mul %8, %10
    %12:f32 = let %11
    %13:ptr<uniform, f32, read> = access %camera, 6u
    %14:f32 = load %13
    %15:ptr<uniform, f32, read> = access %camera, 7u
    %16:f32 = load %15
    %17:f32 = sub %14, %16
    %18:ptr<uniform, f32, read> = access %camera, 7u
    %19:f32 = load %18
    %20:f32 = fma %depthSample, %17, %19
    %21:f32 = div %12, %20
    ret %21
  }
}
%getTile = func(%fragCoord:vec4<f32>):vec3<u32> {
  $B3: {
    %24:ptr<uniform, f32, read> = access %camera, 7u
    %25:f32 = load %24
    %26:ptr<uniform, f32, read> = access %camera, 6u
    %27:f32 = load %26
    %28:f32 = div %25, %27
    %29:f32 = log2 %28
    %30:f32 = div 48.0f, %29
    %sliceScale:f32 = let %30
    %32:ptr<uniform, f32, read> = access %camera, 6u
    %33:f32 = load %32
    %34:f32 = log2 %33
    %35:f32 = mul 48.0f, %34
    %36:f32 = let %35
    %37:ptr<uniform, f32, read> = access %camera, 7u
    %38:f32 = load %37
    %39:ptr<uniform, f32, read> = access %camera, 6u
    %40:f32 = load %39
    %41:f32 = div %38, %40
    %42:f32 = log2 %41
    %43:f32 = div %36, %42
    %44:f32 = negation %43
    %sliceBias:f32 = let %44
    %46:f32 = access %fragCoord, 2u
    %47:f32 = call %linearDepth, %46
    %48:f32 = log2 %47
    %49:f32 = mul %48, %sliceScale
    %50:f32 = add %49, %sliceBias
    %51:f32 = max %50, 0.0f
    %52:u32 = call %tint_f32_to_u32, %51
    %zTile:u32 = let %52
    %55:f32 = access %fragCoord, 0u
    %56:ptr<uniform, vec2<f32>, read> = access %camera, 5u
    %57:f32 = load_vector_element %56, 0u
    %58:f32 = div %57, 32.0f
    %59:f32 = div %55, %58
    %60:u32 = call %tint_f32_to_u32, %59
    %61:u32 = let %60
    %62:f32 = access %fragCoord, 1u
    %63:ptr<uniform, vec2<f32>, read> = access %camera, 5u
    %64:f32 = load_vector_element %63, 1u
    %65:f32 = div %64, 18.0f
    %66:f32 = div %62, %65
    %67:u32 = call %tint_f32_to_u32, %66
    %68:vec3<u32> = construct %61, %67, %zTile
    ret %68
  }
}
%getClusterIndex = func(%fragCoord_1:vec4<f32>):u32 {  # %fragCoord_1: 'fragCoord'
  $B4: {
    %71:vec3<u32> = call %getTile, %fragCoord_1
    %tile:vec3<u32> = let %71
    %73:u32 = access %tile, 0u
    %74:u32 = access %tile, 1u
    %75:u32 = mul %74, 32u
    %76:u32 = add %73, %75
    %77:u32 = access %tile, 2u
    %78:u32 = mul %77, 32u
    %79:u32 = mul %78, 18u
    %80:u32 = add %76, %79
    ret %80
  }
}
%sqDistPointAABB = func(%p:vec3<f32>, %minAABB:vec3<f32>, %maxAABB:vec3<f32>):f32 {
  $B5: {
    %sqDist:ptr<function, f32, read_write> = var, 0.0f
    loop [i: $B6, b: $B7, c: $B8] {  # loop_1
      $B6: {  # initializer
        %i:ptr<function, i32, read_write> = var, 0i
        next_iteration  # -> $B7
      }
      $B7: {  # body
        %87:i32 = load %i
        %88:bool = lt %87, 3i
        if %88 [t: $B9, f: $B10] {  # if_1
          $B9: {  # true
            exit_if  # if_1
          }
          $B10: {  # false
            exit_loop  # loop_1
          }
        }
        %89:i32 = load %i
        %90:f32 = access %p, %89
        %v:f32 = let %90
        %92:i32 = load %i
        %93:f32 = access %minAABB, %92
        %94:bool = lt %v, %93
        if %94 [t: $B11] {  # if_2
          $B11: {  # true
            %95:f32 = load %sqDist
            %96:i32 = load %i
            %97:f32 = access %minAABB, %96
            %98:f32 = sub %97, %v
            %99:i32 = load %i
            %100:f32 = access %minAABB, %99
            %101:f32 = sub %100, %v
            %102:f32 = mul %98, %101
            %103:f32 = add %95, %102
            store %sqDist, %103
            exit_if  # if_2
          }
        }
        %104:i32 = load %i
        %105:f32 = access %maxAABB, %104
        %106:bool = gt %v, %105
        if %106 [t: $B12] {  # if_3
          $B12: {  # true
            %107:f32 = load %sqDist
            %108:i32 = load %i
            %109:f32 = access %maxAABB, %108
            %110:f32 = sub %v, %109
            %111:i32 = load %i
            %112:f32 = access %maxAABB, %111
            %113:f32 = sub %v, %112
            %114:f32 = mul %110, %113
            %115:f32 = add %107, %114
            store %sqDist, %115
            exit_if  # if_3
          }
        }
        continue  # -> $B8
      }
      $B8: {  # continuing
        %116:i32 = load %i
        %117:i32 = add %116, 1i
        store %i, %117
        next_iteration  # -> $B7
      }
    }
    %118:f32 = load %sqDist
    ret %118
  }
}
%computeMain = @compute @workgroup_size(4, 2, 4) func(%global_id:vec3<u32> [@global_invocation_id]):void {
  $B13: {
    %121:u32 = access %global_id, 0u
    %122:u32 = access %global_id, 1u
    %123:u32 = mul %122, 32u
    %124:u32 = add %121, %123
    %125:u32 = access %global_id, 2u
    %126:u32 = mul %125, 32u
    %127:u32 = mul %126, 18u
    %128:u32 = add %124, %127
    %tileIndex:u32 = let %128
    %clusterLightCount:ptr<function, u32, read_write> = var, 0u
    %cluserLightIndices:ptr<function, array<u32, 256>, read_write> = var
    loop [i: $B14, b: $B15, c: $B16] {  # loop_2
      $B14: {  # initializer
        %i_1:ptr<function, u32, read_write> = var, 0u  # %i_1: 'i'
        next_iteration  # -> $B15
      }
      $B15: {  # body
        %133:u32 = load %i_1
        %134:ptr<storage, u32, read> = access %globalLights, 4u
        %135:u32 = load %134
        %136:bool = lt %133, %135
        if %136 [t: $B17, f: $B18] {  # if_4
          $B17: {  # true
            exit_if  # if_4
          }
          $B18: {  # false
            exit_loop  # loop_2
          }
        }
        %137:u32 = load %i_1
        %138:ptr<storage, f32, read> = access %globalLights, 5u, %137, 1u
        %139:f32 = load %138
        %range:f32 = let %139
        %141:bool = lte %range, 0.0f
        %lightInCluster:ptr<function, bool, read_write> = var, %141
        %143:bool = load %lightInCluster
        %144:bool = eq %143, false
        if %144 [t: $B19] {  # if_5
          $B19: {  # true
            %145:ptr<uniform, mat4x4<f32>, read> = access %camera, 2u
            %146:mat4x4<f32> = load %145
            %147:mat4x4<f32> = let %146
            %148:u32 = load %i_1
            %149:ptr<storage, vec3<f32>, read> = access %globalLights, 5u, %148, 0u
            %150:vec3<f32> = load %149
            %151:vec4<f32> = construct %150, 1.0f
            %152:vec4<f32> = mul %147, %151
            %lightViewPos:vec4<f32> = let %152
            %154:vec3<f32> = swizzle %lightViewPos, xyz
            %155:ptr<storage, vec3<f32>, read> = access %clusters, 0u, %tileIndex, 0u
            %156:vec3<f32> = load %155
            %157:ptr<storage, vec3<f32>, read> = access %clusters, 0u, %tileIndex, 1u
            %158:vec3<f32> = load %157
            %159:f32 = call %sqDistPointAABB, %154, %156, %158
            %sqDist_1:f32 = let %159  # %sqDist_1: 'sqDist'
            %161:f32 = mul %range, %range
            %162:bool = lte %sqDist_1, %161
            store %lightInCluster, %162
            exit_if  # if_5
          }
        }
        %163:bool = load %lightInCluster
        if %163 [t: $B20] {  # if_6
          $B20: {  # true
            %164:u32 = load %clusterLightCount
            %165:ptr<function, u32, read_write> = access %cluserLightIndices, %164
            %166:u32 = load %i_1
            store %165, %166
            %167:u32 = load %clusterLightCount
            %168:u32 = add %167, 1u
            store %clusterLightCount, %168
            exit_if  # if_6
          }
        }
        %169:u32 = load %clusterLightCount
        %170:bool = eq %169, 256u
        if %170 [t: $B21] {  # if_7
          $B21: {  # true
            exit_loop  # loop_2
          }
        }
        continue  # -> $B16
      }
      $B16: {  # continuing
        %171:u32 = load %i_1
        %172:u32 = add %171, 1u
        store %i_1, %172
        next_iteration  # -> $B15
      }
    }
    %173:u32 = load %clusterLightCount
    %lightCount:u32 = let %173
    %175:ptr<storage, atomic<u32>, read_write> = access %clusterLights, 0u
    %176:u32 = atomicAdd %175, %lightCount
    %offset:ptr<function, u32, read_write> = var, %176
    %178:u32 = load %offset
    %179:bool = gte %178, 1769472u
    if %179 [t: $B22] {  # if_8
      $B22: {  # true
        ret
      }
    }
    loop [i: $B23, b: $B24, c: $B25] {  # loop_3
      $B23: {  # initializer
        %i_2:ptr<function, u32, read_write> = var, 0u  # %i_2: 'i'
        next_iteration  # -> $B24
      }
      $B24: {  # body
        %181:u32 = load %i_2
        %182:u32 = load %clusterLightCount
        %183:bool = lt %181, %182
        if %183 [t: $B26, f: $B27] {  # if_9
          $B26: {  # true
            exit_if  # if_9
          }
          $B27: {  # false
            exit_loop  # loop_3
          }
        }
        %184:u32 = load %offset
        %185:u32 = load %i_2
        %186:u32 = add %184, %185
        %187:ptr<storage, u32, read_write> = access %clusterLights, 2u, %186
        %188:u32 = load %i_2
        %189:ptr<function, u32, read_write> = access %cluserLightIndices, %188
        %190:u32 = load %189
        store %187, %190
        continue  # -> $B25
      }
      $B25: {  # continuing
        %191:u32 = load %i_2
        %192:u32 = add %191, 1u
        store %i_2, %192
        next_iteration  # -> $B24
      }
    }
    %193:ptr<storage, u32, read_write> = access %clusterLights, 1u, %tileIndex, 0u
    %194:u32 = load %offset
    store %193, %194
    %195:ptr<storage, u32, read_write> = access %clusterLights, 1u, %tileIndex, 1u
    %196:u32 = load %clusterLightCount
    store %195, %196
    ret
  }
}
%tint_f32_to_u32 = func(%value:f32):u32 {
  $B28: {
    %198:u32 = convert %value
    %199:bool = gte %value, 0.0f
    %200:u32 = select 0u, %198, %199
    %201:bool = lte %value, 4294967040.0f
    %202:u32 = select 4294967295u, %200, %201
    ret %202
  }
}

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