[tint] Implement subgroupAll and subgroupAny
Implement subgroupAll and subgroupAny builtins for WGSL, SPIR-V IR
backend (not AST), MSL AST and IR backends and HLSL AST and IR backends.
Subgroup builtins are not yet implemented for the GLSL backend.
Bug: 354738715
Change-Id: I9b3d1401ccca9b56e361f2b893812e3d3a42b4e2
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/201916
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Natalie Chouinard <chouinard@google.com>
diff --git a/src/tint/cmd/fuzz/wgsl/dictionary.txt b/src/tint/cmd/fuzz/wgsl/dictionary.txt
index 3508d1f..222fa52 100644
--- a/src/tint/cmd/fuzz/wgsl/dictionary.txt
+++ b/src/tint/cmd/fuzz/wgsl/dictionary.txt
@@ -368,7 +368,9 @@
"storageBarrier"
"struct"
"subgroupAdd"
+"subgroupAll"
"subgroupAnd"
+"subgroupAny"
"subgroupBallot"
"subgroupBroadcast"
"subgroupExclusiveAdd"
diff --git a/src/tint/lang/core/builtin_fn.cc b/src/tint/lang/core/builtin_fn.cc
index dd3b048..ad8a8a8 100644
--- a/src/tint/lang/core/builtin_fn.cc
+++ b/src/tint/lang/core/builtin_fn.cc
@@ -432,6 +432,12 @@
if (name == "subgroupMax") {
return BuiltinFn::kSubgroupMax;
}
+ if (name == "subgroupAll") {
+ return BuiltinFn::kSubgroupAll;
+ }
+ if (name == "subgroupAny") {
+ return BuiltinFn::kSubgroupAny;
+ }
return BuiltinFn::kNone;
}
@@ -701,6 +707,10 @@
return "subgroupMin";
case BuiltinFn::kSubgroupMax:
return "subgroupMax";
+ case BuiltinFn::kSubgroupAll:
+ return "subgroupAll";
+ case BuiltinFn::kSubgroupAny:
+ return "subgroupAny";
}
return "<unknown>";
}
diff --git a/src/tint/lang/core/builtin_fn.h b/src/tint/lang/core/builtin_fn.h
index 6a9b362..508c3dc 100644
--- a/src/tint/lang/core/builtin_fn.h
+++ b/src/tint/lang/core/builtin_fn.h
@@ -178,6 +178,8 @@
kSubgroupXor,
kSubgroupMin,
kSubgroupMax,
+ kSubgroupAll,
+ kSubgroupAny,
kNone,
};
@@ -331,6 +333,8 @@
BuiltinFn::kSubgroupXor,
BuiltinFn::kSubgroupMin,
BuiltinFn::kSubgroupMax,
+ BuiltinFn::kSubgroupAll,
+ BuiltinFn::kSubgroupAny,
};
/// All builtin function names
@@ -466,6 +470,8 @@
"subgroupXor",
"subgroupMin",
"subgroupMax",
+ "subgroupAll",
+ "subgroupAny",
};
/// Determines if the given `f` is a coarse derivative.
diff --git a/src/tint/lang/core/core.def b/src/tint/lang/core/core.def
index 856b971..d1e1fb4 100644
--- a/src/tint/lang/core/core.def
+++ b/src/tint/lang/core/core.def
@@ -720,6 +720,8 @@
@must_use @stage("fragment", "compute") fn subgroupMin[N: num, T: fiu32_f16](value: vec<N, T>) -> vec<N, T>
@must_use @stage("fragment", "compute") fn subgroupMax[T: fiu32_f16](value: T) -> T
@must_use @stage("fragment", "compute") fn subgroupMax[N: num, T: fiu32_f16](value: vec<N, T>) -> vec<N, T>
+@must_use @stage("fragment", "compute") fn subgroupAll(bool) -> bool
+@must_use @stage("fragment", "compute") fn subgroupAny(bool) -> bool
////////////////////////////////////////////////////////////////////////////////
// Value constructors //
diff --git a/src/tint/lang/core/intrinsic/data.cc b/src/tint/lang/core/intrinsic/data.cc
index e1dab8b..a23c8a5 100644
--- a/src/tint/lang/core/intrinsic/data.cc
+++ b/src/tint/lang/core/intrinsic/data.cc
@@ -9547,6 +9547,17 @@
},
{
/* [469] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(224),
+ /* return_matcher_indices */ MatcherIndicesIndex(43),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [470] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -9557,7 +9568,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(95),
},
{
- /* [470] */
+ /* [471] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -10573,6 +10584,18 @@
/* num overloads */ 2,
/* overloads */ OverloadIndex(407),
},
+ {
+ /* [131] */
+ /* fn subgroupAll(bool) -> bool */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(469),
+ },
+ {
+ /* [132] */
+ /* fn subgroupAny(bool) -> bool */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(469),
+ },
};
constexpr IntrinsicInfo kUnaryOperators[] = {
@@ -10684,13 +10707,13 @@
/* [8] */
/* op &&(bool, bool) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(469),
+ /* overloads */ OverloadIndex(470),
},
{
/* [9] */
/* op ||(bool, bool) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(470),
+ /* overloads */ OverloadIndex(471),
},
{
/* [10] */
diff --git a/src/tint/lang/core/ir/binary/decode.cc b/src/tint/lang/core/ir/binary/decode.cc
index 06b41f5..5bf4b72 100644
--- a/src/tint/lang/core/ir/binary/decode.cc
+++ b/src/tint/lang/core/ir/binary/decode.cc
@@ -1647,6 +1647,10 @@
return core::BuiltinFn::kSubgroupMin;
case pb::BuiltinFn::subgroup_max:
return core::BuiltinFn::kSubgroupMax;
+ case pb::BuiltinFn::subgroup_all:
+ return core::BuiltinFn::kSubgroupAll;
+ case pb::BuiltinFn::subgroup_any:
+ return core::BuiltinFn::kSubgroupAny;
case pb::BuiltinFn::BuiltinFn_INT_MIN_SENTINEL_DO_NOT_USE_:
case pb::BuiltinFn::BuiltinFn_INT_MAX_SENTINEL_DO_NOT_USE_:
diff --git a/src/tint/lang/core/ir/binary/encode.cc b/src/tint/lang/core/ir/binary/encode.cc
index 26fbb9f..741da4d 100644
--- a/src/tint/lang/core/ir/binary/encode.cc
+++ b/src/tint/lang/core/ir/binary/encode.cc
@@ -1159,6 +1159,10 @@
return pb::BuiltinFn::subgroup_min;
case core::BuiltinFn::kSubgroupMax:
return pb::BuiltinFn::subgroup_max;
+ case core::BuiltinFn::kSubgroupAll:
+ return pb::BuiltinFn::subgroup_all;
+ case core::BuiltinFn::kSubgroupAny:
+ return pb::BuiltinFn::subgroup_any;
case core::BuiltinFn::kNone:
break;
}
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
index 4fe6094..8dcd019 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -3116,6 +3116,10 @@
return "WaveActiveMin";
case wgsl::BuiltinFn::kSubgroupMax:
return "WaveActiveMax";
+ case wgsl::BuiltinFn::kSubgroupAll:
+ return "WaveActiveAllTrue";
+ case wgsl::BuiltinFn::kSubgroupAny:
+ return "WaveActiveAnyTrue";
default:
diagnostics_.AddError(Source{}) << "Unknown builtin method: " << builtin->str();
}
diff --git a/src/tint/lang/hlsl/writer/printer/printer.cc b/src/tint/lang/hlsl/writer/printer/printer.cc
index ad315cc..14fd336 100644
--- a/src/tint/lang/hlsl/writer/printer/printer.cc
+++ b/src/tint/lang/hlsl/writer/printer/printer.cc
@@ -1032,6 +1032,12 @@
case core::BuiltinFn::kSubgroupMax:
out << "WaveActiveMax";
break;
+ case core::BuiltinFn::kSubgroupAll:
+ out << "WaveActiveAllTrue";
+ break;
+ case core::BuiltinFn::kSubgroupAny:
+ out << "WaveActiveAnyTrue";
+ break;
default:
TINT_UNREACHABLE() << "unhandled: " << func;
}
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index bc80b2d..a1296cf 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -886,6 +886,24 @@
return true;
}
+ case wgsl::BuiltinFn::kSubgroupAll: {
+ out << "simd_all(";
+ if (!EmitExpression(out, expr->args[0])) {
+ return false;
+ }
+ out << ")";
+ return true;
+ }
+
+ case wgsl::BuiltinFn::kSubgroupAny: {
+ out << "simd_any(";
+ if (!EmitExpression(out, expr->args[0])) {
+ return false;
+ }
+ out << ")";
+ return true;
+ }
+
default:
break;
}
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index d453a25..d89fd9b 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -1006,6 +1006,12 @@
case core::BuiltinFn::kSubgroupMax:
out << "simd_max";
break;
+ case core::BuiltinFn::kSubgroupAll:
+ out << "simd_all";
+ break;
+ case core::BuiltinFn::kSubgroupAny:
+ out << "simd_any";
+ break;
case core::BuiltinFn::kInverseSqrt:
out << "rsqrt";
break;
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.cc b/src/tint/lang/spirv/writer/ast_printer/builder.cc
index 86264ae..9625a86 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.cc
@@ -2573,7 +2573,9 @@
case wgsl::BuiltinFn::kSubgroupOr:
case wgsl::BuiltinFn::kSubgroupXor:
case wgsl::BuiltinFn::kSubgroupMin:
- case wgsl::BuiltinFn::kSubgroupMax: {
+ case wgsl::BuiltinFn::kSubgroupMax:
+ case wgsl::BuiltinFn::kSubgroupAll:
+ case wgsl::BuiltinFn::kSubgroupAny: {
// This file should be removed soon with the SPIR-V AST backend, but in the meantime we
// need this explicit failure to avoid a new ICE that can get caught by the fuzzers.
return 0;
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index 4393248..eb40a7f 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1730,6 +1730,16 @@
operands.push_back(Constant(ir_.constant_values.Get(u32(spv::Scope::Subgroup))));
operands.push_back(U32Operand(u32(spv::GroupOperation::Reduce)));
break;
+ case core::BuiltinFn::kSubgroupAll:
+ module_.PushCapability(SpvCapabilityGroupNonUniformVote);
+ op = spv::Op::OpGroupNonUniformAll;
+ operands.push_back(Constant(ir_.constant_values.Get(u32(spv::Scope::Subgroup))));
+ break;
+ case core::BuiltinFn::kSubgroupAny:
+ module_.PushCapability(SpvCapabilityGroupNonUniformVote);
+ op = spv::Op::OpGroupNonUniformAny;
+ operands.push_back(Constant(ir_.constant_values.Get(u32(spv::Scope::Subgroup))));
+ break;
case core::BuiltinFn::kTan:
glsl_ext_inst(GLSLstd450Tan);
break;
diff --git a/src/tint/lang/wgsl/builtin_fn.cc b/src/tint/lang/wgsl/builtin_fn.cc
index fabc2e9..ed76cfc 100644
--- a/src/tint/lang/wgsl/builtin_fn.cc
+++ b/src/tint/lang/wgsl/builtin_fn.cc
@@ -438,6 +438,12 @@
if (name == "subgroupMax") {
return BuiltinFn::kSubgroupMax;
}
+ if (name == "subgroupAll") {
+ return BuiltinFn::kSubgroupAll;
+ }
+ if (name == "subgroupAny") {
+ return BuiltinFn::kSubgroupAny;
+ }
if (name == "__tint_materialize") {
return BuiltinFn::kTintMaterialize;
}
@@ -714,6 +720,10 @@
return "subgroupMin";
case BuiltinFn::kSubgroupMax:
return "subgroupMax";
+ case BuiltinFn::kSubgroupAll:
+ return "subgroupAll";
+ case BuiltinFn::kSubgroupAny:
+ return "subgroupAny";
case BuiltinFn::kTintMaterialize:
return "__tint_materialize";
}
@@ -799,6 +809,8 @@
case BuiltinFn::kSubgroupXor:
case BuiltinFn::kSubgroupMin:
case BuiltinFn::kSubgroupMax:
+ case BuiltinFn::kSubgroupAll:
+ case BuiltinFn::kSubgroupAny:
return true;
default:
return false;
diff --git a/src/tint/lang/wgsl/builtin_fn.cc.tmpl b/src/tint/lang/wgsl/builtin_fn.cc.tmpl
index 7081457..1f5b736 100644
--- a/src/tint/lang/wgsl/builtin_fn.cc.tmpl
+++ b/src/tint/lang/wgsl/builtin_fn.cc.tmpl
@@ -120,6 +120,8 @@
case BuiltinFn::kSubgroupXor:
case BuiltinFn::kSubgroupMin:
case BuiltinFn::kSubgroupMax:
+ case BuiltinFn::kSubgroupAll:
+ case BuiltinFn::kSubgroupAny:
return true;
default:
return false;
diff --git a/src/tint/lang/wgsl/builtin_fn.h b/src/tint/lang/wgsl/builtin_fn.h
index 449ca60..b3ef461 100644
--- a/src/tint/lang/wgsl/builtin_fn.h
+++ b/src/tint/lang/wgsl/builtin_fn.h
@@ -180,6 +180,8 @@
kSubgroupXor,
kSubgroupMin,
kSubgroupMax,
+ kSubgroupAll,
+ kSubgroupAny,
kTintMaterialize,
kNone,
};
@@ -336,6 +338,8 @@
BuiltinFn::kSubgroupXor,
BuiltinFn::kSubgroupMin,
BuiltinFn::kSubgroupMax,
+ BuiltinFn::kSubgroupAll,
+ BuiltinFn::kSubgroupAny,
BuiltinFn::kTintMaterialize,
};
@@ -474,6 +478,8 @@
"subgroupXor",
"subgroupMin",
"subgroupMax",
+ "subgroupAll",
+ "subgroupAny",
"__tint_materialize",
};
diff --git a/src/tint/lang/wgsl/intrinsic/data.cc b/src/tint/lang/wgsl/intrinsic/data.cc
index 051857d..610515f 100644
--- a/src/tint/lang/wgsl/intrinsic/data.cc
+++ b/src/tint/lang/wgsl/intrinsic/data.cc
@@ -11015,6 +11015,17 @@
},
{
/* [547] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 0,
+ /* templates */ TemplateIndex(/* invalid */),
+ /* parameters */ ParameterIndex(224),
+ /* return_matcher_indices */ MatcherIndicesIndex(55),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [548] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -11025,7 +11036,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
- /* [548] */
+ /* [549] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -11036,7 +11047,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [549] */
+ /* [550] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -11047,7 +11058,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [550] */
+ /* [551] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -11058,7 +11069,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(97),
},
{
- /* [551] */
+ /* [552] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -12115,10 +12126,22 @@
},
{
/* [133] */
- /* fn __tint_materialize[T](T) -> T */
+ /* fn subgroupAll(bool) -> bool */
/* num overloads */ 1,
/* overloads */ OverloadIndex(547),
},
+ {
+ /* [134] */
+ /* fn subgroupAny(bool) -> bool */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(547),
+ },
+ {
+ /* [135] */
+ /* fn __tint_materialize[T](T) -> T */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(548),
+ },
};
constexpr IntrinsicInfo kUnaryOperators[] = {
@@ -12126,13 +12149,13 @@
/* [0] */
/* op &[S : address_space, T, A : access](ref<S, T, A>) -> ptr<S, T, A> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(548),
+ /* overloads */ OverloadIndex(549),
},
{
/* [1] */
/* op *[S : address_space, T, A : access](ptr<S, T, A>) -> ref<S, T, A> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(549),
+ /* overloads */ OverloadIndex(550),
},
{
/* [2] */
@@ -12244,13 +12267,13 @@
/* [8] */
/* op &&(bool, bool) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(550),
+ /* overloads */ OverloadIndex(551),
},
{
/* [9] */
/* op ||(bool, bool) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(551),
+ /* overloads */ OverloadIndex(552),
},
{
/* [10] */
diff --git a/src/tint/lang/wgsl/reader/lower/lower.cc b/src/tint/lang/wgsl/reader/lower/lower.cc
index 7086ca0..c8c2964 100644
--- a/src/tint/lang/wgsl/reader/lower/lower.cc
+++ b/src/tint/lang/wgsl/reader/lower/lower.cc
@@ -177,6 +177,8 @@
CASE(kSubgroupXor)
CASE(kSubgroupMin)
CASE(kSubgroupMax)
+ CASE(kSubgroupAll)
+ CASE(kSubgroupAny)
case tint::wgsl::BuiltinFn::kBitcast: // should lower to ir::Bitcast
case tint::wgsl::BuiltinFn::kWorkgroupUniformLoad: // should be handled in Lower()
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index 2eaf3aa..8e9320a 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -640,6 +640,8 @@
@must_use @stage("fragment", "compute") fn subgroupMin[N: num, T: fiu32_f16](value: vec<N, T>) -> vec<N, T>
@must_use @stage("fragment", "compute") fn subgroupMax[T: fiu32_f16](value: T) -> T
@must_use @stage("fragment", "compute") fn subgroupMax[N: num, T: fiu32_f16](value: vec<N, T>) -> vec<N, T>
+@must_use @stage("fragment", "compute") fn subgroupAll(bool) -> bool
+@must_use @stage("fragment", "compute") fn subgroupAny(bool) -> bool
////////////////////////////////////////////////////////////////////////////////
// Value constructors //
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
index 4e6e8dd..5ea6cf7 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
@@ -655,6 +655,8 @@
case wgsl::BuiltinFn::kSubgroupXor:
case wgsl::BuiltinFn::kSubgroupMin:
case wgsl::BuiltinFn::kSubgroupMax:
+ case wgsl::BuiltinFn::kSubgroupAny:
+ case wgsl::BuiltinFn::kSubgroupAll:
Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
break;
default:
diff --git a/src/tint/lang/wgsl/writer/raise/raise.cc b/src/tint/lang/wgsl/writer/raise/raise.cc
index 0a5d37f..2d9da3c 100644
--- a/src/tint/lang/wgsl/writer/raise/raise.cc
+++ b/src/tint/lang/wgsl/writer/raise/raise.cc
@@ -181,6 +181,8 @@
CASE(kSubgroupXor)
CASE(kSubgroupMin)
CASE(kSubgroupMax)
+ CASE(kSubgroupAny)
+ CASE(kSubgroupAll)
case core::BuiltinFn::kNone:
break;
}
diff --git a/src/tint/utils/protos/ir/ir.proto b/src/tint/utils/protos/ir/ir.proto
index 6a7b2bf..d4a723f 100644
--- a/src/tint/utils/protos/ir/ir.proto
+++ b/src/tint/utils/protos/ir/ir.proto
@@ -641,4 +641,6 @@
subgroup_xor = 128;
subgroup_min = 129;
subgroup_max = 130;
+ subgroup_any = 131;
+ subgroup_all = 132;
}
diff --git a/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl
new file mode 100644
index 0000000..0a8ccde
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl
@@ -0,0 +1,59 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// [hlsl-dxc] flags: --hlsl_shader_model 60
+
+
+enable subgroups;
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+
+
+// fn subgroupAll(bool) -> bool
+fn subgroupAll_c962bd() -> i32{
+ var res: bool = subgroupAll(true);
+ return select(0, 1, all(res == bool()));
+}
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..87d9c8f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.dxc.hlsl
@@ -0,0 +1,17 @@
+RWByteAddressBuffer prevent_dce : register(u0);
+
+int subgroupAll_c962bd() {
+ bool res = WaveActiveAllTrue(true);
+ return (all((res == false)) ? 1 : 0);
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..8df1464
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not support before SM6.0
+
+RWByteAddressBuffer prevent_dce : register(u0);
+
+int subgroupAll_c962bd() {
+ bool res = WaveActiveAllTrue(true);
+ return (all((res == false)) ? 1 : 0);
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.glsl b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.glsl
new file mode 100644
index 0000000..4938fa4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.glsl
@@ -0,0 +1,50 @@
+SKIP: FAILED
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAll_c962bd() -> i32 {
+ var res : bool = subgroupAll(true);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl:41:8 error: GLSL backend does not support extension 'subgroups'
+enable subgroups;
+ ^^^^^^^^^
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAll_c962bd() -> i32 {
+ var res : bool = subgroupAll(true);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl:41:8 error: GLSL backend does not support extension 'subgroups'
+enable subgroups;
+ ^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..133b98b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,16 @@
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int subgroupAll_c962bd() {
+ bool res = WaveActiveAllTrue(true);
+ return ((all((res == false))) ? (1) : (0));
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.ir.msl
new file mode 100644
index 0000000..cf0c117
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.ir.msl
@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ device int* prevent_dce;
+};
+
+int subgroupAll_c962bd() {
+ bool res = simd_all(true);
+ return select(0, 1, all((res == false)));
+}
+
+fragment void fragment_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupAll_c962bd();
+}
+
+kernel void compute_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupAll_c962bd();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.ir.spvasm b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..27385cb
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.ir.spvasm
@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpCapability GroupNonUniformVote
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %subgroupAll_c962bd "subgroupAll_c962bd"
+ OpName %res "res"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ %int = OpTypeInt 32 1
+%tint_symbol_1 = OpTypeStruct %int
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %1 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %6 = OpTypeFunction %int
+ %bool = OpTypeBool
+ %uint = OpTypeInt 32 0
+ %uint_3 = OpConstant %uint 3
+ %true = OpConstantTrue %bool
+%_ptr_Function_bool = OpTypePointer Function %bool
+ %false = OpConstantFalse %bool
+ %int_1 = OpConstant %int 1
+ %int_0 = OpConstant %int 0
+ %void = OpTypeVoid
+ %23 = OpTypeFunction %void
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %uint_0 = OpConstant %uint 0
+%subgroupAll_c962bd = OpFunction %int None %6
+ %7 = OpLabel
+ %res = OpVariable %_ptr_Function_bool Function
+ %8 = OpGroupNonUniformAll %bool %uint_3 %true
+ OpStore %res %8
+ %15 = OpLoad %bool %res
+ %16 = OpLogicalEqual %bool %15 %false
+ %18 = OpSelect %int %16 %int_1 %int_0
+ OpReturnValue %18
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %23
+ %24 = OpLabel
+ %25 = OpFunctionCall %int %subgroupAll_c962bd
+ %26 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0
+ OpStore %26 %25
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %23
+ %30 = OpLabel
+ %31 = OpFunctionCall %int %subgroupAll_c962bd
+ %32 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0
+ OpStore %32 %31
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.msl
new file mode 100644
index 0000000..e80e901
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.msl
@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int subgroupAll_c962bd() {
+ bool res = simd_all(true);
+ return select(0, 1, all((res == false)));
+}
+
+fragment void fragment_main(device int* tint_symbol [[buffer(0)]]) {
+ *(tint_symbol) = subgroupAll_c962bd();
+ return;
+}
+
+kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]]) {
+ *(tint_symbol_1) = subgroupAll_c962bd();
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.spvasm
new file mode 100644
index 0000000..4892627
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAll_c962bd() -> i32 {
+ var res : bool = subgroupAll(true);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+Failed to generate:
diff --git a/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.wgsl
new file mode 100644
index 0000000..c2bb531
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAll/c962bd.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAll_c962bd() -> i32 {
+ var res : bool = subgroupAll(true);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl
new file mode 100644
index 0000000..1892aef
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl
@@ -0,0 +1,59 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// [hlsl-dxc] flags: --hlsl_shader_model 60
+
+
+enable subgroups;
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+
+
+// fn subgroupAny(bool) -> bool
+fn subgroupAny_cddda0() -> i32{
+ var res: bool = subgroupAny(true);
+ return select(0, 1, all(res == bool()));
+}
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..6644f72
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.dxc.hlsl
@@ -0,0 +1,17 @@
+RWByteAddressBuffer prevent_dce : register(u0);
+
+int subgroupAny_cddda0() {
+ bool res = WaveActiveAnyTrue(true);
+ return (all((res == false)) ? 1 : 0);
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..439e998
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.fxc.hlsl
@@ -0,0 +1,19 @@
+SKIP: Wave ops not support before SM6.0
+
+RWByteAddressBuffer prevent_dce : register(u0);
+
+int subgroupAny_cddda0() {
+ bool res = WaveActiveAnyTrue(true);
+ return (all((res == false)) ? 1 : 0);
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.glsl b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.glsl
new file mode 100644
index 0000000..065fa6dc
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.glsl
@@ -0,0 +1,50 @@
+SKIP: FAILED
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAny_cddda0() -> i32 {
+ var res : bool = subgroupAny(true);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl:41:8 error: GLSL backend does not support extension 'subgroups'
+enable subgroups;
+ ^^^^^^^^^
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAny_cddda0() -> i32 {
+ var res : bool = subgroupAny(true);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl:41:8 error: GLSL backend does not support extension 'subgroups'
+enable subgroups;
+ ^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..fb1d0d3
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,16 @@
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int subgroupAny_cddda0() {
+ bool res = WaveActiveAnyTrue(true);
+ return ((all((res == false))) ? (1) : (0));
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.ir.msl
new file mode 100644
index 0000000..8d04db3
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.ir.msl
@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ device int* prevent_dce;
+};
+
+int subgroupAny_cddda0() {
+ bool res = simd_any(true);
+ return select(0, 1, all((res == false)));
+}
+
+fragment void fragment_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupAny_cddda0();
+}
+
+kernel void compute_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupAny_cddda0();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.ir.spvasm b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..0c7e508
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.ir.spvasm
@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpCapability GroupNonUniformVote
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %subgroupAny_cddda0 "subgroupAny_cddda0"
+ OpName %res "res"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ %int = OpTypeInt 32 1
+%tint_symbol_1 = OpTypeStruct %int
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %1 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %6 = OpTypeFunction %int
+ %bool = OpTypeBool
+ %uint = OpTypeInt 32 0
+ %uint_3 = OpConstant %uint 3
+ %true = OpConstantTrue %bool
+%_ptr_Function_bool = OpTypePointer Function %bool
+ %false = OpConstantFalse %bool
+ %int_1 = OpConstant %int 1
+ %int_0 = OpConstant %int 0
+ %void = OpTypeVoid
+ %23 = OpTypeFunction %void
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %uint_0 = OpConstant %uint 0
+%subgroupAny_cddda0 = OpFunction %int None %6
+ %7 = OpLabel
+ %res = OpVariable %_ptr_Function_bool Function
+ %8 = OpGroupNonUniformAny %bool %uint_3 %true
+ OpStore %res %8
+ %15 = OpLoad %bool %res
+ %16 = OpLogicalEqual %bool %15 %false
+ %18 = OpSelect %int %16 %int_1 %int_0
+ OpReturnValue %18
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %23
+ %24 = OpLabel
+ %25 = OpFunctionCall %int %subgroupAny_cddda0
+ %26 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0
+ OpStore %26 %25
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %23
+ %30 = OpLabel
+ %31 = OpFunctionCall %int %subgroupAny_cddda0
+ %32 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0
+ OpStore %32 %31
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.msl
new file mode 100644
index 0000000..85ad2c2
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.msl
@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int subgroupAny_cddda0() {
+ bool res = simd_any(true);
+ return select(0, 1, all((res == false)));
+}
+
+fragment void fragment_main(device int* tint_symbol [[buffer(0)]]) {
+ *(tint_symbol) = subgroupAny_cddda0();
+ return;
+}
+
+kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]]) {
+ *(tint_symbol_1) = subgroupAny_cddda0();
+ return;
+}
+
diff --git a/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.spvasm
new file mode 100644
index 0000000..fde942c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAny_cddda0() -> i32 {
+ var res : bool = subgroupAny(true);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+Failed to generate:
diff --git a/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.wgsl
new file mode 100644
index 0000000..d8b2b7d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupAny/cddda0.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAny_cddda0() -> i32 {
+ var res : bool = subgroupAny(true);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
diff --git a/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl
new file mode 100644
index 0000000..444ee4e
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl
@@ -0,0 +1,60 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// [hlsl-dxc] flags: --hlsl_shader_model 60
+
+
+enable subgroups;
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+
+
+// fn subgroupAll(bool) -> bool
+fn subgroupAll_c962bd() -> i32{
+ var arg_0 = true;
+ var res: bool = subgroupAll(arg_0);
+ return select(0, 1, all(res == bool()));
+}
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
diff --git a/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..8ca14aa
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.dxc.hlsl
@@ -0,0 +1,18 @@
+RWByteAddressBuffer prevent_dce : register(u0);
+
+int subgroupAll_c962bd() {
+ bool arg_0 = true;
+ bool res = WaveActiveAllTrue(arg_0);
+ return (all((res == false)) ? 1 : 0);
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+ return;
+}
diff --git a/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..a195482
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.fxc.hlsl
@@ -0,0 +1,20 @@
+SKIP: Wave ops not support before SM6.0
+
+RWByteAddressBuffer prevent_dce : register(u0);
+
+int subgroupAll_c962bd() {
+ bool arg_0 = true;
+ bool res = WaveActiveAllTrue(arg_0);
+ return (all((res == false)) ? 1 : 0);
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+ return;
+}
diff --git a/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.glsl b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.glsl
new file mode 100644
index 0000000..c1812ca
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.glsl
@@ -0,0 +1,52 @@
+SKIP: FAILED
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAll_c962bd() -> i32 {
+ var arg_0 = true;
+ var res : bool = subgroupAll(arg_0);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl:41:8 error: GLSL backend does not support extension 'subgroups'
+enable subgroups;
+ ^^^^^^^^^
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAll_c962bd() -> i32 {
+ var arg_0 = true;
+ var res : bool = subgroupAll(arg_0);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl:41:8 error: GLSL backend does not support extension 'subgroups'
+enable subgroups;
+ ^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..e1ad543
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,17 @@
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int subgroupAll_c962bd() {
+ bool arg_0 = true;
+ bool res = WaveActiveAllTrue(arg_0);
+ return ((all((res == false))) ? (1) : (0));
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAll_c962bd()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.ir.msl
new file mode 100644
index 0000000..203d32a
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.ir.msl
@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ device int* prevent_dce;
+};
+
+int subgroupAll_c962bd() {
+ bool arg_0 = true;
+ bool res = simd_all(arg_0);
+ return select(0, 1, all((res == false)));
+}
+
+fragment void fragment_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupAll_c962bd();
+}
+
+kernel void compute_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupAll_c962bd();
+}
diff --git a/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.ir.spvasm b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..a812588
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.ir.spvasm
@@ -0,0 +1,67 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 35
+; Schema: 0
+ OpCapability Shader
+ OpCapability GroupNonUniformVote
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %subgroupAll_c962bd "subgroupAll_c962bd"
+ OpName %arg_0 "arg_0"
+ OpName %res "res"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ %int = OpTypeInt 32 1
+%tint_symbol_1 = OpTypeStruct %int
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %1 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %6 = OpTypeFunction %int
+ %bool = OpTypeBool
+%_ptr_Function_bool = OpTypePointer Function %bool
+ %true = OpConstantTrue %bool
+ %uint = OpTypeInt 32 0
+ %uint_3 = OpConstant %uint 3
+ %false = OpConstantFalse %bool
+ %int_1 = OpConstant %int 1
+ %int_0 = OpConstant %int 0
+ %void = OpTypeVoid
+ %25 = OpTypeFunction %void
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %uint_0 = OpConstant %uint 0
+%subgroupAll_c962bd = OpFunction %int None %6
+ %7 = OpLabel
+ %arg_0 = OpVariable %_ptr_Function_bool Function
+ %res = OpVariable %_ptr_Function_bool Function
+ OpStore %arg_0 %true
+ %12 = OpLoad %bool %arg_0
+ %13 = OpGroupNonUniformAll %bool %uint_3 %12
+ OpStore %res %13
+ %17 = OpLoad %bool %res
+ %18 = OpLogicalEqual %bool %17 %false
+ %20 = OpSelect %int %18 %int_1 %int_0
+ OpReturnValue %20
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %25
+ %26 = OpLabel
+ %27 = OpFunctionCall %int %subgroupAll_c962bd
+ %28 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0
+ OpStore %28 %27
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %25
+ %32 = OpLabel
+ %33 = OpFunctionCall %int %subgroupAll_c962bd
+ %34 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0
+ OpStore %34 %33
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.msl
new file mode 100644
index 0000000..61778bb
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.msl
@@ -0,0 +1,19 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int subgroupAll_c962bd() {
+ bool arg_0 = true;
+ bool res = simd_all(arg_0);
+ return select(0, 1, all((res == false)));
+}
+
+fragment void fragment_main(device int* tint_symbol [[buffer(0)]]) {
+ *(tint_symbol) = subgroupAll_c962bd();
+ return;
+}
+
+kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]]) {
+ *(tint_symbol_1) = subgroupAll_c962bd();
+ return;
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.spvasm b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.spvasm
new file mode 100644
index 0000000..48684d2
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.spvasm
@@ -0,0 +1,24 @@
+SKIP: FAILED
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAll_c962bd() -> i32 {
+ var arg_0 = true;
+ var res : bool = subgroupAll(arg_0);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+Failed to generate:
diff --git a/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.wgsl
new file mode 100644
index 0000000..7588f0c
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAll/c962bd.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAll_c962bd() -> i32 {
+ var arg_0 = true;
+ var res : bool = subgroupAll(arg_0);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAll_c962bd();
+}
diff --git a/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl
new file mode 100644
index 0000000..f28250f
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl
@@ -0,0 +1,60 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// [hlsl-dxc] flags: --hlsl_shader_model 60
+
+
+enable subgroups;
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+
+
+// fn subgroupAny(bool) -> bool
+fn subgroupAny_cddda0() -> i32{
+ var arg_0 = true;
+ var res: bool = subgroupAny(arg_0);
+ return select(0, 1, all(res == bool()));
+}
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
diff --git a/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..c66324b
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.dxc.hlsl
@@ -0,0 +1,18 @@
+RWByteAddressBuffer prevent_dce : register(u0);
+
+int subgroupAny_cddda0() {
+ bool arg_0 = true;
+ bool res = WaveActiveAnyTrue(arg_0);
+ return (all((res == false)) ? 1 : 0);
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+ return;
+}
diff --git a/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..5f24ce8
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.fxc.hlsl
@@ -0,0 +1,20 @@
+SKIP: Wave ops not support before SM6.0
+
+RWByteAddressBuffer prevent_dce : register(u0);
+
+int subgroupAny_cddda0() {
+ bool arg_0 = true;
+ bool res = WaveActiveAnyTrue(arg_0);
+ return (all((res == false)) ? 1 : 0);
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+ return;
+}
diff --git a/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.glsl b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.glsl
new file mode 100644
index 0000000..b7bd6bb
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.glsl
@@ -0,0 +1,52 @@
+SKIP: FAILED
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAny_cddda0() -> i32 {
+ var arg_0 = true;
+ var res : bool = subgroupAny(arg_0);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl:41:8 error: GLSL backend does not support extension 'subgroups'
+enable subgroups;
+ ^^^^^^^^^
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAny_cddda0() -> i32 {
+ var arg_0 = true;
+ var res : bool = subgroupAny(arg_0);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl:41:8 error: GLSL backend does not support extension 'subgroups'
+enable subgroups;
+ ^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..f23eee6
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,17 @@
+
+RWByteAddressBuffer prevent_dce : register(u0);
+int subgroupAny_cddda0() {
+ bool arg_0 = true;
+ bool res = WaveActiveAnyTrue(arg_0);
+ return ((all((res == false))) ? (1) : (0));
+}
+
+void fragment_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ prevent_dce.Store(0u, asuint(subgroupAny_cddda0()));
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.ir.msl
new file mode 100644
index 0000000..08076fc
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.ir.msl
@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ device int* prevent_dce;
+};
+
+int subgroupAny_cddda0() {
+ bool arg_0 = true;
+ bool res = simd_any(arg_0);
+ return select(0, 1, all((res == false)));
+}
+
+fragment void fragment_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupAny_cddda0();
+}
+
+kernel void compute_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = subgroupAny_cddda0();
+}
diff --git a/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.ir.spvasm b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..fbf9b15
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.ir.spvasm
@@ -0,0 +1,67 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 35
+; Schema: 0
+ OpCapability Shader
+ OpCapability GroupNonUniformVote
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %subgroupAny_cddda0 "subgroupAny_cddda0"
+ OpName %arg_0 "arg_0"
+ OpName %res "res"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ %int = OpTypeInt 32 1
+%tint_symbol_1 = OpTypeStruct %int
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %1 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %6 = OpTypeFunction %int
+ %bool = OpTypeBool
+%_ptr_Function_bool = OpTypePointer Function %bool
+ %true = OpConstantTrue %bool
+ %uint = OpTypeInt 32 0
+ %uint_3 = OpConstant %uint 3
+ %false = OpConstantFalse %bool
+ %int_1 = OpConstant %int 1
+ %int_0 = OpConstant %int 0
+ %void = OpTypeVoid
+ %25 = OpTypeFunction %void
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %uint_0 = OpConstant %uint 0
+%subgroupAny_cddda0 = OpFunction %int None %6
+ %7 = OpLabel
+ %arg_0 = OpVariable %_ptr_Function_bool Function
+ %res = OpVariable %_ptr_Function_bool Function
+ OpStore %arg_0 %true
+ %12 = OpLoad %bool %arg_0
+ %13 = OpGroupNonUniformAny %bool %uint_3 %12
+ OpStore %res %13
+ %17 = OpLoad %bool %res
+ %18 = OpLogicalEqual %bool %17 %false
+ %20 = OpSelect %int %18 %int_1 %int_0
+ OpReturnValue %20
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %25
+ %26 = OpLabel
+ %27 = OpFunctionCall %int %subgroupAny_cddda0
+ %28 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0
+ OpStore %28 %27
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %25
+ %32 = OpLabel
+ %33 = OpFunctionCall %int %subgroupAny_cddda0
+ %34 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0
+ OpStore %34 %33
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.msl
new file mode 100644
index 0000000..76aa7ce
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.msl
@@ -0,0 +1,19 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int subgroupAny_cddda0() {
+ bool arg_0 = true;
+ bool res = simd_any(arg_0);
+ return select(0, 1, all((res == false)));
+}
+
+fragment void fragment_main(device int* tint_symbol [[buffer(0)]]) {
+ *(tint_symbol) = subgroupAny_cddda0();
+ return;
+}
+
+kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]]) {
+ *(tint_symbol_1) = subgroupAny_cddda0();
+ return;
+}
+
diff --git a/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.spvasm b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.spvasm
new file mode 100644
index 0000000..24548a7
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.spvasm
@@ -0,0 +1,24 @@
+SKIP: FAILED
+
+
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAny_cddda0() -> i32 {
+ var arg_0 = true;
+ var res : bool = subgroupAny(arg_0);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+Failed to generate:
diff --git a/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.wgsl
new file mode 100644
index 0000000..5119940
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupAny/cddda0.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+enable subgroups;
+
+@group(0) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+fn subgroupAny_cddda0() -> i32 {
+ var arg_0 = true;
+ var res : bool = subgroupAny(arg_0);
+ return select(0, 1, all((res == bool())));
+}
+
+@fragment
+fn fragment_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ prevent_dce = subgroupAny_cddda0();
+}
