Import Tint changes from Dawn
Changes:
- f25140fe6ff0babe2fc7712883da6fceb72a3cc7 tint: fix emitting duplicate structs for atomicCompareExc... by Antonio Maiorano <amaiorano@google.com>
- f0c150b01bd832ea2922c49b37860cca5e23cc83 Add parsing of shorter stage attributes. by dan sinclair <dsinclair@chromium.org>
- c0af5c5c9c4163ee166fba14f7ab8e816753b50a tint: Add builtin tests for arguments passed by var by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: f25140fe6ff0babe2fc7712883da6fceb72a3cc7
Change-Id: If6624924e2fa280b061a80c497590e2a06d04e65
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/92580
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 0c7cfb4..5ac5773 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -496,22 +496,22 @@
fn textureDimensions<F: texel_format, A: write_only>(texture: texture_storage_2d_array<F, A>) -> vec2<i32>
fn textureDimensions<F: texel_format, A: write_only>(texture: texture_storage_3d<F, A>) -> vec3<i32>
fn textureDimensions(texture: texture_external) -> vec2<i32>
-fn textureGather<T: fiu32>(component: i32, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>) -> vec4<T>
-fn textureGather<T: fiu32>(component: i32, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<T>
-fn textureGather<T: fiu32>(component: i32, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<T>
-fn textureGather<T: fiu32>(component: i32, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<T>
-fn textureGather<T: fiu32>(component: i32, texture: texture_cube<T>, sampler: sampler, coords: vec3<f32>) -> vec4<T>
-fn textureGather<T: fiu32>(component: i32, texture: texture_cube_array<T>, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<T>
+fn textureGather<T: fiu32>(@const component: i32, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>) -> vec4<T>
+fn textureGather<T: fiu32>(@const component: i32, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<T>
+fn textureGather<T: fiu32>(@const component: i32, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<T>
+fn textureGather<T: fiu32>(@const component: i32, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: i32, @const offset: vec2<i32>) -> vec4<T>
+fn textureGather<T: fiu32>(@const component: i32, texture: texture_cube<T>, sampler: sampler, coords: vec3<f32>) -> vec4<T>
+fn textureGather<T: fiu32>(@const component: i32, texture: texture_cube_array<T>, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<T>
fn textureGather(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
-fn textureGather(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<f32>
+fn textureGather(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32>
fn textureGather(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<f32>
-fn textureGather(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<f32>
+fn textureGather(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, @const offset: vec2<i32>) -> vec4<f32>
fn textureGather(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
fn textureGather(texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<f32>
fn textureGatherCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> vec4<f32>
-fn textureGatherCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, offset: vec2<i32>) -> vec4<f32>
+fn textureGatherCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> vec4<f32>
fn textureGatherCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32) -> vec4<f32>
-fn textureGatherCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, offset: vec2<i32>) -> vec4<f32>
+fn textureGatherCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, @const offset: vec2<i32>) -> vec4<f32>
fn textureGatherCompare(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> vec4<f32>
fn textureGatherCompare(texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: i32, depth_ref: f32) -> vec4<f32>
fn textureNumLayers<T: fiu32>(texture: texture_2d_array<T>) -> i32
@@ -533,59 +533,59 @@
fn textureNumSamples(texture: texture_depth_multisampled_2d) -> i32
@stage("fragment") fn textureSample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32>
@stage("fragment") fn textureSample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
-@stage("fragment") fn textureSample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<f32>
+@stage("fragment") fn textureSample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32>
@stage("fragment") fn textureSample(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<f32>
-@stage("fragment") fn textureSample(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<f32>
+@stage("fragment") fn textureSample(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, @const offset: vec2<i32>) -> vec4<f32>
@stage("fragment") fn textureSample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
-@stage("fragment") fn textureSample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, offset: vec3<i32>) -> vec4<f32>
+@stage("fragment") fn textureSample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32>
@stage("fragment") fn textureSample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
@stage("fragment") fn textureSample(texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<f32>
@stage("fragment") fn textureSample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> f32
-@stage("fragment") fn textureSample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> f32
+@stage("fragment") fn textureSample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> f32
@stage("fragment") fn textureSample(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32) -> f32
-@stage("fragment") fn textureSample(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> f32
+@stage("fragment") fn textureSample(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, @const offset: vec2<i32>) -> f32
@stage("fragment") fn textureSample(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> f32
@stage("fragment") fn textureSample(texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: i32) -> f32
@stage("fragment") fn textureSampleBias(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, bias: f32) -> vec4<f32>
-@stage("fragment") fn textureSampleBias(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, bias: f32, offset: vec2<i32>) -> vec4<f32>
+@stage("fragment") fn textureSampleBias(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, bias: f32, @const offset: vec2<i32>) -> vec4<f32>
@stage("fragment") fn textureSampleBias(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, bias: f32) -> vec4<f32>
-@stage("fragment") fn textureSampleBias(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, bias: f32, offset: vec2<i32>) -> vec4<f32>
+@stage("fragment") fn textureSampleBias(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, bias: f32, @const offset: vec2<i32>) -> vec4<f32>
@stage("fragment") fn textureSampleBias(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, bias: f32) -> vec4<f32>
-@stage("fragment") fn textureSampleBias(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, bias: f32, offset: vec3<i32>) -> vec4<f32>
+@stage("fragment") fn textureSampleBias(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, bias: f32, @const offset: vec3<i32>) -> vec4<f32>
@stage("fragment") fn textureSampleBias(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>, bias: f32) -> vec4<f32>
@stage("fragment") fn textureSampleBias(texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: i32, bias: f32) -> vec4<f32>
@stage("fragment") fn textureSampleCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> f32
-@stage("fragment") fn textureSampleCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, offset: vec2<i32>) -> f32
+@stage("fragment") fn textureSampleCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> f32
@stage("fragment") fn textureSampleCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32) -> f32
-@stage("fragment") fn textureSampleCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, offset: vec2<i32>) -> f32
+@stage("fragment") fn textureSampleCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, @const offset: vec2<i32>) -> f32
@stage("fragment") fn textureSampleCompare(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> f32
@stage("fragment") fn textureSampleCompare(texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: i32, depth_ref: f32) -> f32
fn textureSampleCompareLevel(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> f32
-fn textureSampleCompareLevel(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, offset: vec2<i32>) -> f32
+fn textureSampleCompareLevel(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> f32
fn textureSampleCompareLevel(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32) -> f32
-fn textureSampleCompareLevel(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, offset: vec2<i32>) -> f32
+fn textureSampleCompareLevel(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, @const offset: vec2<i32>) -> f32
fn textureSampleCompareLevel(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> f32
fn textureSampleCompareLevel(texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: i32, depth_ref: f32) -> f32
fn textureSampleGrad(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, ddx: vec2<f32>, ddy: vec2<f32>) -> vec4<f32>
-fn textureSampleGrad(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, ddx: vec2<f32>, ddy: vec2<f32>, offset: vec2<i32>) -> vec4<f32>
+fn textureSampleGrad(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, ddx: vec2<f32>, ddy: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32>
fn textureSampleGrad(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, ddx: vec2<f32>, ddy: vec2<f32>) -> vec4<f32>
-fn textureSampleGrad(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, ddx: vec2<f32>, ddy: vec2<f32>, offset: vec2<i32>) -> vec4<f32>
+fn textureSampleGrad(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, ddx: vec2<f32>, ddy: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32>
fn textureSampleGrad(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>) -> vec4<f32>
-fn textureSampleGrad(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>, offset: vec3<i32>) -> vec4<f32>
+fn textureSampleGrad(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32>
fn textureSampleGrad(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>) -> vec4<f32>
fn textureSampleGrad(texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: i32, ddx: vec3<f32>, ddy: vec3<f32>) -> vec4<f32>
fn textureSampleLevel(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, level: f32) -> vec4<f32>
-fn textureSampleLevel(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, level: f32, offset: vec2<i32>) -> vec4<f32>
+fn textureSampleLevel(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, level: f32, @const offset: vec2<i32>) -> vec4<f32>
fn textureSampleLevel(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, level: f32) -> vec4<f32>
-fn textureSampleLevel(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, level: f32, offset: vec2<i32>) -> vec4<f32>
+fn textureSampleLevel(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, level: f32, @const offset: vec2<i32>) -> vec4<f32>
fn textureSampleLevel(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, level: f32) -> vec4<f32>
-fn textureSampleLevel(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, level: f32, offset: vec3<i32>) -> vec4<f32>
+fn textureSampleLevel(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, level: f32, @const offset: vec3<i32>) -> vec4<f32>
fn textureSampleLevel(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>, level: f32) -> vec4<f32>
fn textureSampleLevel(texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: i32, level: f32) -> vec4<f32>
fn textureSampleLevel(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, level: i32) -> f32
-fn textureSampleLevel(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, level: i32, offset: vec2<i32>) -> f32
+fn textureSampleLevel(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, level: i32, @const offset: vec2<i32>) -> f32
fn textureSampleLevel(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, level: i32) -> f32
-fn textureSampleLevel(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, level: i32, offset: vec2<i32>) -> f32
+fn textureSampleLevel(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, level: i32, @const offset: vec2<i32>) -> f32
fn textureSampleLevel(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>, level: i32) -> f32
fn textureSampleLevel(texture: texture_depth_cube_array,sampler: sampler, coords: vec3<f32>, array_index: i32, level: i32) -> f32
fn textureSampleLevel(texture: texture_external, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index 4ad0f76..a28b798 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -3258,9 +3258,38 @@
return Failure::kErrored;
}
+ // TODO(crbug.com/tint/1503): Enable this once all the Dawn and CTS
+ // tests are updated to use the new format so we can avoid spamming
+ // the log files.
+ if ((false)) {
+ std::string warning = "stage should use @";
+ switch (stage.value) {
+ case ast::PipelineStage::kVertex:
+ warning += "vertex";
+ break;
+ case ast::PipelineStage::kFragment:
+ warning += "fragment";
+ break;
+ case ast::PipelineStage::kCompute:
+ warning += "compute";
+ break;
+ case ast::PipelineStage::kNone:
+ break;
+ }
+ deprecated(t.source(), warning);
+ }
return create<ast::StageAttribute>(t.source(), stage.value);
});
}
+ if (t == kComputeStage) {
+ return create<ast::StageAttribute>(t.source(), ast::PipelineStage::kCompute);
+ }
+ if (t == kVertexStage) {
+ return create<ast::StageAttribute>(t.source(), ast::PipelineStage::kVertex);
+ }
+ if (t == kFragmentStage) {
+ return create<ast::StageAttribute>(t.source(), ast::PipelineStage::kFragment);
+ }
if (t == kSizeAttribute) {
const char* use = "size attribute";
diff --git a/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc b/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc
index 356241b..83ebeac 100644
--- a/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_attribute_list_test.cc
@@ -18,7 +18,7 @@
namespace tint::reader::wgsl {
namespace {
-TEST_F(ParserImplTest, AttributeList_Parses) {
+TEST_F(ParserImplTest, AttributeList_Parses_Stage) {
auto p = parser("@workgroup_size(2) @stage(compute)");
auto attrs = p->attribute_list();
EXPECT_FALSE(p->has_error()) << p->error();
@@ -44,6 +44,32 @@
EXPECT_EQ(attr_1->As<ast::StageAttribute>()->stage, ast::PipelineStage::kCompute);
}
+TEST_F(ParserImplTest, AttributeList_Parses) {
+ auto p = parser("@workgroup_size(2) @compute");
+ auto attrs = p->attribute_list();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(attrs.errored);
+ EXPECT_TRUE(attrs.matched);
+ ASSERT_EQ(attrs.value.size(), 2u);
+
+ auto* attr_0 = attrs.value[0]->As<ast::Attribute>();
+ auto* attr_1 = attrs.value[1]->As<ast::Attribute>();
+ ASSERT_NE(attr_0, nullptr);
+ ASSERT_NE(attr_1, nullptr);
+
+ ASSERT_TRUE(attr_0->Is<ast::WorkgroupAttribute>());
+ const ast::Expression* x = attr_0->As<ast::WorkgroupAttribute>()->x;
+ ASSERT_NE(x, nullptr);
+ auto* x_literal = x->As<ast::LiteralExpression>();
+ ASSERT_NE(x_literal, nullptr);
+ ASSERT_TRUE(x_literal->Is<ast::IntLiteralExpression>());
+ EXPECT_EQ(x_literal->As<ast::IntLiteralExpression>()->value, 2);
+ EXPECT_EQ(x_literal->As<ast::IntLiteralExpression>()->suffix,
+ ast::IntLiteralExpression::Suffix::kNone);
+ ASSERT_TRUE(attr_1->Is<ast::StageAttribute>());
+ EXPECT_EQ(attr_1->As<ast::StageAttribute>()->stage, ast::PipelineStage::kCompute);
+}
+
TEST_F(ParserImplTest, AttributeList_Invalid) {
auto p = parser("@invalid");
auto attrs = p->attribute_list();
diff --git a/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc b/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
index d2585b3..15be222 100644
--- a/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
@@ -265,5 +265,44 @@
EXPECT_EQ(p->error(), "1:14: expected ')' for stage attribute");
}
+TEST_F(ParserImplTest, Attribute_Compute) {
+ auto p = parser("compute");
+ auto attr = p->attribute();
+ EXPECT_TRUE(attr.matched);
+ EXPECT_FALSE(attr.errored);
+ ASSERT_NE(attr.value, nullptr) << p->error();
+ ASSERT_FALSE(p->has_error());
+ auto* func_attr = attr.value->As<ast::Attribute>();
+ ASSERT_NE(func_attr, nullptr);
+ ASSERT_TRUE(func_attr->Is<ast::StageAttribute>());
+ EXPECT_EQ(func_attr->As<ast::StageAttribute>()->stage, ast::PipelineStage::kCompute);
+}
+
+TEST_F(ParserImplTest, Attribute_Vertex) {
+ auto p = parser("vertex");
+ auto attr = p->attribute();
+ EXPECT_TRUE(attr.matched);
+ EXPECT_FALSE(attr.errored);
+ ASSERT_NE(attr.value, nullptr) << p->error();
+ ASSERT_FALSE(p->has_error());
+ auto* func_attr = attr.value->As<ast::Attribute>();
+ ASSERT_NE(func_attr, nullptr);
+ ASSERT_TRUE(func_attr->Is<ast::StageAttribute>());
+ EXPECT_EQ(func_attr->As<ast::StageAttribute>()->stage, ast::PipelineStage::kVertex);
+}
+
+TEST_F(ParserImplTest, Attribute_Fragment) {
+ auto p = parser("fragment");
+ auto attr = p->attribute();
+ EXPECT_TRUE(attr.matched);
+ EXPECT_FALSE(attr.errored);
+ ASSERT_NE(attr.value, nullptr) << p->error();
+ ASSERT_FALSE(p->has_error());
+ auto* func_attr = attr.value->As<ast::Attribute>();
+ ASSERT_NE(func_attr, nullptr);
+ ASSERT_TRUE(func_attr->Is<ast::StageAttribute>());
+ EXPECT_EQ(func_attr->As<ast::StageAttribute>()->stage, ast::PipelineStage::kFragment);
+}
+
} // namespace
} // namespace tint::reader::wgsl
diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl
index fab7783..521e5e1 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -13601,16 +13601,16 @@
},
{
/* [83] */
- /* fn textureGather<T : fiu32>(component: i32, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>) -> vec4<T> */
- /* fn textureGather<T : fiu32>(component: i32, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<T> */
- /* fn textureGather<T : fiu32>(component: i32, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<T> */
- /* fn textureGather<T : fiu32>(component: i32, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<T> */
- /* fn textureGather<T : fiu32>(component: i32, texture: texture_cube<T>, sampler: sampler, coords: vec3<f32>) -> vec4<T> */
- /* fn textureGather<T : fiu32>(component: i32, texture: texture_cube_array<T>, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<T> */
+ /* fn textureGather<T : fiu32>(@const component: i32, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>) -> vec4<T> */
+ /* fn textureGather<T : fiu32>(@const component: i32, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<T> */
+ /* fn textureGather<T : fiu32>(@const component: i32, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<T> */
+ /* fn textureGather<T : fiu32>(@const component: i32, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: i32, @const offset: vec2<i32>) -> vec4<T> */
+ /* fn textureGather<T : fiu32>(@const component: i32, texture: texture_cube<T>, sampler: sampler, coords: vec3<f32>) -> vec4<T> */
+ /* fn textureGather<T : fiu32>(@const component: i32, texture: texture_cube_array<T>, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<T> */
/* fn textureGather(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> vec4<f32> */
- /* fn textureGather(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureGather(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureGather(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<f32> */
- /* fn textureGather(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureGather(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureGather(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> vec4<f32> */
/* fn textureGather(texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<f32> */
/* num overloads */ 12,
@@ -13619,9 +13619,9 @@
{
/* [84] */
/* fn textureGatherCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> vec4<f32> */
- /* fn textureGatherCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureGatherCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureGatherCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32) -> vec4<f32> */
- /* fn textureGatherCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureGatherCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureGatherCompare(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> vec4<f32> */
/* fn textureGatherCompare(texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: i32, depth_ref: f32) -> vec4<f32> */
/* num overloads */ 6,
@@ -13663,17 +13663,17 @@
/* [88] */
/* fn textureSample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32> */
/* fn textureSample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32> */
- /* fn textureSample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureSample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureSample(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<f32> */
- /* fn textureSample(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureSample(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureSample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32> */
- /* fn textureSample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, offset: vec3<i32>) -> vec4<f32> */
+ /* fn textureSample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32> */
/* fn textureSample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32> */
/* fn textureSample(texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<f32> */
/* fn textureSample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> f32 */
- /* fn textureSample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> f32 */
+ /* fn textureSample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> f32 */
/* fn textureSample(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32) -> f32 */
- /* fn textureSample(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> f32 */
+ /* fn textureSample(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, @const offset: vec2<i32>) -> f32 */
/* fn textureSample(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> f32 */
/* fn textureSample(texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: i32) -> f32 */
/* num overloads */ 15,
@@ -13682,11 +13682,11 @@
{
/* [89] */
/* fn textureSampleBias(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, bias: f32) -> vec4<f32> */
- /* fn textureSampleBias(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, bias: f32, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureSampleBias(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, bias: f32, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureSampleBias(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, bias: f32) -> vec4<f32> */
- /* fn textureSampleBias(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, bias: f32, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureSampleBias(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, bias: f32, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureSampleBias(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, bias: f32) -> vec4<f32> */
- /* fn textureSampleBias(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, bias: f32, offset: vec3<i32>) -> vec4<f32> */
+ /* fn textureSampleBias(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, bias: f32, @const offset: vec3<i32>) -> vec4<f32> */
/* fn textureSampleBias(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>, bias: f32) -> vec4<f32> */
/* fn textureSampleBias(texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: i32, bias: f32) -> vec4<f32> */
/* num overloads */ 8,
@@ -13695,9 +13695,9 @@
{
/* [90] */
/* fn textureSampleCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> f32 */
- /* fn textureSampleCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, offset: vec2<i32>) -> f32 */
+ /* fn textureSampleCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> f32 */
/* fn textureSampleCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32) -> f32 */
- /* fn textureSampleCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, offset: vec2<i32>) -> f32 */
+ /* fn textureSampleCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, @const offset: vec2<i32>) -> f32 */
/* fn textureSampleCompare(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> f32 */
/* fn textureSampleCompare(texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: i32, depth_ref: f32) -> f32 */
/* num overloads */ 6,
@@ -13706,9 +13706,9 @@
{
/* [91] */
/* fn textureSampleCompareLevel(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> f32 */
- /* fn textureSampleCompareLevel(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, offset: vec2<i32>) -> f32 */
+ /* fn textureSampleCompareLevel(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> f32 */
/* fn textureSampleCompareLevel(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32) -> f32 */
- /* fn textureSampleCompareLevel(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, offset: vec2<i32>) -> f32 */
+ /* fn textureSampleCompareLevel(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, @const offset: vec2<i32>) -> f32 */
/* fn textureSampleCompareLevel(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> f32 */
/* fn textureSampleCompareLevel(texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: i32, depth_ref: f32) -> f32 */
/* num overloads */ 6,
@@ -13717,11 +13717,11 @@
{
/* [92] */
/* fn textureSampleGrad(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, ddx: vec2<f32>, ddy: vec2<f32>) -> vec4<f32> */
- /* fn textureSampleGrad(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, ddx: vec2<f32>, ddy: vec2<f32>, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureSampleGrad(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, ddx: vec2<f32>, ddy: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureSampleGrad(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, ddx: vec2<f32>, ddy: vec2<f32>) -> vec4<f32> */
- /* fn textureSampleGrad(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, ddx: vec2<f32>, ddy: vec2<f32>, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureSampleGrad(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, ddx: vec2<f32>, ddy: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureSampleGrad(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>) -> vec4<f32> */
- /* fn textureSampleGrad(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>, offset: vec3<i32>) -> vec4<f32> */
+ /* fn textureSampleGrad(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32> */
/* fn textureSampleGrad(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>, ddx: vec3<f32>, ddy: vec3<f32>) -> vec4<f32> */
/* fn textureSampleGrad(texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: i32, ddx: vec3<f32>, ddy: vec3<f32>) -> vec4<f32> */
/* num overloads */ 8,
@@ -13730,17 +13730,17 @@
{
/* [93] */
/* fn textureSampleLevel(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, level: f32) -> vec4<f32> */
- /* fn textureSampleLevel(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, level: f32, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureSampleLevel(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, level: f32, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureSampleLevel(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, level: f32) -> vec4<f32> */
- /* fn textureSampleLevel(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, level: f32, offset: vec2<i32>) -> vec4<f32> */
+ /* fn textureSampleLevel(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, level: f32, @const offset: vec2<i32>) -> vec4<f32> */
/* fn textureSampleLevel(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, level: f32) -> vec4<f32> */
- /* fn textureSampleLevel(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, level: f32, offset: vec3<i32>) -> vec4<f32> */
+ /* fn textureSampleLevel(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, level: f32, @const offset: vec3<i32>) -> vec4<f32> */
/* fn textureSampleLevel(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>, level: f32) -> vec4<f32> */
/* fn textureSampleLevel(texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: i32, level: f32) -> vec4<f32> */
/* fn textureSampleLevel(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, level: i32) -> f32 */
- /* fn textureSampleLevel(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, level: i32, offset: vec2<i32>) -> f32 */
+ /* fn textureSampleLevel(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, level: i32, @const offset: vec2<i32>) -> f32 */
/* fn textureSampleLevel(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, level: i32) -> f32 */
- /* fn textureSampleLevel(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, level: i32, offset: vec2<i32>) -> f32 */
+ /* fn textureSampleLevel(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, level: i32, @const offset: vec2<i32>) -> f32 */
/* fn textureSampleLevel(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>, level: i32) -> f32 */
/* fn textureSampleLevel(texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: i32, level: i32) -> f32 */
/* fn textureSampleLevel(texture: texture_external, sampler: sampler, coords: vec2<f32>) -> vec4<f32> */
diff --git a/src/tint/transform/canonicalize_entry_point_io.h b/src/tint/transform/canonicalize_entry_point_io.h
index 64a10f2..95f8b19 100644
--- a/src/tint/transform/canonicalize_entry_point_io.h
+++ b/src/tint/transform/canonicalize_entry_point_io.h
@@ -34,7 +34,7 @@
/// @location(2) loc2 : vec4<u32>;
/// };
///
-/// @stage(fragment)
+/// @fragment
/// fn frag_main(@builtin(position) coord : vec4<f32>,
/// locations : Locations) -> @location(0) f32 {
/// if (coord.w > 1.0) {
@@ -71,7 +71,7 @@
/// return col;
/// }
///
-/// @stage(fragment)
+/// @fragment
/// fn frag_main(in : frag_main_in) -> frag_main_out {
/// let inner_retval = frag_main_inner(in.coord, Locations(in.loc1, in.loc2));
/// var wrapper_result : frag_main_out;
diff --git a/src/tint/transform/module_scope_var_to_entry_point_param.h b/src/tint/transform/module_scope_var_to_entry_point_param.h
index f268197..e3a50f4 100644
--- a/src/tint/transform/module_scope_var_to_entry_point_param.h
+++ b/src/tint/transform/module_scope_var_to_entry_point_param.h
@@ -43,7 +43,7 @@
/// p = p + f;
/// }
///
-/// @stage(compute) @workgroup_size(1)
+/// @compute @workgroup_size(1)
/// fn main() {
/// foo();
/// }
@@ -55,7 +55,7 @@
/// *p = *p + (*sptr).f;
/// }
///
-/// @stage(compute) @workgroup_size(1)
+/// @compute @workgroup_size(1)
/// fn main(sptr : ptr<storage, S, read>) {
/// var<private> p : f32 = 2.0;
/// foo(&p, sptr);
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index eb44f14..e07139b 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -920,7 +920,7 @@
case sem::BuiltinType::kAtomicCompareExchangeWeak: {
// Emit the builtin return type unique to this overload. This does not
// exist in the AST, so it will not be generated in Generate().
- if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+ if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
return false;
}
@@ -2822,6 +2822,14 @@
return true;
}
+bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
+ auto it = emitted_structs_.emplace(str);
+ if (!it.second) {
+ return true;
+ }
+ return EmitStructType(buffer, str);
+}
+
bool GeneratorImpl::EmitStructMembers(TextBuffer* b, const sem::Struct* str, bool emit_offsets) {
ScopedIndent si(b);
for (auto* mem : str->Members()) {
diff --git a/src/tint/writer/glsl/generator_impl.h b/src/tint/writer/glsl/generator_impl.h
index bcf84b0..819c79b 100644
--- a/src/tint/writer/glsl/generator_impl.h
+++ b/src/tint/writer/glsl/generator_impl.h
@@ -411,6 +411,12 @@
/// @param ty the struct to generate
/// @returns true if the struct is emitted
bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty);
+ /// Handles generating a structure declaration only the first time called. Subsequent calls are
+ /// a no-op and return true.
+ /// @param buffer the text buffer that the type declaration will be written to
+ /// @param ty the struct to generate
+ /// @returns true if the struct is emitted
+ bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
/// Handles generating the members of a structure
/// @param buffer the text buffer that the struct members will be written to
/// @param ty the struct to generate
@@ -503,6 +509,7 @@
std::unordered_map<const sem::Vector*, std::string> dynamic_vector_write_;
std::unordered_map<const sem::Vector*, std::string> int_dot_funcs_;
std::unordered_map<const sem::Type*, std::string> float_modulo_funcs_;
+ std::unordered_set<const sem::Struct*> emitted_structs_;
bool requires_oes_sample_variables_ = false;
bool requires_default_precision_qualifier_ = false;
Version version_;
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 137c466..19af4fa 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -1767,7 +1767,7 @@
case sem::BuiltinType::kAtomicCompareExchangeWeak: {
// Emit the builtin return type unique to this overload. This does not
// exist in the AST, so it will not be generated in Generate().
- if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+ if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
return false;
}
@@ -3921,6 +3921,14 @@
return true;
}
+bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
+ auto it = emitted_structs_.emplace(str);
+ if (!it.second) {
+ return true;
+ }
+ return EmitStructType(buffer, str);
+}
+
bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) {
switch (expr->op) {
case ast::UnaryOp::kIndirection:
diff --git a/src/tint/writer/hlsl/generator_impl.h b/src/tint/writer/hlsl/generator_impl.h
index c58d004..af7e4c9 100644
--- a/src/tint/writer/hlsl/generator_impl.h
+++ b/src/tint/writer/hlsl/generator_impl.h
@@ -411,6 +411,12 @@
/// @param ty the struct to generate
/// @returns true if the struct is emitted
bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty);
+ /// Handles generating a structure declaration only the first time called. Subsequent calls are
+ /// a no-op and return true.
+ /// @param buffer the text buffer that the type declaration will be written to
+ /// @param ty the struct to generate
+ /// @returns true if the struct is emitted
+ bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
/// Handles a unary op expression
/// @param out the output of the expression stream
/// @param expr the expression to emit
@@ -530,6 +536,7 @@
std::unordered_map<const sem::Matrix*, std::string> dynamic_matrix_vector_write_;
std::unordered_map<const sem::Matrix*, std::string> dynamic_matrix_scalar_write_;
std::unordered_map<const sem::Type*, std::string> value_or_one_if_zero_;
+ std::unordered_set<const sem::Struct*> emitted_structs_;
};
} // namespace tint::writer::hlsl
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index cd984b2..3228e86 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -826,46 +826,66 @@
return call("atomic_exchange_explicit", true);
case sem::BuiltinType::kAtomicCompareExchangeWeak: {
- // Emit the builtin return type unique to this overload. This does not
- // exist in the AST, so it will not be generated in Generate().
- if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
- return false;
- }
-
auto* ptr_ty = TypeOf(expr->args[0])->UnwrapRef()->As<sem::Pointer>();
auto sc = ptr_ty->StorageClass();
+ auto* str = builtin->ReturnType()->As<sem::Struct>();
- auto func = utils::GetOrCreate(atomicCompareExchangeWeak_, sc, [&]() -> std::string {
- auto name = UniqueIdentifier("atomicCompareExchangeWeak");
- auto& buf = helpers_;
-
- line(&buf) << "template <typename A, typename T>";
- {
- auto f = line(&buf);
- auto str_name = StructName(builtin->ReturnType()->As<sem::Struct>());
- f << str_name << " " << name << "(";
- if (!EmitStorageClass(f, sc)) {
+ auto func = utils::GetOrCreate(
+ atomicCompareExchangeWeak_, ACEWKeyType{{sc, str}}, [&]() -> std::string {
+ // Emit the builtin return type unique to this overload. This does not
+ // exist in the AST, so it will not be generated in Generate().
+ if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
return "";
}
- f << " A* atomic, T compare, T value) {";
- }
- buf.IncrementIndent();
- TINT_DEFER({
- buf.DecrementIndent();
- line(&buf) << "}";
- line(&buf);
+ auto name = UniqueIdentifier("atomicCompareExchangeWeak");
+ auto& buf = helpers_;
+ auto* atomic_ty = builtin->Parameters()[0]->Type();
+ auto* arg_ty = builtin->Parameters()[1]->Type();
+
+ {
+ auto f = line(&buf);
+ auto str_name = StructName(builtin->ReturnType()->As<sem::Struct>());
+ f << str_name << " " << name << "(";
+ if (!EmitTypeAndName(f, atomic_ty, "atomic")) {
+ return "";
+ }
+ f << ", ";
+ if (!EmitTypeAndName(f, arg_ty, "compare")) {
+ return "";
+ }
+ f << ", ";
+ if (!EmitTypeAndName(f, arg_ty, "value")) {
+ return "";
+ }
+ f << ") {";
+ }
+
+ buf.IncrementIndent();
+ TINT_DEFER({
+ buf.DecrementIndent();
+ line(&buf) << "}";
+ line(&buf);
+ });
+
+ {
+ auto f = line(&buf);
+ if (!EmitTypeAndName(f, arg_ty, "old_value")) {
+ return "";
+ }
+ f << " = compare;";
+ }
+ line(&buf) << "bool exchanged = "
+ "atomic_compare_exchange_weak_explicit(atomic, "
+ "&old_value, value, memory_order_relaxed, "
+ "memory_order_relaxed);";
+ line(&buf) << "return {old_value, exchanged};";
+ return name;
});
- line(&buf) << "T old_value = compare;";
- line(&buf) << "bool exchanged = "
- "atomic_compare_exchange_weak_explicit(atomic, "
- "&old_value, value, memory_order_relaxed, "
- "memory_order_relaxed);";
- line(&buf) << "return {old_value, exchanged};";
- return name;
- });
-
+ if (func.empty()) {
+ return false;
+ }
return call(func, false);
}
@@ -2765,6 +2785,14 @@
return true;
}
+bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
+ auto it = emitted_structs_.emplace(str);
+ if (!it.second) {
+ return true;
+ }
+ return EmitStructType(buffer, str);
+}
+
bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) {
// Handle `-e` when `e` is signed, so that we ensure that if `e` is the
// largest negative value, it returns `e`.
diff --git a/src/tint/writer/msl/generator_impl.h b/src/tint/writer/msl/generator_impl.h
index 21dee28..be98a86 100644
--- a/src/tint/writer/msl/generator_impl.h
+++ b/src/tint/writer/msl/generator_impl.h
@@ -16,6 +16,7 @@
#define SRC_TINT_WRITER_MSL_GENERATOR_IMPL_H_
#include <string>
+#include <tuple>
#include <unordered_map>
#include <unordered_set>
#include <vector>
@@ -332,6 +333,12 @@
/// @param str the struct to generate
/// @returns true if the struct is emitted
bool EmitStructType(TextBuffer* buffer, const sem::Struct* str);
+ /// Handles generating a structure declaration only the first time called. Subsequent calls are
+ /// a no-op and return true.
+ /// @param buffer the text buffer that the type declaration will be written to
+ /// @param ty the struct to generate
+ /// @returns true if the struct is emitted
+ bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
/// Handles a unary op expression
/// @param out the output of the expression stream
/// @param expr the expression to emit
@@ -400,13 +407,13 @@
/// type.
SizeAndAlign MslPackedTypeSizeAndAlign(const sem::Type* ty);
- using StorageClassToString = std::unordered_map<ast::StorageClass, std::string>;
-
std::function<bool()> emit_continuing_;
/// Name of atomicCompareExchangeWeak() helper for the given pointer storage
- /// class.
- StorageClassToString atomicCompareExchangeWeak_;
+ /// class and struct return type
+ using ACEWKeyType =
+ utils::UnorderedKeyWrapper<std::tuple<ast::StorageClass, const sem::Struct*>>;
+ std::unordered_map<ACEWKeyType, std::string> atomicCompareExchangeWeak_;
/// Unique name of the 'TINT_INVARIANT' preprocessor define. Non-empty only if
/// an invariant attribute has been generated.
@@ -423,6 +430,7 @@
std::unordered_map<const sem::Builtin*, std::string> builtins_;
std::unordered_map<const sem::Type*, std::string> unary_minus_funcs_;
std::unordered_map<uint32_t, std::string> int_dot_funcs_;
+ std::unordered_set<const sem::Struct*> emitted_structs_;
};
} // namespace tint::writer::msl
