diff --git a/src/ast/builtin.cc b/src/ast/builtin.cc
index 03a689e4fa..3dac5cdf76 100644
--- a/src/ast/builtin.cc
+++ b/src/ast/builtin.cc
@@ -63,6 +63,10 @@ std::ostream& operator<<(std::ostream& out, Builtin builtin) {
       out << "sample_index";
       break;
     }
+    case Builtin::kSampleMask: {
+      out << "sample_mask";
+      break;
+    }
     case Builtin::kSampleMaskIn: {
       out << "sample_mask_in";
       break;
diff --git a/src/ast/builtin.h b/src/ast/builtin.h
index 3806e4477b..e065d69df8 100644
--- a/src/ast/builtin.h
+++ b/src/ast/builtin.h
@@ -33,8 +33,9 @@ enum class Builtin {
   kLocalInvocationIndex,
   kGlobalInvocationId,
   kSampleIndex,
-  kSampleMaskIn,
-  kSampleMaskOut,
+  kSampleMask,
+  kSampleMaskIn,   // TODO(crbug.com/tint/715): Remove this
+  kSampleMaskOut,  // TODO(crbug.com/tint/715): Remove this
 
   // Below are not currently WGSL builtins, but are included in this enum as
   // they are used by certain backends.
diff --git a/src/inspector/inspector_test.cc b/src/inspector/inspector_test.cc
index e88edbc51a..21555637cd 100644
--- a/src/inspector/inspector_test.cc
+++ b/src/inspector/inspector_test.cc
@@ -1041,7 +1041,7 @@ TEST_F(InspectorGetEntryPointTest, BuiltInsNotStageVariables) {
   auto* in_var1 = Param("in_var1", ty.u32(), {Location(0u)});
   Func("foo", {in_var0, in_var1}, ty.u32(), {Return(Expr("in_var1"))},
        {Stage(ast::PipelineStage::kFragment)},
-       {Builtin(ast::Builtin::kSampleMaskOut)});
+       {Builtin(ast::Builtin::kSampleMask)});
   Inspector& inspector = Build();
 
   auto result = inspector.GetEntryPoints();
diff --git a/src/reader/spirv/enum_converter.cc b/src/reader/spirv/enum_converter.cc
index 15491eb824..364b596669 100644
--- a/src/reader/spirv/enum_converter.cc
+++ b/src/reader/spirv/enum_converter.cc
@@ -66,7 +66,7 @@ ast::StorageClass EnumConverter::ToStorageClass(const SpvStorageClass sc) {
   return ast::StorageClass::kNone;
 }
 
-ast::Builtin EnumConverter::ToBuiltin(SpvBuiltIn b, ast::StorageClass sc) {
+ast::Builtin EnumConverter::ToBuiltin(SpvBuiltIn b) {
   switch (b) {
     case SpvBuiltInPosition:
       return ast::Builtin::kPosition;
@@ -89,8 +89,7 @@ ast::Builtin EnumConverter::ToBuiltin(SpvBuiltIn b, ast::StorageClass sc) {
     case SpvBuiltInSampleId:
       return ast::Builtin::kSampleIndex;
     case SpvBuiltInSampleMask:
-      return sc == ast::StorageClass::kInput ? ast::Builtin::kSampleMaskIn
-                                             : ast::Builtin::kSampleMaskOut;
+      return ast::Builtin::kSampleMask;
     default:
       break;
   }
diff --git a/src/reader/spirv/enum_converter.h b/src/reader/spirv/enum_converter.h
index 7c9d4f77a2..5831a88a73 100644
--- a/src/reader/spirv/enum_converter.h
+++ b/src/reader/spirv/enum_converter.h
@@ -50,9 +50,8 @@ class EnumConverter {
   /// Converts a SPIR-V Builtin value a Tint Builtin.
   /// On failure, logs an error and returns kNone
   /// @param b the SPIR-V builtin
-  /// @param sc the Tint storage class
   /// @returns a Tint AST builtin
-  ast::Builtin ToBuiltin(SpvBuiltIn b, ast::StorageClass sc);
+  ast::Builtin ToBuiltin(SpvBuiltIn b);
 
   /// Converts a possibly arrayed SPIR-V Dim to a Tint texture dimension.
   /// On failure, logs an error and returns kNone
diff --git a/src/reader/spirv/enum_converter_test.cc b/src/reader/spirv/enum_converter_test.cc
index 7011216355..58e69241ef 100644
--- a/src/reader/spirv/enum_converter_test.cc
+++ b/src/reader/spirv/enum_converter_test.cc
@@ -162,12 +162,11 @@ INSTANTIATE_TEST_SUITE_P(EnumConverterBad,
 
 struct BuiltinCase {
   SpvBuiltIn builtin;
-  ast::StorageClass sc;
   bool expect_success;
   ast::Builtin expected;
 };
 inline std::ostream& operator<<(std::ostream& out, BuiltinCase bc) {
-  out << "BuiltinCase{ SpvBuiltIn:" << int(bc.builtin) << " sc:" << int(bc.sc)
+  out << "BuiltinCase{ SpvBuiltIn:" << int(bc.builtin)
       << " expect_success?:" << int(bc.expect_success)
       << " expected:" << int(bc.expected) << "}";
   return out;
@@ -192,7 +191,7 @@ class SpvBuiltinTest : public testing::TestWithParam<BuiltinCase> {
 TEST_P(SpvBuiltinTest, Samples) {
   const auto params = GetParam();
 
-  const auto result = converter_.ToBuiltin(params.builtin, params.sc);
+  const auto result = converter_.ToBuiltin(params.builtin);
   EXPECT_EQ(success_, params.expect_success);
   if (params.expect_success) {
     EXPECT_EQ(result, params.expected);
@@ -207,46 +206,35 @@ INSTANTIATE_TEST_SUITE_P(
     EnumConverterGood_Input,
     SpvBuiltinTest,
     testing::Values(
-        BuiltinCase{SpvBuiltInPosition, ast::StorageClass::kInput, true,
-                    ast::Builtin::kPosition},
-        BuiltinCase{SpvBuiltInInstanceIndex, ast::StorageClass::kInput, true,
+        BuiltinCase{SpvBuiltInPosition, true, ast::Builtin::kPosition},
+        BuiltinCase{SpvBuiltInInstanceIndex, true,
                     ast::Builtin::kInstanceIndex},
-        BuiltinCase{SpvBuiltInFrontFacing, ast::StorageClass::kInput, true,
-                    ast::Builtin::kFrontFacing},
-        BuiltinCase{SpvBuiltInFragCoord, ast::StorageClass::kInput, true,
-                    ast::Builtin::kPosition},
-        BuiltinCase{SpvBuiltInLocalInvocationId, ast::StorageClass::kInput,
-                    true, ast::Builtin::kLocalInvocationId},
-        BuiltinCase{SpvBuiltInLocalInvocationIndex, ast::StorageClass::kInput,
-                    true, ast::Builtin::kLocalInvocationIndex},
-        BuiltinCase{SpvBuiltInGlobalInvocationId, ast::StorageClass::kInput,
-                    true, ast::Builtin::kGlobalInvocationId},
-        BuiltinCase{SpvBuiltInSampleId, ast::StorageClass::kInput, true,
-                    ast::Builtin::kSampleIndex},
-        BuiltinCase{SpvBuiltInSampleMask, ast::StorageClass::kInput, true,
-                    ast::Builtin::kSampleMaskIn}));
+        BuiltinCase{SpvBuiltInFrontFacing, true, ast::Builtin::kFrontFacing},
+        BuiltinCase{SpvBuiltInFragCoord, true, ast::Builtin::kPosition},
+        BuiltinCase{SpvBuiltInLocalInvocationId, true,
+                    ast::Builtin::kLocalInvocationId},
+        BuiltinCase{SpvBuiltInLocalInvocationIndex, true,
+                    ast::Builtin::kLocalInvocationIndex},
+        BuiltinCase{SpvBuiltInGlobalInvocationId, true,
+                    ast::Builtin::kGlobalInvocationId},
+        BuiltinCase{SpvBuiltInSampleId, true, ast::Builtin::kSampleIndex},
+        BuiltinCase{SpvBuiltInSampleMask, true, ast::Builtin::kSampleMask}));
 
 INSTANTIATE_TEST_SUITE_P(
     EnumConverterGood_Output,
     SpvBuiltinTest,
-    testing::Values(BuiltinCase{SpvBuiltInPosition, ast::StorageClass::kOutput,
-                                true, ast::Builtin::kPosition},
-                    BuiltinCase{SpvBuiltInFragDepth, ast::StorageClass::kOutput,
-                                true, ast::Builtin::kFragDepth},
-                    BuiltinCase{SpvBuiltInSampleMask,
-                                ast::StorageClass::kOutput, true,
-                                ast::Builtin::kSampleMaskOut}));
+    testing::Values(
+        BuiltinCase{SpvBuiltInPosition, true, ast::Builtin::kPosition},
+        BuiltinCase{SpvBuiltInFragDepth, true, ast::Builtin::kFragDepth},
+        BuiltinCase{SpvBuiltInSampleMask, true, ast::Builtin::kSampleMask}));
 
 INSTANTIATE_TEST_SUITE_P(
     EnumConverterBad,
     SpvBuiltinTest,
     testing::Values(
-        BuiltinCase{static_cast<SpvBuiltIn>(9999), ast::StorageClass::kInput,
-                    false, ast::Builtin::kNone},
-        BuiltinCase{static_cast<SpvBuiltIn>(9999), ast::StorageClass::kOutput,
-                    false, ast::Builtin::kNone},
-        BuiltinCase{SpvBuiltInNumWorkgroups, ast::StorageClass::kInput, false,
-                    ast::Builtin::kNone}));
+        BuiltinCase{static_cast<SpvBuiltIn>(9999), false, ast::Builtin::kNone},
+        BuiltinCase{static_cast<SpvBuiltIn>(9999), false, ast::Builtin::kNone},
+        BuiltinCase{SpvBuiltInNumWorkgroups, false, ast::Builtin::kNone}));
 
 // Dim
 
diff --git a/src/reader/spirv/parser_impl.cc b/src/reader/spirv/parser_impl.cc
index b5c6eab5da..fd3ada40ff 100644
--- a/src/reader/spirv/parser_impl.cc
+++ b/src/reader/spirv/parser_impl.cc
@@ -1304,7 +1304,7 @@ ast::Variable* ParserImpl::MakeVariable(uint32_t id,
         default:
           break;
       }
-      auto ast_builtin = enum_converter_.ToBuiltin(spv_builtin, sc);
+      auto ast_builtin = enum_converter_.ToBuiltin(spv_builtin);
       if (ast_builtin == ast::Builtin::kNone) {
         return nullptr;
       }
diff --git a/src/reader/spirv/parser_impl_module_var_test.cc b/src/reader/spirv/parser_impl_module_var_test.cc
index 3af4662e04..b70bdef136 100644
--- a/src/reader/spirv/parser_impl_module_var_test.cc
+++ b/src/reader/spirv/parser_impl_module_var_test.cc
@@ -2473,7 +2473,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_In_U32_Direct) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_in}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     in
@@ -2515,7 +2515,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_In_U32_CopyObject) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_in}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     in
@@ -2557,7 +2557,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_In_U32_AccessChain) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_in}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     in
@@ -2598,7 +2598,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_In_I32_Direct) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_in}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     in
@@ -2643,7 +2643,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_In_I32_CopyObject) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_in}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     in
@@ -2688,7 +2688,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_In_I32_AccessChain) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_in}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     in
@@ -2751,7 +2751,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_Out_U32_Direct) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_out}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     out
@@ -2787,7 +2787,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_Out_U32_CopyObject) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_out}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     out
@@ -2823,7 +2823,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_Out_U32_AccessChain) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_out}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     out
@@ -2858,7 +2858,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_Out_I32_Direct) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_out}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     out
@@ -2900,7 +2900,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_Out_I32_CopyObject) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_out}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     out
@@ -2942,7 +2942,7 @@ TEST_F(SpvModuleScopeVarParserTest, SampleMask_Out_I32_AccessChain) {
   EXPECT_THAT(module_str, HasSubstr(R"(
   Variable{
     Decorations{
-      BuiltinDecoration{sample_mask_out}
+      BuiltinDecoration{sample_mask}
     }
     x_1
     out
diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc
index 923053c832..dd0b884648 100644
--- a/src/reader/wgsl/parser_impl.cc
+++ b/src/reader/wgsl/parser_impl.cc
@@ -88,6 +88,9 @@ ast::Builtin ident_to_builtin(const std::string& str) {
   if (str == "sample_index") {
     return ast::Builtin::kSampleIndex;
   }
+  if (str == "sample_mask") {
+    return ast::Builtin::kSampleMask;
+  }
   if (str == "sample_mask_in") {
     return ast::Builtin::kSampleMaskIn;
   }
@@ -1382,6 +1385,12 @@ Expect<ast::Builtin> ParserImpl::expect_builtin() {
   if (builtin == ast::Builtin::kFragCoord) {
     deprecated(ident.source, "use 'position' instead of 'frag_coord'");
   }
+  if (builtin == ast::Builtin::kSampleMaskIn) {
+    deprecated(ident.source, "use 'sample_mask' instead of 'sample_mask_in'");
+  }
+  if (builtin == ast::Builtin::kSampleMaskOut) {
+    deprecated(ident.source, "use 'sample_mask' instead of 'sample_mask_out'");
+  }
 
   return {builtin, ident.source};
 }
diff --git a/src/reader/wgsl/parser_impl_variable_decoration_test.cc b/src/reader/wgsl/parser_impl_variable_decoration_test.cc
index ea44069c51..9430cf1172 100644
--- a/src/reader/wgsl/parser_impl_variable_decoration_test.cc
+++ b/src/reader/wgsl/parser_impl_variable_decoration_test.cc
@@ -122,6 +122,7 @@ INSTANTIATE_TEST_SUITE_P(
                     ast::Builtin::kLocalInvocationIndex},
         BuiltinData{"global_invocation_id", ast::Builtin::kGlobalInvocationId},
         BuiltinData{"sample_index", ast::Builtin::kSampleIndex},
+        BuiltinData{"sample_mask", ast::Builtin::kSampleMask},
         BuiltinData{"sample_mask_in", ast::Builtin::kSampleMaskIn},
         BuiltinData{"sample_mask_out", ast::Builtin::kSampleMaskOut}));
 
@@ -327,6 +328,50 @@ builtin(frag_coord)
 )");
 }
 
+TEST_F(ParserImplTest, Decoration_SampleMaskIn_Deprecated) {
+  auto p = parser("builtin(sample_mask_in)");
+  auto deco = p->decoration();
+  EXPECT_TRUE(deco.matched);
+  EXPECT_FALSE(deco.errored);
+  ASSERT_NE(deco.value, nullptr);
+  auto* var_deco = deco.value->As<ast::Decoration>();
+  ASSERT_NE(var_deco, nullptr);
+  ASSERT_FALSE(p->has_error());
+  ASSERT_TRUE(var_deco->Is<ast::BuiltinDecoration>());
+
+  auto* builtin = var_deco->As<ast::BuiltinDecoration>();
+  EXPECT_EQ(builtin->value(), ast::Builtin::kSampleMaskIn);
+
+  EXPECT_EQ(
+      p->builder().Diagnostics().str(),
+      R"(test.wgsl:1:9 warning: use of deprecated language feature: use 'sample_mask' instead of 'sample_mask_in'
+builtin(sample_mask_in)
+        ^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(ParserImplTest, Decoration_SampleMaskOut_Deprecated) {
+  auto p = parser("builtin(sample_mask_out)");
+  auto deco = p->decoration();
+  EXPECT_TRUE(deco.matched);
+  EXPECT_FALSE(deco.errored);
+  ASSERT_NE(deco.value, nullptr);
+  auto* var_deco = deco.value->As<ast::Decoration>();
+  ASSERT_NE(var_deco, nullptr);
+  ASSERT_FALSE(p->has_error());
+  ASSERT_TRUE(var_deco->Is<ast::BuiltinDecoration>());
+
+  auto* builtin = var_deco->As<ast::BuiltinDecoration>();
+  EXPECT_EQ(builtin->value(), ast::Builtin::kSampleMaskOut);
+
+  EXPECT_EQ(
+      p->builder().Diagnostics().str(),
+      R"(test.wgsl:1:9 warning: use of deprecated language feature: use 'sample_mask' instead of 'sample_mask_out'
+builtin(sample_mask_out)
+        ^^^^^^^^^^^^^^^
+)");
+}
+
 }  // namespace
 }  // namespace wgsl
 }  // namespace reader
