intrinsics: Add scalar overload of all() & any()

Fixed: tint:1253
Change-Id: I0bdc865a9df9e0171c09daa9918b25bba033ba3b
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67061
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>

docs/origin-trial-changes.md

@@ -8,6 +8,10 @@
 * Deprecated texture builtin functions that accepted a `read` access controlled storage texture have been removed.
 * Storage textures must now only use the `write` access control.
 
+### New Features
+
+* `any()` and `all()` now support a `bool` parameter. These simply return the passed argument. [tint:1253](https://crbug.com/tint/1253)
+
 ### Fixes
 
 * Swizzling of `vec3` types in `storage` and `uniform` buffers has been fixed for Metal 1.x. [tint:1249](https://crbug.com/tint/1249)

src/intrinsics.def

@@ -267,7 +267,9 @@ fn abs<T: fiu32>(T) -> T
 fn abs<N: num, T: fiu32>(vec<N, T>) -> vec<N, T>
 fn acos(f32) -> f32
 fn acos<N: num>(vec<N, f32>) -> vec<N, f32>
+fn all(bool) -> bool
 fn all<N: num>(vec<N, bool>) -> bool
+fn any(bool) -> bool
 fn any<N: num>(vec<N, bool>) -> bool
 fn arrayLength<T, A: access>(ptr<storage, array<T>, A>) -> u32
 fn asin(f32) -> f32

src/resolver/intrinsic_test.cc

@@ -107,8 +107,21 @@ INSTANTIATE_TEST_SUITE_P(ResolverTest,
                                          "fwidthCoarse",
                                          "fwidthFine"));
 
-using ResolverIntrinsic = ResolverTestWithParam<std::string>;
-TEST_P(ResolverIntrinsic, Test) {
+using ResolverIntrinsicTest_BoolMethod = ResolverTestWithParam<std::string>;
+TEST_P(ResolverIntrinsicTest_BoolMethod, Scalar) {
+  auto name = GetParam();
+
+  Global("my_var", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* expr = Call(name, "my_var");
+  WrapInFunction(expr);
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+  ASSERT_NE(TypeOf(expr), nullptr);
+  EXPECT_TRUE(TypeOf(expr)->Is<sem::Bool>());
+}
+TEST_P(ResolverIntrinsicTest_BoolMethod, Vector) {
   auto name = GetParam();
 
   Global("my_var", ty.vec3<bool>(), ast::StorageClass::kPrivate);
@@ -122,7 +135,7 @@ TEST_P(ResolverIntrinsic, Test) {
   EXPECT_TRUE(TypeOf(expr)->Is<sem::Bool>());
 }
 INSTANTIATE_TEST_SUITE_P(ResolverTest,
-                         ResolverIntrinsic,
+                         ResolverIntrinsicTest_BoolMethod,
                          testing::Values("any", "all"));
 
 using ResolverIntrinsicTest_FloatMethod = ResolverTestWithParam<std::string>;

src/writer/spirv/builder.cc

@@ -2307,10 +2307,10 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
     return result_id;
   }
 
-  // Generates the SPIR-V ID for the expression for the indexed call parameter,
+  // Generates the SPIR-V ID for the expression for the indexed call argument,
   // and loads it if necessary. Returns 0 on error.
-  auto get_param_as_value_id = [&](size_t i,
-                                   bool generate_load = true) -> uint32_t {
+  auto get_arg_as_value_id = [&](size_t i,
+                                 bool generate_load = true) -> uint32_t {
     auto* arg = call->args[i];
     auto* param = intrinsic->Parameters()[i];
     auto val_id = GenerateExpression(arg);
@@ -2327,7 +2327,7 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
   OperandList params = {Operand::Int(result_type_id), result};
   spv::Op op = spv::Op::OpNop;
 
-  // Pushes the parameters for a GlslStd450 extended instruction, and sets op
+  // Pushes the arguments for a GlslStd450 extended instruction, and sets op
   // to OpExtInst.
   auto glsl_std450 = [&](uint32_t inst_id) {
     auto set_id = GetGLSLstd450Import();
@@ -2338,9 +2338,17 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
 
   switch (intrinsic->Type()) {
     case IntrinsicType::kAny:
+      if (intrinsic->Parameters()[0]->Type()->Is<sem::Bool>()) {
+        // any(v: bool) just resolves to v.
+        return get_arg_as_value_id(0);
+      }
       op = spv::Op::OpAny;
       break;
     case IntrinsicType::kAll:
+      if (intrinsic->Parameters()[0]->Type()->Is<sem::Bool>()) {
+        // all(v: bool) just resolves to v.
+        return get_arg_as_value_id(0);
+      }
       op = spv::Op::OpAll;
       break;
     case IntrinsicType::kArrayLength: {
@@ -2424,7 +2432,7 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
       // Evaluate the single argument, return the non-zero result_id which isn't
       // associated with any op (ignore returns void, so this cannot be used in
       // an expression).
-      if (!get_param_as_value_id(0, false)) {
+      if (!get_arg_as_value_id(0, false)) {
         return 0;
       }
       return result_id;
@@ -2436,7 +2444,7 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
       break;
     case IntrinsicType::kIsFinite: {
       // Implemented as:   not(IsInf or IsNan)
-      auto val_id = get_param_as_value_id(0);
+      auto val_id = get_arg_as_value_id(0);
       if (!val_id) {
         return 0;
       }
@@ -2468,7 +2476,7 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
       //   clamped = uclamp(1,254,exponent_bits);
       //   result = (clamped == exponent_bits);
       //
-      auto val_id = get_param_as_value_id(0);
+      auto val_id = get_arg_as_value_id(0);
       if (!val_id) {
         return 0;
       }
@@ -2541,9 +2549,9 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
     case IntrinsicType::kMix: {
       auto std450 = Operand::Int(GetGLSLstd450Import());
 
-      auto a_id = get_param_as_value_id(0);
-      auto b_id = get_param_as_value_id(1);
-      auto f_id = get_param_as_value_id(2);
+      auto a_id = get_arg_as_value_id(0);
+      auto b_id = get_arg_as_value_id(1);
+      auto f_id = get_arg_as_value_id(2);
       if (!a_id || !b_id || !f_id) {
         return 0;
       }
@@ -2572,9 +2580,9 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
       break;
     case IntrinsicType::kSelect: {
       // Note: Argument order is different in WGSL and SPIR-V
-      auto cond_id = get_param_as_value_id(2);
-      auto true_id = get_param_as_value_id(1);
-      auto false_id = get_param_as_value_id(0);
+      auto cond_id = get_arg_as_value_id(2);
+      auto true_id = get_arg_as_value_id(1);
+      auto false_id = get_arg_as_value_id(0);
       if (!cond_id || !true_id || !false_id) {
         return 0;
       }
@@ -2611,7 +2619,7 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
       if (intrinsic->ReturnType()->is_unsigned_scalar_or_vector()) {
         // abs() only operates on *signed* integers.
         // This is a no-op for unsigned integers.
-        return get_param_as_value_id(0);
+        return get_arg_as_value_id(0);
       }
       if (intrinsic->ReturnType()->is_float_scalar_or_vector()) {
         glsl_std450(GLSLstd450FAbs);
@@ -2637,7 +2645,7 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
   }
 
   for (size_t i = 0; i < call->args.size(); i++) {
-    if (auto val_id = get_param_as_value_id(i)) {
+    if (auto val_id = get_arg_as_value_id(i)) {
       params.emplace_back(Operand::Int(val_id));
     } else {
       return 0;

src/writer/spirv/builder_intrinsic_test.cc

@@ -39,7 +39,32 @@ inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
 }
 
 using IntrinsicBoolTest = IntrinsicBuilderTestWithParam<IntrinsicData>;
-TEST_P(IntrinsicBoolTest, Call_Bool) {
+TEST_P(IntrinsicBoolTest, Call_Bool_Scalar) {
+  auto param = GetParam();
+
+  auto* var = Global("v", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* expr = Call(param.name, "v");
+  WrapInFunction(expr);
+
+  spirv::Builder& b = Build();
+
+  b.push_function(Function{});
+  ASSERT_TRUE(b.GenerateGlobalVariable(var)) << b.error();
+
+  EXPECT_EQ(b.GenerateCallExpression(expr), 6u) << b.error();
+  EXPECT_EQ(DumpInstructions(b.types()), R"(%3 = OpTypeBool
+%2 = OpTypePointer Private %3
+%4 = OpConstantNull %3
+%1 = OpVariable %2 Private %4
+)");
+
+  // both any and all are 'passthrough' for scalar booleans
+  EXPECT_EQ(DumpInstructions(b.functions()[0].instructions()),
+            "%6 = OpLoad %3 %1\n");
+}
+
+TEST_P(IntrinsicBoolTest, Call_Bool_Vector) {
   auto param = GetParam();
 
   auto* var = Global("v", ty.vec3<bool>(), ast::StorageClass::kPrivate);

test/intrinsics/gen/all/353d6a.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn all(bool) -> bool
+fn all_353d6a() {
+  var res: bool = all(bool());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  all_353d6a();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  all_353d6a();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  all_353d6a();
+}

test/intrinsics/gen/all/353d6a.wgsl.expected.hlsl

@@ -0,0 +1,30 @@
+void all_353d6a() {
+  bool res = all(false);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  all_353d6a();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  all_353d6a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  all_353d6a();
+  return;
+}

test/intrinsics/gen/all/353d6a.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void all_353d6a() {
+  bool res = all(bool());
+}
+
+float4 vertex_main_inner() {
+  all_353d6a();
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  all_353d6a();
+  return;
+}
+
+kernel void compute_main() {
+  all_353d6a();
+  return;
+}
+

test/intrinsics/gen/all/353d6a.wgsl.expected.spvasm

@@ -0,0 +1,65 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %all_353d6a "all_353d6a"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+       %bool = OpTypeBool
+         %15 = OpConstantNull %bool
+%_ptr_Function_bool = OpTypePointer Function %bool
+         %18 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %all_353d6a = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_bool Function %15
+               OpStore %res %15
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %18
+         %20 = OpLabel
+         %21 = OpFunctionCall %void %all_353d6a
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %23 = OpLabel
+         %24 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %24
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %all_353d6a
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %all_353d6a
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/all/353d6a.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn all_353d6a() {
+  var res : bool = all(bool());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  all_353d6a();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  all_353d6a();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  all_353d6a();
+}

test/intrinsics/gen/any/2ab91a.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn any(bool) -> bool
+fn any_2ab91a() {
+  var res: bool = any(bool());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  any_2ab91a();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  any_2ab91a();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  any_2ab91a();
+}

test/intrinsics/gen/any/2ab91a.wgsl.expected.hlsl

@@ -0,0 +1,30 @@
+void any_2ab91a() {
+  bool res = any(false);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  any_2ab91a();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  any_2ab91a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  any_2ab91a();
+  return;
+}

test/intrinsics/gen/any/2ab91a.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void any_2ab91a() {
+  bool res = any(bool());
+}
+
+float4 vertex_main_inner() {
+  any_2ab91a();
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  any_2ab91a();
+  return;
+}
+
+kernel void compute_main() {
+  any_2ab91a();
+  return;
+}
+

test/intrinsics/gen/any/2ab91a.wgsl.expected.spvasm

@@ -0,0 +1,65 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %any_2ab91a "any_2ab91a"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+       %bool = OpTypeBool
+         %15 = OpConstantNull %bool
+%_ptr_Function_bool = OpTypePointer Function %bool
+         %18 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %any_2ab91a = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_bool Function %15
+               OpStore %res %15
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %18
+         %20 = OpLabel
+         %21 = OpFunctionCall %void %any_2ab91a
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %23 = OpLabel
+         %24 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %24
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %any_2ab91a
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %any_2ab91a
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/any/2ab91a.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn any_2ab91a() {
+  var res : bool = any(bool());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  any_2ab91a();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  any_2ab91a();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  any_2ab91a();
+}