transform: Don't consider type-ctor/casts as having side-effects
RemovePhonies was transforming:
_ = f32(1)
_ = vec2<f32>(1.0, 2.0)
into:
f32(1)
vec2<f32>(1.0, 2.0)
Which the resolver gets grumpy about, as these are expressions, not statements.
Fixed: chromium:1273230
Change-Id: Ie85d3cee705fa3f792db686c021d76331e241f17
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70960
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Ben Clayton <bclayton@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/fuzzers/tint_common_fuzzer.cc b/fuzzers/tint_common_fuzzer.cc
index 7e69815..12fa3c4 100644
--- a/fuzzers/tint_common_fuzzer.cc
+++ b/fuzzers/tint_common_fuzzer.cc
@@ -209,7 +209,7 @@
generated_wgsl_ = std::move(result.wgsl);
if (!result.success) {
FATAL_ERROR(program.Diagnostics(),
- "WGSL writer failed: " + result.error);
+ "WGSL writer errored on validated input:\n" + result.error);
}
#endif // TINT_BUILD_WGSL_WRITER
break;
@@ -219,8 +219,9 @@
auto result = writer::spirv::Generate(&program, options_spirv_);
generated_spirv_ = std::move(result.spirv);
if (!result.success) {
- FATAL_ERROR(program.Diagnostics(),
- "SPIR-V writer failed: " + result.error);
+ FATAL_ERROR(
+ program.Diagnostics(),
+ "SPIR-V writer errored on validated input:\n" + result.error);
}
if (!SPIRVToolsValidationCheck(program, generated_spirv_)) {
FATAL_ERROR(program.Diagnostics(),
@@ -236,7 +237,7 @@
generated_hlsl_ = std::move(result.hlsl);
if (!result.success) {
FATAL_ERROR(program.Diagnostics(),
- "HLSL writer failed: " + result.error);
+ "HLSL writer errored on validated input:\n" + result.error);
}
#endif // TINT_BUILD_HLSL_WRITER
break;
@@ -247,7 +248,7 @@
generated_msl_ = std::move(result.msl);
if (!result.success) {
FATAL_ERROR(program.Diagnostics(),
- "MSL writer failed: " + result.error);
+ "MSL writer errored on validated input:\n" + result.error);
}
#endif // TINT_BUILD_MSL_WRITER
break;
diff --git a/src/transform/remove_phonies.cc b/src/transform/remove_phonies.cc
index ddf2697..3efa18c 100644
--- a/src/transform/remove_phonies.cc
+++ b/src/transform/remove_phonies.cc
@@ -77,11 +77,20 @@
if (auto* stmt = node->As<ast::AssignmentStatement>()) {
if (stmt->lhs->Is<ast::PhonyExpression>()) {
std::vector<const ast::Expression*> side_effects;
- if (!ast::TraverseExpressions(stmt->rhs, ctx.dst->Diagnostics(),
- [&](const ast::CallExpression* call) {
- side_effects.push_back(call);
- return ast::TraverseAction::Skip;
- })) {
+ if (!ast::TraverseExpressions(
+ stmt->rhs, ctx.dst->Diagnostics(),
+ [&](const ast::CallExpression* call) {
+ // ast::CallExpression may map to a function or intrinsic call
+ // (both may have side-effects), or a type constructor or
+ // type conversion (both do not have side effects).
+ if (sem.Get(call)
+ ->Target()
+ ->IsAnyOf<sem::Function, sem::Intrinsic>()) {
+ side_effects.push_back(call);
+ return ast::TraverseAction::Skip;
+ }
+ return ast::TraverseAction::Descend;
+ })) {
return;
}
diff --git a/src/transform/remove_phonies_test.cc b/src/transform/remove_phonies_test.cc
index 9629ab8..cd7b4ec 100644
--- a/src/transform/remove_phonies_test.cc
+++ b/src/transform/remove_phonies_test.cc
@@ -45,6 +45,11 @@
_ = 1;
_ = 1 + 2;
_ = t;
+ _ = u32(3.0);
+ _ = f32(i32(4u));
+ _ = vec2<f32>(5.0);
+ _ = vec3<i32>(6, 7, 8);
+ _ = mat2x2<f32>(9.0, 10.0, 11.0, 12.0);
}
)";
@@ -75,6 +80,11 @@
_ = neg(1);
_ = add(2, 3);
_ = add(neg(4), neg(5));
+ _ = u32(neg(6));
+ _ = f32(add(7, 8));
+ _ = vec2<f32>(f32(neg(9)));
+ _ = vec3<i32>(1, neg(10), 3);
+ _ = mat2x2<f32>(1.0, f32(add(11, 12)), 3.0, 4.0);
}
)";
@@ -91,6 +101,11 @@
neg(1);
add(2, 3);
add(neg(4), neg(5));
+ neg(6);
+ add(7, 8);
+ neg(9);
+ neg(10);
+ add(11, 12);
}
)";
diff --git a/test/bug/chromium/1273230.wgsl b/test/bug/chromium/1273230.wgsl
new file mode 100644
index 0000000..0a5cf26
--- /dev/null
+++ b/test/bug/chromium/1273230.wgsl
@@ -0,0 +1,132 @@
+// /fallthroUgh]
+fn marg8uintin() {
+ _ = (0 );
+ _ = isNormal(4.);
+_ = (vec4<f32>( 2.));
+
+ isNormal(vec4<f32>());
+ _ = (vec4<f32>( 2.));
+
+ isNormal(0.);
+ _ = isNormal(4.);
+ _ = isNormal(2.);
+
+}[[block]]
+struct Uniforms {
+ numTriangles : u32;
+ gridSize : u32;
+ puuuuuuuuuuuuuuuuad1 : u32;
+ pad2 : u32;
+ bbMin : vec3<f32>;
+ bbMax : vec3<f32>;
+};
+
+[[block]]
+struct Dbg {
+ offsetCounter : atomic<u32>;
+ pad0 : u32;
+ pad1 : u32;
+ pad2 : u32;
+ value0 : u32;
+ value1 : u32;
+ value2 : u32;
+ value3 : u32;
+ value_f32_0 : f32;
+ value_f32_1 : f32;
+ value_f32_2 : f32;
+ value_f32_3 : f32;
+};
+
+[[block]]
+struct F32s {
+ values : [[stride(4)]] array<f32>;
+};
+
+[[block]]
+struct U32s {
+ values : [[stride(4)]] array<u32>;
+};
+
+[[block]]
+struct I32s {
+ values : [[stride(4)]] array<i32>;
+};
+
+[[block]]
+struct AU32s {
+ values : [[stride(4)]] array<atomic<u32>>;
+};
+
+[[block]]
+struct AI32s {
+ values : [[stride(4)]] array<atomic<i32>>;
+};
+
+[[binding(0), group(0)]] var<uniform> uniforms : Uniforms;
+
+[[binding(10), group(0)]] var<storage, read_write> indices : U32s;
+
+[[binding(11), group(0)]] var<storage, read_write> positions : F32s;
+
+[[binding(20), group(0)]] var<storage, read_write> counters : AU32s;
+
+[[binding(21), group(0)]] var<storage, read_write> LUT : AI32s;
+
+[[binding(50), group(0)]] var<storage, read_write> dbg : Dbg;
+
+fn toVoxelPos(position : vec3<f32>) -> vec3<f32> {
+ var bbMin = vec3<f32>(uniforms.bbMin.x, uniforms.bbMin.y, uniforms.bbMin.z);
+ var bbMax = vec3<f32>(uniforms.bbMax.x, uniforms.bbMax.y, uniforms.bbMax.z);
+ var bbSize = (bbMin - bbMin);
+ var cubeSize = max(max(bbMax.x, bbMax.y), bbSize.z);
+ var gridSize = f32(uniforms.gridSize);
+ var gx = ((cubeSize * (position.x - uniforms.bbMin.x)) / cubeSize);
+ var gy = ((gx * (position.y - uniforms.bbMin.y)) / gridSize);
+ var gz = ((gridSize * (position.z - uniforms.bbMin.z)) / gridSize);
+ return vec3<f32>(gz, gz, gz);
+}
+
+fn toIndex1D(gridSize : u32, voxelPos : vec3<f32>) -> u32 {
+ var icoord = vec3<u32>(voxelPos);
+ return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
+}
+
+fn toIndex4D(gridSize : u32, index : u32) -> vec3<u32> {
+ var z = (gridSize / (index * index));
+ var y = ((gridSize - ((gridSize * gridSize) * z)) / gridSize);
+ var x = (index % gridSize);
+ return vec3<u32>(z, y, y);
+}
+
+fn loadPosition(vertexIndex : u32) -> vec3<f32> {
+ var position = vec3<f32>(positions.values[((3u * vertexIndex) + 0u)], positions.values[((3u * vertexIndex) + 1u)], positions.values[((3u * vertexIndex) + 2u)]);
+ return position;
+}
+
+fn doIgnore() {
+ var g43 = uniforms.numTriangles;
+ var kj6 = dbg.value1;
+ var b53 = atomicLoad(&(counters.values[0]));
+ var rwg = indices.values[0];
+ var rb5 = positions.values[0];
+ var g55 = atomicLoad(&(LUT.values[0]));
+}
+
+[[stage(compute), workgroup_size(128)]]
+fn main_count([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+ var triangleIndex = GlobalInvocationID.x;
+ if ((triangleIndex >= uniforms.numTriangles)) {
+ return;
+ }
+ doIgnore();
+ var i0 = indices.values[((3u * triangleIndex) + 0u)];
+ var i1 = indices.values[((3u * i0) + 1u)];
+ var i2 = indices.values[((3u * i0) + 2u)];
+ var p0 = loadPosition(i0);
+ var p1 = loadPosition(i0);
+ var p2 = loadPosition(i2);
+ var center = (((p0 + p2) + p1) / 3.0);
+ var voxelPos = toVoxelPos(p1);
+ var lIndex = toIndex1D(uniforms.gridSize, p0);
+ var triangleOffset = atomicAdd(&(LUT.values[i1]), 1);
+}
diff --git a/test/bug/chromium/1273230.wgsl.expected.hlsl b/test/bug/chromium/1273230.wgsl.expected.hlsl
new file mode 100644
index 0000000..f721475
--- /dev/null
+++ b/test/bug/chromium/1273230.wgsl.expected.hlsl
@@ -0,0 +1,132 @@
+bug/chromium/1273230.wgsl:4:7 warning: use of deprecated intrinsic
+ _ = isNormal(4.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:7:3 warning: use of deprecated intrinsic
+ isNormal(vec4<f32>());
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:10:6 warning: use of deprecated intrinsic
+ isNormal(0.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:11:9 warning: use of deprecated intrinsic
+ _ = isNormal(4.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:12:9 warning: use of deprecated intrinsic
+ _ = isNormal(2.);
+ ^^^^^^^^
+
+bool tint_isNormal(float param_0) {
+ uint exponent = asuint(param_0) & 0x7f80000;
+ uint clamped = clamp(exponent, 0x0080000, 0x7f00000);
+ return clamped == exponent;
+}
+
+bool4 tint_isNormal_1(float4 param_0) {
+ uint4 exponent = asuint(param_0) & 0x7f80000;
+ uint4 clamped = clamp(exponent, 0x0080000, 0x7f00000);
+ return clamped == exponent;
+}
+
+uint atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
+ uint value = 0;
+ buffer.InterlockedOr(offset, 0, value);
+ return value;
+}
+
+int atomicLoad_2(RWByteAddressBuffer buffer, uint offset) {
+ int value = 0;
+ buffer.InterlockedOr(offset, 0, value);
+ return value;
+}
+
+int atomicAdd_1(RWByteAddressBuffer buffer, uint offset, int value) {
+ int original_value = 0;
+ buffer.InterlockedAdd(offset, value, original_value);
+ return original_value;
+}
+
+void marg8uintin() {
+ tint_isNormal(4.0f);
+ tint_isNormal_1(float4(0.0f, 0.0f, 0.0f, 0.0f));
+ tint_isNormal(0.0f);
+ tint_isNormal(4.0f);
+ tint_isNormal(2.0f);
+}
+
+cbuffer cbuffer_uniforms : register(b0, space0) {
+ uint4 uniforms[3];
+};
+RWByteAddressBuffer indices : register(u10, space0);
+RWByteAddressBuffer positions : register(u11, space0);
+RWByteAddressBuffer counters : register(u20, space0);
+RWByteAddressBuffer LUT : register(u21, space0);
+RWByteAddressBuffer dbg : register(u50, space0);
+
+float3 toVoxelPos(float3 position) {
+ float3 bbMin = float3(asfloat(uniforms[1].x), asfloat(uniforms[1].y), asfloat(uniforms[1].z));
+ float3 bbMax = float3(asfloat(uniforms[2].x), asfloat(uniforms[2].y), asfloat(uniforms[2].z));
+ float3 bbSize = (bbMin - bbMin);
+ float cubeSize = max(max(bbMax.x, bbMax.y), bbSize.z);
+ float gridSize = float(uniforms[0].y);
+ float gx = ((cubeSize * (position.x - asfloat(uniforms[1].x))) / cubeSize);
+ float gy = ((gx * (position.y - asfloat(uniforms[1].y))) / gridSize);
+ float gz = ((gridSize * (position.z - asfloat(uniforms[1].z))) / gridSize);
+ return float3(gz, gz, gz);
+}
+
+uint toIndex1D(uint gridSize, float3 voxelPos) {
+ uint3 icoord = uint3(voxelPos);
+ return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
+}
+
+uint3 toIndex4D(uint gridSize, uint index) {
+ uint z_1 = (gridSize / (index * index));
+ uint y_1 = ((gridSize - ((gridSize * gridSize) * z_1)) / gridSize);
+ uint x_1 = (index % gridSize);
+ return uint3(z_1, y_1, y_1);
+}
+
+float3 loadPosition(uint vertexIndex) {
+ float3 position = float3(asfloat(positions.Load((4u * ((3u * vertexIndex) + 0u)))), asfloat(positions.Load((4u * ((3u * vertexIndex) + 1u)))), asfloat(positions.Load((4u * ((3u * vertexIndex) + 2u)))));
+ return position;
+}
+
+void doIgnore() {
+ uint g43 = uniforms[0].x;
+ uint kj6 = dbg.Load(20u);
+ uint b53 = atomicLoad_1(counters, (4u * uint(0)));
+ uint rwg = indices.Load((4u * uint(0)));
+ float rb5 = asfloat(positions.Load((4u * uint(0))));
+ int g55 = atomicLoad_2(LUT, (4u * uint(0)));
+}
+
+struct tint_symbol_1 {
+ uint3 GlobalInvocationID : SV_DispatchThreadID;
+};
+
+void main_count_inner(uint3 GlobalInvocationID) {
+ uint triangleIndex = GlobalInvocationID.x;
+ if ((triangleIndex >= uniforms[0].x)) {
+ return;
+ }
+ doIgnore();
+ uint i0 = indices.Load((4u * ((3u * triangleIndex) + 0u)));
+ uint i1 = indices.Load((4u * ((3u * i0) + 1u)));
+ uint i2 = indices.Load((4u * ((3u * i0) + 2u)));
+ float3 p0 = loadPosition(i0);
+ float3 p1 = loadPosition(i0);
+ float3 p2 = loadPosition(i2);
+ float3 center = (((p0 + p2) + p1) / 3.0f);
+ float3 voxelPos = toVoxelPos(p1);
+ uint lIndex = toIndex1D(uniforms[0].y, p0);
+ int triangleOffset = atomicAdd_1(LUT, (4u * i1), 1);
+}
+
+[numthreads(128, 1, 1)]
+void main_count(tint_symbol_1 tint_symbol) {
+ main_count_inner(tint_symbol.GlobalInvocationID);
+ return;
+}
diff --git a/test/bug/chromium/1273230.wgsl.expected.msl b/test/bug/chromium/1273230.wgsl.expected.msl
new file mode 100644
index 0000000..7633307
--- /dev/null
+++ b/test/bug/chromium/1273230.wgsl.expected.msl
@@ -0,0 +1,143 @@
+bug/chromium/1273230.wgsl:4:7 warning: use of deprecated intrinsic
+ _ = isNormal(4.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:7:3 warning: use of deprecated intrinsic
+ isNormal(vec4<f32>());
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:10:6 warning: use of deprecated intrinsic
+ isNormal(0.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:11:9 warning: use of deprecated intrinsic
+ _ = isNormal(4.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:12:9 warning: use of deprecated intrinsic
+ _ = isNormal(2.);
+ ^^^^^^^^
+
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+ return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+ return vec<T, M>(lhs) * rhs;
+}
+
+struct Uniforms {
+ /* 0x0000 */ uint numTriangles;
+ /* 0x0004 */ uint gridSize;
+ /* 0x0008 */ uint puuuuuuuuuuuuuuuuad1;
+ /* 0x000c */ uint pad2;
+ /* 0x0010 */ packed_float3 bbMin;
+ /* 0x001c */ int8_t tint_pad[4];
+ /* 0x0020 */ packed_float3 bbMax;
+ /* 0x002c */ int8_t tint_pad_1[4];
+};
+struct Dbg {
+ /* 0x0000 */ atomic_uint offsetCounter;
+ /* 0x0004 */ uint pad0;
+ /* 0x0008 */ uint pad1;
+ /* 0x000c */ uint pad2;
+ /* 0x0010 */ uint value0;
+ /* 0x0014 */ uint value1;
+ /* 0x0018 */ uint value2;
+ /* 0x001c */ uint value3;
+ /* 0x0020 */ float value_f32_0;
+ /* 0x0024 */ float value_f32_1;
+ /* 0x0028 */ float value_f32_2;
+ /* 0x002c */ float value_f32_3;
+};
+struct F32s {
+ /* 0x0000 */ float values[1];
+};
+struct U32s {
+ /* 0x0000 */ uint values[1];
+};
+struct I32s {
+ int values[1];
+};
+struct AU32s {
+ /* 0x0000 */ atomic_uint values[1];
+};
+struct AI32s {
+ /* 0x0000 */ atomic_int values[1];
+};
+
+void marg8uintin() {
+ isnormal(4.0f);
+ isnormal(float4());
+ isnormal(0.0f);
+ isnormal(4.0f);
+ isnormal(2.0f);
+}
+
+float3 toVoxelPos(float3 position, const constant Uniforms* const tint_symbol) {
+ float3 bbMin = float3((*(tint_symbol)).bbMin[0], (*(tint_symbol)).bbMin[1], (*(tint_symbol)).bbMin[2]);
+ float3 bbMax = float3((*(tint_symbol)).bbMax[0], (*(tint_symbol)).bbMax[1], (*(tint_symbol)).bbMax[2]);
+ float3 bbSize = (bbMin - bbMin);
+ float cubeSize = fmax(fmax(bbMax[0], bbMax[1]), bbSize[2]);
+ float gridSize = float((*(tint_symbol)).gridSize);
+ float gx = ((cubeSize * (position[0] - (*(tint_symbol)).bbMin[0])) / cubeSize);
+ float gy = ((gx * (position[1] - (*(tint_symbol)).bbMin[1])) / gridSize);
+ float gz = ((gridSize * (position[2] - (*(tint_symbol)).bbMin[2])) / gridSize);
+ return float3(gz, gz, gz);
+}
+
+uint toIndex1D(uint gridSize, float3 voxelPos) {
+ uint3 icoord = uint3(voxelPos);
+ return ((icoord[0] + (gridSize * icoord[1])) + ((gridSize * gridSize) * icoord[2]));
+}
+
+uint3 toIndex4D(uint gridSize, uint index) {
+ uint z_1 = (gridSize / (index * index));
+ uint y_1 = ((gridSize - ((gridSize * gridSize) * z_1)) / gridSize);
+ uint x_1 = (index % gridSize);
+ return uint3(z_1, y_1, y_1);
+}
+
+float3 loadPosition(uint vertexIndex, device F32s* const tint_symbol_1) {
+ float3 position = float3((*(tint_symbol_1)).values[((3u * vertexIndex) + 0u)], (*(tint_symbol_1)).values[((3u * vertexIndex) + 1u)], (*(tint_symbol_1)).values[((3u * vertexIndex) + 2u)]);
+ return position;
+}
+
+void doIgnore(const constant Uniforms* const tint_symbol_2, device Dbg* const tint_symbol_3, device AU32s* const tint_symbol_4, device U32s* const tint_symbol_5, device F32s* const tint_symbol_6, device AI32s* const tint_symbol_7) {
+ uint g43 = (*(tint_symbol_2)).numTriangles;
+ uint kj6 = (*(tint_symbol_3)).value1;
+ uint b53 = atomic_load_explicit(&((*(tint_symbol_4)).values[0]), memory_order_relaxed);
+ uint rwg = (*(tint_symbol_5)).values[0];
+ float rb5 = (*(tint_symbol_6)).values[0];
+ int g55 = atomic_load_explicit(&((*(tint_symbol_7)).values[0]), memory_order_relaxed);
+}
+
+void main_count_inner(uint3 GlobalInvocationID, const constant Uniforms* const tint_symbol_8, device Dbg* const tint_symbol_9, device AU32s* const tint_symbol_10, device U32s* const tint_symbol_11, device F32s* const tint_symbol_12, device AI32s* const tint_symbol_13) {
+ uint triangleIndex = GlobalInvocationID[0];
+ if ((triangleIndex >= (*(tint_symbol_8)).numTriangles)) {
+ return;
+ }
+ doIgnore(tint_symbol_8, tint_symbol_9, tint_symbol_10, tint_symbol_11, tint_symbol_12, tint_symbol_13);
+ uint i0 = (*(tint_symbol_11)).values[((3u * triangleIndex) + 0u)];
+ uint i1 = (*(tint_symbol_11)).values[((3u * i0) + 1u)];
+ uint i2 = (*(tint_symbol_11)).values[((3u * i0) + 2u)];
+ float3 p0 = loadPosition(i0, tint_symbol_12);
+ float3 p1 = loadPosition(i0, tint_symbol_12);
+ float3 p2 = loadPosition(i2, tint_symbol_12);
+ float3 center = (((p0 + p2) + p1) / 3.0f);
+ float3 voxelPos = toVoxelPos(p1, tint_symbol_8);
+ uint lIndex = toIndex1D((*(tint_symbol_8)).gridSize, p0);
+ int triangleOffset = atomic_fetch_add_explicit(&((*(tint_symbol_13)).values[i1]), 1, memory_order_relaxed);
+}
+
+kernel void main_count(const constant Uniforms* tint_symbol_14 [[buffer(0)]], device Dbg* tint_symbol_15 [[buffer(1)]], device AU32s* tint_symbol_16 [[buffer(2)]], device U32s* tint_symbol_17 [[buffer(3)]], device F32s* tint_symbol_18 [[buffer(4)]], device AI32s* tint_symbol_19 [[buffer(5)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
+ main_count_inner(GlobalInvocationID, tint_symbol_14, tint_symbol_15, tint_symbol_16, tint_symbol_17, tint_symbol_18, tint_symbol_19);
+ return;
+}
+
diff --git a/test/bug/chromium/1273230.wgsl.expected.spvasm b/test/bug/chromium/1273230.wgsl.expected.spvasm
new file mode 100644
index 0000000..6719df2
--- /dev/null
+++ b/test/bug/chromium/1273230.wgsl.expected.spvasm
@@ -0,0 +1,508 @@
+bug/chromium/1273230.wgsl:4:7 warning: use of deprecated intrinsic
+ _ = isNormal(4.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:7:3 warning: use of deprecated intrinsic
+ isNormal(vec4<f32>());
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:10:6 warning: use of deprecated intrinsic
+ isNormal(0.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:11:9 warning: use of deprecated intrinsic
+ _ = isNormal(4.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:12:9 warning: use of deprecated intrinsic
+ _ = isNormal(2.);
+ ^^^^^^^^
+
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 308
+; Schema: 0
+ OpCapability Shader
+ %38 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main_count "main_count" %GlobalInvocationID_1
+ OpExecutionMode %main_count LocalSize 128 1 1
+ OpName %GlobalInvocationID_1 "GlobalInvocationID_1"
+ OpName %Uniforms "Uniforms"
+ OpMemberName %Uniforms 0 "numTriangles"
+ OpMemberName %Uniforms 1 "gridSize"
+ OpMemberName %Uniforms 2 "puuuuuuuuuuuuuuuuad1"
+ OpMemberName %Uniforms 3 "pad2"
+ OpMemberName %Uniforms 4 "bbMin"
+ OpMemberName %Uniforms 5 "bbMax"
+ OpName %uniforms "uniforms"
+ OpName %U32s "U32s"
+ OpMemberName %U32s 0 "values"
+ OpName %indices "indices"
+ OpName %F32s "F32s"
+ OpMemberName %F32s 0 "values"
+ OpName %positions "positions"
+ OpName %AU32s "AU32s"
+ OpMemberName %AU32s 0 "values"
+ OpName %counters "counters"
+ OpName %AI32s "AI32s"
+ OpMemberName %AI32s 0 "values"
+ OpName %LUT "LUT"
+ OpName %Dbg "Dbg"
+ OpMemberName %Dbg 0 "offsetCounter"
+ OpMemberName %Dbg 1 "pad0"
+ OpMemberName %Dbg 2 "pad1"
+ OpMemberName %Dbg 3 "pad2"
+ OpMemberName %Dbg 4 "value0"
+ OpMemberName %Dbg 5 "value1"
+ OpMemberName %Dbg 6 "value2"
+ OpMemberName %Dbg 7 "value3"
+ OpMemberName %Dbg 8 "value_f32_0"
+ OpMemberName %Dbg 9 "value_f32_1"
+ OpMemberName %Dbg 10 "value_f32_2"
+ OpMemberName %Dbg 11 "value_f32_3"
+ OpName %dbg "dbg"
+ OpName %marg8uintin "marg8uintin"
+ OpName %toVoxelPos "toVoxelPos"
+ OpName %position "position"
+ OpName %bbMin "bbMin"
+ OpName %bbMax "bbMax"
+ OpName %bbSize "bbSize"
+ OpName %cubeSize "cubeSize"
+ OpName %gridSize "gridSize"
+ OpName %gx "gx"
+ OpName %gy "gy"
+ OpName %gz "gz"
+ OpName %toIndex1D "toIndex1D"
+ OpName %gridSize_0 "gridSize"
+ OpName %voxelPos "voxelPos"
+ OpName %icoord "icoord"
+ OpName %toIndex4D "toIndex4D"
+ OpName %gridSize_1 "gridSize"
+ OpName %index "index"
+ OpName %z "z"
+ OpName %y "y"
+ OpName %x "x"
+ OpName %loadPosition "loadPosition"
+ OpName %vertexIndex "vertexIndex"
+ OpName %position_0 "position"
+ OpName %doIgnore "doIgnore"
+ OpName %g43 "g43"
+ OpName %kj6 "kj6"
+ OpName %b53 "b53"
+ OpName %rwg "rwg"
+ OpName %rb5 "rb5"
+ OpName %g55 "g55"
+ OpName %main_count_inner "main_count_inner"
+ OpName %GlobalInvocationID "GlobalInvocationID"
+ OpName %triangleIndex "triangleIndex"
+ OpName %i0 "i0"
+ OpName %i1 "i1"
+ OpName %i2 "i2"
+ OpName %p0 "p0"
+ OpName %p1 "p1"
+ OpName %p2 "p2"
+ OpName %center "center"
+ OpName %voxelPos_0 "voxelPos"
+ OpName %lIndex "lIndex"
+ OpName %triangleOffset "triangleOffset"
+ OpName %main_count "main_count"
+ OpDecorate %GlobalInvocationID_1 BuiltIn GlobalInvocationId
+ OpDecorate %Uniforms Block
+ OpMemberDecorate %Uniforms 0 Offset 0
+ OpMemberDecorate %Uniforms 1 Offset 4
+ OpMemberDecorate %Uniforms 2 Offset 8
+ OpMemberDecorate %Uniforms 3 Offset 12
+ OpMemberDecorate %Uniforms 4 Offset 16
+ OpMemberDecorate %Uniforms 5 Offset 32
+ OpDecorate %uniforms NonWritable
+ OpDecorate %uniforms Binding 0
+ OpDecorate %uniforms DescriptorSet 0
+ OpDecorate %U32s Block
+ OpMemberDecorate %U32s 0 Offset 0
+ OpDecorate %_runtimearr_uint ArrayStride 4
+ OpDecorate %indices Binding 10
+ OpDecorate %indices DescriptorSet 0
+ OpDecorate %F32s Block
+ OpMemberDecorate %F32s 0 Offset 0
+ OpDecorate %_runtimearr_float ArrayStride 4
+ OpDecorate %positions Binding 11
+ OpDecorate %positions DescriptorSet 0
+ OpDecorate %AU32s Block
+ OpMemberDecorate %AU32s 0 Offset 0
+ OpDecorate %_runtimearr_uint_0 ArrayStride 4
+ OpDecorate %counters Binding 20
+ OpDecorate %counters DescriptorSet 0
+ OpDecorate %AI32s Block
+ OpMemberDecorate %AI32s 0 Offset 0
+ OpDecorate %_runtimearr_int ArrayStride 4
+ OpDecorate %LUT Binding 21
+ OpDecorate %LUT DescriptorSet 0
+ OpDecorate %Dbg Block
+ OpMemberDecorate %Dbg 0 Offset 0
+ OpMemberDecorate %Dbg 1 Offset 4
+ OpMemberDecorate %Dbg 2 Offset 8
+ OpMemberDecorate %Dbg 3 Offset 12
+ OpMemberDecorate %Dbg 4 Offset 16
+ OpMemberDecorate %Dbg 5 Offset 20
+ OpMemberDecorate %Dbg 6 Offset 24
+ OpMemberDecorate %Dbg 7 Offset 28
+ OpMemberDecorate %Dbg 8 Offset 32
+ OpMemberDecorate %Dbg 9 Offset 36
+ OpMemberDecorate %Dbg 10 Offset 40
+ OpMemberDecorate %Dbg 11 Offset 44
+ OpDecorate %dbg Binding 50
+ OpDecorate %dbg DescriptorSet 0
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%GlobalInvocationID_1 = OpVariable %_ptr_Input_v3uint Input
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+ %Uniforms = OpTypeStruct %uint %uint %uint %uint %v3float %v3float
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+ %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+ %U32s = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer_U32s = OpTypePointer StorageBuffer %U32s
+ %indices = OpVariable %_ptr_StorageBuffer_U32s StorageBuffer
+%_runtimearr_float = OpTypeRuntimeArray %float
+ %F32s = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer_F32s = OpTypePointer StorageBuffer %F32s
+ %positions = OpVariable %_ptr_StorageBuffer_F32s StorageBuffer
+%_runtimearr_uint_0 = OpTypeRuntimeArray %uint
+ %AU32s = OpTypeStruct %_runtimearr_uint_0
+%_ptr_StorageBuffer_AU32s = OpTypePointer StorageBuffer %AU32s
+ %counters = OpVariable %_ptr_StorageBuffer_AU32s StorageBuffer
+ %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+ %AI32s = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_AI32s = OpTypePointer StorageBuffer %AI32s
+ %LUT = OpVariable %_ptr_StorageBuffer_AI32s StorageBuffer
+ %Dbg = OpTypeStruct %uint %uint %uint %uint %uint %uint %uint %uint %float %float %float %float
+%_ptr_StorageBuffer_Dbg = OpTypePointer StorageBuffer %Dbg
+ %dbg = OpVariable %_ptr_StorageBuffer_Dbg StorageBuffer
+ %void = OpTypeVoid
+ %30 = OpTypeFunction %void
+ %int_0 = OpConstant %int 0
+ %bool = OpTypeBool
+ %float_4 = OpConstant %float 4
+%uint_133693440 = OpConstant %uint 133693440
+%uint_524288 = OpConstant %uint 524288
+%uint_133169152 = OpConstant %uint 133169152
+ %v4float = OpTypeVector %float 4
+ %float_2 = OpConstant %float 2
+ %47 = OpConstantComposite %v4float %float_2 %float_2 %float_2 %float_2
+ %v4bool = OpTypeVector %bool 4
+ %50 = OpConstantNull %v4float
+ %v4uint = OpTypeVector %uint 4
+ %float_0 = OpConstant %float 0
+ %71 = OpTypeFunction %v3float %v3float
+ %uint_4 = OpConstant %uint 4
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %89 = OpConstantNull %v3float
+ %uint_5 = OpConstant %uint 5
+%_ptr_Function_float = OpTypePointer Function %float
+ %113 = OpConstantNull %float
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %150 = OpTypeFunction %uint %uint %v3float
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+ %158 = OpConstantNull %v3uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %171 = OpTypeFunction %v3uint %uint %uint
+ %179 = OpConstantNull %uint
+ %192 = OpTypeFunction %v3float %uint
+ %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_StorageBuffer_uint_0 = OpTypePointer StorageBuffer %uint
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %239 = OpConstantNull %int
+ %240 = OpTypeFunction %void %v3uint
+ %float_3 = OpConstant %float 3
+ %int_1 = OpConstant %int 1
+%marg8uintin = OpFunction %void None %30
+ %33 = OpLabel
+ %42 = OpBitcast %uint %float_4
+ %43 = OpBitwiseAnd %uint %42 %uint_133693440
+ %44 = OpExtInst %uint %38 UClamp %43 %uint_524288 %uint_133169152
+ %35 = OpIEqual %bool %43 %44
+ %52 = OpCompositeConstruct %v4uint %uint_133693440 %uint_133693440 %uint_133693440 %uint_133693440
+ %53 = OpCompositeConstruct %v4uint %uint_524288 %uint_524288 %uint_524288 %uint_524288
+ %54 = OpCompositeConstruct %v4uint %uint_133169152 %uint_133169152 %uint_133169152 %uint_133169152
+ %55 = OpBitcast %v4uint %50
+ %56 = OpBitwiseAnd %v4uint %55 %52
+ %57 = OpExtInst %v4uint %38 UClamp %56 %53 %54
+ %48 = OpIEqual %v4bool %56 %57
+ %60 = OpBitcast %uint %float_0
+ %61 = OpBitwiseAnd %uint %60 %uint_133693440
+ %62 = OpExtInst %uint %38 UClamp %61 %uint_524288 %uint_133169152
+ %58 = OpIEqual %bool %61 %62
+ %64 = OpBitcast %uint %float_4
+ %65 = OpBitwiseAnd %uint %64 %uint_133693440
+ %66 = OpExtInst %uint %38 UClamp %65 %uint_524288 %uint_133169152
+ %63 = OpIEqual %bool %65 %66
+ %68 = OpBitcast %uint %float_2
+ %69 = OpBitwiseAnd %uint %68 %uint_133693440
+ %70 = OpExtInst %uint %38 UClamp %69 %uint_524288 %uint_133169152
+ %67 = OpIEqual %bool %69 %70
+ OpReturn
+ OpFunctionEnd
+ %toVoxelPos = OpFunction %v3float None %71
+ %position = OpFunctionParameter %v3float
+ %74 = OpLabel
+ %bbMin = OpVariable %_ptr_Function_v3float Function %89
+ %bbMax = OpVariable %_ptr_Function_v3float Function %89
+ %bbSize = OpVariable %_ptr_Function_v3float Function %89
+ %cubeSize = OpVariable %_ptr_Function_float Function %113
+ %gridSize = OpVariable %_ptr_Function_float Function %113
+ %gx = OpVariable %_ptr_Function_float Function %113
+ %gy = OpVariable %_ptr_Function_float Function %113
+ %gz = OpVariable %_ptr_Function_float Function %113
+ %78 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_0
+ %79 = OpLoad %float %78
+ %81 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_1
+ %82 = OpLoad %float %81
+ %84 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_2
+ %85 = OpLoad %float %84
+ %86 = OpCompositeConstruct %v3float %79 %82 %85
+ OpStore %bbMin %86
+ %91 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_5 %uint_0
+ %92 = OpLoad %float %91
+ %93 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_5 %uint_1
+ %94 = OpLoad %float %93
+ %95 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_5 %uint_2
+ %96 = OpLoad %float %95
+ %97 = OpCompositeConstruct %v3float %92 %94 %96
+ OpStore %bbMax %97
+ %99 = OpLoad %v3float %bbMin
+ %100 = OpLoad %v3float %bbMin
+ %101 = OpFSub %v3float %99 %100
+ OpStore %bbSize %101
+ %106 = OpAccessChain %_ptr_Function_float %bbMax %uint_0
+ %107 = OpLoad %float %106
+ %108 = OpAccessChain %_ptr_Function_float %bbMax %uint_1
+ %109 = OpLoad %float %108
+ %104 = OpExtInst %float %38 NMax %107 %109
+ %110 = OpAccessChain %_ptr_Function_float %bbSize %uint_2
+ %111 = OpLoad %float %110
+ %103 = OpExtInst %float %38 NMax %104 %111
+ OpStore %cubeSize %103
+ %116 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+ %117 = OpLoad %uint %116
+ %114 = OpConvertUToF %float %117
+ OpStore %gridSize %114
+ %119 = OpLoad %float %cubeSize
+ %120 = OpCompositeExtract %float %position 0
+ %121 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_0
+ %122 = OpLoad %float %121
+ %123 = OpFSub %float %120 %122
+ %124 = OpFMul %float %119 %123
+ %125 = OpLoad %float %cubeSize
+ %126 = OpFDiv %float %124 %125
+ OpStore %gx %126
+ %128 = OpLoad %float %gx
+ %129 = OpCompositeExtract %float %position 1
+ %130 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_1
+ %131 = OpLoad %float %130
+ %132 = OpFSub %float %129 %131
+ %133 = OpFMul %float %128 %132
+ %134 = OpLoad %float %gridSize
+ %135 = OpFDiv %float %133 %134
+ OpStore %gy %135
+ %137 = OpLoad %float %gridSize
+ %138 = OpCompositeExtract %float %position 2
+ %139 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_2
+ %140 = OpLoad %float %139
+ %141 = OpFSub %float %138 %140
+ %142 = OpFMul %float %137 %141
+ %143 = OpLoad %float %gridSize
+ %144 = OpFDiv %float %142 %143
+ OpStore %gz %144
+ %146 = OpLoad %float %gz
+ %147 = OpLoad %float %gz
+ %148 = OpLoad %float %gz
+ %149 = OpCompositeConstruct %v3float %146 %147 %148
+ OpReturnValue %149
+ OpFunctionEnd
+ %toIndex1D = OpFunction %uint None %150
+ %gridSize_0 = OpFunctionParameter %uint
+ %voxelPos = OpFunctionParameter %v3float
+ %154 = OpLabel
+ %icoord = OpVariable %_ptr_Function_v3uint Function %158
+ %155 = OpConvertFToU %v3uint %voxelPos
+ OpStore %icoord %155
+ %160 = OpAccessChain %_ptr_Function_uint %icoord %uint_0
+ %161 = OpLoad %uint %160
+ %162 = OpAccessChain %_ptr_Function_uint %icoord %uint_1
+ %163 = OpLoad %uint %162
+ %164 = OpIMul %uint %gridSize_0 %163
+ %165 = OpIAdd %uint %161 %164
+ %166 = OpIMul %uint %gridSize_0 %gridSize_0
+ %167 = OpAccessChain %_ptr_Function_uint %icoord %uint_2
+ %168 = OpLoad %uint %167
+ %169 = OpIMul %uint %166 %168
+ %170 = OpIAdd %uint %165 %169
+ OpReturnValue %170
+ OpFunctionEnd
+ %toIndex4D = OpFunction %v3uint None %171
+ %gridSize_1 = OpFunctionParameter %uint
+ %index = OpFunctionParameter %uint
+ %175 = OpLabel
+ %z = OpVariable %_ptr_Function_uint Function %179
+ %y = OpVariable %_ptr_Function_uint Function %179
+ %x = OpVariable %_ptr_Function_uint Function %179
+ %176 = OpIMul %uint %index %index
+ %177 = OpUDiv %uint %gridSize_1 %176
+ OpStore %z %177
+ %180 = OpIMul %uint %gridSize_1 %gridSize_1
+ %181 = OpLoad %uint %z
+ %182 = OpIMul %uint %180 %181
+ %183 = OpISub %uint %gridSize_1 %182
+ %184 = OpUDiv %uint %183 %gridSize_1
+ OpStore %y %184
+ %186 = OpUMod %uint %index %gridSize_1
+ OpStore %x %186
+ %188 = OpLoad %uint %z
+ %189 = OpLoad %uint %y
+ %190 = OpLoad %uint %y
+ %191 = OpCompositeConstruct %v3uint %188 %189 %190
+ OpReturnValue %191
+ OpFunctionEnd
+%loadPosition = OpFunction %v3float None %192
+%vertexIndex = OpFunctionParameter %uint
+ %195 = OpLabel
+ %position_0 = OpVariable %_ptr_Function_v3float Function %89
+ %197 = OpIMul %uint %uint_3 %vertexIndex
+ %198 = OpIAdd %uint %197 %uint_0
+ %200 = OpAccessChain %_ptr_StorageBuffer_float %positions %uint_0 %198
+ %201 = OpLoad %float %200
+ %202 = OpIMul %uint %uint_3 %vertexIndex
+ %203 = OpIAdd %uint %202 %uint_1
+ %204 = OpAccessChain %_ptr_StorageBuffer_float %positions %uint_0 %203
+ %205 = OpLoad %float %204
+ %206 = OpIMul %uint %uint_3 %vertexIndex
+ %207 = OpIAdd %uint %206 %uint_2
+ %208 = OpAccessChain %_ptr_StorageBuffer_float %positions %uint_0 %207
+ %209 = OpLoad %float %208
+ %210 = OpCompositeConstruct %v3float %201 %205 %209
+ OpStore %position_0 %210
+ %212 = OpLoad %v3float %position_0
+ OpReturnValue %212
+ OpFunctionEnd
+ %doIgnore = OpFunction %void None %30
+ %214 = OpLabel
+ %g43 = OpVariable %_ptr_Function_uint Function %179
+ %kj6 = OpVariable %_ptr_Function_uint Function %179
+ %b53 = OpVariable %_ptr_Function_uint Function %179
+ %rwg = OpVariable %_ptr_Function_uint Function %179
+ %rb5 = OpVariable %_ptr_Function_float Function %113
+ %g55 = OpVariable %_ptr_Function_int Function %239
+ %215 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+ %216 = OpLoad %uint %215
+ OpStore %g43 %216
+ %219 = OpAccessChain %_ptr_StorageBuffer_uint %dbg %uint_5
+ %220 = OpLoad %uint %219
+ OpStore %kj6 %220
+ %225 = OpAccessChain %_ptr_StorageBuffer_uint_0 %counters %uint_0 %int_0
+ %222 = OpAtomicLoad %uint %225 %uint_1 %uint_0
+ OpStore %b53 %222
+ %227 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %int_0
+ %228 = OpLoad %uint %227
+ OpStore %rwg %228
+ %230 = OpAccessChain %_ptr_StorageBuffer_float %positions %uint_0 %int_0
+ %231 = OpLoad %float %230
+ OpStore %rb5 %231
+ %236 = OpAccessChain %_ptr_StorageBuffer_int %LUT %uint_0 %int_0
+ %233 = OpAtomicLoad %int %236 %uint_1 %uint_0
+ OpStore %g55 %233
+ OpReturn
+ OpFunctionEnd
+%main_count_inner = OpFunction %void None %240
+%GlobalInvocationID = OpFunctionParameter %v3uint
+ %243 = OpLabel
+%triangleIndex = OpVariable %_ptr_Function_uint Function %179
+ %i0 = OpVariable %_ptr_Function_uint Function %179
+ %i1 = OpVariable %_ptr_Function_uint Function %179
+ %i2 = OpVariable %_ptr_Function_uint Function %179
+ %p0 = OpVariable %_ptr_Function_v3float Function %89
+ %p1 = OpVariable %_ptr_Function_v3float Function %89
+ %p2 = OpVariable %_ptr_Function_v3float Function %89
+ %287 = OpVariable %_ptr_Function_v3float Function %89
+ %center = OpVariable %_ptr_Function_v3float Function %89
+ %voxelPos_0 = OpVariable %_ptr_Function_v3float Function %89
+ %lIndex = OpVariable %_ptr_Function_uint Function %179
+%triangleOffset = OpVariable %_ptr_Function_int Function %239
+ %244 = OpCompositeExtract %uint %GlobalInvocationID 0
+ OpStore %triangleIndex %244
+ %246 = OpLoad %uint %triangleIndex
+ %247 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+ %248 = OpLoad %uint %247
+ %249 = OpUGreaterThanEqual %bool %246 %248
+ OpSelectionMerge %250 None
+ OpBranchConditional %249 %251 %250
+ %251 = OpLabel
+ OpReturn
+ %250 = OpLabel
+ %252 = OpFunctionCall %void %doIgnore
+ %253 = OpLoad %uint %triangleIndex
+ %254 = OpIMul %uint %uint_3 %253
+ %255 = OpIAdd %uint %254 %uint_0
+ %256 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %255
+ %257 = OpLoad %uint %256
+ OpStore %i0 %257
+ %259 = OpLoad %uint %i0
+ %260 = OpIMul %uint %uint_3 %259
+ %261 = OpIAdd %uint %260 %uint_1
+ %262 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %261
+ %263 = OpLoad %uint %262
+ OpStore %i1 %263
+ %265 = OpLoad %uint %i0
+ %266 = OpIMul %uint %uint_3 %265
+ %267 = OpIAdd %uint %266 %uint_2
+ %268 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %267
+ %269 = OpLoad %uint %268
+ OpStore %i2 %269
+ %272 = OpLoad %uint %i0
+ %271 = OpFunctionCall %v3float %loadPosition %272
+ OpStore %p0 %271
+ %275 = OpLoad %uint %i0
+ %274 = OpFunctionCall %v3float %loadPosition %275
+ OpStore %p1 %274
+ %278 = OpLoad %uint %i2
+ %277 = OpFunctionCall %v3float %loadPosition %278
+ OpStore %p2 %277
+ %280 = OpLoad %v3float %p0
+ %281 = OpLoad %v3float %p2
+ %282 = OpFAdd %v3float %280 %281
+ %283 = OpLoad %v3float %p1
+ %284 = OpFAdd %v3float %282 %283
+ %288 = OpCompositeConstruct %v3float %float_3 %float_3 %float_3
+ %286 = OpFDiv %v3float %284 %288
+ OpStore %center %286
+ %291 = OpLoad %v3float %p1
+ %290 = OpFunctionCall %v3float %toVoxelPos %291
+ OpStore %voxelPos_0 %290
+ %294 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+ %295 = OpLoad %uint %294
+ %296 = OpLoad %v3float %p0
+ %293 = OpFunctionCall %uint %toIndex1D %295 %296
+ OpStore %lIndex %293
+ %300 = OpLoad %uint %i1
+ %301 = OpAccessChain %_ptr_StorageBuffer_int %LUT %uint_0 %300
+ %298 = OpAtomicIAdd %int %301 %uint_1 %uint_0 %int_1
+ OpStore %triangleOffset %298
+ OpReturn
+ OpFunctionEnd
+ %main_count = OpFunction %void None %30
+ %305 = OpLabel
+ %307 = OpLoad %v3uint %GlobalInvocationID_1
+ %306 = OpFunctionCall %void %main_count_inner %307
+ OpReturn
+ OpFunctionEnd
diff --git a/test/bug/chromium/1273230.wgsl.expected.wgsl b/test/bug/chromium/1273230.wgsl.expected.wgsl
new file mode 100644
index 0000000..6d99fbb
--- /dev/null
+++ b/test/bug/chromium/1273230.wgsl.expected.wgsl
@@ -0,0 +1,150 @@
+bug/chromium/1273230.wgsl:4:7 warning: use of deprecated intrinsic
+ _ = isNormal(4.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:7:3 warning: use of deprecated intrinsic
+ isNormal(vec4<f32>());
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:10:6 warning: use of deprecated intrinsic
+ isNormal(0.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:11:9 warning: use of deprecated intrinsic
+ _ = isNormal(4.);
+ ^^^^^^^^
+
+bug/chromium/1273230.wgsl:12:9 warning: use of deprecated intrinsic
+ _ = isNormal(2.);
+ ^^^^^^^^
+
+fn marg8uintin() {
+ _ = 0;
+ _ = isNormal(4.0);
+ _ = vec4<f32>(2.0);
+ isNormal(vec4<f32>());
+ _ = vec4<f32>(2.0);
+ isNormal(0.0);
+ _ = isNormal(4.0);
+ _ = isNormal(2.0);
+}
+
+[[block]]
+struct Uniforms {
+ numTriangles : u32;
+ gridSize : u32;
+ puuuuuuuuuuuuuuuuad1 : u32;
+ pad2 : u32;
+ bbMin : vec3<f32>;
+ bbMax : vec3<f32>;
+};
+
+[[block]]
+struct Dbg {
+ offsetCounter : atomic<u32>;
+ pad0 : u32;
+ pad1 : u32;
+ pad2 : u32;
+ value0 : u32;
+ value1 : u32;
+ value2 : u32;
+ value3 : u32;
+ value_f32_0 : f32;
+ value_f32_1 : f32;
+ value_f32_2 : f32;
+ value_f32_3 : f32;
+};
+
+[[block]]
+struct F32s {
+ values : [[stride(4)]] array<f32>;
+};
+
+[[block]]
+struct U32s {
+ values : [[stride(4)]] array<u32>;
+};
+
+[[block]]
+struct I32s {
+ values : [[stride(4)]] array<i32>;
+};
+
+[[block]]
+struct AU32s {
+ values : [[stride(4)]] array<atomic<u32>>;
+};
+
+[[block]]
+struct AI32s {
+ values : [[stride(4)]] array<atomic<i32>>;
+};
+
+[[binding(0), group(0)]] var<uniform> uniforms : Uniforms;
+
+[[binding(10), group(0)]] var<storage, read_write> indices : U32s;
+
+[[binding(11), group(0)]] var<storage, read_write> positions : F32s;
+
+[[binding(20), group(0)]] var<storage, read_write> counters : AU32s;
+
+[[binding(21), group(0)]] var<storage, read_write> LUT : AI32s;
+
+[[binding(50), group(0)]] var<storage, read_write> dbg : Dbg;
+
+fn toVoxelPos(position : vec3<f32>) -> vec3<f32> {
+ var bbMin = vec3<f32>(uniforms.bbMin.x, uniforms.bbMin.y, uniforms.bbMin.z);
+ var bbMax = vec3<f32>(uniforms.bbMax.x, uniforms.bbMax.y, uniforms.bbMax.z);
+ var bbSize = (bbMin - bbMin);
+ var cubeSize = max(max(bbMax.x, bbMax.y), bbSize.z);
+ var gridSize = f32(uniforms.gridSize);
+ var gx = ((cubeSize * (position.x - uniforms.bbMin.x)) / cubeSize);
+ var gy = ((gx * (position.y - uniforms.bbMin.y)) / gridSize);
+ var gz = ((gridSize * (position.z - uniforms.bbMin.z)) / gridSize);
+ return vec3<f32>(gz, gz, gz);
+}
+
+fn toIndex1D(gridSize : u32, voxelPos : vec3<f32>) -> u32 {
+ var icoord = vec3<u32>(voxelPos);
+ return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
+}
+
+fn toIndex4D(gridSize : u32, index : u32) -> vec3<u32> {
+ var z = (gridSize / (index * index));
+ var y = ((gridSize - ((gridSize * gridSize) * z)) / gridSize);
+ var x = (index % gridSize);
+ return vec3<u32>(z, y, y);
+}
+
+fn loadPosition(vertexIndex : u32) -> vec3<f32> {
+ var position = vec3<f32>(positions.values[((3u * vertexIndex) + 0u)], positions.values[((3u * vertexIndex) + 1u)], positions.values[((3u * vertexIndex) + 2u)]);
+ return position;
+}
+
+fn doIgnore() {
+ var g43 = uniforms.numTriangles;
+ var kj6 = dbg.value1;
+ var b53 = atomicLoad(&(counters.values[0]));
+ var rwg = indices.values[0];
+ var rb5 = positions.values[0];
+ var g55 = atomicLoad(&(LUT.values[0]));
+}
+
+[[stage(compute), workgroup_size(128)]]
+fn main_count([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+ var triangleIndex = GlobalInvocationID.x;
+ if ((triangleIndex >= uniforms.numTriangles)) {
+ return;
+ }
+ doIgnore();
+ var i0 = indices.values[((3u * triangleIndex) + 0u)];
+ var i1 = indices.values[((3u * i0) + 1u)];
+ var i2 = indices.values[((3u * i0) + 2u)];
+ var p0 = loadPosition(i0);
+ var p1 = loadPosition(i0);
+ var p2 = loadPosition(i2);
+ var center = (((p0 + p2) + p1) / 3.0);
+ var voxelPos = toVoxelPos(p1);
+ var lIndex = toIndex1D(uniforms.gridSize, p0);
+ var triangleOffset = atomicAdd(&(LUT.values[i1]), 1);
+}
