diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index 9314cae730..2d74b20363 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -2931,6 +2931,7 @@ Expect<const ast::Attribute*> ParserImpl::expect_attribute() {
 //                                   interpolation_sample_name COMMA? PAREN_RIGHT
 //   | ATTR 'invariant'
 //   | ATTR 'location' PAREN_LEFT expression COMMA? PAREN_RIGHT
+//   | ATTR 'must_use'
 //   | ATTR 'size' PAREN_LEFT expression COMMA? PAREN_RIGHT
 //   | ATTR 'workgroup_size' PAREN_LEFT expression COMMA? PAREN_RIGHT
 //   | ATTR 'workgroup_size' PAREN_LEFT expression COMMA expression COMMA? PAREN_RIGHT
@@ -3085,6 +3086,10 @@ Maybe<const ast::Attribute*> ParserImpl::attribute() {
         });
     }
 
+    if (t == "must_use") {
+        return create<ast::MustUseAttribute>(t.source());
+    }
+
     if (t == "size") {
         const char* use = "size attribute";
         return expect_paren_block(use, [&]() -> Result {
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 05151d0d4d..9425cc5661 100644
--- a/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_attribute_test.cc
@@ -435,5 +435,17 @@ TEST_F(ParserImplTest, Attribute_Fragment) {
     EXPECT_EQ(func_attr->As<ast::StageAttribute>()->stage, ast::PipelineStage::kFragment);
 }
 
+TEST_F(ParserImplTest, Attribute_MustUse) {
+    auto p = parser("must_use");
+    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);
+    EXPECT_TRUE(func_attr->Is<ast::MustUseAttribute>());
+}
+
 }  // namespace
 }  // namespace tint::reader::wgsl
diff --git a/src/tint/reader/wgsl/parser_impl_function_decl_test.cc b/src/tint/reader/wgsl/parser_impl_function_decl_test.cc
index 2c26bacfe0..3888326f02 100644
--- a/src/tint/reader/wgsl/parser_impl_function_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_decl_test.cc
@@ -262,6 +262,23 @@ TEST_F(ParserImplTest, FunctionDecl_ReturnTypeAttributeList) {
     EXPECT_TRUE(body->statements[0]->Is<ast::ReturnStatement>());
 }
 
+TEST_F(ParserImplTest, FunctionDecl_MustUse) {
+    auto p = parser("@must_use fn main() { return; }");
+    auto attrs = p->attribute_list();
+    EXPECT_FALSE(p->has_error()) << p->error();
+    ASSERT_FALSE(attrs.errored);
+    ASSERT_TRUE(attrs.matched);
+    auto f = p->function_decl(attrs.value);
+    EXPECT_FALSE(p->has_error()) << p->error();
+    EXPECT_FALSE(f.errored);
+    EXPECT_TRUE(f.matched);
+    ASSERT_NE(f.value, nullptr);
+
+    auto& attributes = f->attributes;
+    ASSERT_EQ(attributes.Length(), 1u);
+    ASSERT_TRUE(attributes[0]->Is<ast::MustUseAttribute>());
+}
+
 TEST_F(ParserImplTest, FunctionDecl_InvalidHeader) {
     auto p = parser("fn main() -> { }");
     auto attrs = p->attribute_list();
diff --git a/src/tint/resolver/attribute_validation_test.cc b/src/tint/resolver/attribute_validation_test.cc
index e0c47d07c6..c6dcad5ae6 100644
--- a/src/tint/resolver/attribute_validation_test.cc
+++ b/src/tint/resolver/attribute_validation_test.cc
@@ -61,6 +61,7 @@ enum class AttributeKind {
     kInterpolate,
     kInvariant,
     kLocation,
+    kMustUse,
     kOffset,
     kSize,
     kStage,
@@ -113,6 +114,8 @@ static utils::Vector<const ast::Attribute*, 2> createAttributes(const Source& so
             return {builder.Location(source, 1_a)};
         case AttributeKind::kOffset:
             return {builder.MemberOffset(source, 4_a)};
+        case AttributeKind::kMustUse:
+            return {builder.MustUse(source)};
         case AttributeKind::kSize:
             return {builder.MemberSize(source, 16_a)};
         case AttributeKind::kStage:
@@ -158,6 +161,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, false},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -195,6 +199,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, false},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -246,6 +251,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, false},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -284,6 +290,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          // kInterpolate tested separately (requires @location)
                                          TestParams{AttributeKind::kInvariant, true},
                                          TestParams{AttributeKind::kLocation, true},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -338,6 +345,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, true},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, true},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -389,6 +397,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, false},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -440,6 +449,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, true},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -489,6 +499,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          // kInterpolate tested separately (requires @location)
                                          TestParams{AttributeKind::kInvariant, true},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -593,6 +604,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, false},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -629,6 +641,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          // kInterpolate tested separately (requires @location)
                                          // kInvariant tested separately (requires position builtin)
                                          TestParams{AttributeKind::kLocation, true},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, true},
                                          TestParams{AttributeKind::kSize, true},
                                          TestParams{AttributeKind::kStage, false},
