From 9943de6813695e831a9e6aa43ae36322e81ae3e6 Mon Sep 17 00:00:00 2001 From: Antonio Maiorano Date: Wed, 22 Dec 2021 15:02:09 +0000 Subject: [PATCH] Fix FXC compile errors on modulo by zero Just like for divide, FXC fails with the exact same error when performing a modulo on a value that FXC determines to be zero. We address it in the same way as we do for divide. This also fixes a couple of the vk-gl-cts tests for which I manually generated expectation files for. Bug: tint:1083 Change-Id: Ia388bf002112afded542adb791d37e88e35a77ff Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/74220 Reviewed-by: James Price Kokoro: Kokoro Commit-Queue: Antonio Maiorano --- src/program_builder.h | 10 ++ src/writer/hlsl/generator_impl.cc | 10 +- src/writer/hlsl/generator_impl_binary_test.cc | 126 ++++++++++++------ test/bug/chromium/1273230.wgsl.expected.hlsl | 2 +- test/bug/tint/1113.wgsl.expected.hlsl | 2 +- .../by_constant/scalar-scalar/f32.wgsl | 6 + .../scalar-scalar/f32.wgsl.expected.hlsl | 5 + .../scalar-scalar/f32.wgsl.expected.msl | 10 ++ .../scalar-scalar/f32.wgsl.expected.spvasm | 20 +++ .../scalar-scalar/f32.wgsl.expected.wgsl | 6 + .../by_constant/scalar-scalar/i32.wgsl | 6 + .../scalar-scalar/i32.wgsl.expected.hlsl | 5 + .../scalar-scalar/i32.wgsl.expected.msl | 10 ++ .../scalar-scalar/i32.wgsl.expected.spvasm | 20 +++ .../scalar-scalar/i32.wgsl.expected.wgsl | 6 + .../by_constant/scalar-scalar/u32.wgsl | 6 + .../scalar-scalar/u32.wgsl.expected.hlsl | 5 + .../scalar-scalar/u32.wgsl.expected.msl | 10 ++ .../scalar-scalar/u32.wgsl.expected.spvasm | 20 +++ .../scalar-scalar/u32.wgsl.expected.wgsl | 6 + .../by_constant/scalar-vec3/i32.wgsl | 6 + .../scalar-vec3/i32.wgsl.expected.hlsl | 7 + .../scalar-vec3/i32.wgsl.expected.msl | 10 ++ .../scalar-vec3/i32.wgsl.expected.spvasm | 27 ++++ .../scalar-vec3/i32.wgsl.expected.wgsl | 6 + .../by_constant/scalar-vec3/u32.wgsl | 6 + .../scalar-vec3/u32.wgsl.expected.hlsl | 7 + .../scalar-vec3/u32.wgsl.expected.msl | 10 ++ .../scalar-vec3/u32.wgsl.expected.spvasm | 27 ++++ .../scalar-vec3/u32.wgsl.expected.wgsl | 6 + .../by_constant/vec3-scalar/i32.wgsl | 6 + .../vec3-scalar/i32.wgsl.expected.hlsl | 6 + .../vec3-scalar/i32.wgsl.expected.msl | 10 ++ .../vec3-scalar/i32.wgsl.expected.spvasm | 28 ++++ .../vec3-scalar/i32.wgsl.expected.wgsl | 6 + .../by_constant/vec3-scalar/u32.wgsl | 6 + .../vec3-scalar/u32.wgsl.expected.hlsl | 6 + .../vec3-scalar/u32.wgsl.expected.msl | 10 ++ .../vec3-scalar/u32.wgsl.expected.spvasm | 28 ++++ .../vec3-scalar/u32.wgsl.expected.wgsl | 6 + .../by_constant/vec3-vec3/f32.wgsl | 6 + .../vec3-vec3/f32.wgsl.expected.hlsl | 7 + .../vec3-vec3/f32.wgsl.expected.msl | 10 ++ .../vec3-vec3/f32.wgsl.expected.spvasm | 26 ++++ .../vec3-vec3/f32.wgsl.expected.wgsl | 6 + .../by_constant/vec3-vec3/i32.wgsl | 6 + .../vec3-vec3/i32.wgsl.expected.hlsl | 7 + .../vec3-vec3/i32.wgsl.expected.msl | 10 ++ .../vec3-vec3/i32.wgsl.expected.spvasm | 26 ++++ .../vec3-vec3/i32.wgsl.expected.wgsl | 6 + .../by_constant/vec3-vec3/u32.wgsl | 6 + .../vec3-vec3/u32.wgsl.expected.hlsl | 7 + .../vec3-vec3/u32.wgsl.expected.msl | 10 ++ .../vec3-vec3/u32.wgsl.expected.spvasm | 26 ++++ .../vec3-vec3/u32.wgsl.expected.wgsl | 6 + .../by_expression/scalar-scalar/f32.wgsl | 6 + .../scalar-scalar/f32.wgsl.expected.hlsl | 7 + .../scalar-scalar/f32.wgsl.expected.msl | 10 ++ .../scalar-scalar/f32.wgsl.expected.spvasm | 32 +++++ .../scalar-scalar/f32.wgsl.expected.wgsl | 6 + .../by_expression/scalar-scalar/i32.wgsl | 6 + .../scalar-scalar/i32.wgsl.expected.hlsl | 11 ++ .../scalar-scalar/i32.wgsl.expected.msl | 10 ++ .../scalar-scalar/i32.wgsl.expected.spvasm | 32 +++++ .../scalar-scalar/i32.wgsl.expected.wgsl | 6 + .../by_expression/scalar-scalar/u32.wgsl | 6 + .../scalar-scalar/u32.wgsl.expected.hlsl | 11 ++ .../scalar-scalar/u32.wgsl.expected.msl | 10 ++ .../scalar-scalar/u32.wgsl.expected.spvasm | 32 +++++ .../scalar-scalar/u32.wgsl.expected.wgsl | 6 + .../by_expression/scalar-vec3/i32.wgsl | 6 + .../scalar-vec3/i32.wgsl.expected.hlsl | 11 ++ .../scalar-vec3/i32.wgsl.expected.msl | 10 ++ .../scalar-vec3/i32.wgsl.expected.spvasm | 39 ++++++ .../scalar-vec3/i32.wgsl.expected.wgsl | 6 + .../by_expression/scalar-vec3/u32.wgsl | 6 + .../scalar-vec3/u32.wgsl.expected.hlsl | 11 ++ .../scalar-vec3/u32.wgsl.expected.msl | 10 ++ .../scalar-vec3/u32.wgsl.expected.spvasm | 39 ++++++ .../scalar-vec3/u32.wgsl.expected.wgsl | 6 + .../by_expression/vec3-scalar/i32.wgsl | 6 + .../vec3-scalar/i32.wgsl.expected.hlsl | 11 ++ .../vec3-scalar/i32.wgsl.expected.msl | 10 ++ .../vec3-scalar/i32.wgsl.expected.spvasm | 40 ++++++ .../vec3-scalar/i32.wgsl.expected.wgsl | 6 + .../by_expression/vec3-scalar/u32.wgsl | 6 + .../vec3-scalar/u32.wgsl.expected.hlsl | 11 ++ .../vec3-scalar/u32.wgsl.expected.msl | 10 ++ .../vec3-scalar/u32.wgsl.expected.spvasm | 40 ++++++ .../vec3-scalar/u32.wgsl.expected.wgsl | 6 + .../by_expression/vec3-vec3/f32.wgsl | 6 + .../vec3-vec3/f32.wgsl.expected.hlsl | 7 + .../vec3-vec3/f32.wgsl.expected.msl | 10 ++ .../vec3-vec3/f32.wgsl.expected.spvasm | 38 ++++++ .../vec3-vec3/f32.wgsl.expected.wgsl | 6 + .../by_expression/vec3-vec3/i32.wgsl | 6 + .../vec3-vec3/i32.wgsl.expected.hlsl | 11 ++ .../vec3-vec3/i32.wgsl.expected.msl | 10 ++ .../vec3-vec3/i32.wgsl.expected.spvasm | 38 ++++++ .../vec3-vec3/i32.wgsl.expected.wgsl | 6 + .../by_expression/vec3-vec3/u32.wgsl | 6 + .../vec3-vec3/u32.wgsl.expected.hlsl | 11 ++ .../vec3-vec3/u32.wgsl.expected.msl | 10 ++ .../vec3-vec3/u32.wgsl.expected.spvasm | 38 ++++++ .../vec3-vec3/u32.wgsl.expected.wgsl | 6 + .../by_identifier/scalar-scalar/f32.wgsl | 6 + .../scalar-scalar/f32.wgsl.expected.hlsl | 7 + .../scalar-scalar/f32.wgsl.expected.msl | 10 ++ .../scalar-scalar/f32.wgsl.expected.spvasm | 30 +++++ .../scalar-scalar/f32.wgsl.expected.wgsl | 6 + .../by_identifier/scalar-scalar/i32.wgsl | 6 + .../scalar-scalar/i32.wgsl.expected.hlsl | 7 + .../scalar-scalar/i32.wgsl.expected.msl | 10 ++ .../scalar-scalar/i32.wgsl.expected.spvasm | 30 +++++ .../scalar-scalar/i32.wgsl.expected.wgsl | 6 + .../by_identifier/scalar-scalar/u32.wgsl | 6 + .../scalar-scalar/u32.wgsl.expected.hlsl | 7 + .../scalar-scalar/u32.wgsl.expected.msl | 10 ++ .../scalar-scalar/u32.wgsl.expected.spvasm | 30 +++++ .../scalar-scalar/u32.wgsl.expected.wgsl | 6 + .../by_identifier/scalar-vec3/i32.wgsl | 6 + .../scalar-vec3/i32.wgsl.expected.hlsl | 7 + .../scalar-vec3/i32.wgsl.expected.msl | 10 ++ .../scalar-vec3/i32.wgsl.expected.spvasm | 37 +++++ .../scalar-vec3/i32.wgsl.expected.wgsl | 6 + .../by_identifier/scalar-vec3/u32.wgsl | 6 + .../scalar-vec3/u32.wgsl.expected.hlsl | 7 + .../scalar-vec3/u32.wgsl.expected.msl | 10 ++ .../scalar-vec3/u32.wgsl.expected.spvasm | 37 +++++ .../scalar-vec3/u32.wgsl.expected.wgsl | 6 + .../by_identifier/vec3-scalar/i32.wgsl | 6 + .../vec3-scalar/i32.wgsl.expected.hlsl | 7 + .../vec3-scalar/i32.wgsl.expected.msl | 10 ++ .../vec3-scalar/i32.wgsl.expected.spvasm | 38 ++++++ .../vec3-scalar/i32.wgsl.expected.wgsl | 6 + .../by_identifier/vec3-scalar/u32.wgsl | 6 + .../vec3-scalar/u32.wgsl.expected.hlsl | 7 + .../vec3-scalar/u32.wgsl.expected.msl | 10 ++ .../vec3-scalar/u32.wgsl.expected.spvasm | 38 ++++++ .../vec3-scalar/u32.wgsl.expected.wgsl | 6 + .../by_identifier/vec3-vec3/f32.wgsl | 6 + .../vec3-vec3/f32.wgsl.expected.hlsl | 7 + .../vec3-vec3/f32.wgsl.expected.msl | 10 ++ .../vec3-vec3/f32.wgsl.expected.spvasm | 36 +++++ .../vec3-vec3/f32.wgsl.expected.wgsl | 6 + .../by_identifier/vec3-vec3/i32.wgsl | 6 + .../vec3-vec3/i32.wgsl.expected.hlsl | 7 + .../vec3-vec3/i32.wgsl.expected.msl | 10 ++ .../vec3-vec3/i32.wgsl.expected.spvasm | 36 +++++ .../vec3-vec3/i32.wgsl.expected.wgsl | 6 + .../by_identifier/vec3-vec3/u32.wgsl | 6 + .../vec3-vec3/u32.wgsl.expected.hlsl | 7 + .../vec3-vec3/u32.wgsl.expected.msl | 10 ++ .../vec3-vec3/u32.wgsl.expected.spvasm | 36 +++++ .../vec3-vec3/u32.wgsl.expected.wgsl | 6 + .../0-opt.spvasm.expected.hlsl | 7 +- .../0-opt.wgsl.expected.hlsl | 7 +- .../0-opt.spvasm.expected.hlsl | 6 +- .../0-opt.wgsl.expected.hlsl | 6 +- 159 files changed, 1965 insertions(+), 69 deletions(-) create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.hlsl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.msl create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.spvasm create mode 100644 test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.wgsl diff --git a/src/program_builder.h b/src/program_builder.h index 53929af772..b724328dfd 100644 --- a/src/program_builder.h +++ b/src/program_builder.h @@ -1665,6 +1665,16 @@ class ProgramBuilder { Expr(std::forward(rhs))); } + /// @param lhs the left hand argument to the modulo operation + /// @param rhs the right hand argument to the modulo operation + /// @returns a `ast::BinaryExpression` applying modulo of `lhs` by `rhs` + template + const ast::Expression* Mod(LHS&& lhs, RHS&& rhs) { + return create(ast::BinaryOp::kModulo, + Expr(std::forward(lhs)), + Expr(std::forward(rhs))); + } + /// @param lhs the left hand argument to the bit shift right operation /// @param rhs the right hand argument to the bit shift right operation /// @returns a `ast::BinaryExpression` bit shifting right `lhs` by `rhs` diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 53188fb92c..9108f9f487 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -863,8 +863,8 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, break; case ast::BinaryOp::kDivide: out << "/"; - // BUG(crbug.com/tint/1083): Integer divide by zero is a FXC compile - // error, and undefined behavior in WGSL. + // BUG(crbug.com/tint/1083): Integer divide/modulo by zero is a FXC + // compile error, and undefined behavior in WGSL. if (TypeOf(expr->rhs)->UnwrapRef()->is_integer_scalar_or_vector()) { out << " "; return EmitExpressionOrOneIfZero(out, expr->rhs); @@ -872,6 +872,12 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, break; case ast::BinaryOp::kModulo: out << "%"; + // BUG(crbug.com/tint/1083): Integer divide/modulo by zero is a FXC + // compile error, and undefined behavior in WGSL. + if (TypeOf(expr->rhs)->UnwrapRef()->is_integer_scalar_or_vector()) { + out << " "; + return EmitExpressionOrOneIfZero(out, expr->rhs); + } break; case ast::BinaryOp::kNone: diagnostics_.add_error(diag::System::Writer, diff --git a/src/writer/hlsl/generator_impl_binary_test.cc b/src/writer/hlsl/generator_impl_binary_test.cc index 6443c48b0d..199f9b1574 100644 --- a/src/writer/hlsl/generator_impl_binary_test.cc +++ b/src/writer/hlsl/generator_impl_binary_test.cc @@ -137,10 +137,12 @@ INSTANTIATE_TEST_SUITE_P( BinaryData{"(left + right)", ast::BinaryOp::kAdd}, BinaryData{"(left - right)", ast::BinaryOp::kSubtract}, BinaryData{"(left * right)", ast::BinaryOp::kMultiply}, - // NOTE: Integer divide covered by DivideBy* tests below + // NOTE: Integer divide covered by DivOrModBy* tests below BinaryData{"(left / right)", ast::BinaryOp::kDivide, BinaryData::Types::Float}, - BinaryData{"(left % right)", ast::BinaryOp::kModulo})); + // NOTE: Integer modulo covered by DivOrModBy* tests below + BinaryData{"(left % right)", ast::BinaryOp::kModulo, + BinaryData::Types::Float})); TEST_F(HlslGeneratorImplTest_Binary, Multiply_VectorScalar) { auto* lhs = vec3(1.f, 1.f, 1.f); @@ -526,11 +528,36 @@ foo((tint_tmp), (tint_tmp_1), (tint_tmp_2)); )"); } -TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_i32) { +namespace HlslGeneratorDivMod { + +struct Params { + enum class Type { Div, Mod }; + Type type; +}; + +struct HlslGeneratorDivModTest : TestParamHelper { + std::string Token() { + return GetParam().type == Params::Type::Div ? "/" : "%"; + } + + template + auto Op(Args... args) { + return GetParam().type == Params::Type::Div + ? Div(std::forward(args)...) + : Mod(std::forward(args)...); + } +}; + +INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest, + HlslGeneratorDivModTest, + testing::Values(Params{Params::Type::Div}, + Params{Params::Type::Mod})); + +TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_i32) { Func("fn", {}, ty.void_(), { Decl(Var("a", ty.i32())), - Decl(Const("r", nullptr, Div("a", 0))), + Decl(Const("r", nullptr, Op("a", 0))), }); GeneratorImpl& gen = Build(); @@ -538,16 +565,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_i32) { ASSERT_TRUE(gen.Generate()); EXPECT_EQ(gen.result(), R"(void fn() { int a = 0; - const int r = (a / 1); + const int r = (a )" + Token() + + R"( 1); } )"); } -TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_u32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_u32) { Func("fn", {}, ty.void_(), { Decl(Var("a", ty.u32())), - Decl(Const("r", nullptr, Div("a", 0u))), + Decl(Const("r", nullptr, Op("a", 0u))), }); GeneratorImpl& gen = Build(); @@ -555,16 +583,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_u32) { ASSERT_TRUE(gen.Generate()); EXPECT_EQ(gen.result(), R"(void fn() { uint a = 0u; - const uint r = (a / 1u); + const uint r = (a )" + Token() + + R"( 1u); } )"); -} +} // namespace HlslGeneratorDivMod -TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_vec_i32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_vec_by_vec_i32) { Func("fn", {}, ty.void_(), { Decl(Var("a", nullptr, vec4(100, 100, 100, 100))), - Decl(Const("r", nullptr, Div("a", vec4(50, 0, 25, 0)))), + Decl(Const("r", nullptr, Op("a", vec4(50, 0, 25, 0)))), }); GeneratorImpl& gen = Build(); @@ -572,16 +601,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_vec_i32) { ASSERT_TRUE(gen.Generate()); EXPECT_EQ(gen.result(), R"(void fn() { int4 a = int4(100, 100, 100, 100); - const int4 r = (a / int4(50, 1, 25, 1)); + const int4 r = (a )" + Token() + + R"( int4(50, 1, 25, 1)); } )"); -} +} // namespace -TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_scalar_i32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_vec_by_scalar_i32) { Func("fn", {}, ty.void_(), { Decl(Var("a", nullptr, vec4(100, 100, 100, 100))), - Decl(Const("r", nullptr, Div("a", 0))), + Decl(Const("r", nullptr, Op("a", 0))), }); GeneratorImpl& gen = Build(); @@ -589,16 +619,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_scalar_i32) { ASSERT_TRUE(gen.Generate()); EXPECT_EQ(gen.result(), R"(void fn() { int4 a = int4(100, 100, 100, 100); - const int4 r = (a / 1); + const int4 r = (a )" + Token() + + R"( 1); } )"); -} +} // namespace hlsl -TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_i32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_i32) { Func("fn", {Param("b", ty.i32())}, ty.void_(), { Decl(Var("a", ty.i32())), - Decl(Const("r", nullptr, Div("a", "b"))), + Decl(Const("r", nullptr, Op("a", "b"))), }); GeneratorImpl& gen = Build(); @@ -606,16 +637,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_i32) { ASSERT_TRUE(gen.Generate()); EXPECT_EQ(gen.result(), R"(void fn(int b) { int a = 0; - const int r = (a / (b == 0 ? 1 : b)); + const int r = (a )" + Token() + + R"( (b == 0 ? 1 : b)); } )"); -} +} // namespace writer -TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_u32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_u32) { Func("fn", {Param("b", ty.u32())}, ty.void_(), { Decl(Var("a", ty.u32())), - Decl(Const("r", nullptr, Div("a", "b"))), + Decl(Const("r", nullptr, Op("a", "b"))), }); GeneratorImpl& gen = Build(); @@ -623,16 +655,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_u32) { ASSERT_TRUE(gen.Generate()); EXPECT_EQ(gen.result(), R"(void fn(uint b) { uint a = 0u; - const uint r = (a / (b == 0u ? 1u : b)); + const uint r = (a )" + Token() + + R"( (b == 0u ? 1u : b)); } )"); -} +} // namespace tint -TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_vec_i32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_vec_by_vec_i32) { Func("fn", {Param("b", ty.vec3())}, ty.void_(), { Decl(Var("a", ty.vec3())), - Decl(Const("r", nullptr, Div("a", "b"))), + Decl(Const("r", nullptr, Op("a", "b"))), }); GeneratorImpl& gen = Build(); @@ -640,16 +673,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_vec_i32) { ASSERT_TRUE(gen.Generate()); EXPECT_EQ(gen.result(), R"(void fn(int3 b) { int3 a = int3(0, 0, 0); - const int3 r = (a / (b == int3(0, 0, 0) ? int3(1, 1, 1) : b)); + const int3 r = (a )" + Token() + + R"( (b == int3(0, 0, 0) ? int3(1, 1, 1) : b)); } )"); } -TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_scalar_i32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_vec_by_scalar_i32) { Func("fn", {Param("b", ty.i32())}, ty.void_(), { Decl(Var("a", ty.vec3())), - Decl(Const("r", nullptr, Div("a", "b"))), + Decl(Const("r", nullptr, Op("a", "b"))), }); GeneratorImpl& gen = Build(); @@ -657,12 +691,13 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_scalar_i32) { ASSERT_TRUE(gen.Generate()); EXPECT_EQ(gen.result(), R"(void fn(int b) { int3 a = int3(0, 0, 0); - const int3 r = (a / (b == 0 ? 1 : b)); + const int3 r = (a )" + Token() + + R"( (b == 0 ? 1 : b)); } )"); } -TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_i32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_i32) { Func("zero", {}, ty.i32(), { Return(Expr(0)), @@ -671,7 +706,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_i32) { Func("fn", {}, ty.void_(), { Decl(Var("a", ty.i32())), - Decl(Const("r", nullptr, Div("a", Call("zero")))), + Decl(Const("r", nullptr, Op("a", Call("zero")))), }); GeneratorImpl& gen = Build(); @@ -687,12 +722,13 @@ int zero() { void fn() { int a = 0; - const int r = (a / value_or_one_if_zero_int(zero())); + const int r = (a )" + Token() + + R"( value_or_one_if_zero_int(zero())); } )"); } -TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_u32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_u32) { Func("zero", {}, ty.u32(), { Return(Expr(0u)), @@ -701,7 +737,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_u32) { Func("fn", {}, ty.void_(), { Decl(Var("a", ty.u32())), - Decl(Const("r", nullptr, Div("a", Call("zero")))), + Decl(Const("r", nullptr, Op("a", Call("zero")))), }); GeneratorImpl& gen = Build(); @@ -717,12 +753,13 @@ uint zero() { void fn() { uint a = 0u; - const uint r = (a / value_or_one_if_zero_uint(zero())); + const uint r = (a )" + Token() + + R"( value_or_one_if_zero_uint(zero())); } )"); } -TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_vec_i32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_vec_by_vec_i32) { Func("zero", {}, ty.vec3(), { Return(vec3(0, 0, 0)), @@ -731,7 +768,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_vec_i32) { Func("fn", {}, ty.void_(), { Decl(Var("a", ty.vec3())), - Decl(Const("r", nullptr, Div("a", Call("zero")))), + Decl(Const("r", nullptr, Op("a", Call("zero")))), }); GeneratorImpl& gen = Build(); @@ -747,12 +784,13 @@ int3 zero() { void fn() { int3 a = int3(0, 0, 0); - const int3 r = (a / value_or_one_if_zero_int3(zero())); + const int3 r = (a )" + Token() + + R"( value_or_one_if_zero_int3(zero())); } )"); } -TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_scalar_i32) { +TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_vec_by_scalar_i32) { Func("zero", {}, ty.i32(), { Return(0), @@ -761,7 +799,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_scalar_i32) { Func("fn", {}, ty.void_(), { Decl(Var("a", ty.vec3())), - Decl(Const("r", nullptr, Div("a", Call("zero")))), + Decl(Const("r", nullptr, Op("a", Call("zero")))), }); GeneratorImpl& gen = Build(); @@ -777,10 +815,12 @@ int zero() { void fn() { int3 a = int3(0, 0, 0); - const int3 r = (a / value_or_one_if_zero_int(zero())); + const int3 r = (a )" + Token() + + R"( value_or_one_if_zero_int(zero())); } )"); } +} // namespace HlslGeneratorDivMod } // namespace } // namespace hlsl diff --git a/test/bug/chromium/1273230.wgsl.expected.hlsl b/test/bug/chromium/1273230.wgsl.expected.hlsl index aef846b7d0..eb84f1ae3e 100644 --- a/test/bug/chromium/1273230.wgsl.expected.hlsl +++ b/test/bug/chromium/1273230.wgsl.expected.hlsl @@ -89,7 +89,7 @@ uint toIndex1D(uint gridSize, float3 voxelPos) { uint3 toIndex4D(uint gridSize, uint index) { uint z_1 = (gridSize / value_or_one_if_zero_uint((index * index))); uint y_1 = ((gridSize - ((gridSize * gridSize) * z_1)) / (gridSize == 0u ? 1u : gridSize)); - uint x_1 = (index % gridSize); + uint x_1 = (index % (gridSize == 0u ? 1u : gridSize)); return uint3(z_1, y_1, y_1); } diff --git a/test/bug/tint/1113.wgsl.expected.hlsl b/test/bug/tint/1113.wgsl.expected.hlsl index 656837ed1d..b2da18f9a2 100644 --- a/test/bug/tint/1113.wgsl.expected.hlsl +++ b/test/bug/tint/1113.wgsl.expected.hlsl @@ -60,7 +60,7 @@ uint toIndex1D(uint gridSize, float3 voxelPos) { uint3 toIndex3D(uint gridSize, uint index) { uint z_1 = (index / value_or_one_if_zero_uint((gridSize * gridSize))); uint y_1 = ((index - ((gridSize * gridSize) * z_1)) / (gridSize == 0u ? 1u : gridSize)); - uint x_1 = (index % gridSize); + uint x_1 = (index % (gridSize == 0u ? 1u : gridSize)); return uint3(x_1, y_1, z_1); } diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl new file mode 100644 index 0000000000..b34099c8e1 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = 1.; + let b = 0.; + let r : f32 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..d3de7dbaa6 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.hlsl @@ -0,0 +1,5 @@ +[numthreads(1, 1, 1)] +void f() { + const float r = (1.0f % 0.0f); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.msl new file mode 100644 index 0000000000..25c2f18e89 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + float const a = 1.0f; + float const b = 0.0f; + float const r = fmod(a, b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..71bf9df9b6 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 9 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %float_1 = OpConstant %float 1 + %float_0 = OpConstant %float 0 + %f = OpFunction %void None %1 + %4 = OpLabel + %8 = OpFRem %float %float_1 %float_0 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..235be9368b --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = 1.0; + let b = 0.0; + let r : f32 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl new file mode 100644 index 0000000000..9900e056aa --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = 1; + let b = 0; + let r : i32 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..b79aad9396 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.hlsl @@ -0,0 +1,5 @@ +[numthreads(1, 1, 1)] +void f() { + const int r = (1 % 1); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.msl new file mode 100644 index 0000000000..98f1698b74 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int const a = 1; + int const b = 0; + int const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..37536a116e --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 9 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %int_1 = OpConstant %int 1 + %int_0 = OpConstant %int 0 + %f = OpFunction %void None %1 + %4 = OpLabel + %8 = OpSMod %int %int_1 %int_0 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..4c6fdf98e7 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = 1; + let b = 0; + let r : i32 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl new file mode 100644 index 0000000000..4fb0f9fb4a --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = 1u; + let b = 0u; + let r : u32 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..241c5f5c2b --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.hlsl @@ -0,0 +1,5 @@ +[numthreads(1, 1, 1)] +void f() { + const uint r = (1u % 1u); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.msl new file mode 100644 index 0000000000..456e3b9caa --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint const a = 1u; + uint const b = 0u; + uint const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..e59fd6eae8 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 9 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 + %uint_0 = OpConstant %uint 0 + %f = OpFunction %void None %1 + %4 = OpLabel + %8 = OpUMod %uint %uint_1 %uint_0 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..62a525b2df --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = 1u; + let b = 0u; + let r : u32 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl new file mode 100644 index 0000000000..9b27df2468 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = 4; + let b = vec3(0, 2, 0); + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..fa4482f682 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + const int a = 4; + const int3 b = int3(0, 2, 0); + const int3 r = (a % int3(1, 2, 1)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.msl new file mode 100644 index 0000000000..95a93f72b8 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int const a = 4; + int3 const b = int3(0, 2, 0); + int3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..dc1cd0582c --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.spvasm @@ -0,0 +1,27 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 16 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %int_4 = OpConstant %int 4 + %v3int = OpTypeVector %int 3 + %int_0 = OpConstant %int 0 + %int_2 = OpConstant %int 2 + %10 = OpConstantComposite %v3int %int_0 %int_2 %int_0 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %14 = OpConstantNull %v3int + %f = OpFunction %void None %1 + %4 = OpLabel + %12 = OpVariable %_ptr_Function_v3int Function %14 + %15 = OpCompositeConstruct %v3int %int_4 %int_4 %int_4 + %11 = OpSMod %v3int %15 %10 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..dca74427df --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = 4; + let b = vec3(0, 2, 0); + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl new file mode 100644 index 0000000000..494a712979 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = 4u; + let b = vec3(0u, 2u, 0u); + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..12aa8c01ee --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + const uint a = 4u; + const uint3 b = uint3(0u, 2u, 0u); + const uint3 r = (a % uint3(1u, 2u, 1u)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.msl new file mode 100644 index 0000000000..16fa02517d --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint const a = 4u; + uint3 const b = uint3(0u, 2u, 0u); + uint3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..019876b5ff --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.spvasm @@ -0,0 +1,27 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 16 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_4 = OpConstant %uint 4 + %v3uint = OpTypeVector %uint 3 + %uint_0 = OpConstant %uint 0 + %uint_2 = OpConstant %uint 2 + %10 = OpConstantComposite %v3uint %uint_0 %uint_2 %uint_0 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %14 = OpConstantNull %v3uint + %f = OpFunction %void None %1 + %4 = OpLabel + %12 = OpVariable %_ptr_Function_v3uint Function %14 + %15 = OpCompositeConstruct %v3uint %uint_4 %uint_4 %uint_4 + %11 = OpUMod %v3uint %15 %10 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..c061d41eee --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = 4u; + let b = vec3(0u, 2u, 0u); + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl new file mode 100644 index 0000000000..2fe5f2de4f --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = vec3(1, 2, 3); + let b = 0; + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..6906266038 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.hlsl @@ -0,0 +1,6 @@ +[numthreads(1, 1, 1)] +void f() { + const int3 a = int3(1, 2, 3); + const int3 r = (a % 1); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.msl new file mode 100644 index 0000000000..cbac4c79f6 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int3 const a = int3(1, 2, 3); + int const b = 0; + int3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..ac1f1d7a1b --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.spvasm @@ -0,0 +1,28 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 17 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 + %int_1 = OpConstant %int 1 + %int_2 = OpConstant %int 2 + %int_3 = OpConstant %int 3 + %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3 + %int_0 = OpConstant %int 0 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %15 = OpConstantNull %v3int + %f = OpFunction %void None %1 + %4 = OpLabel + %13 = OpVariable %_ptr_Function_v3int Function %15 + %16 = OpCompositeConstruct %v3int %int_0 %int_0 %int_0 + %12 = OpSMod %v3int %10 %16 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..0aac1fa8b0 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = vec3(1, 2, 3); + let b = 0; + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl new file mode 100644 index 0000000000..25b884fafc --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = vec3(1u, 2u, 3u); + let b = 0u; + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..4d8b0b4eef --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.hlsl @@ -0,0 +1,6 @@ +[numthreads(1, 1, 1)] +void f() { + const uint3 a = uint3(1u, 2u, 3u); + const uint3 r = (a % 1u); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.msl new file mode 100644 index 0000000000..78398ec53e --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint3 const a = uint3(1u, 2u, 3u); + uint const b = 0u; + uint3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..b8ce74fa9c --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.spvasm @@ -0,0 +1,28 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 17 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %uint_2 = OpConstant %uint 2 + %uint_3 = OpConstant %uint 3 + %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3 + %uint_0 = OpConstant %uint 0 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %15 = OpConstantNull %v3uint + %f = OpFunction %void None %1 + %4 = OpLabel + %13 = OpVariable %_ptr_Function_v3uint Function %15 + %16 = OpCompositeConstruct %v3uint %uint_0 %uint_0 %uint_0 + %12 = OpUMod %v3uint %10 %16 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..7665b70741 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = vec3(1u, 2u, 3u); + let b = 0u; + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl new file mode 100644 index 0000000000..61df86dc51 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = vec3(1., 2., 3.); + let b = vec3(0., 5., 0.); + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..5ea2fbd98f --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + const float3 a = float3(1.0f, 2.0f, 3.0f); + const float3 b = float3(0.0f, 5.0f, 0.0f); + const float3 r = (a % b); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.msl new file mode 100644 index 0000000000..9506b9c7b2 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + float3 const a = float3(1.0f, 2.0f, 3.0f); + float3 const b = float3(0.0f, 5.0f, 0.0f); + float3 const r = fmod(a, b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..6f5c42785e --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.spvasm @@ -0,0 +1,26 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 15 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v3float = OpTypeVector %float 3 + %float_1 = OpConstant %float 1 + %float_2 = OpConstant %float 2 + %float_3 = OpConstant %float 3 + %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3 + %float_0 = OpConstant %float 0 + %float_5 = OpConstant %float 5 + %13 = OpConstantComposite %v3float %float_0 %float_5 %float_0 + %f = OpFunction %void None %1 + %4 = OpLabel + %14 = OpFRem %v3float %10 %13 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..180e7d96bb --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = vec3(1.0, 2.0, 3.0); + let b = vec3(0.0, 5.0, 0.0); + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl new file mode 100644 index 0000000000..23aa96c5a6 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = vec3(1, 2, 3); + let b = vec3(0, 5, 0); + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..ee6067476a --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + const int3 a = int3(1, 2, 3); + const int3 b = int3(0, 5, 0); + const int3 r = (a % int3(1, 5, 1)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.msl new file mode 100644 index 0000000000..7619c9c3eb --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int3 const a = int3(1, 2, 3); + int3 const b = int3(0, 5, 0); + int3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..5205a1957c --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.spvasm @@ -0,0 +1,26 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 15 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 + %int_1 = OpConstant %int 1 + %int_2 = OpConstant %int 2 + %int_3 = OpConstant %int 3 + %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3 + %int_0 = OpConstant %int 0 + %int_5 = OpConstant %int 5 + %13 = OpConstantComposite %v3int %int_0 %int_5 %int_0 + %f = OpFunction %void None %1 + %4 = OpLabel + %14 = OpSMod %v3int %10 %13 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..259840db34 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = vec3(1, 2, 3); + let b = vec3(0, 5, 0); + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl new file mode 100644 index 0000000000..ee9eecccd3 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = vec3(1u, 2u, 3u); + let b = vec3(0u, 5u, 0u); + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..2389839626 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + const uint3 a = uint3(1u, 2u, 3u); + const uint3 b = uint3(0u, 5u, 0u); + const uint3 r = (a % uint3(1u, 5u, 1u)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.msl new file mode 100644 index 0000000000..028424af19 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint3 const a = uint3(1u, 2u, 3u); + uint3 const b = uint3(0u, 5u, 0u); + uint3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..f1fa32537a --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.spvasm @@ -0,0 +1,26 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 15 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %uint_2 = OpConstant %uint 2 + %uint_3 = OpConstant %uint 3 + %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3 + %uint_0 = OpConstant %uint 0 + %uint_5 = OpConstant %uint 5 + %13 = OpConstantComposite %v3uint %uint_0 %uint_5 %uint_0 + %f = OpFunction %void None %1 + %4 = OpLabel + %14 = OpUMod %v3uint %10 %13 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..08943bbe3b --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + let a = vec3(1u, 2u, 3u); + let b = vec3(0u, 5u, 0u); + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl new file mode 100644 index 0000000000..a4e2202d3c --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1.; + var b = 0.; + let r : f32 = a % (b + b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..468c33d586 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + float a = 1.0f; + float b = 0.0f; + const float r = (a % (b + b)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.msl new file mode 100644 index 0000000000..c3d64df07b --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + float a = 1.0f; + float b = 0.0f; + float const r = fmod(a, (b + b)); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..33268b0b64 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.spvasm @@ -0,0 +1,32 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 17 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %float_1 = OpConstant %float 1 +%_ptr_Function_float = OpTypePointer Function %float + %9 = OpConstantNull %float + %float_0 = OpConstant %float 0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_float Function %9 + %b = OpVariable %_ptr_Function_float Function %9 + OpStore %a %float_1 + OpStore %b %float_0 + %12 = OpLoad %float %a + %13 = OpLoad %float %b + %14 = OpLoad %float %b + %15 = OpFAdd %float %13 %14 + %16 = OpFRem %float %12 %15 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..b6c7b67340 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1.0; + var b = 0.0; + let r : f32 = (a % (b + b)); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl new file mode 100644 index 0000000000..7f8d8d0c0d --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1; + var b = 0; + let r : i32 = a % (b + b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..fd4f932ff7 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +int value_or_one_if_zero_int(int value) { + return value == 0 ? 1 : value; +} + +[numthreads(1, 1, 1)] +void f() { + int a = 1; + int b = 0; + const int r = (a % value_or_one_if_zero_int((b + b))); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.msl new file mode 100644 index 0000000000..ead2dcea05 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int a = 1; + int b = 0; + int const r = (a % as_type((as_type(b) + as_type(b)))); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..a37441fe4d --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.spvasm @@ -0,0 +1,32 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 17 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %int_1 = OpConstant %int 1 +%_ptr_Function_int = OpTypePointer Function %int + %9 = OpConstantNull %int + %int_0 = OpConstant %int 0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_int Function %9 + %b = OpVariable %_ptr_Function_int Function %9 + OpStore %a %int_1 + OpStore %b %int_0 + %12 = OpLoad %int %a + %13 = OpLoad %int %b + %14 = OpLoad %int %b + %15 = OpIAdd %int %13 %14 + %16 = OpSMod %int %12 %15 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..581f6dfb98 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1; + var b = 0; + let r : i32 = (a % (b + b)); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl new file mode 100644 index 0000000000..edd134b0fd --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1u; + var b = 0u; + let r : u32 = a % (b + b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..f8f7511f41 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +uint value_or_one_if_zero_uint(uint value) { + return value == 0u ? 1u : value; +} + +[numthreads(1, 1, 1)] +void f() { + uint a = 1u; + uint b = 0u; + const uint r = (a % value_or_one_if_zero_uint((b + b))); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.msl new file mode 100644 index 0000000000..33bf9ab1b4 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint a = 1u; + uint b = 0u; + uint const r = (a % (b + b)); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..ae27d46639 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.spvasm @@ -0,0 +1,32 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 17 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 +%_ptr_Function_uint = OpTypePointer Function %uint + %9 = OpConstantNull %uint + %uint_0 = OpConstant %uint 0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_uint Function %9 + %b = OpVariable %_ptr_Function_uint Function %9 + OpStore %a %uint_1 + OpStore %b %uint_0 + %12 = OpLoad %uint %a + %13 = OpLoad %uint %b + %14 = OpLoad %uint %b + %15 = OpIAdd %uint %13 %14 + %16 = OpUMod %uint %12 %15 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..f894e6d5d8 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1u; + var b = 0u; + let r : u32 = (a % (b + b)); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl new file mode 100644 index 0000000000..b90a40471a --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 4; + var b = vec3(0, 2, 0); + let r : vec3 = a % (b + b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..8e8a7d0533 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +int3 value_or_one_if_zero_int3(int3 value) { + return value == int3(0, 0, 0) ? int3(1, 1, 1) : value; +} + +[numthreads(1, 1, 1)] +void f() { + int a = 4; + int3 b = int3(0, 2, 0); + const int3 r = (a % value_or_one_if_zero_int3((b + b))); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl new file mode 100644 index 0000000000..12f1a72cc2 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int a = 4; + int3 b = int3(0, 2, 0); + int3 const r = (a % as_type((as_type(b) + as_type(b)))); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..b293226629 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.spvasm @@ -0,0 +1,39 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 24 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %int_4 = OpConstant %int 4 +%_ptr_Function_int = OpTypePointer Function %int + %9 = OpConstantNull %int + %v3int = OpTypeVector %int 3 + %int_0 = OpConstant %int 0 + %int_2 = OpConstant %int 2 + %13 = OpConstantComposite %v3int %int_0 %int_2 %int_0 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %16 = OpConstantNull %v3int + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_int Function %9 + %b = OpVariable %_ptr_Function_v3int Function %16 + %22 = OpVariable %_ptr_Function_v3int Function %16 + OpStore %a %int_4 + OpStore %b %13 + %17 = OpLoad %int %a + %18 = OpLoad %v3int %b + %19 = OpLoad %v3int %b + %20 = OpIAdd %v3int %18 %19 + %23 = OpCompositeConstruct %v3int %17 %17 %17 + %21 = OpSMod %v3int %23 %20 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..724d89b48a --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 4; + var b = vec3(0, 2, 0); + let r : vec3 = (a % (b + b)); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl new file mode 100644 index 0000000000..e6626f1081 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 4u; + var b = vec3(0u, 2u, 0u); + let r : vec3 = a % (b + b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..2330363261 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +uint3 value_or_one_if_zero_uint3(uint3 value) { + return value == uint3(0u, 0u, 0u) ? uint3(1u, 1u, 1u) : value; +} + +[numthreads(1, 1, 1)] +void f() { + uint a = 4u; + uint3 b = uint3(0u, 2u, 0u); + const uint3 r = (a % value_or_one_if_zero_uint3((b + b))); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.msl new file mode 100644 index 0000000000..a7511f2241 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint a = 4u; + uint3 b = uint3(0u, 2u, 0u); + uint3 const r = (a % (b + b)); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..ee0be7a97d --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.spvasm @@ -0,0 +1,39 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 24 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_4 = OpConstant %uint 4 +%_ptr_Function_uint = OpTypePointer Function %uint + %9 = OpConstantNull %uint + %v3uint = OpTypeVector %uint 3 + %uint_0 = OpConstant %uint 0 + %uint_2 = OpConstant %uint 2 + %13 = OpConstantComposite %v3uint %uint_0 %uint_2 %uint_0 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %16 = OpConstantNull %v3uint + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_uint Function %9 + %b = OpVariable %_ptr_Function_v3uint Function %16 + %22 = OpVariable %_ptr_Function_v3uint Function %16 + OpStore %a %uint_4 + OpStore %b %13 + %17 = OpLoad %uint %a + %18 = OpLoad %v3uint %b + %19 = OpLoad %v3uint %b + %20 = OpIAdd %v3uint %18 %19 + %23 = OpCompositeConstruct %v3uint %17 %17 %17 + %21 = OpUMod %v3uint %23 %20 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..d820f98035 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 4u; + var b = vec3(0u, 2u, 0u); + let r : vec3 = (a % (b + b)); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl new file mode 100644 index 0000000000..e5dd9839fc --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1, 2, 3); + var b = 0; + let r : vec3 = a % (b + b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..ea6b4cbb42 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +int value_or_one_if_zero_int(int value) { + return value == 0 ? 1 : value; +} + +[numthreads(1, 1, 1)] +void f() { + int3 a = int3(1, 2, 3); + int b = 0; + const int3 r = (a % value_or_one_if_zero_int((b + b))); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.msl new file mode 100644 index 0000000000..ba6626b194 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int3 a = int3(1, 2, 3); + int b = 0; + int3 const r = (a % as_type((as_type(b) + as_type(b)))); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..d9dd3f1a99 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.spvasm @@ -0,0 +1,40 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 25 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 + %int_1 = OpConstant %int 1 + %int_2 = OpConstant %int 2 + %int_3 = OpConstant %int 3 + %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %13 = OpConstantNull %v3int + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %17 = OpConstantNull %int + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_v3int Function %13 + %b = OpVariable %_ptr_Function_int Function %17 + %23 = OpVariable %_ptr_Function_v3int Function %13 + OpStore %a %10 + OpStore %b %int_0 + %18 = OpLoad %v3int %a + %19 = OpLoad %int %b + %20 = OpLoad %int %b + %21 = OpIAdd %int %19 %20 + %24 = OpCompositeConstruct %v3int %21 %21 %21 + %22 = OpSMod %v3int %18 %24 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..c924762b31 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1, 2, 3); + var b = 0; + let r : vec3 = (a % (b + b)); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl new file mode 100644 index 0000000000..b775a66126 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1u, 2u, 3u); + var b = 0u; + let r : vec3 = a % (b + b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..dca1662c8d --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +uint value_or_one_if_zero_uint(uint value) { + return value == 0u ? 1u : value; +} + +[numthreads(1, 1, 1)] +void f() { + uint3 a = uint3(1u, 2u, 3u); + uint b = 0u; + const uint3 r = (a % value_or_one_if_zero_uint((b + b))); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.msl new file mode 100644 index 0000000000..96dba81ec4 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint3 a = uint3(1u, 2u, 3u); + uint b = 0u; + uint3 const r = (a % (b + b)); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..78e1a7eb2e --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.spvasm @@ -0,0 +1,40 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 25 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %uint_2 = OpConstant %uint 2 + %uint_3 = OpConstant %uint 3 + %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %13 = OpConstantNull %v3uint + %uint_0 = OpConstant %uint 0 +%_ptr_Function_uint = OpTypePointer Function %uint + %17 = OpConstantNull %uint + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_v3uint Function %13 + %b = OpVariable %_ptr_Function_uint Function %17 + %23 = OpVariable %_ptr_Function_v3uint Function %13 + OpStore %a %10 + OpStore %b %uint_0 + %18 = OpLoad %v3uint %a + %19 = OpLoad %uint %b + %20 = OpLoad %uint %b + %21 = OpIAdd %uint %19 %20 + %24 = OpCompositeConstruct %v3uint %21 %21 %21 + %22 = OpUMod %v3uint %18 %24 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..d73c1a9070 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1u, 2u, 3u); + var b = 0u; + let r : vec3 = (a % (b + b)); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl new file mode 100644 index 0000000000..92e19cf0cd --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1., 2., 3.); + var b = vec3(0., 5., 0.); + let r : vec3 = a % (b + b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..8d2155076d --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + float3 a = float3(1.0f, 2.0f, 3.0f); + float3 b = float3(0.0f, 5.0f, 0.0f); + const float3 r = (a % (b + b)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.msl new file mode 100644 index 0000000000..6ef4b08902 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + float3 a = float3(1.0f, 2.0f, 3.0f); + float3 b = float3(0.0f, 5.0f, 0.0f); + float3 const r = fmod(a, (b + b)); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..fb8ce28375 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.spvasm @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v3float = OpTypeVector %float 3 + %float_1 = OpConstant %float 1 + %float_2 = OpConstant %float 2 + %float_3 = OpConstant %float 3 + %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3 +%_ptr_Function_v3float = OpTypePointer Function %v3float + %13 = OpConstantNull %v3float + %float_0 = OpConstant %float 0 + %float_5 = OpConstant %float 5 + %16 = OpConstantComposite %v3float %float_0 %float_5 %float_0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_v3float Function %13 + %b = OpVariable %_ptr_Function_v3float Function %13 + OpStore %a %10 + OpStore %b %16 + %18 = OpLoad %v3float %a + %19 = OpLoad %v3float %b + %20 = OpLoad %v3float %b + %21 = OpFAdd %v3float %19 %20 + %22 = OpFRem %v3float %18 %21 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..06e70aef33 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1.0, 2.0, 3.0); + var b = vec3(0.0, 5.0, 0.0); + let r : vec3 = (a % (b + b)); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl new file mode 100644 index 0000000000..e843464983 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1, 2, 3); + var b = vec3(0, 5, 0); + let r : vec3 = a % (b + b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..f33a43d8ef --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +int3 value_or_one_if_zero_int3(int3 value) { + return value == int3(0, 0, 0) ? int3(1, 1, 1) : value; +} + +[numthreads(1, 1, 1)] +void f() { + int3 a = int3(1, 2, 3); + int3 b = int3(0, 5, 0); + const int3 r = (a % value_or_one_if_zero_int3((b + b))); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl new file mode 100644 index 0000000000..61356dc074 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int3 a = int3(1, 2, 3); + int3 b = int3(0, 5, 0); + int3 const r = (a % as_type((as_type(b) + as_type(b)))); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..b96d65ac92 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.spvasm @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 + %int_1 = OpConstant %int 1 + %int_2 = OpConstant %int 2 + %int_3 = OpConstant %int 3 + %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %13 = OpConstantNull %v3int + %int_0 = OpConstant %int 0 + %int_5 = OpConstant %int 5 + %16 = OpConstantComposite %v3int %int_0 %int_5 %int_0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_v3int Function %13 + %b = OpVariable %_ptr_Function_v3int Function %13 + OpStore %a %10 + OpStore %b %16 + %18 = OpLoad %v3int %a + %19 = OpLoad %v3int %b + %20 = OpLoad %v3int %b + %21 = OpIAdd %v3int %19 %20 + %22 = OpSMod %v3int %18 %21 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..78f7040d78 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1, 2, 3); + var b = vec3(0, 5, 0); + let r : vec3 = (a % (b + b)); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl new file mode 100644 index 0000000000..ad503237a7 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1u, 2u, 3u); + var b = vec3(0u, 5u, 0u); + let r : vec3 = a % (b + b); +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..f112cf8458 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +uint3 value_or_one_if_zero_uint3(uint3 value) { + return value == uint3(0u, 0u, 0u) ? uint3(1u, 1u, 1u) : value; +} + +[numthreads(1, 1, 1)] +void f() { + uint3 a = uint3(1u, 2u, 3u); + uint3 b = uint3(0u, 5u, 0u); + const uint3 r = (a % value_or_one_if_zero_uint3((b + b))); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.msl new file mode 100644 index 0000000000..c6d0db5670 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint3 a = uint3(1u, 2u, 3u); + uint3 b = uint3(0u, 5u, 0u); + uint3 const r = (a % (b + b)); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..c4b1f903c6 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.spvasm @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %uint_2 = OpConstant %uint 2 + %uint_3 = OpConstant %uint 3 + %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %13 = OpConstantNull %v3uint + %uint_0 = OpConstant %uint 0 + %uint_5 = OpConstant %uint 5 + %16 = OpConstantComposite %v3uint %uint_0 %uint_5 %uint_0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_v3uint Function %13 + %b = OpVariable %_ptr_Function_v3uint Function %13 + OpStore %a %10 + OpStore %b %16 + %18 = OpLoad %v3uint %a + %19 = OpLoad %v3uint %b + %20 = OpLoad %v3uint %b + %21 = OpIAdd %v3uint %19 %20 + %22 = OpUMod %v3uint %18 %21 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..e7f2ce8aa5 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1u, 2u, 3u); + var b = vec3(0u, 5u, 0u); + let r : vec3 = (a % (b + b)); +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl new file mode 100644 index 0000000000..b80c20109f --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1.; + var b = 0.; + let r : f32 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..3d43aff25b --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + float a = 1.0f; + float b = 0.0f; + const float r = (a % b); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.msl new file mode 100644 index 0000000000..4a82529c71 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + float a = 1.0f; + float b = 0.0f; + float const r = fmod(a, b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..14b6b443af --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.spvasm @@ -0,0 +1,30 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 15 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %float_1 = OpConstant %float 1 +%_ptr_Function_float = OpTypePointer Function %float + %9 = OpConstantNull %float + %float_0 = OpConstant %float 0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_float Function %9 + %b = OpVariable %_ptr_Function_float Function %9 + OpStore %a %float_1 + OpStore %b %float_0 + %12 = OpLoad %float %a + %13 = OpLoad %float %b + %14 = OpFRem %float %12 %13 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..1b3af01b94 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1.0; + var b = 0.0; + let r : f32 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl new file mode 100644 index 0000000000..38874c47de --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1; + var b = 0; + let r : i32 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..9f27720ec7 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + int a = 1; + int b = 0; + const int r = (a % (b == 0 ? 1 : b)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.msl new file mode 100644 index 0000000000..4a51c544cc --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int a = 1; + int b = 0; + int const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..f942f0b9bb --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.spvasm @@ -0,0 +1,30 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 15 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %int_1 = OpConstant %int 1 +%_ptr_Function_int = OpTypePointer Function %int + %9 = OpConstantNull %int + %int_0 = OpConstant %int 0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_int Function %9 + %b = OpVariable %_ptr_Function_int Function %9 + OpStore %a %int_1 + OpStore %b %int_0 + %12 = OpLoad %int %a + %13 = OpLoad %int %b + %14 = OpSMod %int %12 %13 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..1b19397f13 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1; + var b = 0; + let r : i32 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl new file mode 100644 index 0000000000..a8b6555226 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1u; + var b = 0u; + let r : u32 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..cc3d3c11d4 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + uint a = 1u; + uint b = 0u; + const uint r = (a % (b == 0u ? 1u : b)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.msl new file mode 100644 index 0000000000..8849d0d91f --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint a = 1u; + uint b = 0u; + uint const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..8d29bb63e8 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.spvasm @@ -0,0 +1,30 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 15 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 +%_ptr_Function_uint = OpTypePointer Function %uint + %9 = OpConstantNull %uint + %uint_0 = OpConstant %uint 0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_uint Function %9 + %b = OpVariable %_ptr_Function_uint Function %9 + OpStore %a %uint_1 + OpStore %b %uint_0 + %12 = OpLoad %uint %a + %13 = OpLoad %uint %b + %14 = OpUMod %uint %12 %13 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..1a5f7266da --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 1u; + var b = 0u; + let r : u32 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl new file mode 100644 index 0000000000..0f9c970891 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 4; + var b = vec3(0, 2, 0); + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..0b14c77017 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + int a = 4; + int3 b = int3(0, 2, 0); + const int3 r = (a % (b == int3(0, 0, 0) ? int3(1, 1, 1) : b)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.msl new file mode 100644 index 0000000000..ae32a8616b --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int a = 4; + int3 b = int3(0, 2, 0); + int3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..1098fb21ae --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.spvasm @@ -0,0 +1,37 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 22 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %int_4 = OpConstant %int 4 +%_ptr_Function_int = OpTypePointer Function %int + %9 = OpConstantNull %int + %v3int = OpTypeVector %int 3 + %int_0 = OpConstant %int 0 + %int_2 = OpConstant %int 2 + %13 = OpConstantComposite %v3int %int_0 %int_2 %int_0 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %16 = OpConstantNull %v3int + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_int Function %9 + %b = OpVariable %_ptr_Function_v3int Function %16 + %20 = OpVariable %_ptr_Function_v3int Function %16 + OpStore %a %int_4 + OpStore %b %13 + %17 = OpLoad %int %a + %18 = OpLoad %v3int %b + %21 = OpCompositeConstruct %v3int %17 %17 %17 + %19 = OpSMod %v3int %21 %18 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..30f19f4828 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 4; + var b = vec3(0, 2, 0); + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl new file mode 100644 index 0000000000..9afcb3c86a --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 4u; + var b = vec3(0u, 2u, 0u); + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..fa88822f21 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + uint a = 4u; + uint3 b = uint3(0u, 2u, 0u); + const uint3 r = (a % (b == uint3(0u, 0u, 0u) ? uint3(1u, 1u, 1u) : b)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.msl new file mode 100644 index 0000000000..bc4bada4be --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint a = 4u; + uint3 b = uint3(0u, 2u, 0u); + uint3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..45bce24034 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.spvasm @@ -0,0 +1,37 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 22 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_4 = OpConstant %uint 4 +%_ptr_Function_uint = OpTypePointer Function %uint + %9 = OpConstantNull %uint + %v3uint = OpTypeVector %uint 3 + %uint_0 = OpConstant %uint 0 + %uint_2 = OpConstant %uint 2 + %13 = OpConstantComposite %v3uint %uint_0 %uint_2 %uint_0 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %16 = OpConstantNull %v3uint + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_uint Function %9 + %b = OpVariable %_ptr_Function_v3uint Function %16 + %20 = OpVariable %_ptr_Function_v3uint Function %16 + OpStore %a %uint_4 + OpStore %b %13 + %17 = OpLoad %uint %a + %18 = OpLoad %v3uint %b + %21 = OpCompositeConstruct %v3uint %17 %17 %17 + %19 = OpUMod %v3uint %21 %18 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..0e406f7bb7 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = 4u; + var b = vec3(0u, 2u, 0u); + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl new file mode 100644 index 0000000000..41f8c7caf7 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1, 2, 3); + var b = 0; + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..208637ee75 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + int3 a = int3(1, 2, 3); + int b = 0; + const int3 r = (a % (b == 0 ? 1 : b)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.msl new file mode 100644 index 0000000000..39d00f808b --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int3 a = int3(1, 2, 3); + int b = 0; + int3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..f25b33386b --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.spvasm @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 + %int_1 = OpConstant %int 1 + %int_2 = OpConstant %int 2 + %int_3 = OpConstant %int 3 + %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %13 = OpConstantNull %v3int + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %17 = OpConstantNull %int + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_v3int Function %13 + %b = OpVariable %_ptr_Function_int Function %17 + %21 = OpVariable %_ptr_Function_v3int Function %13 + OpStore %a %10 + OpStore %b %int_0 + %18 = OpLoad %v3int %a + %19 = OpLoad %int %b + %22 = OpCompositeConstruct %v3int %19 %19 %19 + %20 = OpSMod %v3int %18 %22 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..dcb69395a7 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1, 2, 3); + var b = 0; + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl new file mode 100644 index 0000000000..5ddae5bdc1 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1u, 2u, 3u); + var b = 0u; + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..0cb79cfea2 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + uint3 a = uint3(1u, 2u, 3u); + uint b = 0u; + const uint3 r = (a % (b == 0u ? 1u : b)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.msl new file mode 100644 index 0000000000..3c8b16e08e --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint3 a = uint3(1u, 2u, 3u); + uint b = 0u; + uint3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..0965b4a729 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.spvasm @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %uint_2 = OpConstant %uint 2 + %uint_3 = OpConstant %uint 3 + %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %13 = OpConstantNull %v3uint + %uint_0 = OpConstant %uint 0 +%_ptr_Function_uint = OpTypePointer Function %uint + %17 = OpConstantNull %uint + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_v3uint Function %13 + %b = OpVariable %_ptr_Function_uint Function %17 + %21 = OpVariable %_ptr_Function_v3uint Function %13 + OpStore %a %10 + OpStore %b %uint_0 + %18 = OpLoad %v3uint %a + %19 = OpLoad %uint %b + %22 = OpCompositeConstruct %v3uint %19 %19 %19 + %20 = OpUMod %v3uint %18 %22 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..eabfbcd88c --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1u, 2u, 3u); + var b = 0u; + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl new file mode 100644 index 0000000000..7a91bd08a5 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1., 2., 3.); + var b = vec3(0., 5., 0.); + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.hlsl new file mode 100644 index 0000000000..636719b6d6 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + float3 a = float3(1.0f, 2.0f, 3.0f); + float3 b = float3(0.0f, 5.0f, 0.0f); + const float3 r = (a % b); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.msl new file mode 100644 index 0000000000..4ebdd54b3d --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + float3 a = float3(1.0f, 2.0f, 3.0f); + float3 b = float3(0.0f, 5.0f, 0.0f); + float3 const r = fmod(a, b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.spvasm new file mode 100644 index 0000000000..1cdd95a931 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.spvasm @@ -0,0 +1,36 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 21 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v3float = OpTypeVector %float 3 + %float_1 = OpConstant %float 1 + %float_2 = OpConstant %float 2 + %float_3 = OpConstant %float 3 + %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3 +%_ptr_Function_v3float = OpTypePointer Function %v3float + %13 = OpConstantNull %v3float + %float_0 = OpConstant %float 0 + %float_5 = OpConstant %float 5 + %16 = OpConstantComposite %v3float %float_0 %float_5 %float_0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_v3float Function %13 + %b = OpVariable %_ptr_Function_v3float Function %13 + OpStore %a %10 + OpStore %b %16 + %18 = OpLoad %v3float %a + %19 = OpLoad %v3float %b + %20 = OpFRem %v3float %18 %19 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.wgsl new file mode 100644 index 0000000000..6d96f8922d --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1.0, 2.0, 3.0); + var b = vec3(0.0, 5.0, 0.0); + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl new file mode 100644 index 0000000000..e9364db7eb --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1, 2, 3); + var b = vec3(0, 5, 0); + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.hlsl new file mode 100644 index 0000000000..97e59e6254 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + int3 a = int3(1, 2, 3); + int3 b = int3(0, 5, 0); + const int3 r = (a % (b == int3(0, 0, 0) ? int3(1, 1, 1) : b)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.msl new file mode 100644 index 0000000000..e6b07ad968 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + int3 a = int3(1, 2, 3); + int3 b = int3(0, 5, 0); + int3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.spvasm new file mode 100644 index 0000000000..29c04a65b9 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.spvasm @@ -0,0 +1,36 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 21 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 + %int_1 = OpConstant %int 1 + %int_2 = OpConstant %int 2 + %int_3 = OpConstant %int 3 + %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3 +%_ptr_Function_v3int = OpTypePointer Function %v3int + %13 = OpConstantNull %v3int + %int_0 = OpConstant %int 0 + %int_5 = OpConstant %int 5 + %16 = OpConstantComposite %v3int %int_0 %int_5 %int_0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_v3int Function %13 + %b = OpVariable %_ptr_Function_v3int Function %13 + OpStore %a %10 + OpStore %b %16 + %18 = OpLoad %v3int %a + %19 = OpLoad %v3int %b + %20 = OpSMod %v3int %18 %19 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.wgsl new file mode 100644 index 0000000000..ebd6b91381 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1, 2, 3); + var b = vec3(0, 5, 0); + let r : vec3 = (a % b); +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl new file mode 100644 index 0000000000..cf77b25e4f --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1u, 2u, 3u); + var b = vec3(0u, 5u, 0u); + let r : vec3 = a % b; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.hlsl new file mode 100644 index 0000000000..b041731b55 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.hlsl @@ -0,0 +1,7 @@ +[numthreads(1, 1, 1)] +void f() { + uint3 a = uint3(1u, 2u, 3u); + uint3 b = uint3(0u, 5u, 0u); + const uint3 r = (a % (b == uint3(0u, 0u, 0u) ? uint3(1u, 1u, 1u) : b)); + return; +} diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.msl new file mode 100644 index 0000000000..46e36b05b4 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + uint3 a = uint3(1u, 2u, 3u); + uint3 b = uint3(0u, 5u, 0u); + uint3 const r = (a % b); + return; +} + diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.spvasm new file mode 100644 index 0000000000..3fbdc22a31 --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.spvasm @@ -0,0 +1,36 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 21 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %a "a" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %uint_2 = OpConstant %uint 2 + %uint_3 = OpConstant %uint 3 + %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %13 = OpConstantNull %v3uint + %uint_0 = OpConstant %uint 0 + %uint_5 = OpConstant %uint 5 + %16 = OpConstantComposite %v3uint %uint_0 %uint_5 %uint_0 + %f = OpFunction %void None %1 + %4 = OpLabel + %a = OpVariable %_ptr_Function_v3uint Function %13 + %b = OpVariable %_ptr_Function_v3uint Function %13 + OpStore %a %10 + OpStore %b %16 + %18 = OpLoad %v3uint %a + %19 = OpLoad %v3uint %b + %20 = OpUMod %v3uint %18 %19 + OpReturn + OpFunctionEnd diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.wgsl new file mode 100644 index 0000000000..615b56c89c --- /dev/null +++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var a = vec3(1u, 2u, 3u); + var b = vec3(0u, 5u, 0u); + let r : vec3 = (a % b); +} diff --git a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.spvasm.expected.hlsl index 535c006879..0d695f6558 100644 --- a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.spvasm.expected.hlsl +++ b/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.spvasm.expected.hlsl @@ -1,5 +1,3 @@ -SKIP: FAILED - cbuffer cbuffer_x_6 : register(b0, space0) { uint4 x_6[4]; }; @@ -18,7 +16,7 @@ void main_1() { if ((x_35 == x_37)) { a = (a + 1); } else { - a = (a / i); + a = (a / (i == 0 ? 1 : i)); } } } @@ -59,6 +57,3 @@ tint_symbol main() { wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; return wrapper_result; } -C:\src\tint\test\Shader@0x00000150DB0FB4A0(19,14-18): error X4010: Unsigned integer divide by zero -C:\src\tint\test\Shader@0x00000150DB0FB4A0(19,14-18): warning X3556: integer divides may be much slower, try using uints if possible. - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.wgsl.expected.hlsl index c8f9a0771f..0d695f6558 100644 --- a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.wgsl.expected.hlsl +++ b/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.wgsl.expected.hlsl @@ -1,5 +1,3 @@ -SKIP: FAILED - cbuffer cbuffer_x_6 : register(b0, space0) { uint4 x_6[4]; }; @@ -18,7 +16,7 @@ void main_1() { if ((x_35 == x_37)) { a = (a + 1); } else { - a = (a / i); + a = (a / (i == 0 ? 1 : i)); } } } @@ -59,6 +57,3 @@ tint_symbol main() { wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; return wrapper_result; } -C:\src\tint\test\Shader@0x000001A53C9E4360(19,14-18): error X4010: Unsigned integer divide by zero -C:\src\tint\test\Shader@0x000001A53C9E4360(19,14-18): warning X3556: integer divides may be much slower, try using uints if possible. - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.spvasm.expected.hlsl index 15894664f5..8455574fea 100644 --- a/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.spvasm.expected.hlsl +++ b/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.spvasm.expected.hlsl @@ -1,5 +1,3 @@ -SKIP: FAILED - cbuffer cbuffer_x_8 : register(b2, space0) { uint4 x_8[2]; }; @@ -24,7 +22,7 @@ void main_1() { if ((x_43 < x_45)) { const uint scalar_offset_1 = ((16u * uint(0))) / 4; const uint x_50 = x_12[scalar_offset_1 / 4][scalar_offset_1 % 4]; - b = asint((x_50 % a)); + b = asint((x_50 % (a == 0u ? 1u : a))); } const int x_54 = b; const int x_56 = asint(x_8[1].x); @@ -68,5 +66,3 @@ tint_symbol_2 main(tint_symbol_1 tint_symbol) { wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; return wrapper_result; } -C:\src\tint\test\Shader@0x000002533EADD8C0(25,16-23): error X4010: Unsigned integer divide by zero - diff --git a/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.wgsl.expected.hlsl index 4cecdf3c03..8455574fea 100644 --- a/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.wgsl.expected.hlsl +++ b/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.wgsl.expected.hlsl @@ -1,5 +1,3 @@ -SKIP: FAILED - cbuffer cbuffer_x_8 : register(b2, space0) { uint4 x_8[2]; }; @@ -24,7 +22,7 @@ void main_1() { if ((x_43 < x_45)) { const uint scalar_offset_1 = ((16u * uint(0))) / 4; const uint x_50 = x_12[scalar_offset_1 / 4][scalar_offset_1 % 4]; - b = asint((x_50 % a)); + b = asint((x_50 % (a == 0u ? 1u : a))); } const int x_54 = b; const int x_56 = asint(x_8[1].x); @@ -68,5 +66,3 @@ tint_symbol_2 main(tint_symbol_1 tint_symbol) { wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; return wrapper_result; } -C:\src\tint\test\Shader@0x0000026AFDC08170(25,16-23): error X4010: Unsigned integer divide by zero -