[tint] Add subgroupBallot to the intrinsics table
Not yet supported by any backend.
Bug: tint:2000
Change-Id: I9292933f9653928de01629e4c23c7c4fe2c4cf38
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/143833
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index 4e9b03a..05ecb45 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -883,6 +883,8 @@
@stage("fragment", "compute") fn atomicExchange<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
@stage("fragment", "compute") fn atomicCompareExchangeWeak<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T>
+@must_use @stage("compute") fn subgroupBallot() -> vec4<u32>
+
////////////////////////////////////////////////////////////////////////////////
// Value constructors //
////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/core/builtin/function.cc b/src/tint/lang/core/builtin/function.cc
index 3d67b07..8d61e8a 100644
--- a/src/tint/lang/core/builtin/function.cc
+++ b/src/tint/lang/core/builtin/function.cc
@@ -364,6 +364,9 @@
if (name == "atomicCompareExchangeWeak") {
return Function::kAtomicCompareExchangeWeak;
}
+ if (name == "subgroupBallot") {
+ return Function::kSubgroupBallot;
+ }
if (name == "_tint_materialize") {
return Function::kTintMaterialize;
}
@@ -600,6 +603,8 @@
return "atomicExchange";
case Function::kAtomicCompareExchangeWeak:
return "atomicCompareExchangeWeak";
+ case Function::kSubgroupBallot:
+ return "subgroupBallot";
case Function::kTintMaterialize:
return "_tint_materialize";
}
@@ -666,6 +671,10 @@
return f == Function::kDot4I8Packed || f == Function::kDot4U8Packed;
}
+bool IsSubgroupBuiltin(Function f) {
+ return f == Function::kSubgroupBallot;
+}
+
bool HasSideEffects(Function f) {
switch (f) {
case Function::kAtomicAdd:
diff --git a/src/tint/lang/core/builtin/function.cc.tmpl b/src/tint/lang/core/builtin/function.cc.tmpl
index 3e765d5..656c541 100644
--- a/src/tint/lang/core/builtin/function.cc.tmpl
+++ b/src/tint/lang/core/builtin/function.cc.tmpl
@@ -102,6 +102,10 @@
return f == Function::kDot4I8Packed || f == Function::kDot4U8Packed;
}
+bool IsSubgroupBuiltin(Function f) {
+ return f == Function::kSubgroupBallot;
+}
+
bool HasSideEffects(Function f) {
switch (f) {
case Function::kAtomicAdd:
diff --git a/src/tint/lang/core/builtin/function.h b/src/tint/lang/core/builtin/function.h
index 1ead560..ee04853 100644
--- a/src/tint/lang/core/builtin/function.h
+++ b/src/tint/lang/core/builtin/function.h
@@ -146,6 +146,7 @@
kAtomicXor,
kAtomicExchange,
kAtomicCompareExchangeWeak,
+ kSubgroupBallot,
kTintMaterialize,
};
@@ -281,6 +282,7 @@
Function::kAtomicXor,
Function::kAtomicExchange,
Function::kAtomicCompareExchangeWeak,
+ Function::kSubgroupBallot,
Function::kTintMaterialize,
};
@@ -399,6 +401,7 @@
"atomicXor",
"atomicExchange",
"atomicCompareExchangeWeak",
+ "subgroupBallot",
"_tint_materialize",
};
@@ -452,6 +455,11 @@
/// @returns true if the given `f` is a DP4a builtin
bool IsDP4aBuiltin(Function f);
+/// Determines if the given `f` is a subgroup builtin.
+/// @param f the builtin type
+/// @returns true if the given `f` is a subgroup builtin
+bool IsSubgroupBuiltin(Function f);
+
/// Determines if the given `f` may have side-effects (i.e. writes to at least one of its inputs)
/// @returns true if intrinsic may have side-effects
bool HasSideEffects(Function f);
diff --git a/src/tint/lang/core/builtin/function.h.tmpl b/src/tint/lang/core/builtin/function.h.tmpl
index d6b8fee..f43ab7a 100644
--- a/src/tint/lang/core/builtin/function.h.tmpl
+++ b/src/tint/lang/core/builtin/function.h.tmpl
@@ -110,6 +110,11 @@
/// @returns true if the given `f` is a DP4a builtin
bool IsDP4aBuiltin(Function f);
+/// Determines if the given `f` is a subgroup builtin.
+/// @param f the builtin type
+/// @returns true if the given `f` is a subgroup builtin
+bool IsSubgroupBuiltin(Function f);
+
/// Determines if the given `f` may have side-effects (i.e. writes to at least one of its inputs)
/// @returns true if intrinsic may have side-effects
bool HasSideEffects(Function f);
diff --git a/src/tint/lang/wgsl/resolver/builtin_validation_test.cc b/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
index 06b3cfe..df1c5cb 100644
--- a/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
@@ -654,5 +654,31 @@
R"(error: workgroupUniformLoad must not be called with an argument that contains an atomic type)");
}
+TEST_F(ResolverBuiltinValidationTest, SubgroupBallotWithoutExtension) {
+ // fn func { return subgroupBallot(); }
+ Func("func", tint::Empty, ty.vec4<u32>(),
+ Vector{
+ Return(Call(Source{Source::Location{12, 34}}, "subgroupBallot")),
+ });
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(
+ r()->error(),
+ R"(12:34 error: cannot call built-in function 'subgroupBallot' without extension chromium_experimental_subgroups)");
+}
+
+TEST_F(ResolverBuiltinValidationTest, SubgroupBallotWithExtension) {
+ // enable chromium_experimental_subgroups;
+ // fn func { return subgroupBallot(); }
+ Enable(builtin::Extension::kChromiumExperimentalSubgroups);
+
+ Func("func", tint::Empty, ty.vec4<u32>(),
+ Vector{
+ Return(Call(Source{Source::Location{12, 34}}, "subgroupBallot")),
+ });
+
+ EXPECT_TRUE(r()->Resolve());
+}
+
} // namespace
} // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/resolver/intrinsic_table.inl b/src/tint/lang/wgsl/resolver/intrinsic_table.inl
index 1e01f78..3df3089 100644
--- a/src/tint/lang/wgsl/resolver/intrinsic_table.inl
+++ b/src/tint/lang/wgsl/resolver/intrinsic_table.inl
@@ -14064,6 +14064,18 @@
},
{
/* [471] */
+ /* num parameters */ 0,
+ /* num template types */ 0,
+ /* num template numbers */ 0,
+ /* template types */ &kTemplateTypes[37],
+ /* template numbers */ &kTemplateNumbers[10],
+ /* parameters */ &kParameters[1014],
+ /* return matcher indices */ &kMatcherIndices[156],
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
+ /* const eval */ nullptr,
+ },
+ {
+ /* [472] */
/* num parameters */ 1,
/* num template types */ 1,
/* num template numbers */ 0,
@@ -14075,7 +14087,7 @@
/* const eval */ &ConstEval::Identity,
},
{
- /* [472] */
+ /* [473] */
/* num parameters */ 2,
/* num template types */ 0,
/* num template numbers */ 0,
@@ -14087,7 +14099,7 @@
/* const eval */ &ConstEval::OpLogicalAnd,
},
{
- /* [473] */
+ /* [474] */
/* num parameters */ 2,
/* num template types */ 0,
/* num template numbers */ 0,
@@ -14099,7 +14111,7 @@
/* const eval */ &ConstEval::OpLogicalOr,
},
{
- /* [474] */
+ /* [475] */
/* num parameters */ 1,
/* num template types */ 1,
/* num template numbers */ 0,
@@ -14985,10 +14997,16 @@
},
{
/* [113] */
- /* fn _tint_materialize<T>(T) -> T */
+ /* fn subgroupBallot() -> vec4<u32> */
/* num overloads */ 1,
/* overloads */ &kOverloads[471],
},
+ {
+ /* [114] */
+ /* fn _tint_materialize<T>(T) -> T */
+ /* num overloads */ 1,
+ /* overloads */ &kOverloads[472],
+ },
};
constexpr IntrinsicInfo kUnaryOperators[] = {
@@ -15100,13 +15118,13 @@
/* [8] */
/* op &&(bool, bool) -> bool */
/* num overloads */ 1,
- /* overloads */ &kOverloads[472],
+ /* overloads */ &kOverloads[473],
},
{
/* [9] */
/* op ||(bool, bool) -> bool */
/* num overloads */ 1,
- /* overloads */ &kOverloads[473],
+ /* overloads */ &kOverloads[474],
},
{
/* [10] */
@@ -15381,7 +15399,7 @@
/* [17] */
/* conv packedVec3<T : concrete_scalar>(vec3<T>) -> packedVec3<T> */
/* num overloads */ 1,
- /* overloads */ &kOverloads[474],
+ /* overloads */ &kOverloads[475],
},
};
diff --git a/src/tint/lang/wgsl/sem/builtin.cc b/src/tint/lang/wgsl/sem/builtin.cc
index ee1fe17..36aa4fc 100644
--- a/src/tint/lang/wgsl/sem/builtin.cc
+++ b/src/tint/lang/wgsl/sem/builtin.cc
@@ -84,6 +84,10 @@
return IsDP4aBuiltin(type_);
}
+bool Builtin::IsSubgroup() const {
+ return IsSubgroupBuiltin(type_);
+}
+
bool Builtin::HasSideEffects() const {
return builtin::HasSideEffects(type_);
}
@@ -92,6 +96,9 @@
if (IsDP4a()) {
return builtin::Extension::kChromiumExperimentalDp4A;
}
+ if (IsSubgroup()) {
+ return builtin::Extension::kChromiumExperimentalSubgroups;
+ }
return builtin::Extension::kUndefined;
}
diff --git a/src/tint/lang/wgsl/sem/builtin.h b/src/tint/lang/wgsl/sem/builtin.h
index dfea716..0a7e8eb 100644
--- a/src/tint/lang/wgsl/sem/builtin.h
+++ b/src/tint/lang/wgsl/sem/builtin.h
@@ -92,6 +92,10 @@
/// chromium_experimental_DP4a)
bool IsDP4a() const;
+ /// @returns true if builtin is a subgroup builtin (defined in the extension
+ /// chromium_experimental_subgroups)
+ bool IsSubgroup() const;
+
/// @returns true if intrinsic may have side-effects (i.e. writes to at least
/// one of its inputs)
bool HasSideEffects() const;
diff --git a/test/tint/builtins/gen/gen.wgsl.tmpl b/test/tint/builtins/gen/gen.wgsl.tmpl
index 4631079..b1e9670 100644
--- a/test/tint/builtins/gen/gen.wgsl.tmpl
+++ b/test/tint/builtins/gen/gen.wgsl.tmpl
@@ -200,6 +200,11 @@
enable chromium_experimental_dp4a;
{{ end -}}
+{{- /* Check and emit chromium_experimental_subgroups */ -}}
+{{- if (eq "subgroupBallot" $builtin_name)}}
+enable chromium_experimental_subgroups;
+{{ end -}}
+
{{- /* Check and emit f16 */ -}}
{{- if OverloadUsesF16 $overload}}
enable f16;
diff --git a/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl
new file mode 100644
index 0000000..b522e95
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl
@@ -0,0 +1,36 @@
+// Copyright 2023 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/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable chromium_experimental_subgroups;
+
+// fn subgroupBallot() -> vec4<u32>
+fn subgroupBallot_7e6d0e() {
+ var res: vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
diff --git a/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..fe24933
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl
@@ -0,0 +1,21 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
+
+Failed to generate: builtins/gen/literal/subgroupBallot/7e6d0e.wgsl:24:8 error: HLSL backend does not support extension 'chromium_experimental_subgroups'
+enable chromium_experimental_subgroups;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..fe24933
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.fxc.hlsl
@@ -0,0 +1,21 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
+
+Failed to generate: builtins/gen/literal/subgroupBallot/7e6d0e.wgsl:24:8 error: HLSL backend does not support extension 'chromium_experimental_subgroups'
+enable chromium_experimental_subgroups;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.glsl b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.glsl
new file mode 100644
index 0000000..bb208e9
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
+
+Failed to generate: error: Unknown builtin method: subgroupBallot
diff --git a/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.msl
new file mode 100644
index 0000000..134107e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.msl
@@ -0,0 +1,21 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
+
+Failed to generate: builtins/gen/literal/subgroupBallot/7e6d0e.wgsl:24:8 error: MSL backend does not support extension 'chromium_experimental_subgroups'
+enable chromium_experimental_subgroups;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.spvasm
new file mode 100644
index 0000000..9aa8242
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.spvasm
@@ -0,0 +1,21 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
+
+Failed to generate: builtins/gen/literal/subgroupBallot/7e6d0e.wgsl:24:8 error: SPIR-V backend does not support extension 'chromium_experimental_subgroups'
+enable chromium_experimental_subgroups;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.wgsl
new file mode 100644
index 0000000..0f3dedb
--- /dev/null
+++ b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.wgsl
@@ -0,0 +1,13 @@
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
diff --git a/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl
new file mode 100644
index 0000000..b522e95
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl
@@ -0,0 +1,36 @@
+// Copyright 2023 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/src/cmd/gen
+// using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable chromium_experimental_subgroups;
+
+// fn subgroupBallot() -> vec4<u32>
+fn subgroupBallot_7e6d0e() {
+ var res: vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
diff --git a/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..dc7c411
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl
@@ -0,0 +1,21 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
+
+Failed to generate: builtins/gen/var/subgroupBallot/7e6d0e.wgsl:24:8 error: HLSL backend does not support extension 'chromium_experimental_subgroups'
+enable chromium_experimental_subgroups;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..dc7c411
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.fxc.hlsl
@@ -0,0 +1,21 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
+
+Failed to generate: builtins/gen/var/subgroupBallot/7e6d0e.wgsl:24:8 error: HLSL backend does not support extension 'chromium_experimental_subgroups'
+enable chromium_experimental_subgroups;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.glsl b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.glsl
new file mode 100644
index 0000000..bb208e9
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
+
+Failed to generate: error: Unknown builtin method: subgroupBallot
diff --git a/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.msl
new file mode 100644
index 0000000..b16ab89
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.msl
@@ -0,0 +1,21 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
+
+Failed to generate: builtins/gen/var/subgroupBallot/7e6d0e.wgsl:24:8 error: MSL backend does not support extension 'chromium_experimental_subgroups'
+enable chromium_experimental_subgroups;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.spvasm b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.spvasm
new file mode 100644
index 0000000..698959f
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.spvasm
@@ -0,0 +1,21 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
+
+Failed to generate: builtins/gen/var/subgroupBallot/7e6d0e.wgsl:24:8 error: SPIR-V backend does not support extension 'chromium_experimental_subgroups'
+enable chromium_experimental_subgroups;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.wgsl
new file mode 100644
index 0000000..0f3dedb
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.wgsl
@@ -0,0 +1,13 @@
+enable chromium_experimental_subgroups;
+
+fn subgroupBallot_7e6d0e() {
+ var res : vec4<u32> = subgroupBallot();
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBallot_7e6d0e();
+}