@@ -871,6 +884,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, false},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -911,6 +925,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, false},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -963,6 +978,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, false},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -1007,6 +1023,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, false},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -1113,6 +1130,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kInterpolate, false},
                                          TestParams{AttributeKind::kInvariant, false},
                                          TestParams{AttributeKind::kLocation, false},
+                                         TestParams{AttributeKind::kMustUse, false},
                                          TestParams{AttributeKind::kOffset, false},
                                          TestParams{AttributeKind::kSize, false},
                                          TestParams{AttributeKind::kStage, false},
@@ -1413,6 +1431,24 @@ TEST_F(InvariantAttributeTests, InvariantWithoutPosition) {
 }  // namespace
 }  // namespace InvariantAttributeTests
 
+namespace MustUseAttributeTests {
+namespace {
+
+using MustUseAttributeTests = ResolverTest;
+TEST_F(MustUseAttributeTests, MustUse) {
+    Func("main", utils::Empty, ty.vec4<f32>(),
+         utils::Vector{
+             Return(Call(ty.vec4<f32>())),
+         },
+         utils::Vector{
+             MustUse(Source{{12, 34}}),
+         });
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+}  // namespace
+}  // namespace MustUseAttributeTests
+
 namespace WorkgroupAttributeTests {
 namespace {
 
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index 4db6c1af50..b558e32dbb 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -1016,7 +1016,7 @@ bool Validator::Function(const sem::Function* func, ast::PipelineStage stage) co
                 return false;
             }
         } else if (!attr->IsAnyOf<ast::DiagnosticAttribute, ast::StageAttribute,
-                                  ast::InternalAttribute>()) {
+                                  ast::MustUseAttribute, ast::InternalAttribute>()) {
             AddError("attribute is not valid for functions", attr->source);
             return false;
         }
@@ -2544,8 +2544,7 @@ bool Validator::CheckTypeAccessAddressSpace(
     }
 
     if (address_space == builtin::AddressSpace::kStorage && access == builtin::Access::kWrite) {
-        // The access mode for the storage address space can only be 'read' or
-        // 'read_write'.
+        // The access mode for the storage address space can only be 'read' or 'read_write'.
         AddError("access mode 'write' is not valid for the 'storage' address space", source);
         return false;
     }
diff --git a/src/tint/writer/wgsl/generator_impl.cc b/src/tint/writer/wgsl/generator_impl.cc
index a52bff094e..4363b62aca 100644
--- a/src/tint/writer/wgsl/generator_impl.cc
+++ b/src/tint/writer/wgsl/generator_impl.cc
@@ -600,6 +600,10 @@ bool GeneratorImpl::EmitAttributes(std::ostream& out,
                 out << ")";
                 return true;
             },
+            [&](const ast::MustUseAttribute*) {
+                out << "must_use";
+                return true;
+            },
             [&](const ast::StructMemberOffsetAttribute* offset) {
                 out << "offset(";
                 if (!EmitExpression(out, offset->expr)) {
diff --git a/src/tint/writer/wgsl/generator_impl_function_test.cc b/src/tint/writer/wgsl/generator_impl_function_test.cc
index 04103ffb4c..68b2191667 100644
--- a/src/tint/writer/wgsl/generator_impl_function_test.cc
+++ b/src/tint/writer/wgsl/generator_impl_function_test.cc
@@ -86,6 +86,27 @@ TEST_F(WgslGeneratorImplTest, Emit_Function_WithAttribute_WorkgroupSize) {
 )");
 }
 
+TEST_F(WgslGeneratorImplTest, Emit_Function_WithAttribute_MustUse) {
+    auto* func = Func("my_func", utils::Empty, ty.void_(),
+                      utils::Vector{
+                          Return(),
+                      },
+                      utils::Vector{
+                          MustUse(),
+                      });
+
+    GeneratorImpl& gen = Build();
+
+    gen.increment_indent();
+
+    ASSERT_TRUE(gen.EmitFunction(func));
+    EXPECT_EQ(gen.result(), R"(  @must_use
+  fn my_func() {
+    return;
+  }
+)");
+}
+
 TEST_F(WgslGeneratorImplTest, Emit_Function_WithAttribute_WorkgroupSize_WithIdent) {
     GlobalConst("height", ty.i32(), Expr(2_i));
     auto* func = Func("my_func", utils::Empty, ty.void_(),