[tint] Add subgroup builtin inputs
Adds parsing and SPIR-V codegen for subgroup_invocation_id and
subgroup_size.
Bug: tint:2000
Change-Id: Idcb9e2011c8c6129c5a0e5f0564048c14f8625a6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/143835
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index 05ecb45..0774687 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -37,6 +37,8 @@
num_workgroups
sample_index
sample_mask
+ subgroup_invocation_id
+ subgroup_size
__point_size
}
diff --git a/src/tint/lang/core/builtin/builtin_value.cc b/src/tint/lang/core/builtin/builtin_value.cc
index 644a217..f11b797 100644
--- a/src/tint/lang/core/builtin/builtin_value.cc
+++ b/src/tint/lang/core/builtin/builtin_value.cc
@@ -61,6 +61,12 @@
if (str == "sample_mask") {
return BuiltinValue::kSampleMask;
}
+ if (str == "subgroup_invocation_id") {
+ return BuiltinValue::kSubgroupInvocationId;
+ }
+ if (str == "subgroup_size") {
+ return BuiltinValue::kSubgroupSize;
+ }
if (str == "vertex_index") {
return BuiltinValue::kVertexIndex;
}
@@ -96,6 +102,10 @@
return "sample_index";
case BuiltinValue::kSampleMask:
return "sample_mask";
+ case BuiltinValue::kSubgroupInvocationId:
+ return "subgroup_invocation_id";
+ case BuiltinValue::kSubgroupSize:
+ return "subgroup_size";
case BuiltinValue::kVertexIndex:
return "vertex_index";
case BuiltinValue::kWorkgroupId:
diff --git a/src/tint/lang/core/builtin/builtin_value.h b/src/tint/lang/core/builtin/builtin_value.h
index ef527f3..96c91b0 100644
--- a/src/tint/lang/core/builtin/builtin_value.h
+++ b/src/tint/lang/core/builtin/builtin_value.h
@@ -41,6 +41,8 @@
kPosition,
kSampleIndex,
kSampleMask,
+ kSubgroupInvocationId,
+ kSubgroupSize,
kVertexIndex,
kWorkgroupId,
};
@@ -66,8 +68,8 @@
"__point_size", "frag_depth", "front_facing",
"global_invocation_id", "instance_index", "local_invocation_id",
"local_invocation_index", "num_workgroups", "position",
- "sample_index", "sample_mask", "vertex_index",
- "workgroup_id",
+ "sample_index", "sample_mask", "subgroup_invocation_id",
+ "subgroup_size", "vertex_index", "workgroup_id",
};
} // namespace tint::builtin
diff --git a/src/tint/lang/core/builtin/builtin_value_bench.cc b/src/tint/lang/core/builtin/builtin_value_bench.cc
index 35f68fd..f7b8d3f 100644
--- a/src/tint/lang/core/builtin/builtin_value_bench.cc
+++ b/src/tint/lang/core/builtin/builtin_value_bench.cc
@@ -108,20 +108,34 @@
"33amOe_mas66",
"samoott6QQmask",
"66mple_mask",
- "verzzx_in6Oxx",
- "vertex_yyndex",
- "vetxHHZnZex",
+ "subroup_invoOaxion_i6zz",
+ "subgroup_inyyocation_id",
+ "subgrup_invcatiZHH_id",
+ "subgroup_invocation_id",
+ "subgroqp_inWWoat44on_id",
+ "subgrou_inOOocation_id",
+ "suhgrup_invoYation_id",
+ "subroup_si",
+ "suFgoup_size",
+ "subgowp_size",
+ "subgroup_size",
+ "suffgKup_sie",
+ "KKubgroqp_size",
+ "subFroup3mmize",
+ "ertex_index",
+ "verteq_inex",
+ "verbx_indbbx",
"vertex_index",
- "vWWteq_in44ex",
- "vrtex_OOndex",
- "hrteYooindx",
- "wogroup_i",
- "wokgrouF_id",
- "worgrwup_id",
+ "iertex_indx",
+ "veOOtexqidex",
+ "vertexvvindTTx",
+ "workFFroup_id",
+ "workgPfpQ00d",
+ "woPkgroup_id",
"workgroup_id",
- "workGKou_if",
- "worKKgrouq_id",
- "w3rkgrommp_id",
+ "wosskgrup_i77",
+ "workgroup_bbRC",
+ "workgroup_iXX",
};
for (auto _ : state) {
for (auto* str : kStrings) {
diff --git a/src/tint/lang/core/builtin/builtin_value_test.cc b/src/tint/lang/core/builtin/builtin_value_test.cc
index dc2c26f..ca74e72 100644
--- a/src/tint/lang/core/builtin/builtin_value_test.cc
+++ b/src/tint/lang/core/builtin/builtin_value_test.cc
@@ -54,6 +54,8 @@
{"position", BuiltinValue::kPosition},
{"sample_index", BuiltinValue::kSampleIndex},
{"sample_mask", BuiltinValue::kSampleMask},
+ {"subgroup_invocation_id", BuiltinValue::kSubgroupInvocationId},
+ {"subgroup_size", BuiltinValue::kSubgroupSize},
{"vertex_index", BuiltinValue::kVertexIndex},
{"workgroup_id", BuiltinValue::kWorkgroupId},
};
@@ -92,12 +94,18 @@
{"sample_XXask", BuiltinValue::kUndefined},
{"samII99l55_mask", BuiltinValue::kUndefined},
{"samaale_SSrHHYk", BuiltinValue::kUndefined},
- {"verkkeH_de", BuiltinValue::kUndefined},
- {"verRg_injex", BuiltinValue::kUndefined},
- {"vrtexinbex", BuiltinValue::kUndefined},
- {"workjroup_id", BuiltinValue::kUndefined},
- {"wrkgroup_id", BuiltinValue::kUndefined},
- {"qorkgro_id", BuiltinValue::kUndefined},
+ {"skkgroup_Hnvocatio_d", BuiltinValue::kUndefined},
+ {"gRbgroup_invocajionid", BuiltinValue::kUndefined},
+ {"sbgroup_nbocation_id", BuiltinValue::kUndefined},
+ {"subgroupjsize", BuiltinValue::kUndefined},
+ {"subgroup_sie", BuiltinValue::kUndefined},
+ {"sgroupqsize", BuiltinValue::kUndefined},
+ {"vertx_NNndex", BuiltinValue::kUndefined},
+ {"vvertex_dex", BuiltinValue::kUndefined},
+ {"ertex_inQQex", BuiltinValue::kUndefined},
+ {"wrkgrup_irf", BuiltinValue::kUndefined},
+ {"workgroup_jd", BuiltinValue::kUndefined},
+ {"w82wNNgrou_id", BuiltinValue::kUndefined},
};
using BuiltinValueParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.cc b/src/tint/lang/spirv/writer/ast_printer/builder.cc
index 1457e51..a147279 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.cc
@@ -4005,6 +4005,12 @@
return SpvBuiltInSampleId;
case builtin::BuiltinValue::kSampleMask:
return SpvBuiltInSampleMask;
+ case builtin::BuiltinValue::kSubgroupInvocationId:
+ module_.PushCapability(SpvCapabilityGroupNonUniform);
+ return SpvBuiltInSubgroupLocalInvocationId;
+ case builtin::BuiltinValue::kSubgroupSize:
+ module_.PushCapability(SpvCapabilityGroupNonUniform);
+ return SpvBuiltInSubgroupSize;
case builtin::BuiltinValue::kUndefined:
break;
}
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index f500bae..ea54789 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -232,6 +232,12 @@
return SpvBuiltInSampleId;
case builtin::BuiltinValue::kSampleMask:
return SpvBuiltInSampleMask;
+ case builtin::BuiltinValue::kSubgroupInvocationId:
+ module_.PushCapability(SpvCapabilityGroupNonUniform);
+ return SpvBuiltInSubgroupLocalInvocationId;
+ case builtin::BuiltinValue::kSubgroupSize:
+ module_.PushCapability(SpvCapabilityGroupNonUniform);
+ return SpvBuiltInSubgroupSize;
case builtin::BuiltinValue::kVertexIndex:
return SpvBuiltInVertexIndex;
case builtin::BuiltinValue::kWorkgroupId:
diff --git a/src/tint/lang/wgsl/resolver/unresolved_identifier_test.cc b/src/tint/lang/wgsl/resolver/unresolved_identifier_test.cc
index 4564d4b..e3036e5 100644
--- a/src/tint/lang/wgsl/resolver/unresolved_identifier_test.cc
+++ b/src/tint/lang/wgsl/resolver/unresolved_identifier_test.cc
@@ -46,7 +46,7 @@
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), R"(12:34 error: unresolved builtin value 'positon'
12:34 note: Did you mean 'position'?
-Possible values: 'frag_depth', 'front_facing', 'global_invocation_id', 'instance_index', 'local_invocation_id', 'local_invocation_index', 'num_workgroups', 'position', 'sample_index', 'sample_mask', 'vertex_index', 'workgroup_id')");
+Possible values: 'frag_depth', 'front_facing', 'global_invocation_id', 'instance_index', 'local_invocation_id', 'local_invocation_index', 'num_workgroups', 'position', 'sample_index', 'sample_mask', 'subgroup_invocation_id', 'subgroup_size', 'vertex_index', 'workgroup_id')");
}
TEST_F(ResolverUnresolvedIdentifierSuggestions, TexelFormat) {
diff --git a/test/tint/shader_io/compute_subgroup_builtins.wgsl b/test/tint/shader_io/compute_subgroup_builtins.wgsl
new file mode 100644
index 0000000..8b85ca2
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins.wgsl
@@ -0,0 +1,10 @@
+@group(0) @binding(0)
+var<storage, read_write> output: array<u32>;
+
+@compute @workgroup_size(1)
+fn main(
+ @builtin(subgroup_invocation_id) subgroup_invocation_id : u32,
+ @builtin(subgroup_size) subgroup_size : u32,
+) {
+ output[subgroup_invocation_id] = subgroup_size;
+}
diff --git a/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.dxc.hlsl b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..24241e6
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.dxc.hlsl
@@ -0,0 +1,11 @@
+SKIP: FAILED
+
+
+@group(0) @binding(0) var<storage, read_write> output : array<u32>;
+
+@compute @workgroup_size(1)
+fn main(@builtin(subgroup_invocation_id) subgroup_invocation_id : u32, @builtin(subgroup_size) subgroup_size : u32) {
+ output[subgroup_invocation_id] = subgroup_size;
+}
+
+Failed to generate: error: unsupported builtin
diff --git a/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.fxc.hlsl b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..24241e6
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.fxc.hlsl
@@ -0,0 +1,11 @@
+SKIP: FAILED
+
+
+@group(0) @binding(0) var<storage, read_write> output : array<u32>;
+
+@compute @workgroup_size(1)
+fn main(@builtin(subgroup_invocation_id) subgroup_invocation_id : u32, @builtin(subgroup_size) subgroup_size : u32) {
+ output[subgroup_invocation_id] = subgroup_size;
+}
+
+Failed to generate: error: unsupported builtin
diff --git a/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.glsl b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.glsl
new file mode 100644
index 0000000..f88e037
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+SKIP: FAILED
+
+#version 310 es
+
+layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
+ uint inner[];
+} tint_symbol;
+
+void tint_symbol_1(uint subgroup_invocation_id, uint subgroup_size) {
+ tint_symbol.inner[subgroup_invocation_id] = subgroup_size;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol_1(tint_symbol_2, tint_symbol_3);
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:13: 'tint_symbol_2' : undeclared identifier
+ERROR: 0:13: '' : compilation terminated
+ERROR: 2 compilation errors. No code generated.
+
+
+
diff --git a/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.msl b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.msl
new file mode 100644
index 0000000..c87173c
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.msl
@@ -0,0 +1,11 @@
+SKIP: FAILED
+
+
+@group(0) @binding(0) var<storage, read_write> output : array<u32>;
+
+@compute @workgroup_size(1)
+fn tint_symbol(@builtin(subgroup_invocation_id) subgroup_invocation_id : u32, @builtin(subgroup_size) subgroup_size : u32) {
+ output[subgroup_invocation_id] = subgroup_size;
+}
+
+Failed to generate: error: unknown builtin
diff --git a/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.spvasm b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.spvasm
new file mode 100644
index 0000000..1190c73
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.spvasm
@@ -0,0 +1,2 @@
+error: line 2: Invalid capability operand: 61
+
diff --git a/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.wgsl b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.wgsl
new file mode 100644
index 0000000..91d8bff
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+@group(0) @binding(0) var<storage, read_write> output : array<u32>;
+
+@compute @workgroup_size(1)
+fn main(@builtin(subgroup_invocation_id) subgroup_invocation_id : u32, @builtin(subgroup_size) subgroup_size : u32) {
+ output[subgroup_invocation_id] = subgroup_size;
+}
diff --git a/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl
new file mode 100644
index 0000000..fac4f13
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl
@@ -0,0 +1,12 @@
+@group(0) @binding(0)
+var<storage, read_write> output: array<u32>;
+
+struct ComputeInputs {
+ @builtin(subgroup_invocation_id) subgroup_invocation_id : u32,
+ @builtin(subgroup_size) subgroup_size : u32,
+};
+
+@compute @workgroup_size(1)
+fn main(inputs : ComputeInputs) {
+ output[inputs.subgroup_invocation_id] = inputs.subgroup_size;
+}
diff --git a/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.dxc.hlsl b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..cab4b43
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.dxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+@group(0) @binding(0) var<storage, read_write> output : array<u32>;
+
+struct ComputeInputs {
+ @builtin(subgroup_invocation_id)
+ subgroup_invocation_id : u32,
+ @builtin(subgroup_size)
+ subgroup_size : u32,
+}
+
+@compute @workgroup_size(1)
+fn main(inputs : ComputeInputs) {
+ output[inputs.subgroup_invocation_id] = inputs.subgroup_size;
+}
+
+Failed to generate: error: unsupported builtin
diff --git a/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.fxc.hlsl b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..cab4b43
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+@group(0) @binding(0) var<storage, read_write> output : array<u32>;
+
+struct ComputeInputs {
+ @builtin(subgroup_invocation_id)
+ subgroup_invocation_id : u32,
+ @builtin(subgroup_size)
+ subgroup_size : u32,
+}
+
+@compute @workgroup_size(1)
+fn main(inputs : ComputeInputs) {
+ output[inputs.subgroup_invocation_id] = inputs.subgroup_size;
+}
+
+Failed to generate: error: unsupported builtin
diff --git a/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.glsl b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.glsl
new file mode 100644
index 0000000..224650c
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.glsl
@@ -0,0 +1,30 @@
+SKIP: FAILED
+
+#version 310 es
+
+layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
+ uint inner[];
+} tint_symbol;
+
+struct ComputeInputs {
+ uint subgroup_invocation_id;
+ uint subgroup_size;
+};
+
+void tint_symbol_1(ComputeInputs inputs) {
+ tint_symbol.inner[inputs.subgroup_invocation_id] = inputs.subgroup_size;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ ComputeInputs tint_symbol_4 = ComputeInputs(tint_symbol_2, tint_symbol_3);
+ tint_symbol_1(tint_symbol_4);
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:18: 'tint_symbol_2' : undeclared identifier
+ERROR: 0:18: '' : compilation terminated
+ERROR: 2 compilation errors. No code generated.
+
+
+
diff --git a/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.msl b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.msl
new file mode 100644
index 0000000..feae4f3
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.msl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+@group(0) @binding(0) var<storage, read_write> output : array<u32>;
+
+struct ComputeInputs {
+ @builtin(subgroup_invocation_id)
+ subgroup_invocation_id : u32,
+ @builtin(subgroup_size)
+ subgroup_size : u32,
+}
+
+@compute @workgroup_size(1)
+fn tint_symbol(inputs : ComputeInputs) {
+ output[inputs.subgroup_invocation_id] = inputs.subgroup_size;
+}
+
+Failed to generate: error: unknown builtin
diff --git a/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.spvasm b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..1190c73
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.spvasm
@@ -0,0 +1,2 @@
+error: line 2: Invalid capability operand: 61
+
diff --git a/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.wgsl b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..0191fdb
--- /dev/null
+++ b/test/tint/shader_io/compute_subgroup_builtins_struct.wgsl.expected.wgsl
@@ -0,0 +1,13 @@
+@group(0) @binding(0) var<storage, read_write> output : array<u32>;
+
+struct ComputeInputs {
+ @builtin(subgroup_invocation_id)
+ subgroup_invocation_id : u32,
+ @builtin(subgroup_size)
+ subgroup_size : u32,
+}
+
+@compute @workgroup_size(1)
+fn main(inputs : ComputeInputs) {
+ output[inputs.subgroup_invocation_id] = inputs.subgroup_size;
+}