diff --git a/src/transform/canonicalize_entry_point_io_test.cc b/src/transform/canonicalize_entry_point_io_test.cc
index 923a06f96f..a7e4bf7009 100644
--- a/src/transform/canonicalize_entry_point_io_test.cc
+++ b/src/transform/canonicalize_entry_point_io_test.cc
@@ -197,7 +197,7 @@ TEST_F(CanonicalizeEntryPointIOTest, Return_Struct) {
   auto* src = R"(
 struct FragOutput {
   [[builtin(frag_depth)]] depth : f32;
-  [[builtin(sample_mask_out)]] mask : u32;
+  [[builtin(sample_mask)]] mask : u32;
   [[location(0)]] color : vec4<f32>;
 };
 
@@ -221,7 +221,7 @@ struct FragOutput {
 struct tint_symbol {
   [[builtin(frag_depth)]]
   depth : f32;
-  [[builtin(sample_mask_out)]]
+  [[builtin(sample_mask)]]
   mask : u32;
   [[location(0)]]
   color : vec4<f32>;
diff --git a/src/transform/spirv.cc b/src/transform/spirv.cc
index 4e0bed9511..2d651e16c9 100644
--- a/src/transform/spirv.cc
+++ b/src/transform/spirv.cc
@@ -59,11 +59,11 @@ void Spirv::HandleEntryPointIOTypes(CloneContext& ctx) const {
   // ```
   // struct FragmentInput {
   //   [[builtin(sample_index)]] sample_index : u32;
-  //   [[builtin(sample_mask_in)]] sample_mask_in : u32;
+  //   [[builtin(sample_mask)]] sample_mask : u32;
   // };
   // struct FragmentOutput {
   //   [[builtin(frag_depth)]] depth: f32;
-  //   [[builtin(sample_mask_out)]] mask_out : u32;
+  //   [[builtin(sample_mask)]] mask_out : u32;
   // };
   //
   // [[stage(fragment)]]
@@ -72,7 +72,7 @@ void Spirv::HandleEntryPointIOTypes(CloneContext& ctx) const {
   //   samples : FragmentInput
   // ) -> FragmentOutput {
   //   var output : FragmentOutput = FragmentOutput(1.0,
-  //                                                samples.sample_mask_in);
+  //                                                samples.sample_mask);
   //   return output;
   // }
   // ```
@@ -81,7 +81,7 @@ void Spirv::HandleEntryPointIOTypes(CloneContext& ctx) const {
   // ```
   // struct FragmentInput {
   //   sample_index : u32;
-  //   sample_mask_in : u32;
+  //   sample_mask : u32;
   // };
   // struct FragmentOutput {
   //   depth: f32;
@@ -90,9 +90,9 @@ void Spirv::HandleEntryPointIOTypes(CloneContext& ctx) const {
   //
   // [[builtin(position)]] var<in> coord : vec4<f32>,
   // [[builtin(sample_index)]] var<in> sample_index : u32,
-  // [[builtin(sample_mask_in)]] var<in> sample_mask_in : u32,
+  // [[builtin(sample_mask)]] var<in> sample_mask : u32,
   // [[builtin(frag_depth)]] var<out> depth: f32;
-  // [[builtin(sample_mask_out)]] var<out> mask_out : u32;
+  // [[builtin(sample_mask)]] var<out> mask_out : u32;
   //
   // fn frag_main_ret(retval : FragmentOutput) {
   //   depth = reval.depth;
@@ -101,9 +101,9 @@ void Spirv::HandleEntryPointIOTypes(CloneContext& ctx) const {
   //
   // [[stage(fragment)]]
   // fn frag_main() {
-  //   let samples : FragmentInput(sample_index, sample_mask_in);
+  //   let samples : FragmentInput(sample_index, sample_mask);
   //   var output : FragmentOutput = FragmentOutput(1.0,
-  //                                                samples.sample_mask_in);
+  //                                                samples.sample_mask);
   //   frag_main_ret(output);
   //   return;
   // }
@@ -189,19 +189,19 @@ void Spirv::HandleEntryPointIOTypes(CloneContext& ctx) const {
 }
 
 void Spirv::HandleSampleMaskBuiltins(CloneContext& ctx) const {
-  // Find global variables decorated with [[builtin(sample_mask_{in,out})]] and
+  // Find global variables decorated with [[builtin(sample_mask)]] and
   // change their type from `u32` to `array<u32, 1>`, as required by Vulkan.
   //
   // Before:
   // ```
-  // [[builtin(sample_mask_out)]] var<out> mask_out : u32;
+  // [[builtin(sample_mask)]] var<out> mask_out : u32;
   // fn main() {
   //   mask_out = 1u;
   // }
   // ```
   // After:
   // ```
-  // [[builtin(sample_mask_out)]] var<out> mask_out : array<u32, 1>;
+  // [[builtin(sample_mask)]] var<out> mask_out : array<u32, 1>;
   // fn main() {
   //   mask_out[0] = 1u;
   // }
@@ -210,7 +210,8 @@ void Spirv::HandleSampleMaskBuiltins(CloneContext& ctx) const {
   for (auto* var : ctx.src->AST().GlobalVariables()) {
     for (auto* deco : var->decorations()) {
       if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
-        if (builtin->value() != ast::Builtin::kSampleMaskIn &&
+        if (builtin->value() != ast::Builtin::kSampleMask &&
+            builtin->value() != ast::Builtin::kSampleMaskIn &&
             builtin->value() != ast::Builtin::kSampleMaskOut) {
           continue;
         }
diff --git a/src/transform/spirv_test.cc b/src/transform/spirv_test.cc
index 45e0128be9..240fae3091 100644
--- a/src/transform/spirv_test.cc
+++ b/src/transform/spirv_test.cc
@@ -487,9 +487,9 @@ TEST_F(SpirvTest, HandleSampleMaskBuiltins_Basic) {
   auto* src = R"(
 [[builtin(sample_index)]] var<in> sample_index : u32;
 
-[[builtin(sample_mask_in)]] var<in> mask_in : u32;
+[[builtin(sample_mask)]] var<in> mask_in : u32;
 
-[[builtin(sample_mask_out)]] var<out> mask_out : u32;
+[[builtin(sample_mask)]] var<out> mask_out : u32;
 
 [[stage(fragment)]]
 fn main() {
@@ -500,9 +500,9 @@ fn main() {
   auto* expect = R"(
 [[builtin(sample_index)]] var<in> sample_index : u32;
 
-[[builtin(sample_mask_in)]] var<in> mask_in : array<u32, 1>;
+[[builtin(sample_mask)]] var<in> mask_in : array<u32, 1>;
 
-[[builtin(sample_mask_out)]] var<out> mask_out : array<u32, 1>;
+[[builtin(sample_mask)]] var<out> mask_out : array<u32, 1>;
 
 [[stage(fragment)]]
 fn main() {
@@ -517,9 +517,9 @@ fn main() {
 
 TEST_F(SpirvTest, HandleSampleMaskBuiltins_FunctionArg) {
   auto* src = R"(
-[[builtin(sample_mask_in)]] var<in> mask_in : u32;
+[[builtin(sample_mask)]] var<in> mask_in : u32;
 
-[[builtin(sample_mask_out)]] var<out> mask_out : u32;
+[[builtin(sample_mask)]] var<out> mask_out : u32;
 
 fn filter(mask: u32) -> u32 {
   return (mask & 3u);
@@ -536,9 +536,9 @@ fn main() {
 )";
 
   auto* expect = R"(
-[[builtin(sample_mask_in)]] var<in> mask_in : array<u32, 1>;
+[[builtin(sample_mask)]] var<in> mask_in : array<u32, 1>;
 
-[[builtin(sample_mask_out)]] var<out> mask_out : array<u32, 1>;
+[[builtin(sample_mask)]] var<out> mask_out : array<u32, 1>;
 
 fn filter(mask : u32) -> u32 {
   return (mask & 3u);
@@ -578,8 +578,8 @@ TEST_F(SpirvTest, MultipleTransforms) {
   auto* src = R"(
 [[stage(fragment)]]
 fn main([[builtin(sample_index)]] sample_index : u32,
-        [[builtin(sample_mask_in)]] mask_in : u32)
-        -> [[builtin(sample_mask_out)]] u32 {
+        [[builtin(sample_mask)]] mask_in : u32)
+        -> [[builtin(sample_mask)]] u32 {
   return mask_in;
 }
 )";
@@ -587,9 +587,9 @@ fn main([[builtin(sample_index)]] sample_index : u32,
   auto* expect = R"(
 [[builtin(sample_index)]] var<in> tint_symbol : u32;
 
-[[builtin(sample_mask_in)]] var<in> tint_symbol_1 : array<u32, 1>;
+[[builtin(sample_mask)]] var<in> tint_symbol_1 : array<u32, 1>;
 
-[[builtin(sample_mask_out)]] var<out> tint_symbol_3 : array<u32, 1>;
+[[builtin(sample_mask)]] var<out> tint_symbol_3 : array<u32, 1>;
 
 fn tint_symbol_4(tint_symbol_2 : u32) {
   tint_symbol_3[0] = tint_symbol_2;
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 3738d695ba..a64758588f 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -1954,6 +1954,8 @@ std::string GeneratorImpl::builtin_to_attribute(ast::Builtin builtin) const {
       return "SV_DispatchThreadID";
     case ast::Builtin::kSampleIndex:
       return "SV_SampleIndex";
+    case ast::Builtin::kSampleMask:
+      return "SV_Coverage";
     case ast::Builtin::kSampleMaskIn:
       return "SV_Coverage";
     case ast::Builtin::kSampleMaskOut:
diff --git a/src/writer/hlsl/generator_impl_test.cc b/src/writer/hlsl/generator_impl_test.cc
index 773c1deb54..908c49f139 100644
--- a/src/writer/hlsl/generator_impl_test.cc
+++ b/src/writer/hlsl/generator_impl_test.cc
@@ -83,6 +83,7 @@ INSTANTIATE_TEST_SUITE_P(
         HlslBuiltinData{ast::Builtin::kGlobalInvocationId,
                         "SV_DispatchThreadID"},
         HlslBuiltinData{ast::Builtin::kSampleIndex, "SV_SampleIndex"},
+        HlslBuiltinData{ast::Builtin::kSampleMask, "SV_Coverage"},
         HlslBuiltinData{ast::Builtin::kSampleMaskIn, "SV_Coverage"},
         HlslBuiltinData{ast::Builtin::kSampleMaskOut, "SV_Coverage"}));
 
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index ddee33e177..ea53d7db20 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -1354,6 +1354,8 @@ std::string GeneratorImpl::builtin_to_attribute(ast::Builtin builtin) const {
       return "thread_position_in_grid";
     case ast::Builtin::kSampleIndex:
       return "sample_id";
+    case ast::Builtin::kSampleMask:
+      return "sample_mask";
     case ast::Builtin::kSampleMaskIn:
       return "sample_mask";
     case ast::Builtin::kSampleMaskOut:
diff --git a/src/writer/msl/generator_impl_test.cc b/src/writer/msl/generator_impl_test.cc
index 1fa7776416..bc2ee5532b 100644
--- a/src/writer/msl/generator_impl_test.cc
+++ b/src/writer/msl/generator_impl_test.cc
@@ -74,6 +74,7 @@ INSTANTIATE_TEST_SUITE_P(
                     MslBuiltinData{ast::Builtin::kGlobalInvocationId,
                                    "thread_position_in_grid"},
                     MslBuiltinData{ast::Builtin::kSampleIndex, "sample_id"},
+                    MslBuiltinData{ast::Builtin::kSampleMask, "sample_mask"},
                     MslBuiltinData{ast::Builtin::kSampleMaskIn, "sample_mask"},
                     MslBuiltinData{ast::Builtin::kSampleMaskOut,
                                    "sample_mask"}));
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index 50278ca0fe..4c8187bf0c 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -3308,6 +3308,8 @@ SpvBuiltIn Builder::ConvertBuiltin(ast::Builtin builtin,
     case ast::Builtin::kSampleIndex:
       push_capability(SpvCapabilitySampleRateShading);
       return SpvBuiltInSampleId;
+    case ast::Builtin::kSampleMask:
+      return SpvBuiltInSampleMask;
     case ast::Builtin::kSampleMaskIn:
       return SpvBuiltInSampleMask;
     case ast::Builtin::kSampleMaskOut:
diff --git a/src/writer/spirv/builder_global_variable_test.cc b/src/writer/spirv/builder_global_variable_test.cc
index 3b636878eb..b6bf64c827 100644
--- a/src/writer/spirv/builder_global_variable_test.cc
+++ b/src/writer/spirv/builder_global_variable_test.cc
@@ -370,6 +370,10 @@ INSTANTIATE_TEST_SUITE_P(
                     ast::StorageClass::kInput, SpvBuiltInGlobalInvocationId},
         BuiltinData{ast::Builtin::kSampleIndex, ast::StorageClass::kInput,
                     SpvBuiltInSampleId},
+        BuiltinData{ast::Builtin::kSampleMask, ast::StorageClass::kInput,
+                    SpvBuiltInSampleMask},
+        BuiltinData{ast::Builtin::kSampleMask, ast::StorageClass::kOutput,
+                    SpvBuiltInSampleMask},
         BuiltinData{ast::Builtin::kSampleMaskIn, ast::StorageClass::kInput,
                     SpvBuiltInSampleMask},
         BuiltinData{ast::Builtin::kSampleMaskOut, ast::StorageClass::kOutput,
@@ -622,16 +626,16 @@ TEST_F(BuilderTest, SampleIndex) {
 
 TEST_F(BuilderTest, SampleMask) {
   // Input:
-  // [[builtin(sample_mask_in)]] var<in> mask_in : u32;
-  // [[builtin(sample_mask_out)]] var<out> mask_out : u32;
+  // [[builtin(sample_mask)]] var<in> mask_in : u32;
+  // [[builtin(sample_mask)]] var<out> mask_out : u32;
   // [[stage(fragment)]]
   // fn main() {
   //   mask_out = mask_in;
   // }
 
   // After sanitization:
-  // [[builtin(sample_mask_in)]] var<in> mask_in : array<u32, 1>;
-  // [[builtin(sample_mask_out)]] var<out> mask_out : array<u32, 1>;
+  // [[builtin(sample_mask)]] var<in> mask_in : array<u32, 1>;
+  // [[builtin(sample_mask)]] var<out> mask_out : array<u32, 1>;
   // [[stage(fragment)]]
   // fn main() {
   //   mask_out[0] = mask_in[0];
@@ -639,11 +643,11 @@ TEST_F(BuilderTest, SampleMask) {
 
   Global("mask_in", ty.u32(), ast::StorageClass::kInput, nullptr,
          ast::DecorationList{
-             create<ast::BuiltinDecoration>(ast::Builtin::kSampleMaskIn),
+             create<ast::BuiltinDecoration>(ast::Builtin::kSampleMask),
          });
   Global("mask_out", ty.u32(), ast::StorageClass::kOutput, nullptr,
          ast::DecorationList{
-             create<ast::BuiltinDecoration>(ast::Builtin::kSampleMaskOut),
+             create<ast::BuiltinDecoration>(ast::Builtin::kSampleMask),
          });
   Func("main", ast::VariableList{}, ty.void_(),
        ast::StatementList{