[hlsl-writer] Add support for subgroupBallot
Generate a call to `WaveActiveBallot` with a `true` predicate.
Bug: tint:2000
Change-Id: I05e8a9fe8a923853bdab60e308ba48dfa32e63ae
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/143836
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
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 78ac0e1..88fd9c6 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -340,6 +340,7 @@
core::Extension::kChromiumExperimentalDp4A,
core::Extension::kChromiumExperimentalFullPtrParameters,
core::Extension::kChromiumExperimentalPushConstant,
+ core::Extension::kChromiumExperimentalSubgroups,
core::Extension::kF16,
core::Extension::kChromiumInternalDualSourceBlending,
})) {
@@ -1198,6 +1199,9 @@
if (builtin->IsDP4a()) {
return EmitDP4aCall(out, expr, builtin);
}
+ if (builtin->IsSubgroup()) {
+ return EmitSubgroupCall(out, expr, builtin);
+ }
auto name = generate_builtin_name(builtin);
if (name.empty()) {
@@ -2485,6 +2489,18 @@
return true;
}
+bool ASTPrinter::EmitSubgroupCall(StringStream& out,
+ [[maybe_unused]] const ast::CallExpression* expr,
+ const sem::Builtin* builtin) {
+ if (builtin->Type() == core::Function::kSubgroupBallot) {
+ out << "WaveActiveBallot(true)";
+ } else {
+ TINT_UNREACHABLE() << "unexpected subgroup builtin type " << core::str(builtin->Type());
+ return false;
+ }
+ return true;
+}
+
bool ASTPrinter::EmitTextureCall(StringStream& out,
const sem::Call* call,
const sem::Builtin* builtin) {
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.h b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.h
index 8ca7e4f..188223c 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.h
@@ -282,6 +282,14 @@
bool EmitDP4aCall(StringStream& out,
const ast::CallExpression* expr,
const sem::Builtin* builtin);
+ /// Handles generating a call to subgroup builtins.
+ /// @param out the output of the expression stream
+ /// @param expr the call expression
+ /// @param builtin the semantic information for the builtin
+ /// @returns true if the call expression is emitted
+ bool EmitSubgroupCall(StringStream& out,
+ const ast::CallExpression* expr,
+ const sem::Builtin* builtin);
/// Handles a case statement
/// @param s the switch statement
/// @param case_idx the index of the switch case in the switch statement
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
index fe24933..3064f17 100644
--- a/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl
@@ -1,21 +1,12 @@
-SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
-
-enable chromium_experimental_subgroups;
-
-fn subgroupBallot_7e6d0e() {
- var res : vec4<u32> = subgroupBallot();
- prevent_dce = res;
+void subgroupBallot_7e6d0e() {
+ uint4 res = WaveActiveBallot(true);
+ prevent_dce.Store4(0u, asuint(res));
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
-
-@compute @workgroup_size(1)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
subgroupBallot_7e6d0e();
+ return;
}
-
-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
index fe24933..bfcf146 100644
--- a/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/subgroupBallot/7e6d0e.wgsl.expected.fxc.hlsl
@@ -1,21 +1,14 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
-enable chromium_experimental_subgroups;
-
-fn subgroupBallot_7e6d0e() {
- var res : vec4<u32> = subgroupBallot();
- prevent_dce = res;
+void subgroupBallot_7e6d0e() {
+ uint4 res = WaveActiveBallot(true);
+ prevent_dce.Store4(0u, asuint(res));
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
-
-@compute @workgroup_size(1)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
subgroupBallot_7e6d0e();
+ return;
}
-
-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/var/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl
index dc7c411..3064f17 100644
--- a/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.dxc.hlsl
@@ -1,21 +1,12 @@
-SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
-
-enable chromium_experimental_subgroups;
-
-fn subgroupBallot_7e6d0e() {
- var res : vec4<u32> = subgroupBallot();
- prevent_dce = res;
+void subgroupBallot_7e6d0e() {
+ uint4 res = WaveActiveBallot(true);
+ prevent_dce.Store4(0u, asuint(res));
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
-
-@compute @workgroup_size(1)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
subgroupBallot_7e6d0e();
+ return;
}
-
-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
index dc7c411..bfcf146 100644
--- a/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/subgroupBallot/7e6d0e.wgsl.expected.fxc.hlsl
@@ -1,21 +1,14 @@
SKIP: FAILED
+RWByteAddressBuffer prevent_dce : register(u0, space2);
-enable chromium_experimental_subgroups;
-
-fn subgroupBallot_7e6d0e() {
- var res : vec4<u32> = subgroupBallot();
- prevent_dce = res;
+void subgroupBallot_7e6d0e() {
+ uint4 res = WaveActiveBallot(true);
+ prevent_dce.Store4(0u, asuint(res));
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
-
-@compute @workgroup_size(1)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
subgroupBallot_7e6d0e();
+ return;
}
-
-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;
- ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-
