From 4ea7dc33816895da76a96e30d0a940339ea747be Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Tue, 12 Apr 2022 15:53:55 +0200 Subject: [PATCH] [spv-out] add support for matrix add and sub --- src/back/spv/block.rs | 98 +++++ tests/in/operators.wgsl | 4 +- tests/out/glsl/operators.main.Compute.glsl | 132 +++---- tests/out/hlsl/operators.hlsl | 132 +++---- tests/out/msl/operators.msl | 132 +++---- tests/out/spv/operators.spvasm | 404 +++++++++++---------- tests/out/wgsl/operators.wgsl | 132 +++---- 7 files changed, 580 insertions(+), 454 deletions(-) diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index 15bd07bea0..b22eb56aea 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -346,6 +346,26 @@ impl<'w> BlockContext<'w> { crate::ScalarKind::Float => spirv::Op::FAdd, _ => spirv::Op::IAdd, }, + crate::TypeInner::Matrix { + columns, + rows, + width, + } => { + self.write_matrix_matrix_column_op( + block, + id, + result_type_id, + left_id, + right_id, + columns, + rows, + width, + spirv::Op::FAdd, + ); + + self.cached[expr_handle] = id; + return Ok(()); + } _ => unimplemented!(), }, crate::BinaryOperator::Subtract => match *left_ty_inner { @@ -354,6 +374,26 @@ impl<'w> BlockContext<'w> { crate::ScalarKind::Float => spirv::Op::FSub, _ => spirv::Op::ISub, }, + crate::TypeInner::Matrix { + columns, + rows, + width, + } => { + self.write_matrix_matrix_column_op( + block, + id, + result_type_id, + left_id, + right_id, + columns, + rows, + width, + spirv::Op::FSub, + ); + + self.cached[expr_handle] = id; + return Ok(()); + } _ => unimplemented!(), }, crate::BinaryOperator::Multiply => match (left_dimension, right_dimension) { @@ -1150,6 +1190,64 @@ impl<'w> BlockContext<'w> { Ok(pointer) } + /// Build the instructions for matrix - matrix column operations + #[allow(clippy::too_many_arguments)] + fn write_matrix_matrix_column_op( + &mut self, + block: &mut Block, + result_id: Word, + result_type_id: Word, + left_id: Word, + right_id: Word, + columns: crate::VectorSize, + rows: crate::VectorSize, + width: u8, + op: spirv::Op, + ) { + self.temp_list.clear(); + + let vector_type_id = self.get_type_id(LookupType::Local(LocalType::Value { + vector_size: Some(rows), + kind: crate::ScalarKind::Float, + width, + pointer_space: None, + })); + + for index in 0..columns as u32 { + let column_id_left = self.gen_id(); + let column_id_right = self.gen_id(); + let column_id_res = self.gen_id(); + + block.body.push(Instruction::composite_extract( + vector_type_id, + column_id_left, + left_id, + &[index], + )); + block.body.push(Instruction::composite_extract( + vector_type_id, + column_id_right, + right_id, + &[index], + )); + block.body.push(Instruction::binary( + op, + vector_type_id, + column_id_res, + column_id_left, + column_id_right, + )); + + self.temp_list.push(column_id_res); + } + + block.body.push(Instruction::composite_construct( + result_type_id, + result_id, + &self.temp_list, + )); + } + /// Build the instructions for the arithmetic expression of a dot product fn write_dot_product( &mut self, diff --git a/tests/in/operators.wgsl b/tests/in/operators.wgsl index 31ca1f7405..205895dd45 100644 --- a/tests/in/operators.wgsl +++ b/tests/in/operators.wgsl @@ -150,8 +150,8 @@ fn arithmetic() { let _ = 2 % vec2(1); // Matrix arithmetic - // let _ = mat3x3() + mat3x3(); - // let _ = mat3x3() - mat3x3(); + let _ = mat3x3() + mat3x3(); + let _ = mat3x3() - mat3x3(); let _ = mat3x3() * 1.0; let _ = 2.0 * mat3x3(); diff --git a/tests/out/glsl/operators.main.Compute.glsl b/tests/out/glsl/operators.main.Compute.glsl index 2aeb893a97..dc741b1ebd 100644 --- a/tests/out/glsl/operators.main.Compute.glsl +++ b/tests/out/glsl/operators.main.Compute.glsl @@ -110,77 +110,79 @@ void arithmetic() { ivec2 unnamed_58 = (ivec2(2) / ivec2(1)); ivec2 unnamed_59 = (ivec2(2) % ivec2(1)); ivec2 unnamed_60 = (ivec2(2) % ivec2(1)); - mat3x3 unnamed_61 = (mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * 1.0); - mat3x3 unnamed_62 = (2.0 * mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); - vec3 unnamed_63 = (mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * vec4(1.0)); - vec4 unnamed_64 = (vec3(2.0) * mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); - mat3x3 unnamed_65 = (mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * mat3x4(vec4(0.0, 0.0, 0.0, 0.0), vec4(0.0, 0.0, 0.0, 0.0), vec4(0.0, 0.0, 0.0, 0.0))); + mat3x3 unnamed_61 = (mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) + mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); + mat3x3 unnamed_62 = (mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) - mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); + mat3x3 unnamed_63 = (mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * 1.0); + mat3x3 unnamed_64 = (2.0 * mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); + vec3 unnamed_65 = (mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * vec4(1.0)); + vec4 unnamed_66 = (vec3(2.0) * mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); + mat3x3 unnamed_67 = (mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * mat3x4(vec4(0.0, 0.0, 0.0, 0.0), vec4(0.0, 0.0, 0.0, 0.0), vec4(0.0, 0.0, 0.0, 0.0))); } void bit() { - int unnamed_66 = (~ 1); - uint unnamed_67 = (~ 1u); - ivec2 unnamed_68 = (~ ivec2(1)); - uvec3 unnamed_69 = (~ uvec3(1u)); - int unnamed_70 = (2 | 1); - uint unnamed_71 = (2u | 1u); - ivec2 unnamed_72 = (ivec2(2) | ivec2(1)); - uvec3 unnamed_73 = (uvec3(2u) | uvec3(1u)); - int unnamed_74 = (2 & 1); - uint unnamed_75 = (2u & 1u); - ivec2 unnamed_76 = (ivec2(2) & ivec2(1)); - uvec3 unnamed_77 = (uvec3(2u) & uvec3(1u)); - int unnamed_78 = (2 ^ 1); - uint unnamed_79 = (2u ^ 1u); - ivec2 unnamed_80 = (ivec2(2) ^ ivec2(1)); - uvec3 unnamed_81 = (uvec3(2u) ^ uvec3(1u)); - int unnamed_82 = (2 << 1u); - uint unnamed_83 = (2u << 1u); - ivec2 unnamed_84 = (ivec2(2) << uvec2(1u)); - uvec3 unnamed_85 = (uvec3(2u) << uvec3(1u)); - int unnamed_86 = (2 >> 1u); - uint unnamed_87 = (2u >> 1u); - ivec2 unnamed_88 = (ivec2(2) >> uvec2(1u)); - uvec3 unnamed_89 = (uvec3(2u) >> uvec3(1u)); + int unnamed_68 = (~ 1); + uint unnamed_69 = (~ 1u); + ivec2 unnamed_70 = (~ ivec2(1)); + uvec3 unnamed_71 = (~ uvec3(1u)); + int unnamed_72 = (2 | 1); + uint unnamed_73 = (2u | 1u); + ivec2 unnamed_74 = (ivec2(2) | ivec2(1)); + uvec3 unnamed_75 = (uvec3(2u) | uvec3(1u)); + int unnamed_76 = (2 & 1); + uint unnamed_77 = (2u & 1u); + ivec2 unnamed_78 = (ivec2(2) & ivec2(1)); + uvec3 unnamed_79 = (uvec3(2u) & uvec3(1u)); + int unnamed_80 = (2 ^ 1); + uint unnamed_81 = (2u ^ 1u); + ivec2 unnamed_82 = (ivec2(2) ^ ivec2(1)); + uvec3 unnamed_83 = (uvec3(2u) ^ uvec3(1u)); + int unnamed_84 = (2 << 1u); + uint unnamed_85 = (2u << 1u); + ivec2 unnamed_86 = (ivec2(2) << uvec2(1u)); + uvec3 unnamed_87 = (uvec3(2u) << uvec3(1u)); + int unnamed_88 = (2 >> 1u); + uint unnamed_89 = (2u >> 1u); + ivec2 unnamed_90 = (ivec2(2) >> uvec2(1u)); + uvec3 unnamed_91 = (uvec3(2u) >> uvec3(1u)); } void comparison() { - bool unnamed_90 = (2 == 1); - bool unnamed_91 = (2u == 1u); - bool unnamed_92 = (2.0 == 1.0); - bvec2 unnamed_93 = equal(ivec2(2), ivec2(1)); - bvec3 unnamed_94 = equal(uvec3(2u), uvec3(1u)); - bvec4 unnamed_95 = equal(vec4(2.0), vec4(1.0)); - bool unnamed_96 = (2 != 1); - bool unnamed_97 = (2u != 1u); - bool unnamed_98 = (2.0 != 1.0); - bvec2 unnamed_99 = notEqual(ivec2(2), ivec2(1)); - bvec3 unnamed_100 = notEqual(uvec3(2u), uvec3(1u)); - bvec4 unnamed_101 = notEqual(vec4(2.0), vec4(1.0)); - bool unnamed_102 = (2 < 1); - bool unnamed_103 = (2u < 1u); - bool unnamed_104 = (2.0 < 1.0); - bvec2 unnamed_105 = lessThan(ivec2(2), ivec2(1)); - bvec3 unnamed_106 = lessThan(uvec3(2u), uvec3(1u)); - bvec4 unnamed_107 = lessThan(vec4(2.0), vec4(1.0)); - bool unnamed_108 = (2 <= 1); - bool unnamed_109 = (2u <= 1u); - bool unnamed_110 = (2.0 <= 1.0); - bvec2 unnamed_111 = lessThanEqual(ivec2(2), ivec2(1)); - bvec3 unnamed_112 = lessThanEqual(uvec3(2u), uvec3(1u)); - bvec4 unnamed_113 = lessThanEqual(vec4(2.0), vec4(1.0)); - bool unnamed_114 = (2 > 1); - bool unnamed_115 = (2u > 1u); - bool unnamed_116 = (2.0 > 1.0); - bvec2 unnamed_117 = greaterThan(ivec2(2), ivec2(1)); - bvec3 unnamed_118 = greaterThan(uvec3(2u), uvec3(1u)); - bvec4 unnamed_119 = greaterThan(vec4(2.0), vec4(1.0)); - bool unnamed_120 = (2 >= 1); - bool unnamed_121 = (2u >= 1u); - bool unnamed_122 = (2.0 >= 1.0); - bvec2 unnamed_123 = greaterThanEqual(ivec2(2), ivec2(1)); - bvec3 unnamed_124 = greaterThanEqual(uvec3(2u), uvec3(1u)); - bvec4 unnamed_125 = greaterThanEqual(vec4(2.0), vec4(1.0)); + bool unnamed_92 = (2 == 1); + bool unnamed_93 = (2u == 1u); + bool unnamed_94 = (2.0 == 1.0); + bvec2 unnamed_95 = equal(ivec2(2), ivec2(1)); + bvec3 unnamed_96 = equal(uvec3(2u), uvec3(1u)); + bvec4 unnamed_97 = equal(vec4(2.0), vec4(1.0)); + bool unnamed_98 = (2 != 1); + bool unnamed_99 = (2u != 1u); + bool unnamed_100 = (2.0 != 1.0); + bvec2 unnamed_101 = notEqual(ivec2(2), ivec2(1)); + bvec3 unnamed_102 = notEqual(uvec3(2u), uvec3(1u)); + bvec4 unnamed_103 = notEqual(vec4(2.0), vec4(1.0)); + bool unnamed_104 = (2 < 1); + bool unnamed_105 = (2u < 1u); + bool unnamed_106 = (2.0 < 1.0); + bvec2 unnamed_107 = lessThan(ivec2(2), ivec2(1)); + bvec3 unnamed_108 = lessThan(uvec3(2u), uvec3(1u)); + bvec4 unnamed_109 = lessThan(vec4(2.0), vec4(1.0)); + bool unnamed_110 = (2 <= 1); + bool unnamed_111 = (2u <= 1u); + bool unnamed_112 = (2.0 <= 1.0); + bvec2 unnamed_113 = lessThanEqual(ivec2(2), ivec2(1)); + bvec3 unnamed_114 = lessThanEqual(uvec3(2u), uvec3(1u)); + bvec4 unnamed_115 = lessThanEqual(vec4(2.0), vec4(1.0)); + bool unnamed_116 = (2 > 1); + bool unnamed_117 = (2u > 1u); + bool unnamed_118 = (2.0 > 1.0); + bvec2 unnamed_119 = greaterThan(ivec2(2), ivec2(1)); + bvec3 unnamed_120 = greaterThan(uvec3(2u), uvec3(1u)); + bvec4 unnamed_121 = greaterThan(vec4(2.0), vec4(1.0)); + bool unnamed_122 = (2 >= 1); + bool unnamed_123 = (2u >= 1u); + bool unnamed_124 = (2.0 >= 1.0); + bvec2 unnamed_125 = greaterThanEqual(ivec2(2), ivec2(1)); + bvec3 unnamed_126 = greaterThanEqual(uvec3(2u), uvec3(1u)); + bvec4 unnamed_127 = greaterThanEqual(vec4(2.0), vec4(1.0)); } void assignment() { diff --git a/tests/out/hlsl/operators.hlsl b/tests/out/hlsl/operators.hlsl index 5275cbe3b9..e640cfaff4 100644 --- a/tests/out/hlsl/operators.hlsl +++ b/tests/out/hlsl/operators.hlsl @@ -138,79 +138,81 @@ void arithmetic() int2 unnamed_58 = (int2(2.xx) / int2(1.xx)); int2 unnamed_59 = (int2(2.xx) % int2(1.xx)); int2 unnamed_60 = (int2(2.xx) % int2(1.xx)); - float3x3 unnamed_61 = mul(1.0, float3x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0))); - float3x3 unnamed_62 = mul(float3x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0)), 2.0); - float3 unnamed_63 = mul(float4(1.0.xxxx), float4x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0))); - float4 unnamed_64 = mul(float4x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0)), float3(2.0.xxx)); - float3x3 unnamed_65 = mul(float3x4(float4(0.0, 0.0, 0.0, 0.0), float4(0.0, 0.0, 0.0, 0.0), float4(0.0, 0.0, 0.0, 0.0)), float4x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0))); + float3x3 unnamed_61 = (float3x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0)) + float3x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0))); + float3x3 unnamed_62 = (float3x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0)) - float3x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0))); + float3x3 unnamed_63 = mul(1.0, float3x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0))); + float3x3 unnamed_64 = mul(float3x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0)), 2.0); + float3 unnamed_65 = mul(float4(1.0.xxxx), float4x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0))); + float4 unnamed_66 = mul(float4x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0)), float3(2.0.xxx)); + float3x3 unnamed_67 = mul(float3x4(float4(0.0, 0.0, 0.0, 0.0), float4(0.0, 0.0, 0.0, 0.0), float4(0.0, 0.0, 0.0, 0.0)), float4x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0))); } void bit() { - int unnamed_66 = ~1; - uint unnamed_67 = ~1u; - int2 unnamed_68 = ~int2(1.xx); - uint3 unnamed_69 = ~uint3(1u.xxx); - int unnamed_70 = (2 | 1); - uint unnamed_71 = (2u | 1u); - int2 unnamed_72 = (int2(2.xx) | int2(1.xx)); - uint3 unnamed_73 = (uint3(2u.xxx) | uint3(1u.xxx)); - int unnamed_74 = (2 & 1); - uint unnamed_75 = (2u & 1u); - int2 unnamed_76 = (int2(2.xx) & int2(1.xx)); - uint3 unnamed_77 = (uint3(2u.xxx) & uint3(1u.xxx)); - int unnamed_78 = (2 ^ 1); - uint unnamed_79 = (2u ^ 1u); - int2 unnamed_80 = (int2(2.xx) ^ int2(1.xx)); - uint3 unnamed_81 = (uint3(2u.xxx) ^ uint3(1u.xxx)); - int unnamed_82 = (2 << 1u); - uint unnamed_83 = (2u << 1u); - int2 unnamed_84 = (int2(2.xx) << uint2(1u.xx)); - uint3 unnamed_85 = (uint3(2u.xxx) << uint3(1u.xxx)); - int unnamed_86 = (2 >> 1u); - uint unnamed_87 = (2u >> 1u); - int2 unnamed_88 = (int2(2.xx) >> uint2(1u.xx)); - uint3 unnamed_89 = (uint3(2u.xxx) >> uint3(1u.xxx)); + int unnamed_68 = ~1; + uint unnamed_69 = ~1u; + int2 unnamed_70 = ~int2(1.xx); + uint3 unnamed_71 = ~uint3(1u.xxx); + int unnamed_72 = (2 | 1); + uint unnamed_73 = (2u | 1u); + int2 unnamed_74 = (int2(2.xx) | int2(1.xx)); + uint3 unnamed_75 = (uint3(2u.xxx) | uint3(1u.xxx)); + int unnamed_76 = (2 & 1); + uint unnamed_77 = (2u & 1u); + int2 unnamed_78 = (int2(2.xx) & int2(1.xx)); + uint3 unnamed_79 = (uint3(2u.xxx) & uint3(1u.xxx)); + int unnamed_80 = (2 ^ 1); + uint unnamed_81 = (2u ^ 1u); + int2 unnamed_82 = (int2(2.xx) ^ int2(1.xx)); + uint3 unnamed_83 = (uint3(2u.xxx) ^ uint3(1u.xxx)); + int unnamed_84 = (2 << 1u); + uint unnamed_85 = (2u << 1u); + int2 unnamed_86 = (int2(2.xx) << uint2(1u.xx)); + uint3 unnamed_87 = (uint3(2u.xxx) << uint3(1u.xxx)); + int unnamed_88 = (2 >> 1u); + uint unnamed_89 = (2u >> 1u); + int2 unnamed_90 = (int2(2.xx) >> uint2(1u.xx)); + uint3 unnamed_91 = (uint3(2u.xxx) >> uint3(1u.xxx)); } void comparison() { - bool unnamed_90 = (2 == 1); - bool unnamed_91 = (2u == 1u); - bool unnamed_92 = (2.0 == 1.0); - bool2 unnamed_93 = (int2(2.xx) == int2(1.xx)); - bool3 unnamed_94 = (uint3(2u.xxx) == uint3(1u.xxx)); - bool4 unnamed_95 = (float4(2.0.xxxx) == float4(1.0.xxxx)); - bool unnamed_96 = (2 != 1); - bool unnamed_97 = (2u != 1u); - bool unnamed_98 = (2.0 != 1.0); - bool2 unnamed_99 = (int2(2.xx) != int2(1.xx)); - bool3 unnamed_100 = (uint3(2u.xxx) != uint3(1u.xxx)); - bool4 unnamed_101 = (float4(2.0.xxxx) != float4(1.0.xxxx)); - bool unnamed_102 = (2 < 1); - bool unnamed_103 = (2u < 1u); - bool unnamed_104 = (2.0 < 1.0); - bool2 unnamed_105 = (int2(2.xx) < int2(1.xx)); - bool3 unnamed_106 = (uint3(2u.xxx) < uint3(1u.xxx)); - bool4 unnamed_107 = (float4(2.0.xxxx) < float4(1.0.xxxx)); - bool unnamed_108 = (2 <= 1); - bool unnamed_109 = (2u <= 1u); - bool unnamed_110 = (2.0 <= 1.0); - bool2 unnamed_111 = (int2(2.xx) <= int2(1.xx)); - bool3 unnamed_112 = (uint3(2u.xxx) <= uint3(1u.xxx)); - bool4 unnamed_113 = (float4(2.0.xxxx) <= float4(1.0.xxxx)); - bool unnamed_114 = (2 > 1); - bool unnamed_115 = (2u > 1u); - bool unnamed_116 = (2.0 > 1.0); - bool2 unnamed_117 = (int2(2.xx) > int2(1.xx)); - bool3 unnamed_118 = (uint3(2u.xxx) > uint3(1u.xxx)); - bool4 unnamed_119 = (float4(2.0.xxxx) > float4(1.0.xxxx)); - bool unnamed_120 = (2 >= 1); - bool unnamed_121 = (2u >= 1u); - bool unnamed_122 = (2.0 >= 1.0); - bool2 unnamed_123 = (int2(2.xx) >= int2(1.xx)); - bool3 unnamed_124 = (uint3(2u.xxx) >= uint3(1u.xxx)); - bool4 unnamed_125 = (float4(2.0.xxxx) >= float4(1.0.xxxx)); + bool unnamed_92 = (2 == 1); + bool unnamed_93 = (2u == 1u); + bool unnamed_94 = (2.0 == 1.0); + bool2 unnamed_95 = (int2(2.xx) == int2(1.xx)); + bool3 unnamed_96 = (uint3(2u.xxx) == uint3(1u.xxx)); + bool4 unnamed_97 = (float4(2.0.xxxx) == float4(1.0.xxxx)); + bool unnamed_98 = (2 != 1); + bool unnamed_99 = (2u != 1u); + bool unnamed_100 = (2.0 != 1.0); + bool2 unnamed_101 = (int2(2.xx) != int2(1.xx)); + bool3 unnamed_102 = (uint3(2u.xxx) != uint3(1u.xxx)); + bool4 unnamed_103 = (float4(2.0.xxxx) != float4(1.0.xxxx)); + bool unnamed_104 = (2 < 1); + bool unnamed_105 = (2u < 1u); + bool unnamed_106 = (2.0 < 1.0); + bool2 unnamed_107 = (int2(2.xx) < int2(1.xx)); + bool3 unnamed_108 = (uint3(2u.xxx) < uint3(1u.xxx)); + bool4 unnamed_109 = (float4(2.0.xxxx) < float4(1.0.xxxx)); + bool unnamed_110 = (2 <= 1); + bool unnamed_111 = (2u <= 1u); + bool unnamed_112 = (2.0 <= 1.0); + bool2 unnamed_113 = (int2(2.xx) <= int2(1.xx)); + bool3 unnamed_114 = (uint3(2u.xxx) <= uint3(1u.xxx)); + bool4 unnamed_115 = (float4(2.0.xxxx) <= float4(1.0.xxxx)); + bool unnamed_116 = (2 > 1); + bool unnamed_117 = (2u > 1u); + bool unnamed_118 = (2.0 > 1.0); + bool2 unnamed_119 = (int2(2.xx) > int2(1.xx)); + bool3 unnamed_120 = (uint3(2u.xxx) > uint3(1u.xxx)); + bool4 unnamed_121 = (float4(2.0.xxxx) > float4(1.0.xxxx)); + bool unnamed_122 = (2 >= 1); + bool unnamed_123 = (2u >= 1u); + bool unnamed_124 = (2.0 >= 1.0); + bool2 unnamed_125 = (int2(2.xx) >= int2(1.xx)); + bool3 unnamed_126 = (uint3(2u.xxx) >= uint3(1u.xxx)); + bool4 unnamed_127 = (float4(2.0.xxxx) >= float4(1.0.xxxx)); } void assignment() diff --git a/tests/out/msl/operators.msl b/tests/out/msl/operators.msl index 87a10ee44d..573bad1821 100644 --- a/tests/out/msl/operators.msl +++ b/tests/out/msl/operators.msl @@ -136,79 +136,81 @@ void arithmetic( metal::int2 unnamed_58 = metal::int2(2) / metal::int2(1); metal::int2 unnamed_59 = metal::int2(2) % metal::int2(1); metal::int2 unnamed_60 = metal::int2(2) % metal::int2(1); - metal::float3x3 unnamed_61 = const_type_14_ * 1.0; - metal::float3x3 unnamed_62 = 2.0 * const_type_14_; - metal::float3 unnamed_63 = const_type_15_ * metal::float4(1.0); - metal::float4 unnamed_64 = metal::float3(2.0) * const_type_15_; - metal::float3x3 unnamed_65 = const_type_15_ * const_type_16_; + metal::float3x3 unnamed_61 = const_type_14_ + const_type_14_; + metal::float3x3 unnamed_62 = const_type_14_ - const_type_14_; + metal::float3x3 unnamed_63 = const_type_14_ * 1.0; + metal::float3x3 unnamed_64 = 2.0 * const_type_14_; + metal::float3 unnamed_65 = const_type_15_ * metal::float4(1.0); + metal::float4 unnamed_66 = metal::float3(2.0) * const_type_15_; + metal::float3x3 unnamed_67 = const_type_15_ * const_type_16_; } void bit( ) { - int unnamed_66 = ~1; - uint unnamed_67 = ~1u; - metal::int2 unnamed_68 = ~metal::int2(1); - metal::uint3 unnamed_69 = ~metal::uint3(1u); - int unnamed_70 = 2 | 1; - uint unnamed_71 = 2u | 1u; - metal::int2 unnamed_72 = metal::int2(2) | metal::int2(1); - metal::uint3 unnamed_73 = metal::uint3(2u) | metal::uint3(1u); - int unnamed_74 = 2 & 1; - uint unnamed_75 = 2u & 1u; - metal::int2 unnamed_76 = metal::int2(2) & metal::int2(1); - metal::uint3 unnamed_77 = metal::uint3(2u) & metal::uint3(1u); - int unnamed_78 = 2 ^ 1; - uint unnamed_79 = 2u ^ 1u; - metal::int2 unnamed_80 = metal::int2(2) ^ metal::int2(1); - metal::uint3 unnamed_81 = metal::uint3(2u) ^ metal::uint3(1u); - int unnamed_82 = 2 << 1u; - uint unnamed_83 = 2u << 1u; - metal::int2 unnamed_84 = metal::int2(2) << metal::uint2(1u); - metal::uint3 unnamed_85 = metal::uint3(2u) << metal::uint3(1u); - int unnamed_86 = 2 >> 1u; - uint unnamed_87 = 2u >> 1u; - metal::int2 unnamed_88 = metal::int2(2) >> metal::uint2(1u); - metal::uint3 unnamed_89 = metal::uint3(2u) >> metal::uint3(1u); + int unnamed_68 = ~1; + uint unnamed_69 = ~1u; + metal::int2 unnamed_70 = ~metal::int2(1); + metal::uint3 unnamed_71 = ~metal::uint3(1u); + int unnamed_72 = 2 | 1; + uint unnamed_73 = 2u | 1u; + metal::int2 unnamed_74 = metal::int2(2) | metal::int2(1); + metal::uint3 unnamed_75 = metal::uint3(2u) | metal::uint3(1u); + int unnamed_76 = 2 & 1; + uint unnamed_77 = 2u & 1u; + metal::int2 unnamed_78 = metal::int2(2) & metal::int2(1); + metal::uint3 unnamed_79 = metal::uint3(2u) & metal::uint3(1u); + int unnamed_80 = 2 ^ 1; + uint unnamed_81 = 2u ^ 1u; + metal::int2 unnamed_82 = metal::int2(2) ^ metal::int2(1); + metal::uint3 unnamed_83 = metal::uint3(2u) ^ metal::uint3(1u); + int unnamed_84 = 2 << 1u; + uint unnamed_85 = 2u << 1u; + metal::int2 unnamed_86 = metal::int2(2) << metal::uint2(1u); + metal::uint3 unnamed_87 = metal::uint3(2u) << metal::uint3(1u); + int unnamed_88 = 2 >> 1u; + uint unnamed_89 = 2u >> 1u; + metal::int2 unnamed_90 = metal::int2(2) >> metal::uint2(1u); + metal::uint3 unnamed_91 = metal::uint3(2u) >> metal::uint3(1u); } void comparison( ) { - bool unnamed_90 = 2 == 1; - bool unnamed_91 = 2u == 1u; - bool unnamed_92 = 2.0 == 1.0; - metal::bool2 unnamed_93 = metal::int2(2) == metal::int2(1); - metal::bool3 unnamed_94 = metal::uint3(2u) == metal::uint3(1u); - metal::bool4 unnamed_95 = metal::float4(2.0) == metal::float4(1.0); - bool unnamed_96 = 2 != 1; - bool unnamed_97 = 2u != 1u; - bool unnamed_98 = 2.0 != 1.0; - metal::bool2 unnamed_99 = metal::int2(2) != metal::int2(1); - metal::bool3 unnamed_100 = metal::uint3(2u) != metal::uint3(1u); - metal::bool4 unnamed_101 = metal::float4(2.0) != metal::float4(1.0); - bool unnamed_102 = 2 < 1; - bool unnamed_103 = 2u < 1u; - bool unnamed_104 = 2.0 < 1.0; - metal::bool2 unnamed_105 = metal::int2(2) < metal::int2(1); - metal::bool3 unnamed_106 = metal::uint3(2u) < metal::uint3(1u); - metal::bool4 unnamed_107 = metal::float4(2.0) < metal::float4(1.0); - bool unnamed_108 = 2 <= 1; - bool unnamed_109 = 2u <= 1u; - bool unnamed_110 = 2.0 <= 1.0; - metal::bool2 unnamed_111 = metal::int2(2) <= metal::int2(1); - metal::bool3 unnamed_112 = metal::uint3(2u) <= metal::uint3(1u); - metal::bool4 unnamed_113 = metal::float4(2.0) <= metal::float4(1.0); - bool unnamed_114 = 2 > 1; - bool unnamed_115 = 2u > 1u; - bool unnamed_116 = 2.0 > 1.0; - metal::bool2 unnamed_117 = metal::int2(2) > metal::int2(1); - metal::bool3 unnamed_118 = metal::uint3(2u) > metal::uint3(1u); - metal::bool4 unnamed_119 = metal::float4(2.0) > metal::float4(1.0); - bool unnamed_120 = 2 >= 1; - bool unnamed_121 = 2u >= 1u; - bool unnamed_122 = 2.0 >= 1.0; - metal::bool2 unnamed_123 = metal::int2(2) >= metal::int2(1); - metal::bool3 unnamed_124 = metal::uint3(2u) >= metal::uint3(1u); - metal::bool4 unnamed_125 = metal::float4(2.0) >= metal::float4(1.0); + bool unnamed_92 = 2 == 1; + bool unnamed_93 = 2u == 1u; + bool unnamed_94 = 2.0 == 1.0; + metal::bool2 unnamed_95 = metal::int2(2) == metal::int2(1); + metal::bool3 unnamed_96 = metal::uint3(2u) == metal::uint3(1u); + metal::bool4 unnamed_97 = metal::float4(2.0) == metal::float4(1.0); + bool unnamed_98 = 2 != 1; + bool unnamed_99 = 2u != 1u; + bool unnamed_100 = 2.0 != 1.0; + metal::bool2 unnamed_101 = metal::int2(2) != metal::int2(1); + metal::bool3 unnamed_102 = metal::uint3(2u) != metal::uint3(1u); + metal::bool4 unnamed_103 = metal::float4(2.0) != metal::float4(1.0); + bool unnamed_104 = 2 < 1; + bool unnamed_105 = 2u < 1u; + bool unnamed_106 = 2.0 < 1.0; + metal::bool2 unnamed_107 = metal::int2(2) < metal::int2(1); + metal::bool3 unnamed_108 = metal::uint3(2u) < metal::uint3(1u); + metal::bool4 unnamed_109 = metal::float4(2.0) < metal::float4(1.0); + bool unnamed_110 = 2 <= 1; + bool unnamed_111 = 2u <= 1u; + bool unnamed_112 = 2.0 <= 1.0; + metal::bool2 unnamed_113 = metal::int2(2) <= metal::int2(1); + metal::bool3 unnamed_114 = metal::uint3(2u) <= metal::uint3(1u); + metal::bool4 unnamed_115 = metal::float4(2.0) <= metal::float4(1.0); + bool unnamed_116 = 2 > 1; + bool unnamed_117 = 2u > 1u; + bool unnamed_118 = 2.0 > 1.0; + metal::bool2 unnamed_119 = metal::int2(2) > metal::int2(1); + metal::bool3 unnamed_120 = metal::uint3(2u) > metal::uint3(1u); + metal::bool4 unnamed_121 = metal::float4(2.0) > metal::float4(1.0); + bool unnamed_122 = 2 >= 1; + bool unnamed_123 = 2u >= 1u; + bool unnamed_124 = 2.0 >= 1.0; + metal::bool2 unnamed_125 = metal::int2(2) >= metal::int2(1); + metal::bool3 unnamed_126 = metal::uint3(2u) >= metal::uint3(1u); + metal::bool4 unnamed_127 = metal::float4(2.0) >= metal::float4(1.0); } void assignment( diff --git a/tests/out/spv/operators.spvasm b/tests/out/spv/operators.spvasm index d3d9919ce3..294f48cc9c 100644 --- a/tests/out/spv/operators.spvasm +++ b/tests/out/spv/operators.spvasm @@ -1,12 +1,12 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 440 +; Bound: 460 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %428 "main" -OpExecutionMode %428 LocalSize 1 1 1 +OpEntryPoint GLCompute %448 "main" +OpExecutionMode %448 LocalSize 1 1 1 OpMemberDecorate %31 0 Offset 0 OpMemberDecorate %31 1 Offset 16 OpDecorate %36 ArrayStride 32 @@ -295,203 +295,223 @@ OpBranch %171 %263 = OpCompositeConstruct %172 %7 %7 %264 = OpCompositeConstruct %172 %18 %18 %265 = OpSMod %172 %264 %263 -%266 = OpMatrixTimesScalar %38 %52 %3 -%267 = OpMatrixTimesScalar %38 %52 %14 -%268 = OpCompositeConstruct %27 %3 %3 %3 %3 -%269 = OpMatrixTimesVector %30 %53 %268 -%270 = OpCompositeConstruct %30 %14 %14 %14 -%271 = OpVectorTimesMatrix %27 %270 %53 -%272 = OpMatrixTimesMatrix %38 %53 %54 +%267 = OpCompositeExtract %30 %52 0 +%268 = OpCompositeExtract %30 %52 0 +%269 = OpFAdd %30 %267 %268 +%270 = OpCompositeExtract %30 %52 1 +%271 = OpCompositeExtract %30 %52 1 +%272 = OpFAdd %30 %270 %271 +%273 = OpCompositeExtract %30 %52 2 +%274 = OpCompositeExtract %30 %52 2 +%275 = OpFAdd %30 %273 %274 +%266 = OpCompositeConstruct %38 %269 %272 %275 +%277 = OpCompositeExtract %30 %52 0 +%278 = OpCompositeExtract %30 %52 0 +%279 = OpFSub %30 %277 %278 +%280 = OpCompositeExtract %30 %52 1 +%281 = OpCompositeExtract %30 %52 1 +%282 = OpFSub %30 %280 %281 +%283 = OpCompositeExtract %30 %52 2 +%284 = OpCompositeExtract %30 %52 2 +%285 = OpFSub %30 %283 %284 +%276 = OpCompositeConstruct %38 %279 %282 %285 +%286 = OpMatrixTimesScalar %38 %52 %3 +%287 = OpMatrixTimesScalar %38 %52 %14 +%288 = OpCompositeConstruct %27 %3 %3 %3 %3 +%289 = OpMatrixTimesVector %30 %53 %288 +%290 = OpCompositeConstruct %30 %14 %14 %14 +%291 = OpVectorTimesMatrix %27 %290 %53 +%292 = OpMatrixTimesMatrix %38 %53 %54 OpReturn OpFunctionEnd -%274 = OpFunction %2 None %153 -%273 = OpLabel -OpBranch %275 -%275 = OpLabel -%276 = OpNot %8 %7 -%277 = OpNot %20 %26 -%278 = OpCompositeConstruct %172 %7 %7 -%279 = OpNot %172 %278 -%280 = OpCompositeConstruct %183 %26 %26 %26 -%281 = OpNot %183 %280 -%282 = OpBitwiseOr %8 %18 %7 -%283 = OpBitwiseOr %20 %25 %26 -%284 = OpCompositeConstruct %172 %18 %18 -%285 = OpCompositeConstruct %172 %7 %7 -%286 = OpBitwiseOr %172 %284 %285 -%287 = OpCompositeConstruct %183 %25 %25 %25 -%288 = OpCompositeConstruct %183 %26 %26 %26 -%289 = OpBitwiseOr %183 %287 %288 -%290 = OpBitwiseAnd %8 %18 %7 -%291 = OpBitwiseAnd %20 %25 %26 -%292 = OpCompositeConstruct %172 %18 %18 -%293 = OpCompositeConstruct %172 %7 %7 -%294 = OpBitwiseAnd %172 %292 %293 -%295 = OpCompositeConstruct %183 %25 %25 %25 -%296 = OpCompositeConstruct %183 %26 %26 %26 -%297 = OpBitwiseAnd %183 %295 %296 -%298 = OpBitwiseXor %8 %18 %7 -%299 = OpBitwiseXor %20 %25 %26 -%300 = OpCompositeConstruct %172 %18 %18 -%301 = OpCompositeConstruct %172 %7 %7 -%302 = OpBitwiseXor %172 %300 %301 -%303 = OpCompositeConstruct %183 %25 %25 %25 -%304 = OpCompositeConstruct %183 %26 %26 %26 -%305 = OpBitwiseXor %183 %303 %304 -%306 = OpShiftLeftLogical %8 %18 %26 -%307 = OpShiftLeftLogical %20 %25 %26 -%308 = OpCompositeConstruct %172 %18 %18 -%309 = OpCompositeConstruct %35 %26 %26 -%310 = OpShiftLeftLogical %172 %308 %309 -%311 = OpCompositeConstruct %183 %25 %25 %25 -%312 = OpCompositeConstruct %183 %26 %26 %26 -%313 = OpShiftLeftLogical %183 %311 %312 -%314 = OpShiftRightArithmetic %8 %18 %26 -%315 = OpShiftRightLogical %20 %25 %26 -%316 = OpCompositeConstruct %172 %18 %18 -%317 = OpCompositeConstruct %35 %26 %26 -%318 = OpShiftRightArithmetic %172 %316 %317 -%319 = OpCompositeConstruct %183 %25 %25 %25 -%320 = OpCompositeConstruct %183 %26 %26 %26 -%321 = OpShiftRightLogical %183 %319 %320 -OpReturn -OpFunctionEnd -%323 = OpFunction %2 None %153 -%322 = OpLabel -OpBranch %324 -%324 = OpLabel -%325 = OpIEqual %10 %18 %7 -%326 = OpIEqual %10 %25 %26 -%327 = OpFOrdEqual %10 %14 %3 +%294 = OpFunction %2 None %153 +%293 = OpLabel +OpBranch %295 +%295 = OpLabel +%296 = OpNot %8 %7 +%297 = OpNot %20 %26 +%298 = OpCompositeConstruct %172 %7 %7 +%299 = OpNot %172 %298 +%300 = OpCompositeConstruct %183 %26 %26 %26 +%301 = OpNot %183 %300 +%302 = OpBitwiseOr %8 %18 %7 +%303 = OpBitwiseOr %20 %25 %26 +%304 = OpCompositeConstruct %172 %18 %18 +%305 = OpCompositeConstruct %172 %7 %7 +%306 = OpBitwiseOr %172 %304 %305 +%307 = OpCompositeConstruct %183 %25 %25 %25 +%308 = OpCompositeConstruct %183 %26 %26 %26 +%309 = OpBitwiseOr %183 %307 %308 +%310 = OpBitwiseAnd %8 %18 %7 +%311 = OpBitwiseAnd %20 %25 %26 +%312 = OpCompositeConstruct %172 %18 %18 +%313 = OpCompositeConstruct %172 %7 %7 +%314 = OpBitwiseAnd %172 %312 %313 +%315 = OpCompositeConstruct %183 %25 %25 %25 +%316 = OpCompositeConstruct %183 %26 %26 %26 +%317 = OpBitwiseAnd %183 %315 %316 +%318 = OpBitwiseXor %8 %18 %7 +%319 = OpBitwiseXor %20 %25 %26 +%320 = OpCompositeConstruct %172 %18 %18 +%321 = OpCompositeConstruct %172 %7 %7 +%322 = OpBitwiseXor %172 %320 %321 +%323 = OpCompositeConstruct %183 %25 %25 %25 +%324 = OpCompositeConstruct %183 %26 %26 %26 +%325 = OpBitwiseXor %183 %323 %324 +%326 = OpShiftLeftLogical %8 %18 %26 +%327 = OpShiftLeftLogical %20 %25 %26 %328 = OpCompositeConstruct %172 %18 %18 -%329 = OpCompositeConstruct %172 %7 %7 -%330 = OpIEqual %156 %328 %329 +%329 = OpCompositeConstruct %35 %26 %26 +%330 = OpShiftLeftLogical %172 %328 %329 %331 = OpCompositeConstruct %183 %25 %25 %25 %332 = OpCompositeConstruct %183 %26 %26 %26 -%333 = OpIEqual %101 %331 %332 -%334 = OpCompositeConstruct %27 %14 %14 %14 %14 -%335 = OpCompositeConstruct %27 %3 %3 %3 %3 -%336 = OpFOrdEqual %29 %334 %335 -%337 = OpINotEqual %10 %18 %7 -%338 = OpINotEqual %10 %25 %26 -%339 = OpFOrdNotEqual %10 %14 %3 -%340 = OpCompositeConstruct %172 %18 %18 -%341 = OpCompositeConstruct %172 %7 %7 -%342 = OpINotEqual %156 %340 %341 -%343 = OpCompositeConstruct %183 %25 %25 %25 -%344 = OpCompositeConstruct %183 %26 %26 %26 -%345 = OpINotEqual %101 %343 %344 -%346 = OpCompositeConstruct %27 %14 %14 %14 %14 -%347 = OpCompositeConstruct %27 %3 %3 %3 %3 -%348 = OpFOrdNotEqual %29 %346 %347 -%349 = OpSLessThan %10 %18 %7 -%350 = OpULessThan %10 %25 %26 -%351 = OpFOrdLessThan %10 %14 %3 -%352 = OpCompositeConstruct %172 %18 %18 -%353 = OpCompositeConstruct %172 %7 %7 -%354 = OpSLessThan %156 %352 %353 -%355 = OpCompositeConstruct %183 %25 %25 %25 -%356 = OpCompositeConstruct %183 %26 %26 %26 -%357 = OpULessThan %101 %355 %356 -%358 = OpCompositeConstruct %27 %14 %14 %14 %14 -%359 = OpCompositeConstruct %27 %3 %3 %3 %3 -%360 = OpFOrdLessThan %29 %358 %359 -%361 = OpSLessThanEqual %10 %18 %7 -%362 = OpULessThanEqual %10 %25 %26 -%363 = OpFOrdLessThanEqual %10 %14 %3 -%364 = OpCompositeConstruct %172 %18 %18 -%365 = OpCompositeConstruct %172 %7 %7 -%366 = OpSLessThanEqual %156 %364 %365 -%367 = OpCompositeConstruct %183 %25 %25 %25 -%368 = OpCompositeConstruct %183 %26 %26 %26 -%369 = OpULessThanEqual %101 %367 %368 -%370 = OpCompositeConstruct %27 %14 %14 %14 %14 -%371 = OpCompositeConstruct %27 %3 %3 %3 %3 -%372 = OpFOrdLessThanEqual %29 %370 %371 -%373 = OpSGreaterThan %10 %18 %7 -%374 = OpUGreaterThan %10 %25 %26 -%375 = OpFOrdGreaterThan %10 %14 %3 -%376 = OpCompositeConstruct %172 %18 %18 -%377 = OpCompositeConstruct %172 %7 %7 -%378 = OpSGreaterThan %156 %376 %377 -%379 = OpCompositeConstruct %183 %25 %25 %25 -%380 = OpCompositeConstruct %183 %26 %26 %26 -%381 = OpUGreaterThan %101 %379 %380 -%382 = OpCompositeConstruct %27 %14 %14 %14 %14 -%383 = OpCompositeConstruct %27 %3 %3 %3 %3 -%384 = OpFOrdGreaterThan %29 %382 %383 -%385 = OpSGreaterThanEqual %10 %18 %7 -%386 = OpUGreaterThanEqual %10 %25 %26 -%387 = OpFOrdGreaterThanEqual %10 %14 %3 -%388 = OpCompositeConstruct %172 %18 %18 -%389 = OpCompositeConstruct %172 %7 %7 -%390 = OpSGreaterThanEqual %156 %388 %389 -%391 = OpCompositeConstruct %183 %25 %25 %25 -%392 = OpCompositeConstruct %183 %26 %26 %26 -%393 = OpUGreaterThanEqual %101 %391 %392 -%394 = OpCompositeConstruct %27 %14 %14 %14 %14 -%395 = OpCompositeConstruct %27 %3 %3 %3 %3 -%396 = OpFOrdGreaterThanEqual %29 %394 %395 +%333 = OpShiftLeftLogical %183 %331 %332 +%334 = OpShiftRightArithmetic %8 %18 %26 +%335 = OpShiftRightLogical %20 %25 %26 +%336 = OpCompositeConstruct %172 %18 %18 +%337 = OpCompositeConstruct %35 %26 %26 +%338 = OpShiftRightArithmetic %172 %336 %337 +%339 = OpCompositeConstruct %183 %25 %25 %25 +%340 = OpCompositeConstruct %183 %26 %26 %26 +%341 = OpShiftRightLogical %183 %339 %340 OpReturn OpFunctionEnd -%399 = OpFunction %2 None %153 -%398 = OpLabel -%397 = OpVariable %112 Function %7 -OpBranch %400 -%400 = OpLabel -%401 = OpLoad %8 %397 -%402 = OpIAdd %8 %401 %7 -OpStore %397 %402 -%403 = OpLoad %8 %397 -%404 = OpISub %8 %403 %7 -OpStore %397 %404 -%405 = OpLoad %8 %397 -%406 = OpLoad %8 %397 -%407 = OpIMul %8 %405 %406 -OpStore %397 %407 -%408 = OpLoad %8 %397 -%409 = OpLoad %8 %397 -%410 = OpSDiv %8 %408 %409 -OpStore %397 %410 -%411 = OpLoad %8 %397 -%412 = OpSMod %8 %411 %7 -OpStore %397 %412 -%413 = OpLoad %8 %397 -%414 = OpBitwiseAnd %8 %413 %11 -OpStore %397 %414 -%415 = OpLoad %8 %397 -%416 = OpBitwiseOr %8 %415 %11 -OpStore %397 %416 -%417 = OpLoad %8 %397 -%418 = OpBitwiseXor %8 %417 %11 -OpStore %397 %418 -%419 = OpLoad %8 %397 -%420 = OpShiftLeftLogical %8 %419 %25 -OpStore %397 %420 -%421 = OpLoad %8 %397 -%422 = OpShiftRightArithmetic %8 %421 %26 -OpStore %397 %422 -%423 = OpLoad %8 %397 -%424 = OpIAdd %8 %423 %7 -OpStore %397 %424 -%425 = OpLoad %8 %397 -%426 = OpISub %8 %425 %7 -OpStore %397 %426 +%343 = OpFunction %2 None %153 +%342 = OpLabel +OpBranch %344 +%344 = OpLabel +%345 = OpIEqual %10 %18 %7 +%346 = OpIEqual %10 %25 %26 +%347 = OpFOrdEqual %10 %14 %3 +%348 = OpCompositeConstruct %172 %18 %18 +%349 = OpCompositeConstruct %172 %7 %7 +%350 = OpIEqual %156 %348 %349 +%351 = OpCompositeConstruct %183 %25 %25 %25 +%352 = OpCompositeConstruct %183 %26 %26 %26 +%353 = OpIEqual %101 %351 %352 +%354 = OpCompositeConstruct %27 %14 %14 %14 %14 +%355 = OpCompositeConstruct %27 %3 %3 %3 %3 +%356 = OpFOrdEqual %29 %354 %355 +%357 = OpINotEqual %10 %18 %7 +%358 = OpINotEqual %10 %25 %26 +%359 = OpFOrdNotEqual %10 %14 %3 +%360 = OpCompositeConstruct %172 %18 %18 +%361 = OpCompositeConstruct %172 %7 %7 +%362 = OpINotEqual %156 %360 %361 +%363 = OpCompositeConstruct %183 %25 %25 %25 +%364 = OpCompositeConstruct %183 %26 %26 %26 +%365 = OpINotEqual %101 %363 %364 +%366 = OpCompositeConstruct %27 %14 %14 %14 %14 +%367 = OpCompositeConstruct %27 %3 %3 %3 %3 +%368 = OpFOrdNotEqual %29 %366 %367 +%369 = OpSLessThan %10 %18 %7 +%370 = OpULessThan %10 %25 %26 +%371 = OpFOrdLessThan %10 %14 %3 +%372 = OpCompositeConstruct %172 %18 %18 +%373 = OpCompositeConstruct %172 %7 %7 +%374 = OpSLessThan %156 %372 %373 +%375 = OpCompositeConstruct %183 %25 %25 %25 +%376 = OpCompositeConstruct %183 %26 %26 %26 +%377 = OpULessThan %101 %375 %376 +%378 = OpCompositeConstruct %27 %14 %14 %14 %14 +%379 = OpCompositeConstruct %27 %3 %3 %3 %3 +%380 = OpFOrdLessThan %29 %378 %379 +%381 = OpSLessThanEqual %10 %18 %7 +%382 = OpULessThanEqual %10 %25 %26 +%383 = OpFOrdLessThanEqual %10 %14 %3 +%384 = OpCompositeConstruct %172 %18 %18 +%385 = OpCompositeConstruct %172 %7 %7 +%386 = OpSLessThanEqual %156 %384 %385 +%387 = OpCompositeConstruct %183 %25 %25 %25 +%388 = OpCompositeConstruct %183 %26 %26 %26 +%389 = OpULessThanEqual %101 %387 %388 +%390 = OpCompositeConstruct %27 %14 %14 %14 %14 +%391 = OpCompositeConstruct %27 %3 %3 %3 %3 +%392 = OpFOrdLessThanEqual %29 %390 %391 +%393 = OpSGreaterThan %10 %18 %7 +%394 = OpUGreaterThan %10 %25 %26 +%395 = OpFOrdGreaterThan %10 %14 %3 +%396 = OpCompositeConstruct %172 %18 %18 +%397 = OpCompositeConstruct %172 %7 %7 +%398 = OpSGreaterThan %156 %396 %397 +%399 = OpCompositeConstruct %183 %25 %25 %25 +%400 = OpCompositeConstruct %183 %26 %26 %26 +%401 = OpUGreaterThan %101 %399 %400 +%402 = OpCompositeConstruct %27 %14 %14 %14 %14 +%403 = OpCompositeConstruct %27 %3 %3 %3 %3 +%404 = OpFOrdGreaterThan %29 %402 %403 +%405 = OpSGreaterThanEqual %10 %18 %7 +%406 = OpUGreaterThanEqual %10 %25 %26 +%407 = OpFOrdGreaterThanEqual %10 %14 %3 +%408 = OpCompositeConstruct %172 %18 %18 +%409 = OpCompositeConstruct %172 %7 %7 +%410 = OpSGreaterThanEqual %156 %408 %409 +%411 = OpCompositeConstruct %183 %25 %25 %25 +%412 = OpCompositeConstruct %183 %26 %26 %26 +%413 = OpUGreaterThanEqual %101 %411 %412 +%414 = OpCompositeConstruct %27 %14 %14 %14 %14 +%415 = OpCompositeConstruct %27 %3 %3 %3 %3 +%416 = OpFOrdGreaterThanEqual %29 %414 %415 OpReturn OpFunctionEnd -%428 = OpFunction %2 None %153 -%427 = OpLabel -OpBranch %429 -%429 = OpLabel -%430 = OpFunctionCall %27 %56 -%431 = OpFunctionCall %27 %81 -%432 = OpVectorShuffle %30 %41 %41 0 1 2 -%433 = OpFunctionCall %30 %98 %432 -%434 = OpFunctionCall %4 %129 -%435 = OpFunctionCall %2 %152 -%436 = OpFunctionCall %2 %170 -%437 = OpFunctionCall %2 %274 -%438 = OpFunctionCall %2 %323 -%439 = OpFunctionCall %2 %399 +%419 = OpFunction %2 None %153 +%418 = OpLabel +%417 = OpVariable %112 Function %7 +OpBranch %420 +%420 = OpLabel +%421 = OpLoad %8 %417 +%422 = OpIAdd %8 %421 %7 +OpStore %417 %422 +%423 = OpLoad %8 %417 +%424 = OpISub %8 %423 %7 +OpStore %417 %424 +%425 = OpLoad %8 %417 +%426 = OpLoad %8 %417 +%427 = OpIMul %8 %425 %426 +OpStore %417 %427 +%428 = OpLoad %8 %417 +%429 = OpLoad %8 %417 +%430 = OpSDiv %8 %428 %429 +OpStore %417 %430 +%431 = OpLoad %8 %417 +%432 = OpSMod %8 %431 %7 +OpStore %417 %432 +%433 = OpLoad %8 %417 +%434 = OpBitwiseAnd %8 %433 %11 +OpStore %417 %434 +%435 = OpLoad %8 %417 +%436 = OpBitwiseOr %8 %435 %11 +OpStore %417 %436 +%437 = OpLoad %8 %417 +%438 = OpBitwiseXor %8 %437 %11 +OpStore %417 %438 +%439 = OpLoad %8 %417 +%440 = OpShiftLeftLogical %8 %439 %25 +OpStore %417 %440 +%441 = OpLoad %8 %417 +%442 = OpShiftRightArithmetic %8 %441 %26 +OpStore %417 %442 +%443 = OpLoad %8 %417 +%444 = OpIAdd %8 %443 %7 +OpStore %417 %444 +%445 = OpLoad %8 %417 +%446 = OpISub %8 %445 %7 +OpStore %417 %446 +OpReturn +OpFunctionEnd +%448 = OpFunction %2 None %153 +%447 = OpLabel +OpBranch %449 +%449 = OpLabel +%450 = OpFunctionCall %27 %56 +%451 = OpFunctionCall %27 %81 +%452 = OpVectorShuffle %30 %41 %41 0 1 2 +%453 = OpFunctionCall %30 %98 %452 +%454 = OpFunctionCall %4 %129 +%455 = OpFunctionCall %2 %152 +%456 = OpFunctionCall %2 %170 +%457 = OpFunctionCall %2 %294 +%458 = OpFunctionCall %2 %343 +%459 = OpFunctionCall %2 %419 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/operators.wgsl b/tests/out/wgsl/operators.wgsl index ca9877b667..c504b2f84b 100644 --- a/tests/out/wgsl/operators.wgsl +++ b/tests/out/wgsl/operators.wgsl @@ -108,77 +108,79 @@ fn arithmetic() { let unnamed_58 = (vec2(2) / vec2(1)); let unnamed_59 = (vec2(2) % vec2(1)); let unnamed_60 = (vec2(2) % vec2(1)); - let unnamed_61 = (mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * 1.0); - let unnamed_62 = (2.0 * mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); - let unnamed_63 = (mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * vec4(1.0)); - let unnamed_64 = (vec3(2.0) * mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); - let unnamed_65 = (mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * mat3x4(vec4(0.0, 0.0, 0.0, 0.0), vec4(0.0, 0.0, 0.0, 0.0), vec4(0.0, 0.0, 0.0, 0.0))); + let unnamed_61 = (mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) + mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); + let unnamed_62 = (mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) - mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); + let unnamed_63 = (mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * 1.0); + let unnamed_64 = (2.0 * mat3x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); + let unnamed_65 = (mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * vec4(1.0)); + let unnamed_66 = (vec3(2.0) * mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0))); + let unnamed_67 = (mat4x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0)) * mat3x4(vec4(0.0, 0.0, 0.0, 0.0), vec4(0.0, 0.0, 0.0, 0.0), vec4(0.0, 0.0, 0.0, 0.0))); } fn bit() { - let unnamed_66 = ~(1); - let unnamed_67 = ~(1u); - let unnamed_68 = !(vec2(1)); - let unnamed_69 = !(vec3(1u)); - let unnamed_70 = (2 | 1); - let unnamed_71 = (2u | 1u); - let unnamed_72 = (vec2(2) | vec2(1)); - let unnamed_73 = (vec3(2u) | vec3(1u)); - let unnamed_74 = (2 & 1); - let unnamed_75 = (2u & 1u); - let unnamed_76 = (vec2(2) & vec2(1)); - let unnamed_77 = (vec3(2u) & vec3(1u)); - let unnamed_78 = (2 ^ 1); - let unnamed_79 = (2u ^ 1u); - let unnamed_80 = (vec2(2) ^ vec2(1)); - let unnamed_81 = (vec3(2u) ^ vec3(1u)); - let unnamed_82 = (2 << 1u); - let unnamed_83 = (2u << 1u); - let unnamed_84 = (vec2(2) << vec2(1u)); - let unnamed_85 = (vec3(2u) << vec3(1u)); - let unnamed_86 = (2 >> 1u); - let unnamed_87 = (2u >> 1u); - let unnamed_88 = (vec2(2) >> vec2(1u)); - let unnamed_89 = (vec3(2u) >> vec3(1u)); + let unnamed_68 = ~(1); + let unnamed_69 = ~(1u); + let unnamed_70 = !(vec2(1)); + let unnamed_71 = !(vec3(1u)); + let unnamed_72 = (2 | 1); + let unnamed_73 = (2u | 1u); + let unnamed_74 = (vec2(2) | vec2(1)); + let unnamed_75 = (vec3(2u) | vec3(1u)); + let unnamed_76 = (2 & 1); + let unnamed_77 = (2u & 1u); + let unnamed_78 = (vec2(2) & vec2(1)); + let unnamed_79 = (vec3(2u) & vec3(1u)); + let unnamed_80 = (2 ^ 1); + let unnamed_81 = (2u ^ 1u); + let unnamed_82 = (vec2(2) ^ vec2(1)); + let unnamed_83 = (vec3(2u) ^ vec3(1u)); + let unnamed_84 = (2 << 1u); + let unnamed_85 = (2u << 1u); + let unnamed_86 = (vec2(2) << vec2(1u)); + let unnamed_87 = (vec3(2u) << vec3(1u)); + let unnamed_88 = (2 >> 1u); + let unnamed_89 = (2u >> 1u); + let unnamed_90 = (vec2(2) >> vec2(1u)); + let unnamed_91 = (vec3(2u) >> vec3(1u)); } fn comparison() { - let unnamed_90 = (2 == 1); - let unnamed_91 = (2u == 1u); - let unnamed_92 = (2.0 == 1.0); - let unnamed_93 = (vec2(2) == vec2(1)); - let unnamed_94 = (vec3(2u) == vec3(1u)); - let unnamed_95 = (vec4(2.0) == vec4(1.0)); - let unnamed_96 = (2 != 1); - let unnamed_97 = (2u != 1u); - let unnamed_98 = (2.0 != 1.0); - let unnamed_99 = (vec2(2) != vec2(1)); - let unnamed_100 = (vec3(2u) != vec3(1u)); - let unnamed_101 = (vec4(2.0) != vec4(1.0)); - let unnamed_102 = (2 < 1); - let unnamed_103 = (2u < 1u); - let unnamed_104 = (2.0 < 1.0); - let unnamed_105 = (vec2(2) < vec2(1)); - let unnamed_106 = (vec3(2u) < vec3(1u)); - let unnamed_107 = (vec4(2.0) < vec4(1.0)); - let unnamed_108 = (2 <= 1); - let unnamed_109 = (2u <= 1u); - let unnamed_110 = (2.0 <= 1.0); - let unnamed_111 = (vec2(2) <= vec2(1)); - let unnamed_112 = (vec3(2u) <= vec3(1u)); - let unnamed_113 = (vec4(2.0) <= vec4(1.0)); - let unnamed_114 = (2 > 1); - let unnamed_115 = (2u > 1u); - let unnamed_116 = (2.0 > 1.0); - let unnamed_117 = (vec2(2) > vec2(1)); - let unnamed_118 = (vec3(2u) > vec3(1u)); - let unnamed_119 = (vec4(2.0) > vec4(1.0)); - let unnamed_120 = (2 >= 1); - let unnamed_121 = (2u >= 1u); - let unnamed_122 = (2.0 >= 1.0); - let unnamed_123 = (vec2(2) >= vec2(1)); - let unnamed_124 = (vec3(2u) >= vec3(1u)); - let unnamed_125 = (vec4(2.0) >= vec4(1.0)); + let unnamed_92 = (2 == 1); + let unnamed_93 = (2u == 1u); + let unnamed_94 = (2.0 == 1.0); + let unnamed_95 = (vec2(2) == vec2(1)); + let unnamed_96 = (vec3(2u) == vec3(1u)); + let unnamed_97 = (vec4(2.0) == vec4(1.0)); + let unnamed_98 = (2 != 1); + let unnamed_99 = (2u != 1u); + let unnamed_100 = (2.0 != 1.0); + let unnamed_101 = (vec2(2) != vec2(1)); + let unnamed_102 = (vec3(2u) != vec3(1u)); + let unnamed_103 = (vec4(2.0) != vec4(1.0)); + let unnamed_104 = (2 < 1); + let unnamed_105 = (2u < 1u); + let unnamed_106 = (2.0 < 1.0); + let unnamed_107 = (vec2(2) < vec2(1)); + let unnamed_108 = (vec3(2u) < vec3(1u)); + let unnamed_109 = (vec4(2.0) < vec4(1.0)); + let unnamed_110 = (2 <= 1); + let unnamed_111 = (2u <= 1u); + let unnamed_112 = (2.0 <= 1.0); + let unnamed_113 = (vec2(2) <= vec2(1)); + let unnamed_114 = (vec3(2u) <= vec3(1u)); + let unnamed_115 = (vec4(2.0) <= vec4(1.0)); + let unnamed_116 = (2 > 1); + let unnamed_117 = (2u > 1u); + let unnamed_118 = (2.0 > 1.0); + let unnamed_119 = (vec2(2) > vec2(1)); + let unnamed_120 = (vec3(2u) > vec3(1u)); + let unnamed_121 = (vec4(2.0) > vec4(1.0)); + let unnamed_122 = (2 >= 1); + let unnamed_123 = (2u >= 1u); + let unnamed_124 = (2.0 >= 1.0); + let unnamed_125 = (vec2(2) >= vec2(1)); + let unnamed_126 = (vec3(2u) >= vec3(1u)); + let unnamed_127 = (vec4(2.0) >= vec4(1.0)); } fn assignment() {