[msl] Add polyfill for dot builtin
The MSL `dot` builtin does not accept integer parameters, so add a
helper function that manually computes an integer dot product.
Add an MSL builtin for `dot` that only accepts floating point
parameters.
Bug: 42251016
Change-Id: Icefd439229c3b4ffc02e5cbddbc2bda4b5cf06eb
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/194341
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
index 8c88765..bda3140 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -92,6 +92,8 @@
return "write";
case BuiltinFn::kDistance:
return "distance";
+ case BuiltinFn::kDot:
+ return "dot";
case BuiltinFn::kLength:
return "length";
case BuiltinFn::kThreadgroupBarrier:
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
index 0dac6ca..b0cac68 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -72,6 +72,7 @@
kSampleCompare,
kWrite,
kDistance,
+ kDot,
kLength,
kThreadgroupBarrier,
kNone,
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index 7978441..62e30c9 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -4797,12 +4797,18 @@
},
{
/* [25] */
+ /* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(160),
+ },
+ {
+ /* [26] */
/* fn length[N : num, T : f32_f16](vec<N, T>) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(161),
},
{
- /* [26] */
+ /* [27] */
/* fn threadgroup_barrier(u32) */
/* num overloads */ 1,
/* overloads */ OverloadIndex(162),
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index f720c99..3894e23 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -322,6 +322,7 @@
@member_function fn write(texture: texture_storage_3d<u32_texel_format, writable>, value: vec4<u32>, coords: vec3<u32>)
fn distance[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> T
+fn dot[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> T
fn length[N: num, T: f32_f16](vec<N, T>) -> T
@stage("compute") fn threadgroup_barrier(u32)
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index a406bee..0b75610 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -75,6 +75,9 @@
/// A map from an atomic pointer type to an atomicCompareExchangeWeak polyfill.
Hashmap<const core::type::Type*, core::ir::Function*, 2> atomic_compare_exchange_polyfills{};
+ /// A map from an integer vector type to a dot polyfill.
+ Hashmap<const core::type::Vector*, core::ir::Function*, 4> integer_dot_polyfills{};
+
/// Process the module.
void Process() {
// Find the builtins that need replacing.
@@ -94,6 +97,7 @@
case core::BuiltinFn::kAtomicSub:
case core::BuiltinFn::kAtomicXor:
case core::BuiltinFn::kDistance:
+ case core::BuiltinFn::kDot:
case core::BuiltinFn::kLength:
case core::BuiltinFn::kTextureDimensions:
case core::BuiltinFn::kTextureGather:
@@ -159,10 +163,13 @@
AtomicCall(builtin, msl::BuiltinFn::kAtomicFetchXorExplicit);
break;
- // Geometric builtins.
+ // Arithmetic builtins.
case core::BuiltinFn::kDistance:
Distance(builtin);
break;
+ case core::BuiltinFn::kDot:
+ Dot(builtin);
+ break;
case core::BuiltinFn::kLength:
Length(builtin);
break;
@@ -304,6 +311,46 @@
builtin->Destroy();
}
+ /// Polyfill a dot call if necessary.
+ /// @param builtin the builtin call instruction
+ void Dot(core::ir::CoreBuiltinCall* builtin) {
+ b.InsertBefore(builtin, [&] {
+ auto* arg0 = builtin->Args()[0];
+ auto* arg1 = builtin->Args()[1];
+ auto* vec = arg0->Type()->As<core::type::Vector>();
+ if (vec->type()->is_integer_scalar()) {
+ // Calls to `dot` with a integer arguments are replaced with helper functions, as
+ // MSL's `dot` builtin only supports floating point arguments.
+ auto* polyfill = integer_dot_polyfills.GetOrAdd(vec, [&] {
+ // Generate a helper function that performs the following:
+ // fn tint_integer_dot(lhs: vec4i, rhs: vec4i) {
+ // let mul = lhs * rhs;
+ // return mul[0] + mul[1] + mul[2] + mul[3];
+ // }
+ auto* el_ty = vec->type();
+ auto* lhs = b.FunctionParam("lhs", vec);
+ auto* rhs = b.FunctionParam("rhs", vec);
+ auto* func = b.Function("tint_dot", el_ty);
+ func->SetParams({lhs, rhs});
+ b.Append(func->Block(), [&] {
+ auto* mul = b.Multiply(vec, lhs, rhs);
+ auto* sum = b.Access(el_ty, mul, u32(0))->Result(0);
+ for (uint32_t i = 1; i < vec->Width(); i++) {
+ sum = b.Add(el_ty, sum, b.Access(el_ty, mul, u32(i)))->Result(0);
+ }
+ b.Return(func, sum);
+ });
+ return func;
+ });
+ b.CallWithResult(builtin->DetachResult(), polyfill, arg0, arg1);
+ } else {
+ b.CallWithResult<msl::ir::BuiltinCall>(builtin->DetachResult(),
+ msl::BuiltinFn::kDot, arg0, arg1);
+ }
+ });
+ builtin->Destroy();
+ }
+
/// Polyfill a length call if necessary.
/// @param builtin the builtin call instruction
void Length(core::ir::CoreBuiltinCall* builtin) {
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
index ed89002..d9bc55b 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -827,6 +827,201 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BuiltinPolyfillTest, Dot_I32) {
+ auto* value0 = b.FunctionParam<vec4<i32>>("value0");
+ auto* value1 = b.FunctionParam<vec4<i32>>("value1");
+ auto* func = b.Function("foo", ty.i32());
+ func->SetParams({value0, value1});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<i32>(core::BuiltinFn::kDot, value0, value1);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value0:vec4<i32>, %value1:vec4<i32>):i32 {
+ $B1: {
+ %4:i32 = dot %value0, %value1
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value0:vec4<i32>, %value1:vec4<i32>):i32 {
+ $B1: {
+ %4:i32 = call %tint_dot, %value0, %value1
+ ret %4
+ }
+}
+%tint_dot = func(%lhs:vec4<i32>, %rhs:vec4<i32>):i32 {
+ $B2: {
+ %8:vec4<i32> = mul %lhs, %rhs
+ %9:i32 = access %8, 0u
+ %10:i32 = access %8, 1u
+ %11:i32 = add %9, %10
+ %12:i32 = access %8, 2u
+ %13:i32 = add %11, %12
+ %14:i32 = access %8, 3u
+ %15:i32 = add %13, %14
+ ret %15
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, Dot_U32) {
+ auto* value0 = b.FunctionParam<vec3<u32>>("value0");
+ auto* value1 = b.FunctionParam<vec3<u32>>("value1");
+ auto* func = b.Function("foo", ty.u32());
+ func->SetParams({value0, value1});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<u32>(core::BuiltinFn::kDot, value0, value1);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value0:vec3<u32>, %value1:vec3<u32>):u32 {
+ $B1: {
+ %4:u32 = dot %value0, %value1
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value0:vec3<u32>, %value1:vec3<u32>):u32 {
+ $B1: {
+ %4:u32 = call %tint_dot, %value0, %value1
+ ret %4
+ }
+}
+%tint_dot = func(%lhs:vec3<u32>, %rhs:vec3<u32>):u32 {
+ $B2: {
+ %8:vec3<u32> = mul %lhs, %rhs
+ %9:u32 = access %8, 0u
+ %10:u32 = access %8, 1u
+ %11:u32 = add %9, %10
+ %12:u32 = access %8, 2u
+ %13:u32 = add %11, %12
+ ret %13
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, Dot_F32) {
+ auto* value0 = b.FunctionParam<vec2<f32>>("value0");
+ auto* value1 = b.FunctionParam<vec2<f32>>("value1");
+ auto* func = b.Function("foo", ty.f32());
+ func->SetParams({value0, value1});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<f32>(core::BuiltinFn::kDot, value0, value1);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value0:vec2<f32>, %value1:vec2<f32>):f32 {
+ $B1: {
+ %4:f32 = dot %value0, %value1
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value0:vec2<f32>, %value1:vec2<f32>):f32 {
+ $B1: {
+ %4:f32 = msl.dot %value0, %value1
+ ret %4
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, Dot_MultipleCalls) {
+ auto* v4i = b.FunctionParam<vec4<i32>>("v4i");
+ auto* v4u = b.FunctionParam<vec4<u32>>("v4u");
+ auto* func = b.Function("foo", ty.void_());
+ func->SetParams({v4i, v4u});
+ b.Append(func->Block(), [&] {
+ b.Call<i32>(core::BuiltinFn::kDot, v4i, v4i);
+ b.Call<i32>(core::BuiltinFn::kDot, v4i, v4i);
+ b.Call<u32>(core::BuiltinFn::kDot, v4u, v4u);
+ b.Call<u32>(core::BuiltinFn::kDot, v4u, v4u);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%foo = func(%v4i:vec4<i32>, %v4u:vec4<u32>):void {
+ $B1: {
+ %4:i32 = dot %v4i, %v4i
+ %5:i32 = dot %v4i, %v4i
+ %6:u32 = dot %v4u, %v4u
+ %7:u32 = dot %v4u, %v4u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%v4i:vec4<i32>, %v4u:vec4<u32>):void {
+ $B1: {
+ %4:i32 = call %tint_dot, %v4i, %v4i
+ %6:i32 = call %tint_dot, %v4i, %v4i
+ %7:u32 = call %tint_dot_1, %v4u, %v4u
+ %9:u32 = call %tint_dot_1, %v4u, %v4u
+ ret
+ }
+}
+%tint_dot = func(%lhs:vec4<i32>, %rhs:vec4<i32>):i32 {
+ $B2: {
+ %12:vec4<i32> = mul %lhs, %rhs
+ %13:i32 = access %12, 0u
+ %14:i32 = access %12, 1u
+ %15:i32 = add %13, %14
+ %16:i32 = access %12, 2u
+ %17:i32 = add %15, %16
+ %18:i32 = access %12, 3u
+ %19:i32 = add %17, %18
+ ret %19
+ }
+}
+%tint_dot_1 = func(%lhs_1:vec4<u32>, %rhs_1:vec4<u32>):u32 { # %tint_dot_1: 'tint_dot', %lhs_1: 'lhs', %rhs_1: 'rhs'
+ $B3: {
+ %22:vec4<u32> = mul %lhs_1, %rhs_1
+ %23:u32 = access %22, 0u
+ %24:u32 = access %22, 1u
+ %25:u32 = add %23, %24
+ %26:u32 = access %22, 2u
+ %27:u32 = add %25, %26
+ %28:u32 = access %22, 3u
+ %29:u32 = add %27, %28
+ ret %29
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(MslWriter_BuiltinPolyfillTest, Length_Scalar) {
auto* value = b.FunctionParam<f32>("value");
auto* func = b.Function("foo", ty.f32());
diff --git a/test/tint/benchmark/shadow-fragment.wgsl.expected.ir.msl b/test/tint/benchmark/shadow-fragment.wgsl.expected.ir.msl
deleted file mode 100644
index 132aa68..0000000
--- a/test/tint/benchmark/shadow-fragment.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,125 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: Scene = struct @align(16) {
- lightViewProjMatrix:mat4x4<f32> @offset(0)
- cameraViewProjMatrix:mat4x4<f32> @offset(64)
- lightPos:vec3<f32> @offset(128)
-}
-
-FragmentInput = struct @align(16) {
- shadowPos:vec3<f32> @offset(0), @location(0)
- fragPos:vec3<f32> @offset(16), @location(1)
- fragNorm:vec3<f32> @offset(32), @location(2)
-}
-
-$B1: { # root
- %scene:ptr<uniform, Scene, read> = var @binding_point(0, 0)
- %shadowMap:ptr<handle, texture_depth_2d, read> = var @binding_point(0, 1)
- %shadowSampler:ptr<handle, sampler_comparison, read> = var @binding_point(0, 2)
-}
-
-%tint_symbol = @fragment func(%input:FragmentInput):vec4<f32> [@location(0)] {
- $B2: {
- %visibility:ptr<function, f32, read_write> = var, 0.0f
- %oneOverShadowDepthTextureSize:f32 = let 0.0009765625f
- loop [i: $B3, b: $B4, c: $B5] { # loop_1
- $B3: { # initializer
- %y:ptr<function, i32, read_write> = var, -1i
- next_iteration # -> $B4
- }
- $B4: { # body
- %9:i32 = load %y
- %10:bool = lte %9, 1i
- if %10 [t: $B6, f: $B7] { # if_1
- $B6: { # true
- exit_if # if_1
- }
- $B7: { # false
- exit_loop # loop_1
- }
- }
- loop [i: $B8, b: $B9, c: $B10] { # loop_2
- $B8: { # initializer
- %x:ptr<function, i32, read_write> = var, -1i
- next_iteration # -> $B9
- }
- $B9: { # body
- %12:i32 = load %x
- %13:bool = lte %12, 1i
- if %13 [t: $B11, f: $B12] { # if_2
- $B11: { # true
- exit_if # if_2
- }
- $B12: { # false
- exit_loop # loop_2
- }
- }
- %14:i32 = load %x
- %15:f32 = convert %14
- %16:f32 = mul %15, %oneOverShadowDepthTextureSize
- %17:f32 = let %16
- %18:i32 = load %y
- %19:f32 = convert %18
- %20:f32 = mul %19, %oneOverShadowDepthTextureSize
- %21:vec2<f32> = construct %17, %20
- %offset:vec2<f32> = let %21
- %23:f32 = load %visibility
- %24:f32 = let %23
- %25:texture_depth_2d = load %shadowMap
- %26:sampler_comparison = load %shadowSampler
- %27:vec3<f32> = access %input, 0u
- %28:vec2<f32> = swizzle %27, xy
- %29:vec2<f32> = add %28, %offset
- %30:f32 = access %input, 0u, 2u
- %31:f32 = sub %30, 0.00700000021606683731f
- %32:f32 = textureSampleCompare %25, %26, %29, %31
- %33:f32 = add %24, %32
- store %visibility, %33
- continue # -> $B10
- }
- $B10: { # continuing
- %34:i32 = load %x
- %35:i32 = add %34, 1i
- store %x, %35
- next_iteration # -> $B9
- }
- }
- continue # -> $B5
- }
- $B5: { # continuing
- %36:i32 = load %y
- %37:i32 = add %36, 1i
- store %y, %37
- next_iteration # -> $B4
- }
- }
- %38:f32 = load %visibility
- %39:f32 = div %38, 9.0f
- store %visibility, %39
- %40:ptr<uniform, vec3<f32>, read> = access %scene, 2u
- %41:vec3<f32> = load %40
- %42:vec3<f32> = access %input, 1u
- %43:vec3<f32> = sub %41, %42
- %44:vec3<f32> = normalize %43
- %45:vec3<f32> = access %input, 2u
- %46:f32 = dot %44, %45
- %47:f32 = max %46, 0.0f
- %lambertFactor:f32 = let %47
- %49:f32 = load %visibility
- %50:f32 = mul %49, %lambertFactor
- %51:f32 = add 0.20000000298023223877f, %50
- %52:f32 = min %51, 1.0f
- %lightingFactor:f32 = let %52
- %54:vec3<f32> = mul %lightingFactor, vec3<f32>(0.89999997615814208984f)
- %55:vec4<f32> = construct %54, 1.0f
- ret %55
- }
-}
-
-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. *
-********************************************************************
diff --git a/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl.expected.ir.msl b/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl.expected.ir.msl
deleted file mode 100644
index 0d672cf..0000000
--- a/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,940 +0,0 @@
-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)
-}
-
-ClusterLights = struct @align(4) {
- offset:u32 @offset(0)
- count:u32 @offset(4)
-}
-
-ClusterLightGroup = struct @align(4) {
- offset: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)
-}
-
-LightShadowTable = struct @align(4) {
- light:array<i32> @offset(0)
-}
-
-ShadowProperties = struct @align(16) {
- viewport:vec4<f32> @offset(0)
- viewProj:mat4x4<f32> @offset(16)
-}
-
-LightShadows = struct @align(16) {
- properties:array<ShadowProperties> @offset(0)
-}
-
-Material = struct @align(16) {
- baseColorFactor:vec4<f32> @offset(0)
- emissiveFactor:vec3<f32> @offset(16)
- occlusionStrength:f32 @offset(28)
- metallicRoughnessFactor:vec2<f32> @offset(32)
- alphaCutoff:f32 @offset(40)
-}
-
-SurfaceInfo = struct @align(16) {
- baseColor:vec4<f32> @offset(0)
- albedo:vec3<f32> @offset(16)
- metallic:f32 @offset(28)
- roughness:f32 @offset(32)
- normal:vec3<f32> @offset(48)
- f0:vec3<f32> @offset(64)
- ao:f32 @offset(76)
- emissive:vec3<f32> @offset(80)
- v:vec3<f32> @offset(96)
-}
-
-VertexOutput = struct @align(16) {
- position:vec4<f32> @offset(0), @builtin(position)
- worldPos:vec3<f32> @offset(16), @location(0)
- view:vec3<f32> @offset(32), @location(1)
- texcoord:vec2<f32> @offset(48), @location(2)
- texcoord2:vec2<f32> @offset(56), @location(3)
- color:vec4<f32> @offset(64), @location(4)
- instanceColor:vec4<f32> @offset(80), @location(5)
- normal:vec3<f32> @offset(96), @location(6)
- tangent:vec3<f32> @offset(112), @location(7)
- bitangent:vec3<f32> @offset(128), @location(8)
-}
-
-PuctualLight = struct @align(16) {
- lightType:u32 @offset(0)
- pointToLight:vec3<f32> @offset(16)
- range:f32 @offset(28)
- color:vec3<f32> @offset(32)
- intensity:f32 @offset(44)
-}
-
-FragmentOutput = struct @align(16) {
- color:vec4<f32> @offset(0), @location(0)
- emissive:vec4<f32> @offset(16), @location(1)
-}
-
-$B1: { # root
- %camera:ptr<uniform, Camera, read> = var @binding_point(0, 0)
- %clusterLights:ptr<storage, ClusterLightGroup, read> = var @binding_point(0, 1)
- %globalLights:ptr<storage, GlobalLights, read> = var @binding_point(0, 2)
- %defaultSampler:ptr<handle, sampler, read> = var @binding_point(0, 3)
- %shadowTexture:ptr<handle, texture_depth_2d, read> = var @binding_point(0, 4)
- %shadowSampler:ptr<handle, sampler_comparison, read> = var @binding_point(0, 5)
- %lightShadowTable:ptr<storage, LightShadowTable, read> = var @binding_point(0, 6)
- %shadowSampleOffsets:ptr<private, array<vec2<f32>, 16>, read_write> = var, array<vec2<f32>, 16>(vec2<f32>(-1.5f), vec2<f32>(-1.5f, -0.5f), vec2<f32>(-1.5f, 0.5f), vec2<f32>(-1.5f, 1.5f), vec2<f32>(-0.5f, -1.5f), vec2<f32>(-0.5f), vec2<f32>(-0.5f, 0.5f), vec2<f32>(-0.5f, 1.5f), vec2<f32>(0.5f, -1.5f), vec2<f32>(0.5f, -0.5f), vec2<f32>(0.5f), vec2<f32>(0.5f, 1.5f), vec2<f32>(1.5f, -1.5f), vec2<f32>(1.5f, -0.5f), vec2<f32>(1.5f, 0.5f), vec2<f32>(1.5f))
- %shadow:ptr<storage, LightShadows, read> = var @binding_point(0, 7)
- %material:ptr<uniform, Material, read> = var @binding_point(0, 8)
- %baseColorTexture:ptr<handle, texture_2d<f32>, read> = var @binding_point(0, 9)
- %baseColorSampler:ptr<handle, sampler, read> = var @binding_point(0, 10)
- %normalTexture:ptr<handle, texture_2d<f32>, read> = var @binding_point(0, 11)
- %normalSampler:ptr<handle, sampler, read> = var @binding_point(0, 12)
- %metallicRoughnessTexture:ptr<handle, texture_2d<f32>, read> = var @binding_point(0, 13)
- %metallicRoughnessSampler:ptr<handle, sampler, read> = var @binding_point(0, 14)
- %occlusionTexture:ptr<handle, texture_2d<f32>, read> = var @binding_point(0, 15)
- %occlusionSampler:ptr<handle, sampler, read> = var @binding_point(0, 16)
- %emissiveTexture:ptr<handle, texture_2d<f32>, read> = var @binding_point(0, 17)
- %emissiveSampler:ptr<handle, sampler, read> = var @binding_point(0, 18)
- %ssaoTexture:ptr<handle, texture_2d<f32>, read> = var @binding_point(0, 19)
-}
-
-%linearTosRGB = func(%linear:vec3<f32>):vec3<f32> {
- $B2: {
- %INV_GAMMA:f32 = let 0.45454543828964233398f
- %25:vec3<f32> = construct %INV_GAMMA
- %26:vec3<f32> = pow %linear, %25
- ret %26
- }
-}
-%sRGBToLinear = func(%srgb:vec3<f32>):vec3<f32> {
- $B3: {
- %29:vec3<f32> = pow %srgb, vec3<f32>(2.20000004768371582031f)
- ret %29
- }
-}
-%linearDepth = func(%depthSample:f32):f32 {
- $B4: {
- %32:ptr<uniform, f32, read> = access %camera, 7u
- %33:f32 = load %32
- %34:ptr<uniform, f32, read> = access %camera, 6u
- %35:f32 = load %34
- %36:f32 = mul %33, %35
- %37:f32 = let %36
- %38:ptr<uniform, f32, read> = access %camera, 6u
- %39:f32 = load %38
- %40:ptr<uniform, f32, read> = access %camera, 7u
- %41:f32 = load %40
- %42:f32 = sub %39, %41
- %43:ptr<uniform, f32, read> = access %camera, 7u
- %44:f32 = load %43
- %45:f32 = fma %depthSample, %42, %44
- %46:f32 = div %37, %45
- ret %46
- }
-}
-%getTile = func(%fragCoord:vec4<f32>):vec3<u32> {
- $B5: {
- %49:ptr<uniform, f32, read> = access %camera, 7u
- %50:f32 = load %49
- %51:ptr<uniform, f32, read> = access %camera, 6u
- %52:f32 = load %51
- %53:f32 = div %50, %52
- %54:f32 = log2 %53
- %55:f32 = div 48.0f, %54
- %sliceScale:f32 = let %55
- %57:ptr<uniform, f32, read> = access %camera, 6u
- %58:f32 = load %57
- %59:f32 = log2 %58
- %60:f32 = mul 48.0f, %59
- %61:f32 = let %60
- %62:ptr<uniform, f32, read> = access %camera, 7u
- %63:f32 = load %62
- %64:ptr<uniform, f32, read> = access %camera, 6u
- %65:f32 = load %64
- %66:f32 = div %63, %65
- %67:f32 = log2 %66
- %68:f32 = div %61, %67
- %69:f32 = negation %68
- %sliceBias:f32 = let %69
- %71:f32 = access %fragCoord, 2u
- %72:f32 = call %linearDepth, %71
- %73:f32 = log2 %72
- %74:f32 = mul %73, %sliceScale
- %75:f32 = add %74, %sliceBias
- %76:f32 = max %75, 0.0f
- %77:u32 = call %tint_f32_to_u32, %76
- %zTile:u32 = let %77
- %80:f32 = access %fragCoord, 0u
- %81:ptr<uniform, vec2<f32>, read> = access %camera, 5u
- %82:f32 = load_vector_element %81, 0u
- %83:f32 = div %82, 32.0f
- %84:f32 = div %80, %83
- %85:u32 = call %tint_f32_to_u32, %84
- %86:u32 = let %85
- %87:f32 = access %fragCoord, 1u
- %88:ptr<uniform, vec2<f32>, read> = access %camera, 5u
- %89:f32 = load_vector_element %88, 1u
- %90:f32 = div %89, 18.0f
- %91:f32 = div %87, %90
- %92:u32 = call %tint_f32_to_u32, %91
- %93:vec3<u32> = construct %86, %92, %zTile
- ret %93
- }
-}
-%getClusterIndex = func(%fragCoord_1:vec4<f32>):u32 { # %fragCoord_1: 'fragCoord'
- $B6: {
- %96:vec3<u32> = call %getTile, %fragCoord_1
- %tile:vec3<u32> = let %96
- %98:u32 = access %tile, 0u
- %99:u32 = access %tile, 1u
- %100:u32 = mul %99, 32u
- %101:u32 = add %98, %100
- %102:u32 = access %tile, 2u
- %103:u32 = mul %102, 32u
- %104:u32 = mul %103, 18u
- %105:u32 = add %101, %104
- ret %105
- }
-}
-%dirLightVisibility = func(%worldPos:vec3<f32>):f32 {
- $B7: {
- %108:ptr<storage, i32, read> = access %lightShadowTable, 0u, 0u
- %109:i32 = load %108
- %shadowIndex:i32 = let %109
- %111:bool = eq %shadowIndex, -1i
- if %111 [t: $B8] { # if_1
- $B8: { # true
- ret 1.0f
- }
- }
- %112:ptr<storage, vec4<f32>, read> = access %shadow, 0u, %shadowIndex, 0u
- %113:vec4<f32> = load %112
- %viewport:vec4<f32> = let %113
- %115:ptr<storage, mat4x4<f32>, read> = access %shadow, 0u, %shadowIndex, 1u
- %116:mat4x4<f32> = load %115
- %117:mat4x4<f32> = let %116
- %118:vec4<f32> = construct %worldPos, 1.0f
- %119:vec4<f32> = mul %117, %118
- %lightPos:vec4<f32> = let %119
- %121:vec2<f32> = swizzle %lightPos, xy
- %122:f32 = access %lightPos, 3u
- %123:vec2<f32> = div %121, %122
- %124:vec2<f32> = mul %123, vec2<f32>(0.5f, -0.5f)
- %125:vec2<f32> = add %124, vec2<f32>(0.5f)
- %126:f32 = access %lightPos, 2u
- %127:f32 = access %lightPos, 3u
- %128:f32 = div %126, %127
- %129:vec3<f32> = construct %125, %128
- %shadowPos:vec3<f32> = let %129
- %131:vec2<f32> = swizzle %viewport, xy
- %132:vec2<f32> = swizzle %shadowPos, xy
- %133:vec2<f32> = swizzle %viewport, zw
- %134:vec2<f32> = mul %132, %133
- %135:vec2<f32> = add %131, %134
- %136:vec2<f32> = construct %135
- %viewportPos:vec2<f32> = let %136
- %138:texture_depth_2d = load %shadowTexture
- %139:vec2<u32> = textureDimensions %138, 0i
- %140:vec2<f32> = convert %139
- %141:vec2<f32> = div 1.0f, %140
- %texelSize:vec2<f32> = let %141
- %143:vec2<f32> = swizzle %viewport, xy
- %144:vec2<f32> = sub %143, %texelSize
- %145:vec2<f32> = swizzle %viewport, xy
- %146:vec2<f32> = swizzle %viewport, zw
- %147:vec2<f32> = add %145, %146
- %148:vec2<f32> = add %147, %texelSize
- %149:vec4<f32> = construct %144, %148
- %clampRect:vec4<f32> = let %149
- %visibility:ptr<function, f32, read_write> = var, 0.0f
- loop [i: $B9, b: $B10, c: $B11] { # loop_1
- $B9: { # initializer
- %i:ptr<function, u32, read_write> = var, 0u
- next_iteration # -> $B10
- }
- $B10: { # body
- %153:u32 = load %i
- %154:bool = lt %153, 16u
- if %154 [t: $B12, f: $B13] { # if_2
- $B12: { # true
- exit_if # if_2
- }
- $B13: { # false
- exit_loop # loop_1
- }
- }
- %155:f32 = load %visibility
- %156:f32 = let %155
- %157:texture_depth_2d = load %shadowTexture
- %158:texture_depth_2d = let %157
- %159:sampler_comparison = load %shadowSampler
- %160:sampler_comparison = let %159
- %161:u32 = load %i
- %162:ptr<private, vec2<f32>, read_write> = access %shadowSampleOffsets, %161
- %163:vec2<f32> = load %162
- %164:vec2<f32> = mul %163, %texelSize
- %165:vec2<f32> = add %viewportPos, %164
- %166:vec2<f32> = swizzle %clampRect, xy
- %167:vec2<f32> = swizzle %clampRect, zw
- %168:vec2<f32> = clamp %165, %166, %167
- %169:f32 = access %shadowPos, 2u
- %170:f32 = sub %169, 0.00300000002607703209f
- %171:f32 = textureSampleCompareLevel %158, %160, %168, %170
- %172:f32 = add %156, %171
- store %visibility, %172
- continue # -> $B11
- }
- $B11: { # continuing
- %173:u32 = load %i
- %174:u32 = add %173, 1u
- store %i, %174
- next_iteration # -> $B10
- }
- }
- %175:f32 = load %visibility
- %176:f32 = div %175, 16.0f
- ret %176
- }
-}
-%getCubeFace = func(%v:vec3<f32>):i32 {
- $B14: {
- %179:vec3<f32> = abs %v
- %vAbs:vec3<f32> = let %179
- %181:f32 = access %vAbs, 2u
- %182:f32 = access %vAbs, 0u
- %183:bool = gte %181, %182
- %184:bool = if %183 [t: $B15, f: $B16] { # if_3
- $B15: { # true
- %185:f32 = access %vAbs, 2u
- %186:f32 = access %vAbs, 1u
- %187:bool = gte %185, %186
- exit_if %187 # if_3
- }
- $B16: { # false
- exit_if false # if_3
- }
- }
- if %184 [t: $B17] { # if_4
- $B17: { # true
- %188:f32 = access %v, 2u
- %189:bool = lt %188, 0.0f
- if %189 [t: $B18] { # if_5
- $B18: { # true
- ret 5i
- }
- }
- ret 4i
- }
- }
- %190:f32 = access %vAbs, 1u
- %191:f32 = access %vAbs, 0u
- %192:bool = gte %190, %191
- if %192 [t: $B19] { # if_6
- $B19: { # true
- %193:f32 = access %v, 1u
- %194:bool = lt %193, 0.0f
- if %194 [t: $B20] { # if_7
- $B20: { # true
- ret 3i
- }
- }
- ret 2i
- }
- }
- %195:f32 = access %v, 0u
- %196:bool = lt %195, 0.0f
- if %196 [t: $B21] { # if_8
- $B21: { # true
- ret 1i
- }
- }
- ret 0i
- }
-}
-%pointLightVisibility = func(%lightIndex:u32, %worldPos_1:vec3<f32>, %pointToLight:vec3<f32>):f32 { # %worldPos_1: 'worldPos'
- $B22: {
- %201:u32 = add %lightIndex, 1u
- %202:ptr<storage, i32, read> = access %lightShadowTable, 0u, %201
- %203:i32 = load %202
- %shadowIndex_1:ptr<function, i32, read_write> = var, %203 # %shadowIndex_1: 'shadowIndex'
- %205:i32 = load %shadowIndex_1
- %206:bool = eq %205, -1i
- if %206 [t: $B23] { # if_9
- $B23: { # true
- ret 1.0f
- }
- }
- %207:i32 = load %shadowIndex_1
- %208:i32 = let %207
- %209:vec3<f32> = mul %pointToLight, -1.0f
- %210:i32 = call %getCubeFace, %209
- %211:i32 = add %208, %210
- store %shadowIndex_1, %211
- %212:i32 = load %shadowIndex_1
- %213:ptr<storage, vec4<f32>, read> = access %shadow, 0u, %212, 0u
- %214:vec4<f32> = load %213
- %viewport_1:vec4<f32> = let %214 # %viewport_1: 'viewport'
- %216:i32 = load %shadowIndex_1
- %217:ptr<storage, mat4x4<f32>, read> = access %shadow, 0u, %216, 1u
- %218:mat4x4<f32> = load %217
- %219:mat4x4<f32> = let %218
- %220:vec4<f32> = construct %worldPos_1, 1.0f
- %221:vec4<f32> = mul %219, %220
- %lightPos_1:vec4<f32> = let %221 # %lightPos_1: 'lightPos'
- %223:vec2<f32> = swizzle %lightPos_1, xy
- %224:f32 = access %lightPos_1, 3u
- %225:vec2<f32> = div %223, %224
- %226:vec2<f32> = mul %225, vec2<f32>(0.5f, -0.5f)
- %227:vec2<f32> = add %226, vec2<f32>(0.5f)
- %228:f32 = access %lightPos_1, 2u
- %229:f32 = access %lightPos_1, 3u
- %230:f32 = div %228, %229
- %231:vec3<f32> = construct %227, %230
- %shadowPos_1:vec3<f32> = let %231 # %shadowPos_1: 'shadowPos'
- %233:vec2<f32> = swizzle %viewport_1, xy
- %234:vec2<f32> = swizzle %shadowPos_1, xy
- %235:vec2<f32> = swizzle %viewport_1, zw
- %236:vec2<f32> = mul %234, %235
- %237:vec2<f32> = add %233, %236
- %238:vec2<f32> = construct %237
- %viewportPos_1:vec2<f32> = let %238 # %viewportPos_1: 'viewportPos'
- %240:texture_depth_2d = load %shadowTexture
- %241:vec2<u32> = textureDimensions %240, 0i
- %242:vec2<f32> = convert %241
- %243:vec2<f32> = div 1.0f, %242
- %texelSize_1:vec2<f32> = let %243 # %texelSize_1: 'texelSize'
- %245:vec2<f32> = swizzle %viewport_1, xy
- %246:vec2<f32> = swizzle %viewport_1, xy
- %247:vec2<f32> = swizzle %viewport_1, zw
- %248:vec2<f32> = add %246, %247
- %249:vec4<f32> = construct %245, %248
- %clampRect_1:vec4<f32> = let %249 # %clampRect_1: 'clampRect'
- %visibility_1:ptr<function, f32, read_write> = var, 0.0f # %visibility_1: 'visibility'
- loop [i: $B24, b: $B25, c: $B26] { # loop_2
- $B24: { # initializer
- %i_1:ptr<function, u32, read_write> = var, 0u # %i_1: 'i'
- next_iteration # -> $B25
- }
- $B25: { # body
- %253:u32 = load %i_1
- %254:bool = lt %253, 16u
- if %254 [t: $B27, f: $B28] { # if_10
- $B27: { # true
- exit_if # if_10
- }
- $B28: { # false
- exit_loop # loop_2
- }
- }
- %255:f32 = load %visibility_1
- %256:f32 = let %255
- %257:texture_depth_2d = load %shadowTexture
- %258:texture_depth_2d = let %257
- %259:sampler_comparison = load %shadowSampler
- %260:sampler_comparison = let %259
- %261:u32 = load %i_1
- %262:ptr<private, vec2<f32>, read_write> = access %shadowSampleOffsets, %261
- %263:vec2<f32> = load %262
- %264:vec2<f32> = mul %263, %texelSize_1
- %265:vec2<f32> = add %viewportPos_1, %264
- %266:vec2<f32> = swizzle %clampRect_1, xy
- %267:vec2<f32> = swizzle %clampRect_1, zw
- %268:vec2<f32> = clamp %265, %266, %267
- %269:f32 = access %shadowPos_1, 2u
- %270:f32 = sub %269, 0.00999999977648258209f
- %271:f32 = textureSampleCompareLevel %258, %260, %268, %270
- %272:f32 = add %256, %271
- store %visibility_1, %272
- continue # -> $B26
- }
- $B26: { # continuing
- %273:u32 = load %i_1
- %274:u32 = add %273, 1u
- store %i_1, %274
- next_iteration # -> $B25
- }
- }
- %275:f32 = load %visibility_1
- %276:f32 = div %275, 16.0f
- ret %276
- }
-}
-%GetSurfaceInfo = func(%input:VertexOutput):SurfaceInfo {
- $B29: {
- %surface:ptr<function, SurfaceInfo, read_write> = var
- %280:ptr<function, vec3<f32>, read_write> = access %surface, 8u
- %281:vec3<f32> = access %input, 2u
- %282:vec3<f32> = normalize %281
- store %280, %282
- %283:vec3<f32> = access %input, 8u
- %284:vec3<f32> = access %input, 9u
- %285:vec3<f32> = access %input, 7u
- %286:mat3x3<f32> = construct %283, %284, %285
- %tbn:mat3x3<f32> = let %286
- %288:texture_2d<f32> = load %normalTexture
- %289:sampler = load %normalSampler
- %290:vec2<f32> = access %input, 3u
- %291:vec4<f32> = textureSample %288, %289, %290
- %292:vec3<f32> = swizzle %291, xyz
- %normalMap:vec3<f32> = let %292
- %294:ptr<function, vec3<f32>, read_write> = access %surface, 4u
- %295:vec3<f32> = mul 2.0f, %normalMap
- %296:vec3<f32> = sub %295, vec3<f32>(1.0f)
- %297:vec3<f32> = mul %tbn, %296
- %298:vec3<f32> = normalize %297
- store %294, %298
- %299:texture_2d<f32> = load %baseColorTexture
- %300:sampler = load %baseColorSampler
- %301:vec2<f32> = access %input, 3u
- %302:vec4<f32> = textureSample %299, %300, %301
- %baseColorMap:vec4<f32> = let %302
- %304:ptr<function, vec4<f32>, read_write> = access %surface, 0u
- %305:vec4<f32> = access %input, 5u
- %306:ptr<uniform, vec4<f32>, read> = access %material, 0u
- %307:vec4<f32> = load %306
- %308:vec4<f32> = mul %305, %307
- %309:vec4<f32> = mul %308, %baseColorMap
- store %304, %309
- %310:ptr<function, vec4<f32>, read_write> = access %surface, 0u
- %311:f32 = load_vector_element %310, 3u
- %312:ptr<uniform, f32, read> = access %material, 4u
- %313:f32 = load %312
- %314:bool = lt %311, %313
- if %314 [t: $B30] { # if_11
- $B30: { # true
- exit_if # if_11
- }
- }
- %315:ptr<function, vec3<f32>, read_write> = access %surface, 1u
- %316:ptr<function, vec4<f32>, read_write> = access %surface, 0u
- %317:vec4<f32> = load %316
- %318:vec3<f32> = swizzle %317, xyz
- store %315, %318
- %319:texture_2d<f32> = load %metallicRoughnessTexture
- %320:sampler = load %metallicRoughnessSampler
- %321:vec2<f32> = access %input, 3u
- %322:vec4<f32> = textureSample %319, %320, %321
- %metallicRoughnessMap:vec4<f32> = let %322
- %324:ptr<function, f32, read_write> = access %surface, 2u
- %325:ptr<uniform, vec2<f32>, read> = access %material, 3u
- %326:f32 = load_vector_element %325, 0u
- %327:f32 = access %metallicRoughnessMap, 2u
- %328:f32 = mul %326, %327
- store %324, %328
- %329:ptr<function, f32, read_write> = access %surface, 3u
- %330:ptr<uniform, vec2<f32>, read> = access %material, 3u
- %331:f32 = load_vector_element %330, 1u
- %332:f32 = access %metallicRoughnessMap, 1u
- %333:f32 = mul %331, %332
- store %329, %333
- %dielectricSpec:vec3<f32> = let vec3<f32>(0.03999999910593032837f)
- %335:ptr<function, vec3<f32>, read_write> = access %surface, 5u
- %336:ptr<function, vec3<f32>, read_write> = access %surface, 1u
- %337:vec3<f32> = load %336
- %338:vec3<f32> = let %337
- %339:ptr<function, f32, read_write> = access %surface, 2u
- %340:f32 = load %339
- %341:vec3<f32> = construct %340
- %342:vec3<f32> = mix %dielectricSpec, %338, %341
- store %335, %342
- %343:texture_2d<f32> = load %occlusionTexture
- %344:sampler = load %occlusionSampler
- %345:vec2<f32> = access %input, 3u
- %346:vec4<f32> = textureSample %343, %344, %345
- %occlusionMap:vec4<f32> = let %346
- %348:ptr<function, f32, read_write> = access %surface, 6u
- %349:ptr<uniform, f32, read> = access %material, 2u
- %350:f32 = load %349
- %351:f32 = access %occlusionMap, 0u
- %352:f32 = mul %350, %351
- store %348, %352
- %353:texture_2d<f32> = load %emissiveTexture
- %354:sampler = load %emissiveSampler
- %355:vec2<f32> = access %input, 3u
- %356:vec4<f32> = textureSample %353, %354, %355
- %emissiveMap:vec4<f32> = let %356
- %358:ptr<function, vec3<f32>, read_write> = access %surface, 7u
- %359:ptr<uniform, vec3<f32>, read> = access %material, 1u
- %360:vec3<f32> = load %359
- %361:vec3<f32> = swizzle %emissiveMap, xyz
- %362:vec3<f32> = mul %360, %361
- store %358, %362
- %363:f32 = access %input, 6u, 3u
- %364:bool = eq %363, 0.0f
- if %364 [t: $B31, f: $B32] { # if_12
- $B31: { # true
- %365:ptr<function, vec3<f32>, read_write> = access %surface, 1u
- %366:ptr<function, vec3<f32>, read_write> = access %surface, 1u
- %367:vec3<f32> = load %366
- %368:vec4<f32> = access %input, 6u
- %369:vec3<f32> = swizzle %368, xyz
- %370:vec3<f32> = add %367, %369
- store %365, %370
- exit_if # if_12
- }
- $B32: { # false
- %371:ptr<function, vec3<f32>, read_write> = access %surface, 1u
- %372:ptr<function, vec3<f32>, read_write> = access %surface, 1u
- %373:vec3<f32> = load %372
- %374:vec4<f32> = access %input, 6u
- %375:vec3<f32> = swizzle %374, xyz
- %376:vec3<f32> = mul %373, %375
- store %371, %376
- exit_if # if_12
- }
- }
- %377:SurfaceInfo = load %surface
- ret %377
- }
-}
-%FresnelSchlick = func(%cosTheta:f32, %F0:vec3<f32>):vec3<f32> {
- $B33: {
- %381:vec3<f32> = sub vec3<f32>(1.0f), %F0
- %382:f32 = sub 1.0f, %cosTheta
- %383:f32 = pow %382, 5.0f
- %384:vec3<f32> = mul %381, %383
- %385:vec3<f32> = add %F0, %384
- ret %385
- }
-}
-%DistributionGGX = func(%N:vec3<f32>, %H:vec3<f32>, %roughness:f32):f32 {
- $B34: {
- %390:f32 = mul %roughness, %roughness
- %a:f32 = let %390
- %392:f32 = mul %a, %a
- %a2:f32 = let %392
- %394:f32 = dot %N, %H
- %395:f32 = max %394, 0.0f
- %NdotH:f32 = let %395
- %397:f32 = mul %NdotH, %NdotH
- %NdotH2:f32 = let %397
- %num:f32 = let %a2
- %400:f32 = sub %a2, 1.0f
- %401:f32 = mul %NdotH2, %400
- %402:f32 = add %401, 1.0f
- %denom:f32 = let %402
- %404:f32 = mul 3.14159274101257324219f, %denom
- %405:f32 = mul %404, %denom
- %406:f32 = div %num, %405
- ret %406
- }
-}
-%GeometrySchlickGGX = func(%NdotV:f32, %roughness_1:f32):f32 { # %roughness_1: 'roughness'
- $B35: {
- %410:f32 = add %roughness_1, 1.0f
- %r:f32 = let %410
- %412:f32 = mul %r, %r
- %413:f32 = div %412, 8.0f
- %k:f32 = let %413
- %num_1:f32 = let %NdotV # %num_1: 'num'
- %416:f32 = sub 1.0f, %k
- %417:f32 = mul %NdotV, %416
- %418:f32 = add %417, %k
- %denom_1:f32 = let %418 # %denom_1: 'denom'
- %420:f32 = div %num_1, %denom_1
- ret %420
- }
-}
-%GeometrySmith = func(%N_1:vec3<f32>, %V:vec3<f32>, %L:vec3<f32>, %roughness_2:f32):f32 { # %N_1: 'N', %roughness_2: 'roughness'
- $B36: {
- %426:f32 = dot %N_1, %V
- %427:f32 = max %426, 0.0f
- %NdotV_1:f32 = let %427 # %NdotV_1: 'NdotV'
- %429:f32 = dot %N_1, %L
- %430:f32 = max %429, 0.0f
- %NdotL:f32 = let %430
- %432:f32 = call %GeometrySchlickGGX, %NdotV_1, %roughness_2
- %ggx2:f32 = let %432
- %434:f32 = call %GeometrySchlickGGX, %NdotL, %roughness_2
- %ggx1:f32 = let %434
- %436:f32 = mul %ggx1, %ggx2
- ret %436
- }
-}
-%lightAttenuation = func(%light:PuctualLight):f32 {
- $B37: {
- %439:u32 = access %light, 0u
- %440:bool = eq %439, 2u
- if %440 [t: $B38] { # if_13
- $B38: { # true
- ret 1.0f
- }
- }
- %441:vec3<f32> = access %light, 1u
- %442:f32 = length %441
- %distance:f32 = let %442
- %444:f32 = access %light, 2u
- %445:bool = lte %444, 0.0f
- if %445 [t: $B39] { # if_14
- $B39: { # true
- %446:f32 = pow %distance, 2.0f
- %447:f32 = div 1.0f, %446
- ret %447
- }
- }
- %448:f32 = access %light, 2u
- %449:f32 = div %distance, %448
- %450:f32 = pow %449, 4.0f
- %451:f32 = sub 1.0f, %450
- %452:f32 = clamp %451, 0.0f, 1.0f
- %453:f32 = let %452
- %454:f32 = pow %distance, 2.0f
- %455:f32 = div %453, %454
- ret %455
- }
-}
-%lightRadiance = func(%light_1:PuctualLight, %surface_1:SurfaceInfo):vec3<f32> { # %light_1: 'light', %surface_1: 'surface'
- $B40: {
- %459:vec3<f32> = access %light_1, 1u
- %460:vec3<f32> = normalize %459
- %L_1:vec3<f32> = let %460 # %L_1: 'L'
- %462:vec3<f32> = access %surface_1, 8u
- %463:vec3<f32> = add %462, %L_1
- %464:vec3<f32> = normalize %463
- %H_1:vec3<f32> = let %464 # %H_1: 'H'
- %466:vec3<f32> = access %surface_1, 4u
- %467:f32 = access %surface_1, 3u
- %468:f32 = call %DistributionGGX, %466, %H_1, %467
- %NDF:f32 = let %468
- %470:vec3<f32> = access %surface_1, 4u
- %471:vec3<f32> = access %surface_1, 8u
- %472:f32 = access %surface_1, 3u
- %473:f32 = call %GeometrySmith, %470, %471, %L_1, %472
- %G:f32 = let %473
- %475:vec3<f32> = access %surface_1, 8u
- %476:f32 = dot %H_1, %475
- %477:f32 = max %476, 0.0f
- %478:vec3<f32> = access %surface_1, 5u
- %479:vec3<f32> = call %FresnelSchlick, %477, %478
- %F:vec3<f32> = let %479
- %481:vec3<f32> = sub vec3<f32>(1.0f), %F
- %482:f32 = access %surface_1, 2u
- %483:f32 = sub 1.0f, %482
- %484:vec3<f32> = mul %481, %483
- %kD:vec3<f32> = let %484
- %486:vec3<f32> = access %surface_1, 4u
- %487:f32 = dot %486, %L_1
- %488:f32 = max %487, 0.0f
- %NdotL_1:f32 = let %488 # %NdotL_1: 'NdotL'
- %490:f32 = mul %NDF, %G
- %491:vec3<f32> = mul %490, %F
- %numerator:vec3<f32> = let %491
- %493:vec3<f32> = access %surface_1, 4u
- %494:vec3<f32> = access %surface_1, 8u
- %495:f32 = dot %493, %494
- %496:f32 = max %495, 0.0f
- %497:f32 = mul 4.0f, %496
- %498:f32 = mul %497, %NdotL_1
- %499:f32 = max %498, 0.00100000004749745131f
- %denominator:f32 = let %499
- %501:vec3<f32> = construct %denominator
- %502:vec3<f32> = div %numerator, %501
- %specular:vec3<f32> = let %502
- %504:vec3<f32> = access %light_1, 3u
- %505:f32 = access %light_1, 4u
- %506:vec3<f32> = mul %504, %505
- %507:f32 = call %lightAttenuation, %light_1
- %508:vec3<f32> = mul %506, %507
- %radiance:vec3<f32> = let %508
- %510:vec3<f32> = access %surface_1, 1u
- %511:vec3<f32> = mul %kD, %510
- %512:vec3<f32> = div %511, vec3<f32>(3.14159274101257324219f)
- %513:vec3<f32> = add %512, %specular
- %514:vec3<f32> = mul %513, %radiance
- %515:vec3<f32> = mul %514, %NdotL_1
- ret %515
- }
-}
-%fragmentMain = @fragment func(%input_1:VertexOutput):FragmentOutput { # %input_1: 'input'
- $B41: {
- %518:SurfaceInfo = call %GetSurfaceInfo, %input_1
- %surface_2:SurfaceInfo = let %518 # %surface_2: 'surface'
- %Lo:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(0.0f)
- %521:ptr<storage, f32, read> = access %globalLights, 2u
- %522:f32 = load %521
- %523:bool = gt %522, 0.0f
- if %523 [t: $B42] { # if_15
- $B42: { # true
- %light_2:ptr<function, PuctualLight, read_write> = var # %light_2: 'light'
- %525:ptr<function, u32, read_write> = access %light_2, 0u
- store %525, 2u
- %526:ptr<function, vec3<f32>, read_write> = access %light_2, 1u
- %527:ptr<storage, vec3<f32>, read> = access %globalLights, 3u
- %528:vec3<f32> = load %527
- store %526, %528
- %529:ptr<function, vec3<f32>, read_write> = access %light_2, 3u
- %530:ptr<storage, vec3<f32>, read> = access %globalLights, 1u
- %531:vec3<f32> = load %530
- store %529, %531
- %532:ptr<function, f32, read_write> = access %light_2, 4u
- %533:ptr<storage, f32, read> = access %globalLights, 2u
- %534:f32 = load %533
- store %532, %534
- %535:vec3<f32> = access %input_1, 1u
- %536:f32 = call %dirLightVisibility, %535
- %lightVis:f32 = let %536
- %538:vec3<f32> = load %Lo
- %539:vec3<f32> = let %538
- %540:PuctualLight = load %light_2
- %541:vec3<f32> = call %lightRadiance, %540, %surface_2
- %542:vec3<f32> = mul %541, %lightVis
- %543:vec3<f32> = add %539, %542
- store %Lo, %543
- exit_if # if_15
- }
- }
- %544:vec4<f32> = access %input_1, 0u
- %545:u32 = call %getClusterIndex, %544
- %clusterIndex:u32 = let %545
- %547:ptr<storage, u32, read> = access %clusterLights, 1u, %clusterIndex, 0u
- %548:u32 = load %547
- %lightOffset:u32 = let %548
- %550:ptr<storage, u32, read> = access %clusterLights, 1u, %clusterIndex, 1u
- %551:u32 = load %550
- %lightCount:u32 = let %551
- loop [i: $B43, b: $B44, c: $B45] { # loop_3
- $B43: { # initializer
- %lightIndex_1:ptr<function, u32, read_write> = var, 0u # %lightIndex_1: 'lightIndex'
- next_iteration # -> $B44
- }
- $B44: { # body
- %554:u32 = load %lightIndex_1
- %555:bool = lt %554, %lightCount
- if %555 [t: $B46, f: $B47] { # if_16
- $B46: { # true
- exit_if # if_16
- }
- $B47: { # false
- exit_loop # loop_3
- }
- }
- %556:u32 = load %lightIndex_1
- %557:u32 = add %lightOffset, %556
- %558:ptr<storage, u32, read> = access %clusterLights, 2u, %557
- %559:u32 = load %558
- %i_2:u32 = let %559 # %i_2: 'i'
- %light_3:ptr<function, PuctualLight, read_write> = var # %light_3: 'light'
- %562:ptr<function, u32, read_write> = access %light_3, 0u
- store %562, 0u
- %563:ptr<function, vec3<f32>, read_write> = access %light_3, 1u
- %564:ptr<storage, vec3<f32>, read> = access %globalLights, 5u, %i_2, 0u
- %565:vec3<f32> = load %564
- %566:vec3<f32> = swizzle %565, xyz
- %567:vec3<f32> = access %input_1, 1u
- %568:vec3<f32> = sub %566, %567
- store %563, %568
- %569:ptr<function, f32, read_write> = access %light_3, 2u
- %570:ptr<storage, f32, read> = access %globalLights, 5u, %i_2, 1u
- %571:f32 = load %570
- store %569, %571
- %572:ptr<function, vec3<f32>, read_write> = access %light_3, 3u
- %573:ptr<storage, vec3<f32>, read> = access %globalLights, 5u, %i_2, 2u
- %574:vec3<f32> = load %573
- store %572, %574
- %575:ptr<function, f32, read_write> = access %light_3, 4u
- %576:ptr<storage, f32, read> = access %globalLights, 5u, %i_2, 3u
- %577:f32 = load %576
- store %575, %577
- %578:vec3<f32> = access %input_1, 1u
- %579:ptr<function, vec3<f32>, read_write> = access %light_3, 1u
- %580:vec3<f32> = load %579
- %581:f32 = call %pointLightVisibility, %i_2, %578, %580
- %lightVis_1:f32 = let %581 # %lightVis_1: 'lightVis'
- %583:vec3<f32> = load %Lo
- %584:vec3<f32> = let %583
- %585:PuctualLight = load %light_3
- %586:vec3<f32> = call %lightRadiance, %585, %surface_2
- %587:vec3<f32> = mul %586, %lightVis_1
- %588:vec3<f32> = add %584, %587
- store %Lo, %588
- continue # -> $B45
- }
- $B45: { # continuing
- %589:u32 = load %lightIndex_1
- %590:u32 = add %589, 1u
- store %lightIndex_1, %590
- next_iteration # -> $B44
- }
- }
- %591:vec4<f32> = access %input_1, 0u
- %592:vec2<f32> = swizzle %591, xy
- %593:texture_2d<f32> = load %ssaoTexture
- %594:vec2<u32> = textureDimensions %593
- %595:vec2<u32> = swizzle %594, xy
- %596:vec2<f32> = convert %595
- %597:vec2<f32> = div %592, %596
- %ssaoCoord:vec2<f32> = let %597
- %599:texture_2d<f32> = load %ssaoTexture
- %600:sampler = load %defaultSampler
- %601:vec4<f32> = textureSample %599, %600, %ssaoCoord
- %602:f32 = access %601, 0u
- %ssaoFactor:f32 = let %602
- %604:ptr<storage, vec3<f32>, read> = access %globalLights, 0u
- %605:vec3<f32> = load %604
- %606:vec3<f32> = access %surface_2, 1u
- %607:vec3<f32> = mul %605, %606
- %608:f32 = access %surface_2, 6u
- %609:vec3<f32> = mul %607, %608
- %610:vec3<f32> = mul %609, %ssaoFactor
- %ambient:vec3<f32> = let %610
- %612:vec3<f32> = load %Lo
- %613:vec3<f32> = add %612, %ambient
- %614:vec3<f32> = access %surface_2, 7u
- %615:vec3<f32> = add %613, %614
- %616:vec3<f32> = call %linearTosRGB, %615
- %color:vec3<f32> = let %616
- %out:ptr<function, FragmentOutput, read_write> = var
- %619:ptr<function, vec4<f32>, read_write> = access %out, 0u
- %620:f32 = access %surface_2, 0u, 3u
- %621:vec4<f32> = construct %color, %620
- store %619, %621
- %622:ptr<function, vec4<f32>, read_write> = access %out, 1u
- %623:vec3<f32> = access %surface_2, 7u
- %624:f32 = access %surface_2, 0u, 3u
- %625:vec4<f32> = construct %623, %624
- store %622, %625
- %626:FragmentOutput = load %out
- ret %626
- }
-}
-%tint_f32_to_u32 = func(%value:f32):u32 {
- $B48: {
- %628:u32 = convert %value
- %629:bool = gte %value, 0.0f
- %630:u32 = select 0u, %628, %629
- %631:bool = lte %value, 4294967040.0f
- %632:u32 = select 4294967295u, %630, %631
- ret %632
- }
-}
-
-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. *
-********************************************************************
diff --git a/test/tint/bug/tint/1121.wgsl.expected.ir.msl b/test/tint/bug/tint/1121.wgsl.expected.ir.msl
index 02a6258..ad9cbfb 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.ir.msl
@@ -1,459 +1,179 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: LightData = struct @align(16) {
- position:vec4<f32> @offset(0)
- color:vec3<f32> @offset(16)
- radius:f32 @offset(28)
-}
+struct LightData {
+ float4 position;
+ float3 color;
+ float radius;
+};
+struct LightsBuffer {
+ tint_array<LightData, 1> lights;
+};
+struct TileLightIdData {
+ atomic_uint count;
+ tint_array<uint, 64> lightId;
+};
+struct Tiles {
+ tint_array<TileLightIdData, 4> data;
+};
+struct Config {
+ uint numLights;
+ uint numTiles;
+ uint tileCountX;
+ uint tileCountY;
+ uint numTileLightSlot;
+ uint tileSize;
+};
+struct Uniforms {
+ float4 min;
+ float4 max;
+ float4x4 viewMatrix;
+ float4x4 projectionMatrix;
+ float4 fullScreenSize;
+};
+struct tint_module_vars_struct {
+ device LightsBuffer* lightsBuffer;
+ device Tiles* tileLightId;
+ const constant Config* config;
+ const constant Uniforms* uniforms;
+};
-LightsBuffer = struct @align(16) {
- lights:array<LightData> @offset(0)
-}
-
-TileLightIdData = struct @align(4) {
- count:atomic<u32> @offset(0)
- lightId:array<u32, 64> @offset(4)
-}
-
-Tiles = struct @align(4) {
- data:array<TileLightIdData, 4> @offset(0)
-}
-
-Config = struct @align(4) {
- numLights:u32 @offset(0)
- numTiles:u32 @offset(4)
- tileCountX:u32 @offset(8)
- tileCountY:u32 @offset(12)
- numTileLightSlot:u32 @offset(16)
- tileSize:u32 @offset(20)
-}
-
-Uniforms = struct @align(16) {
- min:vec4<f32> @offset(0)
- max:vec4<f32> @offset(16)
- viewMatrix:mat4x4<f32> @offset(32)
- projectionMatrix:mat4x4<f32> @offset(96)
- fullScreenSize:vec4<f32> @offset(160)
-}
-
-$B1: { # root
- %lightsBuffer:ptr<storage, LightsBuffer, read_write> = var @binding_point(0, 0)
- %tileLightId:ptr<storage, Tiles, read_write> = var @binding_point(1, 0)
- %config:ptr<uniform, Config, read> = var @binding_point(2, 0)
- %uniforms:ptr<uniform, Uniforms, read> = var @binding_point(3, 0)
-}
-
-%tint_symbol = @compute @workgroup_size(64, 1, 1) func(%GlobalInvocationID:vec3<u32> [@global_invocation_id]):void {
- $B2: {
- %7:u32 = access %GlobalInvocationID, 0u
- %index:ptr<function, u32, read_write> = var, %7
- %9:u32 = load %index
- %10:ptr<uniform, u32, read> = access %config, 0u
- %11:u32 = load %10
- %12:bool = gte %9, %11
- if %12 [t: $B3] { # if_1
- $B3: { # true
- ret
+void tint_symbol_inner(uint3 GlobalInvocationID, tint_module_vars_struct tint_module_vars) {
+ uint index = GlobalInvocationID[0u];
+ if ((index >= (*tint_module_vars.config).numLights)) {
+ return;
+ }
+ device float4* const v = (&(*tint_module_vars.lightsBuffer).lights[index].position);
+ float const v_1 = ((*tint_module_vars.lightsBuffer).lights[index].position[1u] - 0.10000000149011611938f);
+ float const v_2 = float(index);
+ (*v)[1u] = (v_1 + (0.00100000004749745131f * (v_2 - (64.0f * floor((float(index) / 64.0f))))));
+ if (((*tint_module_vars.lightsBuffer).lights[index].position[1u] < (*tint_module_vars.uniforms).min[1u])) {
+ (*tint_module_vars.lightsBuffer).lights[index].position[1u] = (*tint_module_vars.uniforms).max[1u];
+ }
+ float4x4 M = (*tint_module_vars.uniforms).projectionMatrix;
+ float viewNear = (-(M[3][2]) / (-1.0f + M[2][2]));
+ float viewFar = (-(M[3][2]) / (1.0f + M[2][2]));
+ float4 lightPos = (*tint_module_vars.lightsBuffer).lights[index].position;
+ lightPos = ((*tint_module_vars.uniforms).viewMatrix * lightPos);
+ lightPos = (lightPos / lightPos[3u]);
+ float lightRadius = (*tint_module_vars.lightsBuffer).lights[index].radius;
+ float4 const v_3 = lightPos;
+ float4 boxMin = (v_3 - float4(float3(lightRadius), 0.0f));
+ float4 const v_4 = lightPos;
+ float4 boxMax = (v_4 + float4(float3(lightRadius), 0.0f));
+ tint_array<float4, 6> frustumPlanes = {};
+ frustumPlanes[4] = float4(0.0f, 0.0f, -1.0f, viewNear);
+ frustumPlanes[5] = float4(0.0f, 0.0f, 1.0f, -(viewFar));
+ int const TILE_SIZE = 16;
+ int const TILE_COUNT_X = 2;
+ int const TILE_COUNT_Y = 2;
+ {
+ int y = 0;
+ while(true) {
+ if ((y < TILE_COUNT_Y)) {
+ } else {
+ break;
}
- }
- %13:u32 = load %index
- %14:ptr<storage, vec4<f32>, read_write> = access %lightsBuffer, 0u, %13, 0u
- %15:ptr<storage, vec4<f32>, read_write> = let %14
- %16:u32 = load %index
- %17:ptr<storage, vec4<f32>, read_write> = access %lightsBuffer, 0u, %16, 0u
- %18:f32 = load_vector_element %17, 1u
- %19:f32 = sub %18, 0.10000000149011611938f
- %20:f32 = let %19
- %21:u32 = load %index
- %22:f32 = convert %21
- %23:f32 = let %22
- %24:u32 = load %index
- %25:f32 = convert %24
- %26:f32 = div %25, 64.0f
- %27:f32 = floor %26
- %28:f32 = mul 64.0f, %27
- %29:f32 = sub %23, %28
- %30:f32 = mul 0.00100000004749745131f, %29
- %31:f32 = add %20, %30
- store_vector_element %15, 1u, %31
- %32:u32 = load %index
- %33:ptr<storage, vec4<f32>, read_write> = access %lightsBuffer, 0u, %32, 0u
- %34:f32 = load_vector_element %33, 1u
- %35:ptr<uniform, vec4<f32>, read> = access %uniforms, 0u
- %36:f32 = load_vector_element %35, 1u
- %37:bool = lt %34, %36
- if %37 [t: $B4] { # if_2
- $B4: { # true
- %38:u32 = load %index
- %39:ptr<storage, vec4<f32>, read_write> = access %lightsBuffer, 0u, %38, 0u
- %40:ptr<uniform, vec4<f32>, read> = access %uniforms, 1u
- %41:f32 = load_vector_element %40, 1u
- store_vector_element %39, 1u, %41
- exit_if # if_2
- }
- }
- %42:ptr<uniform, mat4x4<f32>, read> = access %uniforms, 3u
- %43:mat4x4<f32> = load %42
- %M:ptr<function, mat4x4<f32>, read_write> = var, %43
- %45:ptr<function, vec4<f32>, read_write> = access %M, 3i
- %46:f32 = load_vector_element %45, 2i
- %47:f32 = negation %46
- %48:ptr<function, vec4<f32>, read_write> = access %M, 2i
- %49:f32 = load_vector_element %48, 2i
- %50:f32 = add -1.0f, %49
- %51:f32 = div %47, %50
- %viewNear:ptr<function, f32, read_write> = var, %51
- %53:ptr<function, vec4<f32>, read_write> = access %M, 3i
- %54:f32 = load_vector_element %53, 2i
- %55:f32 = negation %54
- %56:ptr<function, vec4<f32>, read_write> = access %M, 2i
- %57:f32 = load_vector_element %56, 2i
- %58:f32 = add 1.0f, %57
- %59:f32 = div %55, %58
- %viewFar:ptr<function, f32, read_write> = var, %59
- %61:u32 = load %index
- %62:ptr<storage, vec4<f32>, read_write> = access %lightsBuffer, 0u, %61, 0u
- %63:vec4<f32> = load %62
- %lightPos:ptr<function, vec4<f32>, read_write> = var, %63
- %65:ptr<uniform, mat4x4<f32>, read> = access %uniforms, 2u
- %66:mat4x4<f32> = load %65
- %67:vec4<f32> = load %lightPos
- %68:vec4<f32> = mul %66, %67
- store %lightPos, %68
- %69:vec4<f32> = load %lightPos
- %70:f32 = load_vector_element %lightPos, 3u
- %71:vec4<f32> = div %69, %70
- store %lightPos, %71
- %72:u32 = load %index
- %73:ptr<storage, f32, read_write> = access %lightsBuffer, 0u, %72, 2u
- %74:f32 = load %73
- %lightRadius:ptr<function, f32, read_write> = var, %74
- %76:vec4<f32> = load %lightPos
- %77:vec4<f32> = let %76
- %78:f32 = load %lightRadius
- %79:vec3<f32> = construct %78
- %80:vec4<f32> = construct %79, 0.0f
- %81:vec4<f32> = sub %77, %80
- %boxMin:ptr<function, vec4<f32>, read_write> = var, %81
- %83:vec4<f32> = load %lightPos
- %84:vec4<f32> = let %83
- %85:f32 = load %lightRadius
- %86:vec3<f32> = construct %85
- %87:vec4<f32> = construct %86, 0.0f
- %88:vec4<f32> = add %84, %87
- %boxMax:ptr<function, vec4<f32>, read_write> = var, %88
- %frustumPlanes:ptr<function, array<vec4<f32>, 6>, read_write> = var
- %91:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 4i
- %92:f32 = load %viewNear
- %93:vec4<f32> = construct 0.0f, 0.0f, -1.0f, %92
- store %91, %93
- %94:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 5i
- %95:f32 = load %viewFar
- %96:f32 = negation %95
- %97:vec4<f32> = construct 0.0f, 0.0f, 1.0f, %96
- store %94, %97
- %TILE_SIZE:i32 = let 16i
- %TILE_COUNT_X:i32 = let 2i
- %TILE_COUNT_Y:i32 = let 2i
- loop [i: $B5, b: $B6, c: $B7] { # loop_1
- $B5: { # initializer
- %y:ptr<function, i32, read_write> = var, 0i
- next_iteration # -> $B6
- }
- $B6: { # body
- %102:i32 = load %y
- %103:bool = lt %102, %TILE_COUNT_Y
- if %103 [t: $B8, f: $B9] { # if_3
- $B8: { # true
- exit_if # if_3
+ {
+ int x = 0;
+ while(true) {
+ if ((x < TILE_COUNT_X)) {
+ } else {
+ break;
}
- $B9: { # false
- exit_loop # loop_1
+ int2 tilePixel0Idx = int2((x * TILE_SIZE), (y * TILE_SIZE));
+ float2 const v_5 = (2.0f * float2(tilePixel0Idx));
+ float2 floorCoord = ((v_5 / (*tint_module_vars.uniforms).fullScreenSize.xy) - float2(1.0f));
+ int2 const v_6 = tilePixel0Idx;
+ float2 const v_7 = (2.0f * float2((v_6 + int2(TILE_SIZE))));
+ float2 ceilCoord = ((v_7 / (*tint_module_vars.uniforms).fullScreenSize.xy) - float2(1.0f));
+ float2 viewFloorCoord = float2((((-(viewNear) * floorCoord[0u]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord[1u]) - (M[2][1] * viewNear)) / M[1][1]));
+ float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord[0u]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord[1u]) - (M[2][1] * viewNear)) / M[1][1]));
+ frustumPlanes[0] = float4(1.0f, 0.0f, (-(viewFloorCoord[0u]) / viewNear), 0.0f);
+ frustumPlanes[1] = float4(-1.0f, 0.0f, (viewCeilCoord[0u] / viewNear), 0.0f);
+ frustumPlanes[2] = float4(0.0f, 1.0f, (-(viewFloorCoord[1u]) / viewNear), 0.0f);
+ frustumPlanes[3] = float4(0.0f, -1.0f, (viewCeilCoord[1u] / viewNear), 0.0f);
+ float dp = 0.0f;
+ {
+ uint i = 0u;
+ while(true) {
+ if ((i < 6u)) {
+ } else {
+ break;
+ }
+ float4 p = 0.0f;
+ if ((frustumPlanes[i][0u] > 0.0f)) {
+ p[0u] = boxMax[0u];
+ } else {
+ p[0u] = boxMin[0u];
+ }
+ if ((frustumPlanes[i][1u] > 0.0f)) {
+ p[1u] = boxMax[1u];
+ } else {
+ p[1u] = boxMin[1u];
+ }
+ if ((frustumPlanes[i][2u] > 0.0f)) {
+ p[2u] = boxMax[2u];
+ } else {
+ p[2u] = boxMin[2u];
+ }
+ p[3u] = 1.0f;
+ float const v_8 = dp;
+ dp = (v_8 + min(0.0f, dot(p, frustumPlanes[i])));
+ {
+ i = (i + 1u);
+ }
+ continue;
+ }
}
+ if ((dp >= 0.0f)) {
+ uint tileId = uint((x + (y * TILE_COUNT_X)));
+ bool v_9 = false;
+ if ((tileId < 0u)) {
+ v_9 = true;
+ } else {
+ v_9 = (tileId >= (*tint_module_vars.config).numTiles);
+ }
+ if (v_9) {
+ {
+ x = (x + 1);
+ }
+ continue;
+ }
+ uint offset = atomic_fetch_add_explicit((&(*tint_module_vars.tileLightId).data[tileId].count), 1u, memory_order_relaxed);
+ if ((offset >= (*tint_module_vars.config).numTileLightSlot)) {
+ {
+ x = (x + 1);
+ }
+ continue;
+ }
+ (*tint_module_vars.tileLightId).data[tileId].lightId[offset] = GlobalInvocationID[0u];
+ }
+ {
+ x = (x + 1);
+ }
+ continue;
}
- loop [i: $B10, b: $B11, c: $B12] { # loop_2
- $B10: { # initializer
- %x:ptr<function, i32, read_write> = var, 0i
- next_iteration # -> $B11
- }
- $B11: { # body
- %105:i32 = load %x
- %106:bool = lt %105, %TILE_COUNT_X
- if %106 [t: $B13, f: $B14] { # if_4
- $B13: { # true
- exit_if # if_4
- }
- $B14: { # false
- exit_loop # loop_2
- }
- }
- %107:i32 = load %x
- %108:i32 = mul %107, %TILE_SIZE
- %109:i32 = load %y
- %110:i32 = mul %109, %TILE_SIZE
- %111:vec2<i32> = construct %108, %110
- %tilePixel0Idx:ptr<function, vec2<i32>, read_write> = var, %111
- %113:vec2<i32> = load %tilePixel0Idx
- %114:vec2<f32> = convert %113
- %115:vec2<f32> = mul 2.0f, %114
- %116:ptr<uniform, vec4<f32>, read> = access %uniforms, 4u
- %117:vec4<f32> = load %116
- %118:vec2<f32> = swizzle %117, xy
- %119:vec2<f32> = div %115, %118
- %120:vec2<f32> = sub %119, vec2<f32>(1.0f)
- %floorCoord:ptr<function, vec2<f32>, read_write> = var, %120
- %122:vec2<i32> = load %tilePixel0Idx
- %123:vec2<i32> = let %122
- %124:vec2<i32> = construct %TILE_SIZE
- %125:vec2<i32> = add %123, %124
- %126:vec2<f32> = convert %125
- %127:vec2<f32> = mul 2.0f, %126
- %128:ptr<uniform, vec4<f32>, read> = access %uniforms, 4u
- %129:vec4<f32> = load %128
- %130:vec2<f32> = swizzle %129, xy
- %131:vec2<f32> = div %127, %130
- %132:vec2<f32> = sub %131, vec2<f32>(1.0f)
- %ceilCoord:ptr<function, vec2<f32>, read_write> = var, %132
- %134:f32 = load %viewNear
- %135:f32 = negation %134
- %136:f32 = load_vector_element %floorCoord, 0u
- %137:f32 = mul %135, %136
- %138:ptr<function, vec4<f32>, read_write> = access %M, 2i
- %139:f32 = load_vector_element %138, 0i
- %140:f32 = load %viewNear
- %141:f32 = mul %139, %140
- %142:f32 = sub %137, %141
- %143:ptr<function, vec4<f32>, read_write> = access %M, 0i
- %144:f32 = load_vector_element %143, 0i
- %145:f32 = div %142, %144
- %146:f32 = load %viewNear
- %147:f32 = negation %146
- %148:f32 = load_vector_element %floorCoord, 1u
- %149:f32 = mul %147, %148
- %150:ptr<function, vec4<f32>, read_write> = access %M, 2i
- %151:f32 = load_vector_element %150, 1i
- %152:f32 = load %viewNear
- %153:f32 = mul %151, %152
- %154:f32 = sub %149, %153
- %155:ptr<function, vec4<f32>, read_write> = access %M, 1i
- %156:f32 = load_vector_element %155, 1i
- %157:f32 = div %154, %156
- %158:vec2<f32> = construct %145, %157
- %viewFloorCoord:ptr<function, vec2<f32>, read_write> = var, %158
- %160:f32 = load %viewNear
- %161:f32 = negation %160
- %162:f32 = load_vector_element %ceilCoord, 0u
- %163:f32 = mul %161, %162
- %164:ptr<function, vec4<f32>, read_write> = access %M, 2i
- %165:f32 = load_vector_element %164, 0i
- %166:f32 = load %viewNear
- %167:f32 = mul %165, %166
- %168:f32 = sub %163, %167
- %169:ptr<function, vec4<f32>, read_write> = access %M, 0i
- %170:f32 = load_vector_element %169, 0i
- %171:f32 = div %168, %170
- %172:f32 = load %viewNear
- %173:f32 = negation %172
- %174:f32 = load_vector_element %ceilCoord, 1u
- %175:f32 = mul %173, %174
- %176:ptr<function, vec4<f32>, read_write> = access %M, 2i
- %177:f32 = load_vector_element %176, 1i
- %178:f32 = load %viewNear
- %179:f32 = mul %177, %178
- %180:f32 = sub %175, %179
- %181:ptr<function, vec4<f32>, read_write> = access %M, 1i
- %182:f32 = load_vector_element %181, 1i
- %183:f32 = div %180, %182
- %184:vec2<f32> = construct %171, %183
- %viewCeilCoord:ptr<function, vec2<f32>, read_write> = var, %184
- %186:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 0i
- %187:f32 = load_vector_element %viewFloorCoord, 0u
- %188:f32 = negation %187
- %189:f32 = load %viewNear
- %190:f32 = div %188, %189
- %191:vec4<f32> = construct 1.0f, 0.0f, %190, 0.0f
- store %186, %191
- %192:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 1i
- %193:f32 = load_vector_element %viewCeilCoord, 0u
- %194:f32 = load %viewNear
- %195:f32 = div %193, %194
- %196:vec4<f32> = construct -1.0f, 0.0f, %195, 0.0f
- store %192, %196
- %197:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 2i
- %198:f32 = load_vector_element %viewFloorCoord, 1u
- %199:f32 = negation %198
- %200:f32 = load %viewNear
- %201:f32 = div %199, %200
- %202:vec4<f32> = construct 0.0f, 1.0f, %201, 0.0f
- store %197, %202
- %203:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, 3i
- %204:f32 = load_vector_element %viewCeilCoord, 1u
- %205:f32 = load %viewNear
- %206:f32 = div %204, %205
- %207:vec4<f32> = construct 0.0f, -1.0f, %206, 0.0f
- store %203, %207
- %dp:ptr<function, f32, read_write> = var, 0.0f
- loop [i: $B15, b: $B16, c: $B17] { # loop_3
- $B15: { # initializer
- %i:ptr<function, u32, read_write> = var, 0u
- next_iteration # -> $B16
- }
- $B16: { # body
- %210:u32 = load %i
- %211:bool = lt %210, 6u
- if %211 [t: $B18, f: $B19] { # if_5
- $B18: { # true
- exit_if # if_5
- }
- $B19: { # false
- exit_loop # loop_3
- }
- }
- %p:ptr<function, vec4<f32>, read_write> = var
- %213:u32 = load %i
- %214:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, %213
- %215:f32 = load_vector_element %214, 0u
- %216:bool = gt %215, 0.0f
- if %216 [t: $B20, f: $B21] { # if_6
- $B20: { # true
- %217:f32 = load_vector_element %boxMax, 0u
- store_vector_element %p, 0u, %217
- exit_if # if_6
- }
- $B21: { # false
- %218:f32 = load_vector_element %boxMin, 0u
- store_vector_element %p, 0u, %218
- exit_if # if_6
- }
- }
- %219:u32 = load %i
- %220:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, %219
- %221:f32 = load_vector_element %220, 1u
- %222:bool = gt %221, 0.0f
- if %222 [t: $B22, f: $B23] { # if_7
- $B22: { # true
- %223:f32 = load_vector_element %boxMax, 1u
- store_vector_element %p, 1u, %223
- exit_if # if_7
- }
- $B23: { # false
- %224:f32 = load_vector_element %boxMin, 1u
- store_vector_element %p, 1u, %224
- exit_if # if_7
- }
- }
- %225:u32 = load %i
- %226:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, %225
- %227:f32 = load_vector_element %226, 2u
- %228:bool = gt %227, 0.0f
- if %228 [t: $B24, f: $B25] { # if_8
- $B24: { # true
- %229:f32 = load_vector_element %boxMax, 2u
- store_vector_element %p, 2u, %229
- exit_if # if_8
- }
- $B25: { # false
- %230:f32 = load_vector_element %boxMin, 2u
- store_vector_element %p, 2u, %230
- exit_if # if_8
- }
- }
- store_vector_element %p, 3u, 1.0f
- %231:f32 = load %dp
- %232:f32 = let %231
- %233:vec4<f32> = load %p
- %234:u32 = load %i
- %235:ptr<function, vec4<f32>, read_write> = access %frustumPlanes, %234
- %236:vec4<f32> = load %235
- %237:f32 = dot %233, %236
- %238:f32 = min 0.0f, %237
- %239:f32 = add %232, %238
- store %dp, %239
- continue # -> $B17
- }
- $B17: { # continuing
- %240:u32 = load %i
- %241:u32 = add %240, 1u
- store %i, %241
- next_iteration # -> $B16
- }
- }
- %242:f32 = load %dp
- %243:bool = gte %242, 0.0f
- if %243 [t: $B26] { # if_9
- $B26: { # true
- %244:i32 = load %x
- %245:i32 = load %y
- %246:i32 = mul %245, %TILE_COUNT_X
- %247:i32 = add %244, %246
- %248:u32 = convert %247
- %tileId:ptr<function, u32, read_write> = var, %248
- %250:u32 = load %tileId
- %251:bool = lt %250, 0u
- %252:bool = if %251 [t: $B27, f: $B28] { # if_10
- $B27: { # true
- exit_if true # if_10
- }
- $B28: { # false
- %253:u32 = load %tileId
- %254:ptr<uniform, u32, read> = access %config, 1u
- %255:u32 = load %254
- %256:bool = gte %253, %255
- exit_if %256 # if_10
- }
- }
- if %252 [t: $B29] { # if_11
- $B29: { # true
- continue # -> $B12
- }
- }
- %257:u32 = load %tileId
- %258:ptr<storage, atomic<u32>, read_write> = access %tileLightId, 0u, %257, 0u
- %259:u32 = atomicAdd %258, 1u
- %offset:ptr<function, u32, read_write> = var, %259
- %261:u32 = load %offset
- %262:ptr<uniform, u32, read> = access %config, 4u
- %263:u32 = load %262
- %264:bool = gte %261, %263
- if %264 [t: $B30] { # if_12
- $B30: { # true
- continue # -> $B12
- }
- }
- %265:u32 = load %tileId
- %266:u32 = load %offset
- %267:ptr<storage, u32, read_write> = access %tileLightId, 0u, %265, 1u, %266
- %268:u32 = access %GlobalInvocationID, 0u
- store %267, %268
- exit_if # if_9
- }
- }
- continue # -> $B12
- }
- $B12: { # continuing
- %269:i32 = load %x
- %270:i32 = add %269, 1i
- store %x, %270
- next_iteration # -> $B11
- }
- }
- continue # -> $B7
}
- $B7: { # continuing
- %271:i32 = load %y
- %272:i32 = add %271, 1i
- store %y, %272
- next_iteration # -> $B6
+ {
+ y = (y + 1);
}
+ continue;
}
- 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. *
-********************************************************************
+kernel void tint_symbol(uint3 GlobalInvocationID [[thread_position_in_grid]], device LightsBuffer* lightsBuffer [[buffer(2)]], device Tiles* tileLightId [[buffer(3)]], const constant Config* config [[buffer(0)]], const constant Uniforms* uniforms [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.lightsBuffer=lightsBuffer, .tileLightId=tileLightId, .config=config, .uniforms=uniforms};
+ tint_symbol_inner(GlobalInvocationID, tint_module_vars);
+}
diff --git a/test/tint/bug/tint/1534.wgsl.expected.ir.msl b/test/tint/bug/tint/1534.wgsl.expected.ir.msl
index d9246b1..4b5c318 100644
--- a/test/tint/bug/tint/1534.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1534.wgsl.expected.ir.msl
@@ -1,38 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct g {
+ uint3 a;
+};
+struct h {
+ uint a;
+};
+struct tint_module_vars_struct {
+ const constant g* i;
+ device h* j;
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: g = struct @align(16) {
- a:vec3<u32> @offset(0)
+uint tint_dot(uint3 lhs, uint3 rhs) {
+ return (((lhs * rhs)[0u] + (lhs * rhs)[1u]) + (lhs * rhs)[2u]);
}
-
-h = struct @align(4) {
- a:u32 @offset(0)
+kernel void tint_symbol(const constant g* i [[buffer(0)]], device h* j [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.i=i, .j=j};
+ uint const l = tint_dot((*tint_module_vars.i).a, (*tint_module_vars.i).a);
+ (*tint_module_vars.j).a = (*tint_module_vars.i).a[0u];
}
-
-$B1: { # root
- %i:ptr<uniform, g, read> = var @binding_point(0, 0)
- %j:ptr<storage, h, read_write> = var @binding_point(0, 1)
-}
-
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %4:ptr<uniform, vec3<u32>, read> = access %i, 0u
- %5:vec3<u32> = load %4
- %6:ptr<uniform, vec3<u32>, read> = access %i, 0u
- %7:vec3<u32> = load %6
- %8:u32 = dot %5, %7
- %l:u32 = let %8
- %10:ptr<storage, u32, read_write> = access %j, 0u
- %11:ptr<uniform, vec3<u32>, read> = access %i, 0u
- %12:u32 = load_vector_element %11, 0u
- store %10, %12
- 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. *
-********************************************************************
diff --git a/test/tint/bug/tint/922.wgsl.expected.ir.msl b/test/tint/bug/tint/922.wgsl.expected.ir.msl
index edb93e9..2cbdc5b 100644
--- a/test/tint/bug/tint/922.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/922.wgsl.expected.ir.msl
@@ -1,509 +1,309 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct Mat4x3_ {
+ float4 mx;
+ float4 my;
+ float4 mz;
+};
+struct Mat4x4_ {
+ float4 mx;
+ float4 my;
+ float4 mz;
+ float4 mw;
+};
+struct Mat4x2_ {
+ float4 mx;
+ float4 my;
+};
+struct ub_SceneParams {
+ Mat4x4_ u_Projection;
+};
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: Mat4x4_ = struct @align(16) {
- mx:vec4<f32> @offset(0)
- my:vec4<f32> @offset(16)
- mz:vec4<f32> @offset(32)
- mw:vec4<f32> @offset(48)
-}
+struct ub_MaterialParams {
+ tint_array<Mat4x2_, 1> u_TexMtx;
+ float4 u_Misc0_;
+};
+struct ub_PacketParams {
+ tint_array<Mat4x3_, 32> u_PosMtx;
+};
+struct tint_module_vars_struct {
+ const constant ub_SceneParams* global;
+ const constant ub_MaterialParams* global1;
+ const constant ub_PacketParams* global2;
+ thread float3* a_Position1;
+ thread float2* a_UV1;
+ thread float4* a_Color1;
+ thread float3* a_Normal1;
+ thread float* a_PosMtxIdx1;
+ thread float4* v_Color;
+ thread float2* v_TexCoord;
+ thread float4* gl_Position;
+};
+struct VertexOutput {
+ float4 v_Color;
+ float2 v_TexCoord;
+ float4 member;
+};
+struct tint_symbol_outputs {
+ float4 VertexOutput_v_Color [[user(locn0)]];
+ float2 VertexOutput_v_TexCoord [[user(locn1)]];
+ float4 VertexOutput_member [[position]];
+};
+struct tint_symbol_inputs {
+ float3 a_Position [[attribute(0)]];
+ float2 a_UV [[attribute(1)]];
+ float4 a_Color [[attribute(2)]];
+ float3 a_Normal [[attribute(3)]];
+ float a_PosMtxIdx [[attribute(4)]];
+};
-ub_SceneParams = struct @align(16) {
- u_Projection:Mat4x4_ @offset(0)
+float3 Mat4x3GetCol0_(Mat4x3_ m) {
+ Mat4x3_ m1 = {};
+ m1 = m;
+ Mat4x3_ const x_e2 = m1;
+ Mat4x3_ const x_e5 = m1;
+ Mat4x3_ const x_e8 = m1;
+ return float3(x_e2.mx[0u], x_e5.my[0u], x_e8.mz[0u]);
}
-
-Mat4x2_ = struct @align(16) {
- mx:vec4<f32> @offset(0)
- my:vec4<f32> @offset(16)
+float3 Mat4x3GetCol1_(Mat4x3_ m2) {
+ Mat4x3_ m3 = {};
+ m3 = m2;
+ Mat4x3_ const x_e2 = m3;
+ Mat4x3_ const x_e5 = m3;
+ Mat4x3_ const x_e8 = m3;
+ return float3(x_e2.mx[1u], x_e5.my[1u], x_e8.mz[1u]);
}
-
-ub_MaterialParams = struct @align(16) {
- u_TexMtx:array<Mat4x2_, 1> @offset(0)
- u_Misc0_:vec4<f32> @offset(32)
+float3 Mat4x3GetCol2_(Mat4x3_ m4) {
+ Mat4x3_ m5 = {};
+ m5 = m4;
+ Mat4x3_ const x_e2 = m5;
+ Mat4x3_ const x_e5 = m5;
+ Mat4x3_ const x_e8 = m5;
+ return float3(x_e2.mx[2u], x_e5.my[2u], x_e8.mz[2u]);
}
-
-Mat4x3_ = struct @align(16) {
- mx:vec4<f32> @offset(0)
- my:vec4<f32> @offset(16)
- mz:vec4<f32> @offset(32)
+float3 Mat4x3GetCol3_(Mat4x3_ m6) {
+ Mat4x3_ m7 = {};
+ m7 = m6;
+ Mat4x3_ const x_e2 = m7;
+ Mat4x3_ const x_e5 = m7;
+ Mat4x3_ const x_e8 = m7;
+ return float3(x_e2.mx[3u], x_e5.my[3u], x_e8.mz[3u]);
}
-
-ub_PacketParams = struct @align(16) {
- u_PosMtx:array<Mat4x3_, 32> @offset(0)
+float4 Mul(Mat4x4_ m8, float4 v) {
+ Mat4x4_ m9 = {};
+ float4 v1 = 0.0f;
+ m9 = m8;
+ v1 = v;
+ Mat4x4_ const x_e4 = m9;
+ float4 const x_e6 = v1;
+ Mat4x4_ const x_e8 = m9;
+ float4 const x_e10 = v1;
+ Mat4x4_ const x_e12 = m9;
+ float4 const x_e14 = v1;
+ Mat4x4_ const x_e16 = m9;
+ float4 const x_e18 = v1;
+ float const v_1 = dot(x_e4.mx, x_e6);
+ float const v_2 = dot(x_e8.my, x_e10);
+ float const v_3 = dot(x_e12.mz, x_e14);
+ return float4(v_1, v_2, v_3, dot(x_e16.mw, x_e18));
}
-
-VertexOutput = struct @align(16) {
- v_Color:vec4<f32> @offset(0), @location(0)
- v_TexCoord:vec2<f32> @offset(16), @location(1)
- member:vec4<f32> @offset(32), @builtin(position)
+float3 Mul1(Mat4x3_ m10, float4 v2) {
+ Mat4x3_ m11 = {};
+ float4 v3 = 0.0f;
+ m11 = m10;
+ v3 = v2;
+ Mat4x3_ const x_e4 = m11;
+ float4 const x_e6 = v3;
+ Mat4x3_ const x_e8 = m11;
+ float4 const x_e10 = v3;
+ Mat4x3_ const x_e12 = m11;
+ float4 const x_e14 = v3;
+ float const v_4 = dot(x_e4.mx, x_e6);
+ float const v_5 = dot(x_e8.my, x_e10);
+ return float3(v_4, v_5, dot(x_e12.mz, x_e14));
}
-
-$B1: { # root
- %global:ptr<uniform, ub_SceneParams, read> = var @binding_point(0, 0)
- %global1:ptr<uniform, ub_MaterialParams, read> = var @binding_point(0, 1)
- %global2:ptr<uniform, ub_PacketParams, read> = var @binding_point(0, 2)
- %a_Position1:ptr<private, vec3<f32>, read_write> = var
- %a_UV1:ptr<private, vec2<f32>, read_write> = var
- %a_Color1:ptr<private, vec4<f32>, read_write> = var
- %a_Normal1:ptr<private, vec3<f32>, read_write> = var
- %a_PosMtxIdx1:ptr<private, f32, read_write> = var
- %v_Color:ptr<private, vec4<f32>, read_write> = var
- %v_TexCoord:ptr<private, vec2<f32>, read_write> = var
- %gl_Position:ptr<private, vec4<f32>, read_write> = var
+float2 Mul2(Mat4x2_ m12, float4 v4) {
+ Mat4x2_ m13 = {};
+ float4 v5 = 0.0f;
+ m13 = m12;
+ v5 = v4;
+ Mat4x2_ const x_e4 = m13;
+ float4 const x_e6 = v5;
+ Mat4x2_ const x_e8 = m13;
+ float4 const x_e10 = v5;
+ float const v_6 = dot(x_e4.mx, x_e6);
+ return float2(v_6, dot(x_e8.my, x_e10));
}
-
-%Mat4x3GetCol0_ = func(%m:Mat4x3_):vec3<f32> {
- $B2: {
- %m1:ptr<function, Mat4x3_, read_write> = var
- store %m1, %m
- %15:Mat4x3_ = load %m1
- %x_e2:Mat4x3_ = let %15
- %17:Mat4x3_ = load %m1
- %x_e5:Mat4x3_ = let %17
- %19:Mat4x3_ = load %m1
- %x_e8:Mat4x3_ = let %19
- %21:f32 = access %x_e2, 0u, 0u
- %22:f32 = access %x_e5, 1u, 0u
- %23:f32 = access %x_e8, 2u, 0u
- %24:vec3<f32> = construct %21, %22, %23
- ret %24
+float4 Mul3(float3 v6, Mat4x3_ m14) {
+ float3 v7 = 0.0f;
+ Mat4x3_ m15 = {};
+ v7 = v6;
+ m15 = m14;
+ Mat4x3_ const x_e5 = m15;
+ float3 const x_e6 = Mat4x3GetCol0_(x_e5);
+ float3 const x_e7 = v7;
+ Mat4x3_ const x_e10 = m15;
+ float3 const x_e11 = Mat4x3GetCol1_(x_e10);
+ float3 const x_e12 = v7;
+ Mat4x3_ const x_e15 = m15;
+ float3 const x_e16 = Mat4x3GetCol2_(x_e15);
+ float3 const x_e17 = v7;
+ Mat4x3_ const x_e20 = m15;
+ float3 const x_e21 = Mat4x3GetCol3_(x_e20);
+ float3 const x_e22 = v7;
+ float const v_7 = dot(x_e6, x_e7);
+ float const v_8 = dot(x_e11, x_e12);
+ float const v_9 = dot(x_e16, x_e17);
+ return float4(v_7, v_8, v_9, dot(x_e21, x_e22));
+}
+Mat4x4_ x_Mat4x4_(float n) {
+ float n1 = 0.0f;
+ Mat4x4_ o = {};
+ n1 = n;
+ float const x_e4 = n1;
+ o.mx = float4(x_e4, 0.0f, 0.0f, 0.0f);
+ float const x_e11 = n1;
+ o.my = float4(0.0f, x_e11, 0.0f, 0.0f);
+ float const x_e18 = n1;
+ o.mz = float4(0.0f, 0.0f, x_e18, 0.0f);
+ float const x_e25 = n1;
+ o.mw = float4(0.0f, 0.0f, 0.0f, x_e25);
+ Mat4x4_ const x_e27 = o;
+ return x_e27;
+}
+Mat4x4_ x_Mat4x4_1(Mat4x3_ m16) {
+ Mat4x3_ m17 = {};
+ Mat4x4_ o1 = {};
+ m17 = m16;
+ Mat4x4_ const x_e4 = x_Mat4x4_(1.0f);
+ o1 = x_e4;
+ Mat4x3_ const x_e7 = m17;
+ o1.mx = x_e7.mx;
+ Mat4x3_ const x_e10 = m17;
+ o1.my = x_e10.my;
+ Mat4x3_ const x_e13 = m17;
+ o1.mz = x_e13.mz;
+ Mat4x4_ const x_e15 = o1;
+ return x_e15;
+}
+Mat4x4_ x_Mat4x4_2(Mat4x2_ m18) {
+ Mat4x2_ m19 = {};
+ Mat4x4_ o2 = {};
+ m19 = m18;
+ Mat4x4_ const x_e4 = x_Mat4x4_(1.0f);
+ o2 = x_e4;
+ Mat4x2_ const x_e7 = m19;
+ o2.mx = x_e7.mx;
+ Mat4x2_ const x_e10 = m19;
+ o2.my = x_e10.my;
+ Mat4x4_ const x_e12 = o2;
+ return x_e12;
+}
+Mat4x3_ x_Mat4x3_(float n2) {
+ float n3 = 0.0f;
+ Mat4x3_ o3 = {};
+ n3 = n2;
+ float const x_e4 = n3;
+ o3.mx = float4(x_e4, 0.0f, 0.0f, 0.0f);
+ float const x_e11 = n3;
+ o3.my = float4(0.0f, x_e11, 0.0f, 0.0f);
+ float const x_e18 = n3;
+ o3.mz = float4(0.0f, 0.0f, x_e18, 0.0f);
+ Mat4x3_ const x_e21 = o3;
+ return x_e21;
+}
+Mat4x3_ x_Mat4x3_1(Mat4x4_ m20) {
+ Mat4x4_ m21 = {};
+ Mat4x3_ o4 = {};
+ m21 = m20;
+ Mat4x4_ const x_e4 = m21;
+ o4.mx = x_e4.mx;
+ Mat4x4_ const x_e7 = m21;
+ o4.my = x_e7.my;
+ Mat4x4_ const x_e10 = m21;
+ o4.mz = x_e10.mz;
+ Mat4x3_ const x_e12 = o4;
+ return x_e12;
+}
+int tint_f32_to_i32(float value) {
+ return select(2147483647, select((-2147483647 - 1), int(value), (value >= -2147483648.0f)), (value <= 2147483520.0f));
+}
+void main1(tint_module_vars_struct tint_module_vars) {
+ Mat4x3_ t_PosMtx = {};
+ float2 t_TexSpaceCoord = 0.0f;
+ float const x_e15 = (*tint_module_vars.a_PosMtxIdx1);
+ Mat4x3_ const x_e18 = (*tint_module_vars.global2).u_PosMtx[tint_f32_to_i32(x_e15)];
+ t_PosMtx = x_e18;
+ Mat4x3_ const x_e23 = t_PosMtx;
+ Mat4x4_ const x_e24 = x_Mat4x4_1(x_e23);
+ float3 const x_e25 = (*tint_module_vars.a_Position1);
+ Mat4x3_ const x_e29 = t_PosMtx;
+ Mat4x4_ const x_e30 = x_Mat4x4_1(x_e29);
+ float3 const x_e31 = (*tint_module_vars.a_Position1);
+ float4 const x_e34 = Mul(x_e30, float4(x_e31, 1.0f));
+ Mat4x4_ const x_e35 = (*tint_module_vars.global).u_Projection;
+ Mat4x3_ const x_e37 = t_PosMtx;
+ Mat4x4_ const x_e38 = x_Mat4x4_1(x_e37);
+ float3 const x_e39 = (*tint_module_vars.a_Position1);
+ Mat4x3_ const x_e43 = t_PosMtx;
+ Mat4x4_ const x_e44 = x_Mat4x4_1(x_e43);
+ float3 const x_e45 = (*tint_module_vars.a_Position1);
+ float4 const x_e48 = Mul(x_e44, float4(x_e45, 1.0f));
+ float4 const x_e49 = Mul(x_e35, x_e48);
+ (*tint_module_vars.gl_Position) = x_e49;
+ float4 const x_e50 = (*tint_module_vars.a_Color1);
+ (*tint_module_vars.v_Color) = x_e50;
+ float4 const x_e52 = (*tint_module_vars.global1).u_Misc0_;
+ if ((x_e52[0u] == 2.0f)) {
+ float3 const x_e59 = (*tint_module_vars.a_Normal1);
+ Mat4x2_ const x_e64 = (*tint_module_vars.global1).u_TexMtx[0];
+ float3 const x_e65 = (*tint_module_vars.a_Normal1);
+ float2 const x_e68 = Mul2(x_e64, float4(x_e65, 1.0f));
+ (*tint_module_vars.v_TexCoord) = x_e68.xy;
+ return;
+ } else {
+ float2 const x_e73 = (*tint_module_vars.a_UV1);
+ Mat4x2_ const x_e79 = (*tint_module_vars.global1).u_TexMtx[0];
+ float2 const x_e80 = (*tint_module_vars.a_UV1);
+ float2 const x_e84 = Mul2(x_e79, float4(x_e80, 1.0f, 1.0f));
+ (*tint_module_vars.v_TexCoord) = x_e84.xy;
+ return;
}
+ /* unreachable */
}
-%Mat4x3GetCol1_ = func(%m2:Mat4x3_):vec3<f32> {
- $B3: {
- %m3:ptr<function, Mat4x3_, read_write> = var
- store %m3, %m2
- %28:Mat4x3_ = load %m3
- %x_e2_1:Mat4x3_ = let %28 # %x_e2_1: 'x_e2'
- %30:Mat4x3_ = load %m3
- %x_e5_1:Mat4x3_ = let %30 # %x_e5_1: 'x_e5'
- %32:Mat4x3_ = load %m3
- %x_e8_1:Mat4x3_ = let %32 # %x_e8_1: 'x_e8'
- %34:f32 = access %x_e2_1, 0u, 1u
- %35:f32 = access %x_e5_1, 1u, 1u
- %36:f32 = access %x_e8_1, 2u, 1u
- %37:vec3<f32> = construct %34, %35, %36
- ret %37
- }
+VertexOutput tint_symbol_inner(float3 a_Position, float2 a_UV, float4 a_Color, float3 a_Normal, float a_PosMtxIdx, tint_module_vars_struct tint_module_vars) {
+ (*tint_module_vars.a_Position1) = a_Position;
+ (*tint_module_vars.a_UV1) = a_UV;
+ (*tint_module_vars.a_Color1) = a_Color;
+ (*tint_module_vars.a_Normal1) = a_Normal;
+ (*tint_module_vars.a_PosMtxIdx1) = a_PosMtxIdx;
+ main1(tint_module_vars);
+ float4 const x_e11 = (*tint_module_vars.v_Color);
+ float2 const x_e13 = (*tint_module_vars.v_TexCoord);
+ float4 const x_e15 = (*tint_module_vars.gl_Position);
+ return VertexOutput{.v_Color=x_e11, .v_TexCoord=x_e13, .member=x_e15};
}
-%Mat4x3GetCol2_ = func(%m4:Mat4x3_):vec3<f32> {
- $B4: {
- %m5:ptr<function, Mat4x3_, read_write> = var
- store %m5, %m4
- %41:Mat4x3_ = load %m5
- %x_e2_2:Mat4x3_ = let %41 # %x_e2_2: 'x_e2'
- %43:Mat4x3_ = load %m5
- %x_e5_2:Mat4x3_ = let %43 # %x_e5_2: 'x_e5'
- %45:Mat4x3_ = load %m5
- %x_e8_2:Mat4x3_ = let %45 # %x_e8_2: 'x_e8'
- %47:f32 = access %x_e2_2, 0u, 2u
- %48:f32 = access %x_e5_2, 1u, 2u
- %49:f32 = access %x_e8_2, 2u, 2u
- %50:vec3<f32> = construct %47, %48, %49
- ret %50
- }
+vertex tint_symbol_outputs tint_symbol(tint_symbol_inputs inputs [[stage_in]], const constant ub_SceneParams* global [[buffer(1)]], const constant ub_MaterialParams* global1 [[buffer(2)]], const constant ub_PacketParams* global2 [[buffer(0)]]) {
+ thread float3 a_Position1 = 0.0f;
+ thread float2 a_UV1 = 0.0f;
+ thread float4 a_Color1 = 0.0f;
+ thread float3 a_Normal1 = 0.0f;
+ thread float a_PosMtxIdx1 = 0.0f;
+ thread float4 v_Color = 0.0f;
+ thread float2 v_TexCoord = 0.0f;
+ thread float4 gl_Position = 0.0f;
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.global=global, .global1=global1, .global2=global2, .a_Position1=(&a_Position1), .a_UV1=(&a_UV1), .a_Color1=(&a_Color1), .a_Normal1=(&a_Normal1), .a_PosMtxIdx1=(&a_PosMtxIdx1), .v_Color=(&v_Color), .v_TexCoord=(&v_TexCoord), .gl_Position=(&gl_Position)};
+ VertexOutput const v_10 = tint_symbol_inner(inputs.a_Position, inputs.a_UV, inputs.a_Color, inputs.a_Normal, inputs.a_PosMtxIdx, tint_module_vars);
+ return tint_symbol_outputs{.VertexOutput_v_Color=v_10.v_Color, .VertexOutput_v_TexCoord=v_10.v_TexCoord, .VertexOutput_member=v_10.member};
}
-%Mat4x3GetCol3_ = func(%m6:Mat4x3_):vec3<f32> {
- $B5: {
- %m7:ptr<function, Mat4x3_, read_write> = var
- store %m7, %m6
- %54:Mat4x3_ = load %m7
- %x_e2_3:Mat4x3_ = let %54 # %x_e2_3: 'x_e2'
- %56:Mat4x3_ = load %m7
- %x_e5_3:Mat4x3_ = let %56 # %x_e5_3: 'x_e5'
- %58:Mat4x3_ = load %m7
- %x_e8_3:Mat4x3_ = let %58 # %x_e8_3: 'x_e8'
- %60:f32 = access %x_e2_3, 0u, 3u
- %61:f32 = access %x_e5_3, 1u, 3u
- %62:f32 = access %x_e8_3, 2u, 3u
- %63:vec3<f32> = construct %60, %61, %62
- ret %63
- }
-}
-%Mul = func(%m8:Mat4x4_, %v:vec4<f32>):vec4<f32> {
- $B6: {
- %m9:ptr<function, Mat4x4_, read_write> = var
- %v1:ptr<function, vec4<f32>, read_write> = var
- store %m9, %m8
- store %v1, %v
- %69:Mat4x4_ = load %m9
- %x_e4:Mat4x4_ = let %69
- %71:vec4<f32> = load %v1
- %x_e6:vec4<f32> = let %71
- %73:Mat4x4_ = load %m9
- %x_e8_4:Mat4x4_ = let %73 # %x_e8_4: 'x_e8'
- %75:vec4<f32> = load %v1
- %x_e10:vec4<f32> = let %75
- %77:Mat4x4_ = load %m9
- %x_e12:Mat4x4_ = let %77
- %79:vec4<f32> = load %v1
- %x_e14:vec4<f32> = let %79
- %81:Mat4x4_ = load %m9
- %x_e16:Mat4x4_ = let %81
- %83:vec4<f32> = load %v1
- %x_e18:vec4<f32> = let %83
- %85:vec4<f32> = access %x_e4, 0u
- %86:f32 = dot %85, %x_e6
- %87:f32 = let %86
- %88:vec4<f32> = access %x_e8_4, 1u
- %89:f32 = dot %88, %x_e10
- %90:f32 = let %89
- %91:vec4<f32> = access %x_e12, 2u
- %92:f32 = dot %91, %x_e14
- %93:f32 = let %92
- %94:vec4<f32> = access %x_e16, 3u
- %95:f32 = dot %94, %x_e18
- %96:vec4<f32> = construct %87, %90, %93, %95
- ret %96
- }
-}
-%Mul1 = func(%m10:Mat4x3_, %v2:vec4<f32>):vec3<f32> {
- $B7: {
- %m11:ptr<function, Mat4x3_, read_write> = var
- %v3:ptr<function, vec4<f32>, read_write> = var
- store %m11, %m10
- store %v3, %v2
- %102:Mat4x3_ = load %m11
- %x_e4_1:Mat4x3_ = let %102 # %x_e4_1: 'x_e4'
- %104:vec4<f32> = load %v3
- %x_e6_1:vec4<f32> = let %104 # %x_e6_1: 'x_e6'
- %106:Mat4x3_ = load %m11
- %x_e8_5:Mat4x3_ = let %106 # %x_e8_5: 'x_e8'
- %108:vec4<f32> = load %v3
- %x_e10_1:vec4<f32> = let %108 # %x_e10_1: 'x_e10'
- %110:Mat4x3_ = load %m11
- %x_e12_1:Mat4x3_ = let %110 # %x_e12_1: 'x_e12'
- %112:vec4<f32> = load %v3
- %x_e14_1:vec4<f32> = let %112 # %x_e14_1: 'x_e14'
- %114:vec4<f32> = access %x_e4_1, 0u
- %115:f32 = dot %114, %x_e6_1
- %116:f32 = let %115
- %117:vec4<f32> = access %x_e8_5, 1u
- %118:f32 = dot %117, %x_e10_1
- %119:f32 = let %118
- %120:vec4<f32> = access %x_e12_1, 2u
- %121:f32 = dot %120, %x_e14_1
- %122:vec3<f32> = construct %116, %119, %121
- ret %122
- }
-}
-%Mul2 = func(%m12:Mat4x2_, %v4:vec4<f32>):vec2<f32> {
- $B8: {
- %m13:ptr<function, Mat4x2_, read_write> = var
- %v5:ptr<function, vec4<f32>, read_write> = var
- store %m13, %m12
- store %v5, %v4
- %128:Mat4x2_ = load %m13
- %x_e4_2:Mat4x2_ = let %128 # %x_e4_2: 'x_e4'
- %130:vec4<f32> = load %v5
- %x_e6_2:vec4<f32> = let %130 # %x_e6_2: 'x_e6'
- %132:Mat4x2_ = load %m13
- %x_e8_6:Mat4x2_ = let %132 # %x_e8_6: 'x_e8'
- %134:vec4<f32> = load %v5
- %x_e10_2:vec4<f32> = let %134 # %x_e10_2: 'x_e10'
- %136:vec4<f32> = access %x_e4_2, 0u
- %137:f32 = dot %136, %x_e6_2
- %138:f32 = let %137
- %139:vec4<f32> = access %x_e8_6, 1u
- %140:f32 = dot %139, %x_e10_2
- %141:vec2<f32> = construct %138, %140
- ret %141
- }
-}
-%Mul3 = func(%v6:vec3<f32>, %m14:Mat4x3_):vec4<f32> {
- $B9: {
- %v7:ptr<function, vec3<f32>, read_write> = var
- %m15:ptr<function, Mat4x3_, read_write> = var
- store %v7, %v6
- store %m15, %m14
- %147:Mat4x3_ = load %m15
- %x_e5_4:Mat4x3_ = let %147 # %x_e5_4: 'x_e5'
- %149:vec3<f32> = call %Mat4x3GetCol0_, %x_e5_4
- %x_e6_3:vec3<f32> = let %149 # %x_e6_3: 'x_e6'
- %151:vec3<f32> = load %v7
- %x_e7:vec3<f32> = let %151
- %153:Mat4x3_ = load %m15
- %x_e10_3:Mat4x3_ = let %153 # %x_e10_3: 'x_e10'
- %155:vec3<f32> = call %Mat4x3GetCol1_, %x_e10_3
- %x_e11:vec3<f32> = let %155
- %157:vec3<f32> = load %v7
- %x_e12_2:vec3<f32> = let %157 # %x_e12_2: 'x_e12'
- %159:Mat4x3_ = load %m15
- %x_e15:Mat4x3_ = let %159
- %161:vec3<f32> = call %Mat4x3GetCol2_, %x_e15
- %x_e16_1:vec3<f32> = let %161 # %x_e16_1: 'x_e16'
- %163:vec3<f32> = load %v7
- %x_e17:vec3<f32> = let %163
- %165:Mat4x3_ = load %m15
- %x_e20:Mat4x3_ = let %165
- %167:vec3<f32> = call %Mat4x3GetCol3_, %x_e20
- %x_e21:vec3<f32> = let %167
- %169:vec3<f32> = load %v7
- %x_e22:vec3<f32> = let %169
- %171:f32 = dot %x_e6_3, %x_e7
- %172:f32 = let %171
- %173:f32 = dot %x_e11, %x_e12_2
- %174:f32 = let %173
- %175:f32 = dot %x_e16_1, %x_e17
- %176:f32 = let %175
- %177:f32 = dot %x_e21, %x_e22
- %178:vec4<f32> = construct %172, %174, %176, %177
- ret %178
- }
-}
-%x_Mat4x4_ = func(%n:f32):Mat4x4_ {
- $B10: {
- %n1:ptr<function, f32, read_write> = var
- %o:ptr<function, Mat4x4_, read_write> = var
- store %n1, %n
- %183:f32 = load %n1
- %x_e4_3:f32 = let %183 # %x_e4_3: 'x_e4'
- %185:ptr<function, vec4<f32>, read_write> = access %o, 0u
- %186:vec4<f32> = construct %x_e4_3, 0.0f, 0.0f, 0.0f
- store %185, %186
- %187:f32 = load %n1
- %x_e11_1:f32 = let %187 # %x_e11_1: 'x_e11'
- %189:ptr<function, vec4<f32>, read_write> = access %o, 1u
- %190:vec4<f32> = construct 0.0f, %x_e11_1, 0.0f, 0.0f
- store %189, %190
- %191:f32 = load %n1
- %x_e18_1:f32 = let %191 # %x_e18_1: 'x_e18'
- %193:ptr<function, vec4<f32>, read_write> = access %o, 2u
- %194:vec4<f32> = construct 0.0f, 0.0f, %x_e18_1, 0.0f
- store %193, %194
- %195:f32 = load %n1
- %x_e25:f32 = let %195
- %197:ptr<function, vec4<f32>, read_write> = access %o, 3u
- %198:vec4<f32> = construct 0.0f, 0.0f, 0.0f, %x_e25
- store %197, %198
- %199:Mat4x4_ = load %o
- %x_e27:Mat4x4_ = let %199
- ret %x_e27
- }
-}
-%x_Mat4x4_1 = func(%m16:Mat4x3_):Mat4x4_ {
- $B11: {
- %m17:ptr<function, Mat4x3_, read_write> = var
- %o1:ptr<function, Mat4x4_, read_write> = var
- store %m17, %m16
- %205:Mat4x4_ = call %x_Mat4x4_, 1.0f
- %x_e4_4:Mat4x4_ = let %205 # %x_e4_4: 'x_e4'
- store %o1, %x_e4_4
- %207:Mat4x3_ = load %m17
- %x_e7_1:Mat4x3_ = let %207 # %x_e7_1: 'x_e7'
- %209:ptr<function, vec4<f32>, read_write> = access %o1, 0u
- %210:vec4<f32> = access %x_e7_1, 0u
- store %209, %210
- %211:Mat4x3_ = load %m17
- %x_e10_4:Mat4x3_ = let %211 # %x_e10_4: 'x_e10'
- %213:ptr<function, vec4<f32>, read_write> = access %o1, 1u
- %214:vec4<f32> = access %x_e10_4, 1u
- store %213, %214
- %215:Mat4x3_ = load %m17
- %x_e13:Mat4x3_ = let %215
- %217:ptr<function, vec4<f32>, read_write> = access %o1, 2u
- %218:vec4<f32> = access %x_e13, 2u
- store %217, %218
- %219:Mat4x4_ = load %o1
- %x_e15_1:Mat4x4_ = let %219 # %x_e15_1: 'x_e15'
- ret %x_e15_1
- }
-}
-%x_Mat4x4_2 = func(%m18:Mat4x2_):Mat4x4_ {
- $B12: {
- %m19:ptr<function, Mat4x2_, read_write> = var
- %o2:ptr<function, Mat4x4_, read_write> = var
- store %m19, %m18
- %225:Mat4x4_ = call %x_Mat4x4_, 1.0f
- %x_e4_5:Mat4x4_ = let %225 # %x_e4_5: 'x_e4'
- store %o2, %x_e4_5
- %227:Mat4x2_ = load %m19
- %x_e7_2:Mat4x2_ = let %227 # %x_e7_2: 'x_e7'
- %229:ptr<function, vec4<f32>, read_write> = access %o2, 0u
- %230:vec4<f32> = access %x_e7_2, 0u
- store %229, %230
- %231:Mat4x2_ = load %m19
- %x_e10_5:Mat4x2_ = let %231 # %x_e10_5: 'x_e10'
- %233:ptr<function, vec4<f32>, read_write> = access %o2, 1u
- %234:vec4<f32> = access %x_e10_5, 1u
- store %233, %234
- %235:Mat4x4_ = load %o2
- %x_e12_3:Mat4x4_ = let %235 # %x_e12_3: 'x_e12'
- ret %x_e12_3
- }
-}
-%x_Mat4x3_ = func(%n2:f32):Mat4x3_ {
- $B13: {
- %n3:ptr<function, f32, read_write> = var
- %o3:ptr<function, Mat4x3_, read_write> = var
- store %n3, %n2
- %241:f32 = load %n3
- %x_e4_6:f32 = let %241 # %x_e4_6: 'x_e4'
- %243:ptr<function, vec4<f32>, read_write> = access %o3, 0u
- %244:vec4<f32> = construct %x_e4_6, 0.0f, 0.0f, 0.0f
- store %243, %244
- %245:f32 = load %n3
- %x_e11_2:f32 = let %245 # %x_e11_2: 'x_e11'
- %247:ptr<function, vec4<f32>, read_write> = access %o3, 1u
- %248:vec4<f32> = construct 0.0f, %x_e11_2, 0.0f, 0.0f
- store %247, %248
- %249:f32 = load %n3
- %x_e18_2:f32 = let %249 # %x_e18_2: 'x_e18'
- %251:ptr<function, vec4<f32>, read_write> = access %o3, 2u
- %252:vec4<f32> = construct 0.0f, 0.0f, %x_e18_2, 0.0f
- store %251, %252
- %253:Mat4x3_ = load %o3
- %x_e21_1:Mat4x3_ = let %253 # %x_e21_1: 'x_e21'
- ret %x_e21_1
- }
-}
-%x_Mat4x3_1 = func(%m20:Mat4x4_):Mat4x3_ {
- $B14: {
- %m21:ptr<function, Mat4x4_, read_write> = var
- %o4:ptr<function, Mat4x3_, read_write> = var
- store %m21, %m20
- %259:Mat4x4_ = load %m21
- %x_e4_7:Mat4x4_ = let %259 # %x_e4_7: 'x_e4'
- %261:ptr<function, vec4<f32>, read_write> = access %o4, 0u
- %262:vec4<f32> = access %x_e4_7, 0u
- store %261, %262
- %263:Mat4x4_ = load %m21
- %x_e7_3:Mat4x4_ = let %263 # %x_e7_3: 'x_e7'
- %265:ptr<function, vec4<f32>, read_write> = access %o4, 1u
- %266:vec4<f32> = access %x_e7_3, 1u
- store %265, %266
- %267:Mat4x4_ = load %m21
- %x_e10_6:Mat4x4_ = let %267 # %x_e10_6: 'x_e10'
- %269:ptr<function, vec4<f32>, read_write> = access %o4, 2u
- %270:vec4<f32> = access %x_e10_6, 2u
- store %269, %270
- %271:Mat4x3_ = load %o4
- %x_e12_4:Mat4x3_ = let %271 # %x_e12_4: 'x_e12'
- ret %x_e12_4
- }
-}
-%main1 = func():void {
- $B15: {
- %t_PosMtx:ptr<function, Mat4x3_, read_write> = var
- %t_TexSpaceCoord:ptr<function, vec2<f32>, read_write> = var
- %276:f32 = load %a_PosMtxIdx1
- %x_e15_2:f32 = let %276 # %x_e15_2: 'x_e15'
- %278:i32 = call %tint_f32_to_i32, %x_e15_2
- %280:ptr<uniform, Mat4x3_, read> = access %global2, 0u, %278
- %281:Mat4x3_ = load %280
- %x_e18_3:Mat4x3_ = let %281 # %x_e18_3: 'x_e18'
- store %t_PosMtx, %x_e18_3
- %283:Mat4x3_ = load %t_PosMtx
- %x_e23:Mat4x3_ = let %283
- %285:Mat4x4_ = call %x_Mat4x4_1, %x_e23
- %x_e24:Mat4x4_ = let %285
- %287:vec3<f32> = load %a_Position1
- %x_e25_1:vec3<f32> = let %287 # %x_e25_1: 'x_e25'
- %289:Mat4x3_ = load %t_PosMtx
- %x_e29:Mat4x3_ = let %289
- %291:Mat4x4_ = call %x_Mat4x4_1, %x_e29
- %x_e30:Mat4x4_ = let %291
- %293:vec3<f32> = load %a_Position1
- %x_e31:vec3<f32> = let %293
- %295:vec4<f32> = construct %x_e31, 1.0f
- %296:vec4<f32> = call %Mul, %x_e30, %295
- %x_e34:vec4<f32> = let %296
- %298:ptr<uniform, Mat4x4_, read> = access %global, 0u
- %299:Mat4x4_ = load %298
- %x_e35:Mat4x4_ = let %299
- %301:Mat4x3_ = load %t_PosMtx
- %x_e37:Mat4x3_ = let %301
- %303:Mat4x4_ = call %x_Mat4x4_1, %x_e37
- %x_e38:Mat4x4_ = let %303
- %305:vec3<f32> = load %a_Position1
- %x_e39:vec3<f32> = let %305
- %307:Mat4x3_ = load %t_PosMtx
- %x_e43:Mat4x3_ = let %307
- %309:Mat4x4_ = call %x_Mat4x4_1, %x_e43
- %x_e44:Mat4x4_ = let %309
- %311:vec3<f32> = load %a_Position1
- %x_e45:vec3<f32> = let %311
- %313:vec4<f32> = construct %x_e45, 1.0f
- %314:vec4<f32> = call %Mul, %x_e44, %313
- %x_e48:vec4<f32> = let %314
- %316:vec4<f32> = call %Mul, %x_e35, %x_e48
- %x_e49:vec4<f32> = let %316
- store %gl_Position, %x_e49
- %318:vec4<f32> = load %a_Color1
- %x_e50:vec4<f32> = let %318
- store %v_Color, %x_e50
- %320:ptr<uniform, vec4<f32>, read> = access %global1, 1u
- %321:vec4<f32> = load %320
- %x_e52:vec4<f32> = let %321
- %323:f32 = access %x_e52, 0u
- %324:bool = eq %323, 2.0f
- if %324 [t: $B16, f: $B17] { # if_1
- $B16: { # true
- %325:vec3<f32> = load %a_Normal1
- %x_e59:vec3<f32> = let %325
- %327:ptr<uniform, Mat4x2_, read> = access %global1, 0u, 0i
- %328:Mat4x2_ = load %327
- %x_e64:Mat4x2_ = let %328
- %330:vec3<f32> = load %a_Normal1
- %x_e65:vec3<f32> = let %330
- %332:vec4<f32> = construct %x_e65, 1.0f
- %333:vec2<f32> = call %Mul2, %x_e64, %332
- %x_e68:vec2<f32> = let %333
- %335:vec2<f32> = swizzle %x_e68, xy
- store %v_TexCoord, %335
- ret
- }
- $B17: { # false
- %336:vec2<f32> = load %a_UV1
- %x_e73:vec2<f32> = let %336
- %338:ptr<uniform, Mat4x2_, read> = access %global1, 0u, 0i
- %339:Mat4x2_ = load %338
- %x_e79:Mat4x2_ = let %339
- %341:vec2<f32> = load %a_UV1
- %x_e80:vec2<f32> = let %341
- %343:vec4<f32> = construct %x_e80, 1.0f, 1.0f
- %344:vec2<f32> = call %Mul2, %x_e79, %343
- %x_e84:vec2<f32> = let %344
- %346:vec2<f32> = swizzle %x_e84, xy
- store %v_TexCoord, %346
- ret
- }
- }
- unreachable
- }
-}
-%tint_symbol = @vertex func(%a_Position:vec3<f32> [@location(0)], %a_UV:vec2<f32> [@location(1)], %a_Color:vec4<f32> [@location(2)], %a_Normal:vec3<f32> [@location(3)], %a_PosMtxIdx:f32 [@location(4)]):VertexOutput {
- $B18: {
- store %a_Position1, %a_Position
- store %a_UV1, %a_UV
- store %a_Color1, %a_Color
- store %a_Normal1, %a_Normal
- store %a_PosMtxIdx1, %a_PosMtxIdx
- %353:void = call %main1
- %354:vec4<f32> = load %v_Color
- %x_e11_3:vec4<f32> = let %354 # %x_e11_3: 'x_e11'
- %356:vec2<f32> = load %v_TexCoord
- %x_e13_1:vec2<f32> = let %356 # %x_e13_1: 'x_e13'
- %358:vec4<f32> = load %gl_Position
- %x_e15_3:vec4<f32> = let %358 # %x_e15_3: 'x_e15'
- %360:VertexOutput = construct %x_e11_3, %x_e13_1, %x_e15_3
- ret %360
- }
-}
-%tint_f32_to_i32 = func(%value:f32):i32 {
- $B19: {
- %362:i32 = convert %value
- %363:bool = gte %value, -2147483648.0f
- %364:i32 = select -2147483648i, %362, %363
- %365:bool = lte %value, 2147483520.0f
- %366:i32 = select 2147483647i, %364, %365
- ret %366
- }
-}
-
-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. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/dot/0c577b.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/0c577b.wgsl.expected.ir.msl
index 98d607a..b0f2c87 100644
--- a/test/tint/builtins/gen/var/dot/0c577b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/0c577b.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float dot_0c577b() {
+ float4 arg_0 = float4(1.0f);
+ float4 arg_1 = float4(1.0f);
+ float res = dot(arg_0, arg_1);
+ return res;
}
-
-%dot_0c577b = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f32>, read_write> = var, vec4<f32>(1.0f)
- %arg_1:ptr<function, vec4<f32>, read_write> = var, vec4<f32>(1.0f)
- %5:vec4<f32> = load %arg_0
- %6:vec4<f32> = load %arg_1
- %7:f32 = dot %5, %6
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_0c577b();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_0c577b
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_0c577b();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_0c577b
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_0c577b();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_0c577b
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/dot/7548a0.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/7548a0.wgsl.expected.ir.msl
index 0d58a72..af4a46d 100644
--- a/test/tint/builtins/gen/var/dot/7548a0.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/7548a0.wgsl.expected.ir.msl
@@ -1,45 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device uint* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ uint prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ uint VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+uint tint_dot(uint3 lhs, uint3 rhs) {
+ return (((lhs * rhs)[0u] + (lhs * rhs)[1u]) + (lhs * rhs)[2u]);
}
-
-%dot_7548a0 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<u32>, read_write> = var, vec3<u32>(1u)
- %arg_1:ptr<function, vec3<u32>, read_write> = var, vec3<u32>(1u)
- %5:vec3<u32> = load %arg_0
- %6:vec3<u32> = load %arg_1
- %7:u32 = dot %5, %6
- %res:ptr<function, u32, read_write> = var, %7
- %9:u32 = load %res
- store %prevent_dce, %9
- ret
- }
+uint dot_7548a0() {
+ uint3 arg_0 = uint3(1u);
+ uint3 arg_1 = uint3(1u);
+ uint res = tint_dot(arg_0, arg_1);
+ return res;
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_7548a0
- ret vec4<f32>(0.0f)
- }
+fragment void fragment_main(device uint* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_7548a0();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_7548a0
- ret
- }
+kernel void compute_main(device uint* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_7548a0();
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_7548a0
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_7548a0();
+ return out;
}
-
-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. *
-********************************************************************
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
+}
diff --git a/test/tint/builtins/gen/var/dot/883f0e.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/883f0e.wgsl.expected.ir.msl
index 4ba0669..ef12b12 100644
--- a/test/tint/builtins/gen/var/dot/883f0e.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/883f0e.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float dot_883f0e() {
+ float2 arg_0 = float2(1.0f);
+ float2 arg_1 = float2(1.0f);
+ float res = dot(arg_0, arg_1);
+ return res;
}
-
-%dot_883f0e = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %arg_1:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %5:vec2<f32> = load %arg_0
- %6:vec2<f32> = load %arg_1
- %7:f32 = dot %5, %6
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_883f0e();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_883f0e
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_883f0e();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_883f0e
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_883f0e();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_883f0e
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/dot/8e40f1.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/8e40f1.wgsl.expected.ir.msl
index 59ba8e4..170eb3b 100644
--- a/test/tint/builtins/gen/var/dot/8e40f1.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/8e40f1.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half dot_8e40f1() {
+ half3 arg_0 = half3(1.0h);
+ half3 arg_1 = half3(1.0h);
+ half res = dot(arg_0, arg_1);
+ return res;
}
-
-%dot_8e40f1 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f16>, read_write> = var, vec3<f16>(1.0h)
- %arg_1:ptr<function, vec3<f16>, read_write> = var, vec3<f16>(1.0h)
- %5:vec3<f16> = load %arg_0
- %6:vec3<f16> = load %arg_1
- %7:f16 = dot %5, %6
- %res:ptr<function, f16, read_write> = var, %7
- %9:f16 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_8e40f1();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_8e40f1
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_8e40f1();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_8e40f1
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_8e40f1();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_8e40f1
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/dot/97c7ee.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/97c7ee.wgsl.expected.ir.msl
index e50a72d..6d1bc7c 100644
--- a/test/tint/builtins/gen/var/dot/97c7ee.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/97c7ee.wgsl.expected.ir.msl
@@ -1,45 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device uint* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ uint prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ uint VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+uint tint_dot(uint2 lhs, uint2 rhs) {
+ return ((lhs * rhs)[0u] + (lhs * rhs)[1u]);
}
-
-%dot_97c7ee = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<u32>, read_write> = var, vec2<u32>(1u)
- %arg_1:ptr<function, vec2<u32>, read_write> = var, vec2<u32>(1u)
- %5:vec2<u32> = load %arg_0
- %6:vec2<u32> = load %arg_1
- %7:u32 = dot %5, %6
- %res:ptr<function, u32, read_write> = var, %7
- %9:u32 = load %res
- store %prevent_dce, %9
- ret
- }
+uint dot_97c7ee() {
+ uint2 arg_0 = uint2(1u);
+ uint2 arg_1 = uint2(1u);
+ uint res = tint_dot(arg_0, arg_1);
+ return res;
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_97c7ee
- ret vec4<f32>(0.0f)
- }
+fragment void fragment_main(device uint* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_97c7ee();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_97c7ee
- ret
- }
+kernel void compute_main(device uint* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_97c7ee();
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_97c7ee
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_97c7ee();
+ return out;
}
-
-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. *
-********************************************************************
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
+}
diff --git a/test/tint/builtins/gen/var/dot/ba4246.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/ba4246.wgsl.expected.ir.msl
index 95efaf3..2ab76da 100644
--- a/test/tint/builtins/gen/var/dot/ba4246.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/ba4246.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+float dot_ba4246() {
+ float3 arg_0 = float3(1.0f);
+ float3 arg_1 = float3(1.0f);
+ float res = dot(arg_0, arg_1);
+ return res;
}
-
-%dot_ba4246 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %arg_1:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %5:vec3<f32> = load %arg_0
- %6:vec3<f32> = load %arg_1
- %7:f32 = dot %5, %6
- %res:ptr<function, f32, read_write> = var, %7
- %9:f32 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_ba4246();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_ba4246
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device float* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_ba4246();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_ba4246
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_ba4246();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_ba4246
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/dot/cd5a04.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/cd5a04.wgsl.expected.ir.msl
index c06f542..ed5d545 100644
--- a/test/tint/builtins/gen/var/dot/cd5a04.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/cd5a04.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half dot_cd5a04() {
+ half2 arg_0 = half2(1.0h);
+ half2 arg_1 = half2(1.0h);
+ half res = dot(arg_0, arg_1);
+ return res;
}
-
-%dot_cd5a04 = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f16>, read_write> = var, vec2<f16>(1.0h)
- %arg_1:ptr<function, vec2<f16>, read_write> = var, vec2<f16>(1.0h)
- %5:vec2<f16> = load %arg_0
- %6:vec2<f16> = load %arg_1
- %7:f16 = dot %5, %6
- %res:ptr<function, f16, read_write> = var, %7
- %9:f16 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_cd5a04();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_cd5a04
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_cd5a04();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_cd5a04
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_cd5a04();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_cd5a04
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/dot/d0d179.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/d0d179.wgsl.expected.ir.msl
index 5f9317e..2ce04af 100644
--- a/test/tint/builtins/gen/var/dot/d0d179.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/d0d179.wgsl.expected.ir.msl
@@ -1,45 +1,38 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device half* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ half prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ half VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f16, read_write> = var @binding_point(2, 0)
+half dot_d0d179() {
+ half4 arg_0 = half4(1.0h);
+ half4 arg_1 = half4(1.0h);
+ half res = dot(arg_0, arg_1);
+ return res;
}
-
-%dot_d0d179 = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f16>, read_write> = var, vec4<f16>(1.0h)
- %arg_1:ptr<function, vec4<f16>, read_write> = var, vec4<f16>(1.0h)
- %5:vec4<f16> = load %arg_0
- %6:vec4<f16> = load %arg_1
- %7:f16 = dot %5, %6
- %res:ptr<function, f16, read_write> = var, %7
- %9:f16 = load %res
- store %prevent_dce, %9
- ret
- }
+fragment void fragment_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_d0d179();
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_d0d179
- ret vec4<f32>(0.0f)
- }
+kernel void compute_main(device half* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_d0d179();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_d0d179
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_d0d179();
+ return out;
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_d0d179
- ret
- }
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
}
-
-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. *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/dot/e994c7.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/e994c7.wgsl.expected.ir.msl
index b54bb20..551a309 100644
--- a/test/tint/builtins/gen/var/dot/e994c7.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/e994c7.wgsl.expected.ir.msl
@@ -1,45 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device uint* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ uint prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ uint VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+uint tint_dot(uint4 lhs, uint4 rhs) {
+ return ((((lhs * rhs)[0u] + (lhs * rhs)[1u]) + (lhs * rhs)[2u]) + (lhs * rhs)[3u]);
}
-
-%dot_e994c7 = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<u32>, read_write> = var, vec4<u32>(1u)
- %arg_1:ptr<function, vec4<u32>, read_write> = var, vec4<u32>(1u)
- %5:vec4<u32> = load %arg_0
- %6:vec4<u32> = load %arg_1
- %7:u32 = dot %5, %6
- %res:ptr<function, u32, read_write> = var, %7
- %9:u32 = load %res
- store %prevent_dce, %9
- ret
- }
+uint dot_e994c7() {
+ uint4 arg_0 = uint4(1u);
+ uint4 arg_1 = uint4(1u);
+ uint res = tint_dot(arg_0, arg_1);
+ return res;
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_e994c7
- ret vec4<f32>(0.0f)
- }
+fragment void fragment_main(device uint* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_e994c7();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_e994c7
- ret
- }
+kernel void compute_main(device uint* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_e994c7();
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_e994c7
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_e994c7();
+ return out;
}
-
-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. *
-********************************************************************
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
+}
diff --git a/test/tint/builtins/gen/var/dot/ef6b1d.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/ef6b1d.wgsl.expected.ir.msl
index c791587..04a2355 100644
--- a/test/tint/builtins/gen/var/dot/ef6b1d.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/ef6b1d.wgsl.expected.ir.msl
@@ -1,45 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device int* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ int prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ int VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, i32, read_write> = var @binding_point(2, 0)
+int tint_dot(int4 lhs, int4 rhs) {
+ return ((((lhs * rhs)[0u] + (lhs * rhs)[1u]) + (lhs * rhs)[2u]) + (lhs * rhs)[3u]);
}
-
-%dot_ef6b1d = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<i32>, read_write> = var, vec4<i32>(1i)
- %arg_1:ptr<function, vec4<i32>, read_write> = var, vec4<i32>(1i)
- %5:vec4<i32> = load %arg_0
- %6:vec4<i32> = load %arg_1
- %7:i32 = dot %5, %6
- %res:ptr<function, i32, read_write> = var, %7
- %9:i32 = load %res
- store %prevent_dce, %9
- ret
- }
+int dot_ef6b1d() {
+ int4 arg_0 = int4(1);
+ int4 arg_1 = int4(1);
+ int res = tint_dot(arg_0, arg_1);
+ return res;
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_ef6b1d
- ret vec4<f32>(0.0f)
- }
+fragment void fragment_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_ef6b1d();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_ef6b1d
- ret
- }
+kernel void compute_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_ef6b1d();
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_ef6b1d
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_ef6b1d();
+ return out;
}
-
-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. *
-********************************************************************
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
+}
diff --git a/test/tint/builtins/gen/var/dot/f1312c.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/f1312c.wgsl.expected.ir.msl
index 4796545..1532a7a 100644
--- a/test/tint/builtins/gen/var/dot/f1312c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/f1312c.wgsl.expected.ir.msl
@@ -1,45 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device int* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ int prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ int VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, i32, read_write> = var @binding_point(2, 0)
+int tint_dot(int3 lhs, int3 rhs) {
+ return (((lhs * rhs)[0u] + (lhs * rhs)[1u]) + (lhs * rhs)[2u]);
}
-
-%dot_f1312c = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<i32>, read_write> = var, vec3<i32>(1i)
- %arg_1:ptr<function, vec3<i32>, read_write> = var, vec3<i32>(1i)
- %5:vec3<i32> = load %arg_0
- %6:vec3<i32> = load %arg_1
- %7:i32 = dot %5, %6
- %res:ptr<function, i32, read_write> = var, %7
- %9:i32 = load %res
- store %prevent_dce, %9
- ret
- }
+int dot_f1312c() {
+ int3 arg_0 = int3(1);
+ int3 arg_1 = int3(1);
+ int res = tint_dot(arg_0, arg_1);
+ return res;
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_f1312c
- ret vec4<f32>(0.0f)
- }
+fragment void fragment_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_f1312c();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_f1312c
- ret
- }
+kernel void compute_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_f1312c();
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_f1312c
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_f1312c();
+ return out;
}
-
-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. *
-********************************************************************
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
+}
diff --git a/test/tint/builtins/gen/var/dot/fc5f7c.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot/fc5f7c.wgsl.expected.ir.msl
index 596c52e..cfff3a1 100644
--- a/test/tint/builtins/gen/var/dot/fc5f7c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot/fc5f7c.wgsl.expected.ir.msl
@@ -1,45 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device int* prevent_dce;
+};
+struct VertexOutput {
+ float4 pos;
+ int prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ int VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, i32, read_write> = var @binding_point(2, 0)
+int tint_dot(int2 lhs, int2 rhs) {
+ return ((lhs * rhs)[0u] + (lhs * rhs)[1u]);
}
-
-%dot_fc5f7c = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<i32>, read_write> = var, vec2<i32>(1i)
- %arg_1:ptr<function, vec2<i32>, read_write> = var, vec2<i32>(1i)
- %5:vec2<i32> = load %arg_0
- %6:vec2<i32> = load %arg_1
- %7:i32 = dot %5, %6
- %res:ptr<function, i32, read_write> = var, %7
- %9:i32 = load %res
- store %prevent_dce, %9
- ret
- }
+int dot_fc5f7c() {
+ int2 arg_0 = int2(1);
+ int2 arg_1 = int2(1);
+ int res = tint_dot(arg_0, arg_1);
+ return res;
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %11:void = call %dot_fc5f7c
- ret vec4<f32>(0.0f)
- }
+fragment void fragment_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_fc5f7c();
}
-%fragment_main = @fragment func():void {
- $B4: {
- %13:void = call %dot_fc5f7c
- ret
- }
+kernel void compute_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot_fc5f7c();
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %15:void = call %dot_fc5f7c
- ret
- }
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot_fc5f7c();
+ return out;
}
-
-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. *
-********************************************************************
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
+}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-access-array-dot/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-access-array-dot/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index c627809..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-access-array-dot/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,129 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: strided_arr = struct @align(4) {
- el:i32 @offset(0)
-}
-
-buf1 = struct @align(4) {
- x_GLF_uniform_int_values:array<strided_arr, 3> @offset(0)
-}
-
-strided_arr_1 = struct @align(4) {
- el:f32 @offset(0)
-}
-
-buf0 = struct @align(4) {
- x_GLF_uniform_float_values:array<strided_arr_1, 3> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf1, read> = var @binding_point(0, 1)
- %x_9:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %numbers:ptr<function, array<i32, 3>, read_write> = var
- %a:ptr<function, vec2<f32>, read_write> = var
- %b:ptr<function, f32, read_write> = var
- %8:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %9:i32 = load %8
- %x_38:i32 = let %9
- %11:ptr<function, i32, read_write> = access %numbers, %x_38
- %12:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %13:i32 = load %12
- store %11, %13
- %14:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %15:i32 = load %14
- %x_43:i32 = let %15
- %17:ptr<function, i32, read_write> = access %numbers, %x_43
- %18:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %19:i32 = load %18
- store %17, %19
- %20:ptr<uniform, i32, read> = access %x_6, 0u, 2i, 0u
- %21:i32 = load %20
- %x_48:i32 = let %21
- %23:ptr<function, i32, read_write> = access %numbers, %x_48
- %24:ptr<uniform, i32, read> = access %x_6, 0u, 2i, 0u
- %25:i32 = load %24
- store %23, %25
- %26:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %27:i32 = load %26
- %28:f32 = convert %27
- %29:f32 = let %28
- %30:ptr<uniform, f32, read> = access %x_9, 0u, 2i, 0u
- %31:f32 = load %30
- %32:bool = lt 0.0f, %31
- %33:i32 = select 2i, 1i, %32
- %34:ptr<function, i32, read_write> = access %numbers, %33
- %35:i32 = load %34
- %36:f32 = convert %35
- %37:vec2<f32> = construct %29, %36
- store %a, %37
- %38:vec2<f32> = load %a
- %39:vec2<f32> = let %38
- %40:ptr<uniform, f32, read> = access %x_9, 0u, 1i, 0u
- %41:f32 = load %40
- %42:ptr<uniform, f32, read> = access %x_9, 0u, 1i, 0u
- %43:f32 = load %42
- %44:vec2<f32> = construct %41, %43
- %45:f32 = dot %39, %44
- store %b, %45
- %46:f32 = load %b
- %47:ptr<uniform, f32, read> = access %x_9, 0u, 0i, 0u
- %48:f32 = load %47
- %49:bool = eq %46, %48
- if %49 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %50:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %51:i32 = load %50
- %52:f32 = convert %51
- %53:f32 = let %52
- %54:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %55:i32 = load %54
- %56:f32 = convert %55
- %57:f32 = let %56
- %58:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %59:i32 = load %58
- %60:f32 = convert %59
- %61:f32 = let %60
- %62:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %63:i32 = load %62
- %64:f32 = convert %63
- %65:vec4<f32> = construct %53, %57, %61, %64
- store %x_GLF_color, %65
- exit_if # if_1
- }
- $B4: { # false
- %66:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %67:i32 = load %66
- %68:f32 = convert %67
- %69:vec4<f32> = construct %68
- store %x_GLF_color, %69
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %71:void = call %main_1
- %72:vec4<f32> = load %x_GLF_color
- %73:main_out = construct %72
- ret %73
- }
-}
-
-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. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-access-array-dot/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-access-array-dot/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index bdac3c1..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-access-array-dot/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,145 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: strided_arr = struct @align(4) {
- el:i32 @offset(0)
-}
-
-buf1 = struct @align(4) {
- x_GLF_uniform_int_values:array<strided_arr, 3> @offset(0)
-}
-
-strided_arr_1 = struct @align(4) {
- el:f32 @offset(0)
-}
-
-buf0 = struct @align(4) {
- x_GLF_uniform_float_values:array<strided_arr_1, 3> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf1, read> = var @binding_point(0, 1)
- %x_9:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %numbers:ptr<function, array<i32, 3>, read_write> = var
- %a:ptr<function, vec2<f32>, read_write> = var
- %b:ptr<function, f32, read_write> = var
- %8:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %9:i32 = load %8
- %x_38:i32 = let %9
- %11:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %12:i32 = load %11
- %x_40:i32 = let %12
- %14:ptr<function, i32, read_write> = access %numbers, %x_38
- store %14, %x_40
- %15:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %16:i32 = load %15
- %x_43:i32 = let %16
- %18:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %19:i32 = load %18
- %x_45:i32 = let %19
- %21:ptr<function, i32, read_write> = access %numbers, %x_43
- store %21, %x_45
- %22:ptr<uniform, i32, read> = access %x_6, 0u, 2i, 0u
- %23:i32 = load %22
- %x_48:i32 = let %23
- %25:ptr<uniform, i32, read> = access %x_6, 0u, 2i, 0u
- %26:i32 = load %25
- %x_50:i32 = let %26
- %28:ptr<function, i32, read_write> = access %numbers, %x_48
- store %28, %x_50
- %29:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %30:i32 = load %29
- %x_53:i32 = let %30
- %32:ptr<uniform, f32, read> = access %x_9, 0u, 2i, 0u
- %33:f32 = load %32
- %x_56:f32 = let %33
- %35:bool = lt 0.0f, %x_56
- %36:i32 = select 2i, 1i, %35
- %37:ptr<function, i32, read_write> = access %numbers, %36
- %38:i32 = load %37
- %x_60:i32 = let %38
- %40:f32 = convert %x_53
- %41:f32 = let %40
- %42:f32 = convert %x_60
- %43:vec2<f32> = construct %41, %42
- store %a, %43
- %44:vec2<f32> = load %a
- %x_63:vec2<f32> = let %44
- %46:ptr<uniform, f32, read> = access %x_9, 0u, 1i, 0u
- %47:f32 = load %46
- %x_65:f32 = let %47
- %49:ptr<uniform, f32, read> = access %x_9, 0u, 1i, 0u
- %50:f32 = load %49
- %x_67:f32 = let %50
- %52:vec2<f32> = construct %x_65, %x_67
- %53:f32 = dot %x_63, %52
- store %b, %53
- %54:f32 = load %b
- %x_70:f32 = let %54
- %56:ptr<uniform, f32, read> = access %x_9, 0u, 0i, 0u
- %57:f32 = load %56
- %x_72:f32 = let %57
- %59:bool = eq %x_70, %x_72
- if %59 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %60:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %61:i32 = load %60
- %x_78:i32 = let %61
- %63:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %64:i32 = load %63
- %x_81:i32 = let %64
- %66:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %67:i32 = load %66
- %x_84:i32 = let %67
- %69:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %70:i32 = load %69
- %x_87:i32 = let %70
- %72:f32 = convert %x_78
- %73:f32 = let %72
- %74:f32 = convert %x_81
- %75:f32 = let %74
- %76:f32 = convert %x_84
- %77:f32 = let %76
- %78:f32 = convert %x_87
- %79:vec4<f32> = construct %73, %75, %77, %78
- store %x_GLF_color, %79
- exit_if # if_1
- }
- $B4: { # false
- %80:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %81:i32 = load %80
- %x_91:i32 = let %81
- %83:f32 = convert %x_91
- %x_92:f32 = let %83
- %85:vec4<f32> = construct %x_92, %x_92, %x_92, %x_92
- store %x_GLF_color, %85
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %87:void = call %main_1
- %88:vec4<f32> = load %x_GLF_color
- %89:main_out = construct %88
- ret %89
- }
-}
-
-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. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-reflect-denorm/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-reflect-denorm/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index 81acf81..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-reflect-denorm/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,124 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: strided_arr = struct @align(4) {
- el:i32 @offset(0)
-}
-
-buf0 = struct @align(4) {
- x_GLF_uniform_int_values:array<strided_arr, 5> @offset(0)
-}
-
-strided_arr_1 = struct @align(4) {
- el:f32 @offset(0)
-}
-
-buf1 = struct @align(4) {
- x_GLF_uniform_float_values:array<strided_arr_1, 3> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_9:ptr<uniform, buf1, read> = var @binding_point(0, 1)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %I:ptr<function, vec4<f32>, read_write> = var
- %N:ptr<function, vec4<f32>, read_write> = var
- %R:ptr<function, vec4<f32>, read_write> = var
- %r:ptr<function, vec4<f32>, read_write> = var
- %9:ptr<uniform, i32, read> = access %x_6, 0u, 2i, 0u
- %10:i32 = load %9
- %11:u32 = bitcast %10
- %12:u32 = let %11
- %13:ptr<uniform, i32, read> = access %x_6, 0u, 3i, 0u
- %14:i32 = load %13
- %15:u32 = bitcast %14
- %16:u32 = let %15
- %17:ptr<uniform, i32, read> = access %x_6, 0u, 4i, 0u
- %18:i32 = load %17
- %19:u32 = bitcast %18
- %20:vec4<u32> = construct %12, %16, %19, 92985u
- %21:vec4<f32> = bitcast %20
- store %I, %21
- %22:ptr<uniform, f32, read> = access %x_9, 0u, 1i, 0u
- %23:f32 = load %22
- %24:vec4<f32> = construct %23
- store %N, %24
- %25:vec4<f32> = load %I
- %26:vec4<f32> = reflect %25, vec4<f32>(0.5f)
- store %R, %26
- %27:vec4<f32> = load %I
- %28:vec4<f32> = let %27
- %29:vec4<f32> = load %N
- %30:vec4<f32> = let %29
- %31:ptr<uniform, f32, read> = access %x_9, 0u, 2i, 0u
- %32:f32 = load %31
- %33:f32 = let %32
- %34:vec4<f32> = load %N
- %35:vec4<f32> = load %I
- %36:f32 = dot %34, %35
- %37:f32 = mul %33, %36
- %38:vec4<f32> = mul %30, %37
- %39:vec4<f32> = sub %28, %38
- store %r, %39
- %40:vec4<f32> = load %R
- %41:vec4<f32> = load %r
- %42:f32 = distance %40, %41
- %43:ptr<uniform, f32, read> = access %x_9, 0u, 0i, 0u
- %44:f32 = load %43
- %45:bool = lt %42, %44
- if %45 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %46:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %47:i32 = load %46
- %48:f32 = convert %47
- %49:f32 = let %48
- %50:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %51:i32 = load %50
- %52:f32 = convert %51
- %53:f32 = let %52
- %54:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %55:i32 = load %54
- %56:f32 = convert %55
- %57:f32 = let %56
- %58:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %59:i32 = load %58
- %60:f32 = convert %59
- %61:vec4<f32> = construct %49, %53, %57, %60
- store %x_GLF_color, %61
- exit_if # if_1
- }
- $B4: { # false
- %62:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %63:i32 = load %62
- %64:f32 = convert %63
- %65:vec4<f32> = construct %64
- store %x_GLF_color, %65
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %67:void = call %main_1
- %68:vec4<f32> = load %x_GLF_color
- %69:main_out = construct %68
- ret %69
- }
-}
-
-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. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-reflect-denorm/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-reflect-denorm/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index 59492a1..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-apfloat-reflect-denorm/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,140 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: strided_arr = struct @align(4) {
- el:i32 @offset(0)
-}
-
-buf0 = struct @align(4) {
- x_GLF_uniform_int_values:array<strided_arr, 5> @offset(0)
-}
-
-strided_arr_1 = struct @align(4) {
- el:f32 @offset(0)
-}
-
-buf1 = struct @align(4) {
- x_GLF_uniform_float_values:array<strided_arr_1, 3> @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_6:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_9:ptr<uniform, buf1, read> = var @binding_point(0, 1)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %I:ptr<function, vec4<f32>, read_write> = var
- %N:ptr<function, vec4<f32>, read_write> = var
- %R:ptr<function, vec4<f32>, read_write> = var
- %r:ptr<function, vec4<f32>, read_write> = var
- %9:ptr<uniform, i32, read> = access %x_6, 0u, 2i, 0u
- %10:i32 = load %9
- %x_40:i32 = let %10
- %12:ptr<uniform, i32, read> = access %x_6, 0u, 3i, 0u
- %13:i32 = load %12
- %x_43:i32 = let %13
- %15:ptr<uniform, i32, read> = access %x_6, 0u, 4i, 0u
- %16:i32 = load %15
- %x_46:i32 = let %16
- %18:u32 = bitcast %x_40
- %19:u32 = let %18
- %20:u32 = bitcast %x_43
- %21:u32 = let %20
- %22:u32 = bitcast %x_46
- %23:vec4<u32> = construct %19, %21, %22, 92985u
- %24:vec4<f32> = bitcast %23
- store %I, %24
- %25:ptr<uniform, f32, read> = access %x_9, 0u, 1i, 0u
- %26:f32 = load %25
- %x_51:f32 = let %26
- %28:vec4<f32> = construct %x_51, %x_51, %x_51, %x_51
- store %N, %28
- %29:vec4<f32> = load %I
- %x_53:vec4<f32> = let %29
- %31:vec4<f32> = reflect %x_53, vec4<f32>(0.5f)
- store %R, %31
- %32:vec4<f32> = load %I
- %x_55:vec4<f32> = let %32
- %34:ptr<uniform, f32, read> = access %x_9, 0u, 2i, 0u
- %35:f32 = load %34
- %x_57:f32 = let %35
- %37:vec4<f32> = load %N
- %x_58:vec4<f32> = let %37
- %39:vec4<f32> = load %I
- %x_59:vec4<f32> = let %39
- %41:vec4<f32> = load %N
- %x_62:vec4<f32> = let %41
- %43:f32 = dot %x_58, %x_59
- %44:f32 = mul %x_57, %43
- %45:vec4<f32> = mul %x_62, %44
- %46:vec4<f32> = sub %x_55, %45
- store %r, %46
- %47:vec4<f32> = load %R
- %x_65:vec4<f32> = let %47
- %49:vec4<f32> = load %r
- %x_66:vec4<f32> = let %49
- %51:ptr<uniform, f32, read> = access %x_9, 0u, 0i, 0u
- %52:f32 = load %51
- %x_69:f32 = let %52
- %54:f32 = distance %x_65, %x_66
- %55:bool = lt %54, %x_69
- if %55 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- %56:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %57:i32 = load %56
- %x_75:i32 = let %57
- %59:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %60:i32 = load %59
- %x_78:i32 = let %60
- %62:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %63:i32 = load %62
- %x_81:i32 = let %63
- %65:ptr<uniform, i32, read> = access %x_6, 0u, 0i, 0u
- %66:i32 = load %65
- %x_84:i32 = let %66
- %68:f32 = convert %x_75
- %69:f32 = let %68
- %70:f32 = convert %x_78
- %71:f32 = let %70
- %72:f32 = convert %x_81
- %73:f32 = let %72
- %74:f32 = convert %x_84
- %75:vec4<f32> = construct %69, %71, %73, %74
- store %x_GLF_color, %75
- exit_if # if_1
- }
- $B4: { # false
- %76:ptr<uniform, i32, read> = access %x_6, 0u, 1i, 0u
- %77:i32 = load %76
- %x_88:i32 = let %77
- %79:f32 = convert %x_88
- %x_89:f32 = let %79
- %81:vec4<f32> = construct %x_89, %x_89, %x_89, %x_89
- store %x_GLF_color, %81
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %83:void = call %main_1
- %84:vec4<f32> = load %x_GLF_color
- %85:main_out = construct %84
- ret %85
- }
-}
-
-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. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-const-folding-dot-condition-true/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-const-folding-dot-condition-true/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index 0bfdeba..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-const-folding-dot-condition-true/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,9 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: dot
-********************************************************************
-* 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. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-const-folding-dot-condition-true/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-const-folding-dot-condition-true/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index 0bfdeba..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-const-folding-dot-condition-true/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,9 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: dot
-********************************************************************
-* 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. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-extract/0-opt.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-extract/0-opt.spvasm.expected.ir.msl
deleted file mode 100644
index 00c893c..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-extract/0-opt.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,51 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(4) {
- three:f32 @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_5:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %4:ptr<uniform, f32, read> = access %x_5, 0u
- %5:f32 = load %4
- %6:vec4<f32> = construct 1.0f, 2.0f, %5, 4.0f
- %7:f32 = dot %6, vec4<f32>(0.0f, 1.0f, 0.0f, 0.0f)
- %8:bool = eq %7, 2.0f
- if %8 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- store %x_GLF_color, vec4<f32>(1.0f, 0.0f, 0.0f, 1.0f)
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %10:void = call %main_1
- %11:vec4<f32> = load %x_GLF_color
- %12:main_out = construct %11
- ret %12
- }
-}
-
-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. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-extract/0-opt.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-extract/0-opt.wgsl.expected.ir.msl
deleted file mode 100644
index ea147f4..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-extract/0-opt.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,52 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(4) {
- three:f32 @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_5:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %4:ptr<uniform, f32, read> = access %x_5, 0u
- %5:f32 = load %4
- %x_25:f32 = let %5
- %7:vec4<f32> = construct 1.0f, 2.0f, %x_25, 4.0f
- %8:f32 = dot %7, vec4<f32>(0.0f, 1.0f, 0.0f, 0.0f)
- %9:bool = eq %8, 2.0f
- if %9 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- store %x_GLF_color, vec4<f32>(1.0f, 0.0f, 0.0f, 1.0f)
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %11:void = call %main_1
- %12:vec4<f32> = load %x_GLF_color
- %13:main_out = construct %12
- ret %13
- }
-}
-
-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. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-no-extract/0.spvasm.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-no-extract/0.spvasm.expected.ir.msl
deleted file mode 100644
index 58496e6..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-no-extract/0.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,51 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(4) {
- three:f32 @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_5:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %4:ptr<uniform, f32, read> = access %x_5, 0u
- %5:f32 = load %4
- %6:vec2<f32> = construct 2.0f, %5
- %7:f32 = dot %6, vec2<f32>(0.0f, 2.0f)
- %8:bool = eq %7, 6.0f
- if %8 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- store %x_GLF_color, vec4<f32>(1.0f, 0.0f, 0.0f, 1.0f)
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %10:void = call %main_1
- %11:vec4<f32> = load %x_GLF_color
- %12:main_out = construct %11
- ret %12
- }
-}
-
-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. *
-********************************************************************
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-no-extract/0.wgsl.expected.ir.msl b/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-no-extract/0.wgsl.expected.ir.msl
deleted file mode 100644
index 5eeba9e..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-folding-rules-dot-no-extract/0.wgsl.expected.ir.msl
+++ /dev/null
@@ -1,52 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: buf0 = struct @align(4) {
- three:f32 @offset(0)
-}
-
-main_out = struct @align(16) {
- x_GLF_color_1:vec4<f32> @offset(0), @location(0)
-}
-
-$B1: { # root
- %x_5:ptr<uniform, buf0, read> = var @binding_point(0, 0)
- %x_GLF_color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%main_1 = func():void {
- $B2: {
- %4:ptr<uniform, f32, read> = access %x_5, 0u
- %5:f32 = load %4
- %x_26:f32 = let %5
- %7:vec2<f32> = construct 2.0f, %x_26
- %8:f32 = dot %7, vec2<f32>(0.0f, 2.0f)
- %9:bool = eq %8, 6.0f
- if %9 [t: $B3, f: $B4] { # if_1
- $B3: { # true
- store %x_GLF_color, vec4<f32>(1.0f, 0.0f, 0.0f, 1.0f)
- exit_if # if_1
- }
- $B4: { # false
- store %x_GLF_color, vec4<f32>(0.0f)
- exit_if # if_1
- }
- }
- ret
- }
-}
-%tint_symbol = @fragment func():main_out {
- $B5: {
- %11:void = call %main_1
- %12:vec4<f32> = load %x_GLF_color
- %13:main_out = construct %12
- ret %13
- }
-}
-
-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. *
-********************************************************************
