Add support for [[builtin(workgroup_id)]]
This is a trivial mapping to/from WGSL in all cases.
Bug: tint:478
Change-Id: I7f21a2392543a880906b54fddbdb8bbd149a526e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/48140
Commit-Queue: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/ast/builtin.cc b/src/ast/builtin.cc
index 3dac5cd..bdfaf88 100644
--- a/src/ast/builtin.cc
+++ b/src/ast/builtin.cc
@@ -59,6 +59,10 @@
out << "global_invocation_id";
break;
}
+ case Builtin::kWorkgroupId: {
+ out << "workgroup_id";
+ break;
+ }
case Builtin::kSampleIndex: {
out << "sample_index";
break;
diff --git a/src/ast/builtin.h b/src/ast/builtin.h
index e065d69..5b37aeb 100644
--- a/src/ast/builtin.h
+++ b/src/ast/builtin.h
@@ -32,6 +32,7 @@
kLocalInvocationId,
kLocalInvocationIndex,
kGlobalInvocationId,
+ kWorkgroupId,
kSampleIndex,
kSampleMask,
kSampleMaskIn, // TODO(crbug.com/tint/715): Remove this
diff --git a/src/reader/spirv/enum_converter.cc b/src/reader/spirv/enum_converter.cc
index 364b596..b07da26 100644
--- a/src/reader/spirv/enum_converter.cc
+++ b/src/reader/spirv/enum_converter.cc
@@ -86,6 +86,8 @@
return ast::Builtin::kLocalInvocationIndex;
case SpvBuiltInGlobalInvocationId:
return ast::Builtin::kGlobalInvocationId;
+ case SpvBuiltInWorkgroupId:
+ return ast::Builtin::kWorkgroupId;
case SpvBuiltInSampleId:
return ast::Builtin::kSampleIndex;
case SpvBuiltInSampleMask:
diff --git a/src/reader/spirv/enum_converter_test.cc b/src/reader/spirv/enum_converter_test.cc
index 58e6924..19f6ad2 100644
--- a/src/reader/spirv/enum_converter_test.cc
+++ b/src/reader/spirv/enum_converter_test.cc
@@ -217,6 +217,7 @@
ast::Builtin::kLocalInvocationIndex},
BuiltinCase{SpvBuiltInGlobalInvocationId, true,
ast::Builtin::kGlobalInvocationId},
+ BuiltinCase{SpvBuiltInWorkgroupId, true, ast::Builtin::kWorkgroupId},
BuiltinCase{SpvBuiltInSampleId, true, ast::Builtin::kSampleIndex},
BuiltinCase{SpvBuiltInSampleMask, true, ast::Builtin::kSampleMask}));
diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc
index dd0b884..794cad8 100644
--- a/src/reader/wgsl/parser_impl.cc
+++ b/src/reader/wgsl/parser_impl.cc
@@ -85,6 +85,9 @@
if (str == "global_invocation_id") {
return ast::Builtin::kGlobalInvocationId;
}
+ if (str == "workgroup_id") {
+ return ast::Builtin::kWorkgroupId;
+ }
if (str == "sample_index") {
return ast::Builtin::kSampleIndex;
}
diff --git a/src/reader/wgsl/parser_impl_variable_decoration_test.cc b/src/reader/wgsl/parser_impl_variable_decoration_test.cc
index 9430cf1..c324483 100644
--- a/src/reader/wgsl/parser_impl_variable_decoration_test.cc
+++ b/src/reader/wgsl/parser_impl_variable_decoration_test.cc
@@ -121,6 +121,7 @@
BuiltinData{"local_invocation_index",
ast::Builtin::kLocalInvocationIndex},
BuiltinData{"global_invocation_id", ast::Builtin::kGlobalInvocationId},
+ BuiltinData{"workgroup_id", ast::Builtin::kWorkgroupId},
BuiltinData{"sample_index", ast::Builtin::kSampleIndex},
BuiltinData{"sample_mask", ast::Builtin::kSampleMask},
BuiltinData{"sample_mask_in", ast::Builtin::kSampleMaskIn},
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 5550534..b7dd0d2 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -1949,6 +1949,8 @@
return "SV_GroupIndex";
case ast::Builtin::kGlobalInvocationId:
return "SV_DispatchThreadID";
+ case ast::Builtin::kWorkgroupId:
+ return "SV_GroupID";
case ast::Builtin::kSampleIndex:
return "SV_SampleIndex";
case ast::Builtin::kSampleMask:
diff --git a/src/writer/hlsl/generator_impl_test.cc b/src/writer/hlsl/generator_impl_test.cc
index d37ff53..b04d496 100644
--- a/src/writer/hlsl/generator_impl_test.cc
+++ b/src/writer/hlsl/generator_impl_test.cc
@@ -82,6 +82,7 @@
HlslBuiltinData{ast::Builtin::kLocalInvocationIndex, "SV_GroupIndex"},
HlslBuiltinData{ast::Builtin::kGlobalInvocationId,
"SV_DispatchThreadID"},
+ HlslBuiltinData{ast::Builtin::kWorkgroupId, "SV_GroupID"},
HlslBuiltinData{ast::Builtin::kSampleIndex, "SV_SampleIndex"},
HlslBuiltinData{ast::Builtin::kSampleMask, "SV_Coverage"},
HlslBuiltinData{ast::Builtin::kSampleMaskIn, "SV_Coverage"},
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index ffca978..74e05e4 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -1352,6 +1352,8 @@
return "thread_index_in_threadgroup";
case ast::Builtin::kGlobalInvocationId:
return "thread_position_in_grid";
+ case ast::Builtin::kWorkgroupId:
+ return "threadgroup_position_in_grid";
case ast::Builtin::kSampleIndex:
return "sample_id";
case ast::Builtin::kSampleMask:
diff --git a/src/writer/msl/generator_impl_test.cc b/src/writer/msl/generator_impl_test.cc
index bc2ee55..9cdf4b3 100644
--- a/src/writer/msl/generator_impl_test.cc
+++ b/src/writer/msl/generator_impl_test.cc
@@ -73,6 +73,8 @@
"thread_index_in_threadgroup"},
MslBuiltinData{ast::Builtin::kGlobalInvocationId,
"thread_position_in_grid"},
+ MslBuiltinData{ast::Builtin::kWorkgroupId,
+ "threadgroup_position_in_grid"},
MslBuiltinData{ast::Builtin::kSampleIndex, "sample_id"},
MslBuiltinData{ast::Builtin::kSampleMask, "sample_mask"},
MslBuiltinData{ast::Builtin::kSampleMaskIn, "sample_mask"},
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index 77e3f11..4482b7d 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -3304,6 +3304,8 @@
return SpvBuiltInGlobalInvocationId;
case ast::Builtin::kPointSize:
return SpvBuiltInPointSize;
+ case ast::Builtin::kWorkgroupId:
+ return SpvBuiltInWorkgroupId;
case ast::Builtin::kSampleIndex:
push_capability(SpvCapabilitySampleRateShading);
return SpvBuiltInSampleId;
diff --git a/src/writer/spirv/builder_global_variable_test.cc b/src/writer/spirv/builder_global_variable_test.cc
index b6bf64c..eb0ed78 100644
--- a/src/writer/spirv/builder_global_variable_test.cc
+++ b/src/writer/spirv/builder_global_variable_test.cc
@@ -368,6 +368,8 @@
ast::StorageClass::kInput, SpvBuiltInLocalInvocationIndex},
BuiltinData{ast::Builtin::kGlobalInvocationId,
ast::StorageClass::kInput, SpvBuiltInGlobalInvocationId},
+ BuiltinData{ast::Builtin::kWorkgroupId, ast::StorageClass::kInput,
+ SpvBuiltInWorkgroupId},
BuiltinData{ast::Builtin::kSampleIndex, ast::StorageClass::kInput,
SpvBuiltInSampleId},
BuiltinData{ast::Builtin::kSampleMask, ast::StorageClass::kInput,
diff --git a/src/writer/wgsl/generator_impl_test.cc b/src/writer/wgsl/generator_impl_test.cc
index 4a6e8c1..287a6a7 100644
--- a/src/writer/wgsl/generator_impl_test.cc
+++ b/src/writer/wgsl/generator_impl_test.cc
@@ -74,6 +74,7 @@
"local_invocation_index"},
WgslBuiltinData{ast::Builtin::kGlobalInvocationId,
"global_invocation_id"},
+ WgslBuiltinData{ast::Builtin::kWorkgroupId, "workgroup_id"},
WgslBuiltinData{ast::Builtin::kSampleIndex, "sample_index"},
WgslBuiltinData{ast::Builtin::kSampleMaskIn, "sample_mask_in"},
WgslBuiltinData{ast::Builtin::kSampleMaskOut, "sample_mask_out"}));
