diff --git a/fuzzers/tint_common_fuzzer.cc b/fuzzers/tint_common_fuzzer.cc index 7e6981501e..12fa3c4a9b 100644 --- a/fuzzers/tint_common_fuzzer.cc +++ b/fuzzers/tint_common_fuzzer.cc @@ -209,7 +209,7 @@ int CommonFuzzer::Run(const uint8_t* data, size_t size) { 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 @@ int CommonFuzzer::Run(const uint8_t* data, size_t size) { 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 @@ int CommonFuzzer::Run(const uint8_t* data, size_t size) { 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 @@ int CommonFuzzer::Run(const uint8_t* data, size_t size) { 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 ddf2697282..3efa18cd10 100644 --- a/src/transform/remove_phonies.cc +++ b/src/transform/remove_phonies.cc @@ -77,11 +77,20 @@ void RemovePhonies::Run(CloneContext& ctx, const DataMap&, DataMap&) { if (auto* stmt = node->As()) { if (stmt->lhs->Is()) { std::vector 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()) { + 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 9629ab84fc..cd7b4ec9f8 100644 --- a/src/transform/remove_phonies_test.cc +++ b/src/transform/remove_phonies_test.cc @@ -45,6 +45,11 @@ fn f() { _ = 1; _ = 1 + 2; _ = t; + _ = u32(3.0); + _ = f32(i32(4u)); + _ = vec2(5.0); + _ = vec3(6, 7, 8); + _ = mat2x2(9.0, 10.0, 11.0, 12.0); } )"; @@ -75,6 +80,11 @@ fn f() { _ = neg(1); _ = add(2, 3); _ = add(neg(4), neg(5)); + _ = u32(neg(6)); + _ = f32(add(7, 8)); + _ = vec2(f32(neg(9))); + _ = vec3(1, neg(10), 3); + _ = mat2x2(1.0, f32(add(11, 12)), 3.0, 4.0); } )"; @@ -91,6 +101,11 @@ fn f() { 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 0000000000..0a5cf26d08 --- /dev/null +++ b/test/bug/chromium/1273230.wgsl @@ -0,0 +1,132 @@ +// /fallthroUgh] +fn marg8uintin() { + _ = (0 ); + _ = isNormal(4.); +_ = (vec4( 2.)); + + isNormal(vec4()); + _ = (vec4( 2.)); + + isNormal(0.); + _ = isNormal(4.); + _ = isNormal(2.); + +}[[block]] +struct Uniforms { + numTriangles : u32; + gridSize : u32; + puuuuuuuuuuuuuuuuad1 : u32; + pad2 : u32; + bbMin : vec3; + bbMax : vec3; +}; + +[[block]] +struct Dbg { + offsetCounter : atomic; + 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; +}; + +[[block]] +struct U32s { + values : [[stride(4)]] array; +}; + +[[block]] +struct I32s { + values : [[stride(4)]] array; +}; + +[[block]] +struct AU32s { + values : [[stride(4)]] array>; +}; + +[[block]] +struct AI32s { + values : [[stride(4)]] array>; +}; + +[[binding(0), group(0)]] var uniforms : Uniforms; + +[[binding(10), group(0)]] var indices : U32s; + +[[binding(11), group(0)]] var positions : F32s; + +[[binding(20), group(0)]] var counters : AU32s; + +[[binding(21), group(0)]] var LUT : AI32s; + +[[binding(50), group(0)]] var dbg : Dbg; + +fn toVoxelPos(position : vec3) -> vec3 { + var bbMin = vec3(uniforms.bbMin.x, uniforms.bbMin.y, uniforms.bbMin.z); + var bbMax = vec3(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(gz, gz, gz); +} + +fn toIndex1D(gridSize : u32, voxelPos : vec3) -> u32 { + var icoord = vec3(voxelPos); + return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z)); +} + +fn toIndex4D(gridSize : u32, index : u32) -> vec3 { + var z = (gridSize / (index * index)); + var y = ((gridSize - ((gridSize * gridSize) * z)) / gridSize); + var x = (index % gridSize); + return vec3(z, y, y); +} + +fn loadPosition(vertexIndex : u32) -> vec3 { + var position = vec3(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) { + 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 0000000000..f721475126 --- /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()); + ^^^^^^^^ + +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 0000000000..76333070eb --- /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()); + ^^^^^^^^ + +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 + +using namespace metal; + +template +inline auto operator*(matrix lhs, packed_vec rhs) { + return lhs * vec(rhs); +} + +template +inline auto operator*(packed_vec lhs, matrix rhs) { + return vec(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 0000000000..6719df277f --- /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()); + ^^^^^^^^ + +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 0000000000..6d99fbb371 --- /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()); + ^^^^^^^^ + +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(2.0); + isNormal(vec4()); + _ = vec4(2.0); + isNormal(0.0); + _ = isNormal(4.0); + _ = isNormal(2.0); +} + +[[block]] +struct Uniforms { + numTriangles : u32; + gridSize : u32; + puuuuuuuuuuuuuuuuad1 : u32; + pad2 : u32; + bbMin : vec3; + bbMax : vec3; +}; + +[[block]] +struct Dbg { + offsetCounter : atomic; + 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; +}; + +[[block]] +struct U32s { + values : [[stride(4)]] array; +}; + +[[block]] +struct I32s { + values : [[stride(4)]] array; +}; + +[[block]] +struct AU32s { + values : [[stride(4)]] array>; +}; + +[[block]] +struct AI32s { + values : [[stride(4)]] array>; +}; + +[[binding(0), group(0)]] var uniforms : Uniforms; + +[[binding(10), group(0)]] var indices : U32s; + +[[binding(11), group(0)]] var positions : F32s; + +[[binding(20), group(0)]] var counters : AU32s; + +[[binding(21), group(0)]] var LUT : AI32s; + +[[binding(50), group(0)]] var dbg : Dbg; + +fn toVoxelPos(position : vec3) -> vec3 { + var bbMin = vec3(uniforms.bbMin.x, uniforms.bbMin.y, uniforms.bbMin.z); + var bbMax = vec3(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(gz, gz, gz); +} + +fn toIndex1D(gridSize : u32, voxelPos : vec3) -> u32 { + var icoord = vec3(voxelPos); + return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z)); +} + +fn toIndex4D(gridSize : u32, index : u32) -> vec3 { + var z = (gridSize / (index * index)); + var y = ((gridSize - ((gridSize * gridSize) * z)) / gridSize); + var x = (index % gridSize); + return vec3(z, y, y); +} + +fn loadPosition(vertexIndex : u32) -> vec3 { + var position = vec3(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) { + 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); +}