diff --git a/docs/tint/origin-trial-changes.md b/docs/tint/origin-trial-changes.md
index ef741c9..2371a83 100644
--- a/docs/tint/origin-trial-changes.md
+++ b/docs/tint/origin-trial-changes.md
@@ -5,6 +5,7 @@
 ### Breaking changes
 
 * The `@block` attribute has been removed. [tint:1324](crbug.com/tint/1324)
+* Attributes using `[[attribute]]` syntax are no longer supported. [tint:1382](crbug.com/tint/1382)
 
 ## Changes for M101
 
diff --git a/src/tint/reader/wgsl/lexer.cc b/src/tint/reader/wgsl/lexer.cc
index d5263aa..1d01375 100644
--- a/src/tint/reader/wgsl/lexer.cc
+++ b/src/tint/reader/wgsl/lexer.cc
@@ -801,14 +801,6 @@
     type = Token::Type::kAttr;
     pos_ += 1;
     location_.column += 1;
-  } else if (matches(pos_, "[[")) {
-    type = Token::Type::kAttrLeft;
-    pos_ += 2;
-    location_.column += 2;
-  } else if (matches(pos_, "]]")) {
-    type = Token::Type::kAttrRight;
-    pos_ += 2;
-    location_.column += 2;
   } else if (matches(pos_, "(")) {
     type = Token::Type::kParenLeft;
     pos_ += 1;
diff --git a/src/tint/reader/wgsl/lexer_test.cc b/src/tint/reader/wgsl/lexer_test.cc
index 93cbf0b..f9354e4 100644
--- a/src/tint/reader/wgsl/lexer_test.cc
+++ b/src/tint/reader/wgsl/lexer_test.cc
@@ -801,8 +801,6 @@
                     TokenData{"&&", Token::Type::kAndAnd},
                     TokenData{"->", Token::Type::kArrow},
                     TokenData{"@", Token::Type::kAttr},
-                    TokenData{"[[", Token::Type::kAttrLeft},
-                    TokenData{"]]", Token::Type::kAttrRight},
                     TokenData{"/", Token::Type::kForwardSlash},
                     TokenData{"!", Token::Type::kBang},
                     TokenData{"[", Token::Type::kBracketLeft},
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index 55abfcf..eaf5fb7 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -122,14 +122,6 @@
 const char kStrideAttribute[] = "stride";
 const char kWorkgroupSizeAttribute[] = "workgroup_size";
 
-bool is_attribute(Token t) {
-  return t == kAlignAttribute || t == kBindingAttribute ||
-         t == kBuiltinAttribute || t == kGroupAttribute || t == kIdAttribute ||
-         t == kInterpolateAttribute || t == kLocationAttribute ||
-         t == kSizeAttribute || t == kStageAttribute || t == kStrideAttribute ||
-         t == kWorkgroupSizeAttribute;
-}
-
 // https://gpuweb.github.io/gpuweb/wgsl.html#reserved-keywords
 bool is_reserved(Token t) {
   return t == "asm" || t == "bf16" || t == "const" || t == "do" ||
@@ -144,7 +136,6 @@
 /// Used by sync_to() to skip over closing block tokens that were opened during
 /// the forward scan.
 struct BlockCounters {
-  int attrs = 0;    // [[ ]]
   int brace = 0;    // {   }
   int bracket = 0;  // [   ]
   int paren = 0;    // (   )
@@ -152,10 +143,6 @@
   /// @return the current enter-exit depth for the given block token type. If
   /// `t` is not a block token type, then 0 is always returned.
   int consume(const Token& t) {
-    if (t.Is(Token::Type::kAttrLeft))  // [DEPRECATED]
-      return attrs++;
-    if (t.Is(Token::Type::kAttrRight))  // [DEPRECATED]
-      return attrs--;
     if (t.Is(Token::Type::kBraceLeft))
       return brace++;
     if (t.Is(Token::Type::kBraceRight))
@@ -1374,8 +1361,7 @@
   while (continue_parsing()) {
     // Check for the end of the list.
     auto t = peek();
-    if (!t.IsIdentifier() && !t.Is(Token::Type::kAttr) &&
-        !t.Is(Token::Type::kAttrLeft)) {
+    if (!t.IsIdentifier() && !t.Is(Token::Type::kAttr)) {
       break;
     }
 
@@ -2819,7 +2805,6 @@
 
 Maybe<ast::AttributeList> ParserImpl::attribute_list() {
   bool errored = false;
-  bool matched = false;
   ast::AttributeList attrs;
 
   while (continue_parsing()) {
@@ -2829,80 +2814,20 @@
       } else {
         attrs.emplace_back(attr.value);
       }
-    } else {  // [DEPRECATED] - old [[attribute]] style
-      auto list = attribute_bracketed_list(attrs);
-      if (list.errored) {
-        errored = true;
-      }
-      if (!list.matched) {
-        break;
-      }
+    } else {
+      break;
     }
-
-    matched = true;
   }
 
   if (errored)
     return Failure::kErrored;
 
-  if (!matched)
+  if (attrs.empty())
     return Failure::kNoMatch;
 
   return attrs;
 }
 
-Maybe<bool> ParserImpl::attribute_bracketed_list(ast::AttributeList& attrs) {
-  const char* use = "attribute list";
-
-  Source source;
-  if (!match(Token::Type::kAttrLeft, &source)) {
-    return Failure::kNoMatch;
-  }
-
-  deprecated(source,
-             "[[attribute]] style attributes have been replaced with "
-             "@attribute style");
-
-  if (match(Token::Type::kAttrRight, &source))
-    return add_error(source, "empty attribute list");
-
-  return sync(Token::Type::kAttrRight, [&]() -> Expect<bool> {
-    bool errored = false;
-
-    while (continue_parsing()) {
-      auto attr = expect_attribute();
-      if (attr.errored) {
-        errored = true;
-      }
-      attrs.emplace_back(attr.value);
-
-      if (match(Token::Type::kComma)) {
-        continue;
-      }
-
-      if (is_attribute(peek())) {
-        // We have two attributes in a bracket without a separating comma.
-        // e.g. @location(1) group(2)
-        //                    ^^^ expected comma
-        expect(use, Token::Type::kComma);
-        return Failure::kErrored;
-      }
-
-      break;
-    }
-
-    if (errored) {
-      return Failure::kErrored;
-    }
-
-    if (!expect(use, Token::Type::kAttrRight)) {
-      return Failure::kErrored;
-    }
-
-    return true;
-  });
-}
-
 Expect<const ast::Attribute*> ParserImpl::expect_attribute() {
   auto t = peek();
   auto attr = attribute();
@@ -3146,17 +3071,6 @@
     return true;
   }
 
-  // Handle the case when `]` is expected but the actual token is `]]`.
-  // For example, in `arr1[arr2[0]]`.
-  if (tok == Token::Type::kBracketRight && t.Is(Token::Type::kAttrRight)) {
-    next();
-    auto source = t.source();
-    source.range.begin.column++;
-    token_queue_.push_front({Token::Type::kBracketRight, source});
-    synchronized_ = true;
-    return true;
-  }
-
   // Error cases
   synchronized_ = false;
   if (handle_error(t)) {
diff --git a/src/tint/reader/wgsl/parser_impl.h b/src/tint/reader/wgsl/parser_impl.h
index 8c99f0a..50344e7 100644
--- a/src/tint/reader/wgsl/parser_impl.h
+++ b/src/tint/reader/wgsl/parser_impl.h
@@ -671,11 +671,6 @@
   /// Parses one or more attribute lists.
   /// @return the parsed attribute list, or an empty list on error.
   Maybe<ast::AttributeList> attribute_list();
-  /// Parses a list of attributes between `ATTR_LEFT` and `ATTR_RIGHT`
-  /// brackets.
-  /// @param attrs the list to append newly parsed attributes to.
-  /// @return true if any attributes were be parsed, otherwise false.
-  Maybe<bool> attribute_bracketed_list(ast::AttributeList& attrs);
   /// Parses a single attribute of the following types:
   /// * `struct_attribute`
   /// * `struct_member_attribute`
diff --git a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
index f8c6136..7c21634 100644
--- a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
@@ -61,20 +61,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_AliasDeclInvalidAttribute) {
-  EXPECT(
-      "[[invariant]]type e=u32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[invariant]]type e=u32;
-^^
-
-test.wgsl:1:3 error: unexpected attributes
-[[invariant]]type e=u32;
-  ^^^^^^^^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, IndexExprInvalidExpr) {
   EXPECT("fn f() { x = y[^]; }",
          R"(test.wgsl:1:16 error: unable to parse expression inside []
@@ -338,34 +324,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_FunctionDeclStageMissingLParen) {
-  EXPECT(
-      "[[stage vertex]] fn f() {}",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[stage vertex]] fn f() {}
-^^
-
-test.wgsl:1:9 error: expected '(' for stage attribute
-[[stage vertex]] fn f() {}
-        ^^^^^^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_FunctionDeclStageMissingRParen) {
-  EXPECT(
-      "[[stage(vertex]] fn f() {}",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[stage(vertex]] fn f() {}
-^^
-
-test.wgsl:1:15 error: expected ')' for stage attribute
-[[stage(vertex]] fn f() {}
-              ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, FunctionDeclStageInvalid) {
   EXPECT("@stage(x) fn f() {}",
          R"(test.wgsl:1:8 error: invalid value for stage attribute
@@ -386,34 +344,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_FunctionDeclWorkgroupSizeMissingLParen) {
-  EXPECT(
-      "[[workgroup_size 1]] fn f() {}",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[workgroup_size 1]] fn f() {}
-^^
-
-test.wgsl:1:18 error: expected '(' for workgroup_size attribute
-[[workgroup_size 1]] fn f() {}
-                 ^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_FunctionDeclWorkgroupSizeMissingRParen) {
-  EXPECT(
-      "[[workgroup_size(1]] fn f() {}",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[workgroup_size(1]] fn f() {}
-^^
-
-test.wgsl:1:19 error: expected ')' for workgroup_size attribute
-[[workgroup_size(1]] fn f() {}
-                  ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, FunctionDeclWorkgroupSizeXInvalid) {
   EXPECT("@workgroup_size() fn f() {}",
          R"(test.wgsl:1:17 error: expected workgroup_size x parameter
@@ -749,34 +679,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclStructMemberAttrEmpty) {
-  EXPECT(
-      "struct S { [[]] i : i32; };",
-      R"(test.wgsl:1:12 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-struct S { [[]] i : i32; };
-           ^^
-
-test.wgsl:1:14 error: empty attribute list
-struct S { [[]] i : i32; };
-             ^^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclStructMemberAttrMissingEnd) {
-  EXPECT(
-      "struct S { [[ i : i32; };",
-      R"(test.wgsl:1:12 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-struct S { [[ i : i32; };
-           ^^
-
-test.wgsl:1:15 error: expected attribute
-struct S { [[ i : i32; };
-              ^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclStructMemberInvalidIdentifier) {
   EXPECT("struct S { 1 : i32; };",
          R"(test.wgsl:1:12 error: expected identifier for struct member
@@ -793,36 +695,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest,
-       DEPRECATED_GlobalDeclStructMemberAlignMissingLParen) {
-  EXPECT(
-      "struct S { [[align 1)]] i : i32; };",
-      R"(test.wgsl:1:12 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-struct S { [[align 1)]] i : i32; };
-           ^^
-
-test.wgsl:1:20 error: expected '(' for align attribute
-struct S { [[align 1)]] i : i32; };
-                   ^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest,
-       DEPRECATED_GlobalDeclStructMemberAlignMissingRParen) {
-  EXPECT(
-      "struct S { [[align(1]] i : i32; };",
-      R"(test.wgsl:1:12 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-struct S { [[align(1]] i : i32; };
-           ^^
-
-test.wgsl:1:21 error: expected ')' for align attribute
-struct S { [[align(1]] i : i32; };
-                    ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclStructMemberAlignInvaldValue) {
   EXPECT(
       "struct S { @align(x) i : i32; };",
@@ -840,36 +712,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest,
-       DEPRECATED_GlobalDeclStructMemberSizeMissingLParen) {
-  EXPECT(
-      "struct S { [[size 1)]] i : i32; };",
-      R"(test.wgsl:1:12 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-struct S { [[size 1)]] i : i32; };
-           ^^
-
-test.wgsl:1:19 error: expected '(' for size attribute
-struct S { [[size 1)]] i : i32; };
-                  ^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest,
-       DEPRECATED_GlobalDeclStructMemberSizeMissingRParen) {
-  EXPECT(
-      "struct S { [[size(1]] i : i32; };",
-      R"(test.wgsl:1:12 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-struct S { [[size(1]] i : i32; };
-           ^^
-
-test.wgsl:1:20 error: expected ')' for size attribute
-struct S { [[size(1]] i : i32; };
-                   ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclStructMemberSizeInvaldValue) {
   EXPECT(
       "struct S { @size(x) i : i32; };",
@@ -916,20 +758,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclTypeAttrInvalid) {
-  EXPECT(
-      "var x : [[]] i32;",
-      R"(test.wgsl:1:9 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-var x : [[]] i32;
-        ^^
-
-test.wgsl:1:11 error: empty attribute list
-var x : [[]] i32;
-          ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclVarArrayMissingLessThan) {
   EXPECT("var i : array;",
          R"(test.wgsl:1:14 error: expected '<' for array declaration
@@ -954,48 +782,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarArrayAttrMissingEnd) {
-  EXPECT(
-      "var i : [[location(1) array<i32>;",
-      R"(test.wgsl:1:9 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-var i : [[location(1) array<i32>;
-        ^^
-
-test.wgsl:1:23 error: expected ']]' for attribute list
-var i : [[location(1) array<i32>;
-                      ^^^^^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarArrayStrideMissingLParen) {
-  EXPECT(
-      "var i : [[stride 1)]] array<i32>;",
-      R"(test.wgsl:1:9 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-var i : [[stride 1)]] array<i32>;
-        ^^
-
-test.wgsl:1:18 error: expected '(' for stride attribute
-var i : [[stride 1)]] array<i32>;
-                 ^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarArrayStrideMissingRParen) {
-  EXPECT(
-      "var i : [[location(1]] array<i32>;",
-      R"(test.wgsl:1:9 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-var i : [[location(1]] array<i32>;
-        ^^
-
-test.wgsl:1:21 error: expected ')' for location attribute
-var i : [[location(1]] array<i32>;
-                    ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclVarArrayStrideInvalid) {
   EXPECT(
       "var i : @stride(x) array<i32>;",
@@ -1037,34 +823,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrListEmpty) {
-  EXPECT(
-      "[[]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[]] var i : i32;
-^^
-
-test.wgsl:1:3 error: empty attribute list
-[[]] var i : i32;
-  ^^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrListInvalid) {
-  EXPECT(
-      "[[location(1), meow]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[location(1), meow]] var i : i32;
-^^
-
-test.wgsl:1:16 error: expected attribute
-[[location(1), meow]] var i : i32;
-               ^^^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclVarAttrListMissingComma) {
   EXPECT("@location(1) group(2) var i : i32;",
          R"(test.wgsl:1:14 error: expected declaration after attributes
@@ -1077,20 +835,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrListMissingEnd) {
-  EXPECT(
-      "[[location(1) meow]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[location(1) meow]] var i : i32;
-^^
-
-test.wgsl:1:15 error: expected ']]' for attribute list
-[[location(1) meow]] var i : i32;
-              ^^^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclVarAttrLocationMissingLParen) {
   EXPECT("@location 1) var i : i32;",
          R"(test.wgsl:1:11 error: expected '(' for location attribute
@@ -1107,34 +851,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrLocationMissingLParen) {
-  EXPECT(
-      "[[location 1]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[location 1]] var i : i32;
-^^
-
-test.wgsl:1:12 error: expected '(' for location attribute
-[[location 1]] var i : i32;
-           ^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrLocationMissingRParen) {
-  EXPECT(
-      "[[location (1]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[location (1]] var i : i32;
-^^
-
-test.wgsl:1:14 error: expected ')' for location attribute
-[[location (1]] var i : i32;
-             ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclVarAttrLocationInvalidValue) {
   EXPECT(
       "@location(x) var i : i32;",
@@ -1160,34 +876,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrBuiltinMissingLParen) {
-  EXPECT(
-      "[[builtin position]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[builtin position]] var i : i32;
-^^
-
-test.wgsl:1:11 error: expected '(' for builtin attribute
-[[builtin position]] var i : i32;
-          ^^^^^^^^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrBuiltinMissingRParen) {
-  EXPECT(
-      "[[builtin(position]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[builtin(position]] var i : i32;
-^^
-
-test.wgsl:1:19 error: expected ')' for builtin attribute
-[[builtin(position]] var i : i32;
-                  ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclVarAttrBuiltinInvalidIdentifer) {
   EXPECT("@builtin(1) var i : i32;",
          R"(test.wgsl:1:10 error: expected identifier for builtin
@@ -1220,34 +908,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrBindingMissingLParen) {
-  EXPECT(
-      "[[binding 1]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[binding 1]] var i : i32;
-^^
-
-test.wgsl:1:11 error: expected '(' for binding attribute
-[[binding 1]] var i : i32;
-          ^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrBindingMissingRParen) {
-  EXPECT(
-      "[[binding(1]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[binding(1]] var i : i32;
-^^
-
-test.wgsl:1:12 error: expected ')' for binding attribute
-[[binding(1]] var i : i32;
-           ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclVarAttrBindingInvalidValue) {
   EXPECT(
       "@binding(x) var i : i32;",
@@ -1273,34 +933,6 @@
 )");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrGroupMissingLParen) {
-  EXPECT(
-      "[[group 1]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[group 1]] var i : i32;
-^^
-
-test.wgsl:1:9 error: expected '(' for group attribute
-[[group 1]] var i : i32;
-        ^
-)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarAttrGroupMissingRParen) {
-  EXPECT(
-      "[[group(1]] var i : i32;",
-      R"(test.wgsl:1:1 warning: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-[[group(1]] var i : i32;
-^^
-
-test.wgsl:1:10 error: expected ')' for group attribute
-[[group(1]] var i : i32;
-         ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclVarAttrBindingGroupValue) {
   EXPECT(
       "@group(x) var i : i32;",
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 5ca1dac..1d275dc 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
@@ -56,18 +56,6 @@
   EXPECT_EQ(p->error(), "1:2: expected attribute");
 }
 
-TEST_F(ParserImplTest, AttributeList_ExtraComma) {
-  auto p = parser("[[workgroup_size(2), ]]");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:22: expected attribute)");
-}
-
 TEST_F(ParserImplTest, AttributeList_BadAttribute) {
   auto p = parser("@stage()");
   auto attrs = p->attribute_list();
@@ -77,85 +65,6 @@
   EXPECT_EQ(p->error(), "1:8: invalid value for stage attribute");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_Empty) {
-  auto p = parser("[[]]");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:3: empty attribute list)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_Invalid) {
-  auto p = parser("[[invalid]]");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_TRUE(attrs.value.empty());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:3: expected attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_ExtraComma) {
-  auto p = parser("[[workgroup_size(2), ]]");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:22: expected attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_MissingComma) {
-  auto p = parser("[[workgroup_size(2) workgroup_size(2)]]");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:21: expected ',' for attribute list)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_BadAttribute) {
-  auto p = parser("[[stage()]]");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:9: invalid value for stage attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_MissingRightAttr) {
-  auto p = parser("[[workgroup_size(2), workgroup_size(3, 4, 5)");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:45: expected ']]' for attribute list)");
-}
-
 }  // namespace
 }  // namespace wgsl
 }  // namespace reader
diff --git a/src/tint/reader/wgsl/parser_impl_function_header_test.cc b/src/tint/reader/wgsl/parser_impl_function_header_test.cc
index 6416aca..45e4d5d 100644
--- a/src/tint/reader/wgsl/parser_impl_function_header_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_header_test.cc
@@ -62,7 +62,7 @@
 }
 
 TEST_F(ParserImplTest, FunctionHeader_InvariantReturnType) {
-  auto p = parser("fn main() -> [[invariant]] f32");
+  auto p = parser("fn main() -> @invariant f32");
   auto f = p->function_header();
   ASSERT_FALSE(p->has_error()) << p->error();
   EXPECT_TRUE(f.matched);
diff --git a/src/tint/reader/wgsl/parser_impl_struct_attribute_decl_test.cc b/src/tint/reader/wgsl/parser_impl_struct_attribute_decl_test.cc
index d7ad8b1..98adb42 100644
--- a/src/tint/reader/wgsl/parser_impl_struct_attribute_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_struct_attribute_decl_test.cc
@@ -71,42 +71,6 @@
   EXPECT_TRUE(attrs.value.empty());
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_attributeDecl_Parses) {
-  auto p = parser("[[invariant]]");
-  auto attrs = p->attribute_list();
-  EXPECT_FALSE(p->has_error());
-  EXPECT_FALSE(attrs.errored);
-  EXPECT_TRUE(attrs.matched);
-  ASSERT_EQ(attrs.value.size(), 1u);
-  auto* invariant_attr = attrs.value[0]->As<ast::Attribute>();
-  EXPECT_TRUE(invariant_attr->Is<ast::InvariantAttribute>());
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_attributeDecl_MissingAttrRight) {
-  auto p = parser("[[invariant");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_TRUE(attrs.value.empty());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:12: expected ']]' for attribute list)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_attributeDecl_Invalidattribute) {
-  auto p = parser("[[invalid]]");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_TRUE(attrs.value.empty());
-}
-
 }  // namespace
 }  // namespace wgsl
 }  // namespace reader
diff --git a/src/tint/reader/wgsl/parser_impl_struct_member_attribute_decl_test.cc b/src/tint/reader/wgsl/parser_impl_struct_member_attribute_decl_test.cc
index 0c2140f..50e1cb4 100644
--- a/src/tint/reader/wgsl/parser_impl_struct_member_attribute_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_struct_member_attribute_decl_test.cc
@@ -28,20 +28,6 @@
   EXPECT_EQ(attrs.value.size(), 0u);
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeDecl_EmptyBlock) {
-  auto p = parser("[[]]");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_EQ(attrs.value.size(), 0u);
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:3: empty attribute list)");
-}
-
 TEST_F(ParserImplTest, AttributeDecl_Single) {
   auto p = parser("@size(4)");
   auto attrs = p->attribute_list();
@@ -64,30 +50,6 @@
             "1:7: expected signed integer literal for size attribute");
 }
 
-TEST_F(ParserImplTest, AttributeDecl_MissingClose) {
-  auto p = parser("[[size(4)");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error()) << p->error();
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:10: expected ']]' for attribute list)");
-}
-
-TEST_F(ParserImplTest, StructMemberAttributeDecl_SizeMissingClose) {
-  auto p = parser("[[size(4)");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error()) << p->error();
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:10: expected ']]' for attribute list)");
-}
-
 }  // namespace
 }  // namespace wgsl
 }  // namespace reader
diff --git a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
index dbb408f..d74ddf6 100644
--- a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
@@ -605,92 +605,6 @@
   EXPECT_EQ(p->error(), "1:9: stride attribute must be greater than 0");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_TypeDecl_Array_Attribute_MissingClosingAttr) {
-  auto p = parser("[[stride(16) array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:3: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
-1:14: expected ']]' for attribute list)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_TypeDecl_Array_Stride_MissingLeftParen) {
-  auto p = parser("[[stride 4)]] array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:10: expected '(' for stride attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_TypeDecl_Array_Stride_MissingRightParen) {
-  auto p = parser("[[stride(4]] array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:3: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
-1:11: expected ')' for stride attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_TypeDecl_Array_Stride_MissingValue) {
-  auto p = parser("[[stride()]] array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:10: expected signed integer literal for stride attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_TypeDecl_Array_Stride_InvalidValue) {
-  auto p = parser("[[stride(invalid)]] array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:10: expected signed integer literal for stride attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_TypeDecl_Array_Stride_InvalidValue_Negative) {
-  auto p = parser("[[stride(-1)]] array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:10: stride attribute must be greater than 0)");
-}
-
 TEST_F(ParserImplTest, TypeDecl_Array_Runtime) {
   auto p = parser("array<u32>");
   auto t = p->type_decl();
diff --git a/src/tint/reader/wgsl/parser_impl_variable_attribute_list_test.cc b/src/tint/reader/wgsl/parser_impl_variable_attribute_list_test.cc
index b751f87..253451c 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_attribute_list_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_attribute_list_test.cc
@@ -59,76 +59,6 @@
   EXPECT_EQ(p->error(), "1:10: invalid value for builtin attribute");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_Empty) {
-  auto p = parser(R"([[]])");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_TRUE(attrs.value.empty());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:3: empty attribute list)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_Invalid) {
-  auto p = parser(R"([[invalid]])");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_TRUE(attrs.value.empty());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:3: expected attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_ExtraComma) {
-  auto p = parser(R"([[builtin(position), ]])");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_TRUE(attrs.value.empty());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:22: expected attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_MissingComma) {
-  auto p = parser(R"([[binding(4) location(5)]])");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_TRUE(attrs.value.empty());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:14: expected ',' for attribute list)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_AttributeList_InvalidValue) {
-  auto p = parser("[[builtin(invalid)]]");
-  auto attrs = p->attribute_list();
-  EXPECT_TRUE(p->has_error());
-  EXPECT_TRUE(attrs.errored);
-  EXPECT_FALSE(attrs.matched);
-  EXPECT_TRUE(attrs.value.empty());
-  EXPECT_EQ(
-      p->error(),
-      R"(1:1: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:11: invalid value for builtin attribute)");
-}
-
 }  // namespace
 }  // namespace wgsl
 }  // namespace reader
diff --git a/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc b/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
index f41ec2d..976284e 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
@@ -104,56 +104,6 @@
   ASSERT_EQ(p->error(), "1:18: expected '(' for stride attribute");
 }
 
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest,
-       DEPRECATED_VariableIdentDecl_AttributeMissingRightBlock) {
-  auto p = parser("my_var : [[location(4) S");
-  auto decl = p->expect_variable_ident_decl("test");
-  ASSERT_TRUE(p->has_error());
-  ASSERT_TRUE(decl.errored);
-  ASSERT_EQ(
-      p->error(),
-      R"(1:10: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:24: expected ']]' for attribute list)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest,
-       DEPRECATED_VariableIdentDecl_AttributeMissingRightParen) {
-  auto p = parser("my_var : [[location(4]] S");
-  auto decl = p->expect_variable_ident_decl("test");
-  ASSERT_TRUE(p->has_error());
-  ASSERT_TRUE(decl.errored);
-  ASSERT_EQ(
-      p->error(),
-      R"(1:10: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:22: expected ')' for location attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_VariableIdentDecl_AttributeMissingLeftParen) {
-  auto p = parser("my_var : [[stride 4)]] S");
-  auto decl = p->expect_variable_ident_decl("test");
-  ASSERT_TRUE(p->has_error());
-  ASSERT_TRUE(decl.errored);
-  ASSERT_EQ(
-      p->error(),
-      R"(1:10: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:19: expected '(' for stride attribute)");
-}
-
-// TODO(crbug.com/tint/1382): Remove
-TEST_F(ParserImplTest, DEPRECATED_VariableIdentDecl_AttributeEmpty) {
-  auto p = parser("my_var : [[]] S");
-  auto decl = p->expect_variable_ident_decl("test");
-  ASSERT_TRUE(p->has_error());
-  ASSERT_TRUE(decl.errored);
-  ASSERT_EQ(
-      p->error(),
-      R"(1:10: use of deprecated language feature: [[attribute]] style attributes have been replaced with @attribute style
-1:12: empty attribute list)");
-}
-
 }  // namespace
 }  // namespace wgsl
 }  // namespace reader
diff --git a/src/tint/reader/wgsl/token.cc b/src/tint/reader/wgsl/token.cc
index abc0dc1..266907e 100644
--- a/src/tint/reader/wgsl/token.cc
+++ b/src/tint/reader/wgsl/token.cc
@@ -44,10 +44,6 @@
       return "->";
     case Token::Type::kAttr:
       return "@";
-    case Token::Type::kAttrLeft:
-      return "[[";
-    case Token::Type::kAttrRight:
-      return "]]";
     case Token::Type::kForwardSlash:
       return "/";
     case Token::Type::kBang:
diff --git a/src/tint/reader/wgsl/token.h b/src/tint/reader/wgsl/token.h
index 58bb92a..dcc0167 100644
--- a/src/tint/reader/wgsl/token.h
+++ b/src/tint/reader/wgsl/token.h
@@ -54,10 +54,6 @@
     kArrow,
     /// A '@'
     kAttr,
-    /// A '[[' - [DEPRECATED] now '@'
-    kAttrLeft,
-    /// A ']]' - [DEPRECATED] now '@'
-    kAttrRight,
     /// A '/'
     kForwardSlash,
     /// A '!'
diff --git a/src/tint/resolver/attribute_validation_test.cc b/src/tint/resolver/attribute_validation_test.cc
index 71224c5..1ad15f9 100644
--- a/src/tint/resolver/attribute_validation_test.cc
+++ b/src/tint/resolver/attribute_validation_test.cc
@@ -271,7 +271,7 @@
                     TestParams{AttributeKind::kBuiltin, true},
                     TestParams{AttributeKind::kGroup, false},
                     TestParams{AttributeKind::kId, false},
-                    // kInterpolate tested separately (requires [[location]])
+                    // kInterpolate tested separately (requires @location)
                     TestParams{AttributeKind::kInvariant, true},
                     TestParams{AttributeKind::kLocation, true},
                     TestParams{AttributeKind::kOffset, false},
@@ -461,7 +461,7 @@
                     TestParams{AttributeKind::kBuiltin, true},
                     TestParams{AttributeKind::kGroup, false},
                     TestParams{AttributeKind::kId, false},
-                    // kInterpolate tested separately (requires [[location]])
+                    // kInterpolate tested separately (requires @location)
                     TestParams{AttributeKind::kInvariant, true},
                     TestParams{AttributeKind::kLocation, false},
                     TestParams{AttributeKind::kOffset, false},
@@ -594,7 +594,7 @@
                     TestParams{AttributeKind::kBuiltin, true},
                     TestParams{AttributeKind::kGroup, false},
                     TestParams{AttributeKind::kId, false},
-                    // kInterpolate tested separately (requires [[location]])
+                    // kInterpolate tested separately (requires @location)
                     // kInvariant tested separately (requires position builtin)
                     TestParams{AttributeKind::kLocation, true},
                     TestParams{AttributeKind::kOffset, true},
diff --git a/src/tint/resolver/function_validation_test.cc b/src/tint/resolver/function_validation_test.cc
index 4efd00b..cb33381 100644
--- a/src/tint/resolver/function_validation_test.cc
+++ b/src/tint/resolver/function_validation_test.cc
@@ -463,7 +463,7 @@
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_U32) {
-  // [[stage(compute), workgroup_size(1u, 2u, 3u)]
+  // @stage(compute) @workgroup_size(1u, 2u, 3u)
   // fn main() {}
 
   Func("main", {}, ty.void_(), {},
@@ -474,7 +474,7 @@
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_MismatchTypeU32) {
-  // [[stage(compute), workgroup_size(1u, 2u, 3)]
+  // @stage(compute) @workgroup_size(1u, 2u, 3)
   // fn main() {}
 
   Func("main", {}, ty.void_(), {},
@@ -488,7 +488,7 @@
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_MismatchTypeI32) {
-  // [[stage(compute), workgroup_size(1, 2u, 3)]
+  // @stage(compute) @workgroup_size(1, 2u, 3)
   // fn main() {}
 
   Func("main", {}, ty.void_(), {},
@@ -503,7 +503,7 @@
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_TypeMismatch) {
   // let x = 64u;
-  // [[stage(compute), workgroup_size(1, x)]
+  // @stage(compute) @workgroup_size(1, x)
   // fn main() {}
   GlobalConst("x", ty.u32(), Expr(64u));
   Func("main", {}, ty.void_(), {},
@@ -519,7 +519,7 @@
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_TypeMismatch2) {
   // let x = 64u;
   // let y = 32;
-  // [[stage(compute), workgroup_size(x, y)]
+  // @stage(compute) @workgroup_size(x, y)
   // fn main() {}
   GlobalConst("x", ty.u32(), Expr(64u));
   GlobalConst("y", ty.i32(), Expr(32));
@@ -535,7 +535,7 @@
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Mismatch_ConstU32) {
   // let x = 4u;
   // let x = 8u;
-  // [[stage(compute), workgroup_size(x, y, 16]
+  // @stage(compute) @workgroup_size(x, y, 16
   // fn main() {}
   GlobalConst("x", ty.u32(), Expr(4u));
   GlobalConst("y", ty.u32(), Expr(8u));
@@ -550,7 +550,7 @@
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_BadType) {
-  // [[stage(compute), workgroup_size(64.0)]
+  // @stage(compute) @workgroup_size(64.0)
   // fn main() {}
 
   Func("main", {}, ty.void_(), {},
@@ -564,7 +564,7 @@
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_Negative) {
-  // [[stage(compute), workgroup_size(-2)]
+  // @stage(compute) @workgroup_size(-2)
   // fn main() {}
 
   Func("main", {}, ty.void_(), {},
@@ -577,7 +577,7 @@
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_Zero) {
-  // [[stage(compute), workgroup_size(0)]
+  // @stage(compute) @workgroup_size(0)
   // fn main() {}
 
   Func("main", {}, ty.void_(), {},
@@ -591,7 +591,7 @@
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_BadType) {
   // let x = 64.0;
-  // [[stage(compute), workgroup_size(x)]
+  // @stage(compute) @workgroup_size(x)
   // fn main() {}
   GlobalConst("x", ty.f32(), Expr(64.f));
   Func("main", {}, ty.void_(), {},
@@ -606,7 +606,7 @@
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Negative) {
   // let x = -2;
-  // [[stage(compute), workgroup_size(x)]
+  // @stage(compute) @workgroup_size(x)
   // fn main() {}
   GlobalConst("x", ty.i32(), Expr(-2));
   Func("main", {}, ty.void_(), {},
@@ -620,7 +620,7 @@
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Zero) {
   // let x = 0;
-  // [[stage(compute), workgroup_size(x)]
+  // @stage(compute) @workgroup_size(x)
   // fn main() {}
   GlobalConst("x", ty.i32(), Expr(0));
   Func("main", {}, ty.void_(), {},
@@ -635,7 +635,7 @@
 TEST_F(ResolverFunctionValidationTest,
        WorkgroupSize_Const_NestedZeroValueConstructor) {
   // let x = i32(i32(i32()));
-  // [[stage(compute), workgroup_size(x)]
+  // @stage(compute) @workgroup_size(x)
   // fn main() {}
   GlobalConst("x", ty.i32(),
               Construct(ty.i32(), Construct(ty.i32(), Construct(ty.i32()))));
@@ -650,7 +650,7 @@
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_NonConst) {
   // var<private> x = 0;
-  // [[stage(compute), workgroup_size(x)]
+  // @stage(compute) @workgroup_size(x)
   // fn main() {}
   Global("x", ty.i32(), ast::StorageClass::kPrivate, Expr(64));
   Func("main", {}, ty.void_(), {},
@@ -664,7 +664,7 @@
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_InvalidExpr) {
-  // [[stage(compute), workgroup_size(i32(1))]
+  // @stage(compute) @workgroup_size(i32(1))
   // fn main() {}
   Func("main", {}, ty.void_(), {},
        {Stage(ast::PipelineStage::kCompute),
diff --git a/src/tint/sem/binding_point.h b/src/tint/sem/binding_point.h
index ed968e2..3f6d307 100644
--- a/src/tint/sem/binding_point.h
+++ b/src/tint/sem/binding_point.h
@@ -26,9 +26,9 @@
 
 /// BindingPoint holds a group and binding index.
 struct BindingPoint {
-  /// The `[[group]]` part of the binding point
+  /// The `@group` part of the binding point
   uint32_t group = 0;
-  /// The `[[binding]]` part of the binding point
+  /// The `@binding` part of the binding point
   uint32_t binding = 0;
 
   /// Equality operator
diff --git a/src/tint/transform/add_empty_entry_point_test.cc b/src/tint/transform/add_empty_entry_point_test.cc
index b5ab815..e8c1c72 100644
--- a/src/tint/transform/add_empty_entry_point_test.cc
+++ b/src/tint/transform/add_empty_entry_point_test.cc
@@ -32,7 +32,7 @@
 
 TEST_F(AddEmptyEntryPointTest, ShouldRunExistingEntryPoint) {
   auto* src = R"(
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn existing() {}
 )";
 
diff --git a/src/tint/transform/array_length_from_uniform_test.cc b/src/tint/transform/array_length_from_uniform_test.cc
index 9170d65..ffb63af 100644
--- a/src/tint/transform/array_length_from_uniform_test.cc
+++ b/src/tint/transform/array_length_from_uniform_test.cc
@@ -39,9 +39,9 @@
   arr : array<i32>;
 };
 
-[[group(0), binding(0)]] var<storage, read> sb : SB;
+@group(0) @binding(0) var<storage, read> sb : SB;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
 }
 )";
@@ -56,9 +56,9 @@
   arr : array<i32>;
 };
 
-[[group(0), binding(0)]] var<storage, read> sb : SB;
+@group(0) @binding(0) var<storage, read> sb : SB;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   var len : u32 = arrayLength(&sb.arr);
 }
@@ -74,9 +74,9 @@
   arr : array<i32>;
 };
 
-[[group(0), binding(0)]] var<storage, read> sb : SB;
+@group(0) @binding(0) var<storage, read> sb : SB;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   var len : u32 = arrayLength(&sb.arr);
 }
diff --git a/src/tint/transform/calculate_array_length_test.cc b/src/tint/transform/calculate_array_length_test.cc
index 9bd7eb0..64589a0 100644
--- a/src/tint/transform/calculate_array_length_test.cc
+++ b/src/tint/transform/calculate_array_length_test.cc
@@ -37,9 +37,9 @@
   arr : array<i32>;
 };
 
-[[group(0), binding(0)]] var<storage, read> sb : SB;
+@group(0) @binding(0) var<storage, read> sb : SB;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
 }
 )";
@@ -54,9 +54,9 @@
   arr : array<i32>;
 };
 
-[[group(0), binding(0)]] var<storage, read> sb : SB;
+@group(0) @binding(0) var<storage, read> sb : SB;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   var len : u32 = arrayLength(&sb.arr);
 }
diff --git a/src/tint/transform/canonicalize_entry_point_io_test.cc b/src/tint/transform/canonicalize_entry_point_io_test.cc
index b4bcfb8..8866f13 100644
--- a/src/tint/transform/canonicalize_entry_point_io_test.cc
+++ b/src/tint/transform/canonicalize_entry_point_io_test.cc
@@ -2219,7 +2219,7 @@
 TEST_F(CanonicalizeEntryPointIOTest, InvariantAttributes) {
   auto* src = R"(
 struct VertexOut {
-  [[builtin(position), invariant]] pos : vec4<f32>;
+  @builtin(position) @invariant pos : vec4<f32>;
 };
 
 @stage(vertex)
@@ -2228,7 +2228,7 @@
 }
 
 @stage(vertex)
-fn main2() -> [[builtin(position), invariant]] vec4<f32> {
+fn main2() -> @builtin(position) @invariant vec4<f32> {
   return vec4<f32>();
 }
 )";
@@ -2289,12 +2289,12 @@
 }
 
 @stage(vertex)
-fn main2() -> [[builtin(position), invariant]] vec4<f32> {
+fn main2() -> @builtin(position) @invariant vec4<f32> {
   return vec4<f32>();
 }
 
 struct VertexOut {
-  [[builtin(position), invariant]] pos : vec4<f32>;
+  @builtin(position) @invariant pos : vec4<f32>;
 };
 )";
 
diff --git a/src/tint/transform/decompose_memory_access_test.cc b/src/tint/transform/decompose_memory_access_test.cc
index 91476f3..8ca3c2b 100644
--- a/src/tint/transform/decompose_memory_access_test.cc
+++ b/src/tint/transform/decompose_memory_access_test.cc
@@ -33,7 +33,7 @@
 struct Buffer {
   i : i32;
 };
-[[group(0), binding(0)]] var<storage, read_write> sb : Buffer;
+@group(0) @binding(0) var<storage, read_write> sb : Buffer;
 )";
 
   EXPECT_TRUE(ShouldRun<DecomposeMemoryAccess>(src));
@@ -44,7 +44,7 @@
 struct Buffer {
   i : i32;
 };
-[[group(0), binding(0)]] var<uniform> ub : Buffer;
+@group(0) @binding(0) var<uniform> ub : Buffer;
 )";
 
   EXPECT_TRUE(ShouldRun<DecomposeMemoryAccess>(src));
diff --git a/src/tint/transform/decompose_strided_matrix.h b/src/tint/transform/decompose_strided_matrix.h
index f223178..9a4c7b7 100644
--- a/src/tint/transform/decompose_strided_matrix.h
+++ b/src/tint/transform/decompose_strided_matrix.h
@@ -21,7 +21,7 @@
 namespace transform {
 
 /// DecomposeStridedMatrix transforms replaces matrix members of storage or
-/// uniform buffer structures, that have a [[stride]] attribute, into an array
+/// uniform buffer structures, that have a stride attribute, into an array
 /// of N column vectors.
 /// This transform is used by the SPIR-V reader to handle the SPIR-V
 /// MatrixStride attribute.
diff --git a/src/tint/transform/first_index_offset_test.cc b/src/tint/transform/first_index_offset_test.cc
index d482114..9de9880 100644
--- a/src/tint/transform/first_index_offset_test.cc
+++ b/src/tint/transform/first_index_offset_test.cc
@@ -34,7 +34,7 @@
 
 TEST_F(FirstIndexOffsetTest, ShouldRunFragmentStage) {
   auto* src = R"(
-[[stage(fragment)]]
+@stage(fragment)
 fn entry() {
   return;
 }
@@ -45,8 +45,8 @@
 
 TEST_F(FirstIndexOffsetTest, ShouldRunVertexStage) {
   auto* src = R"(
-[[stage(vertex)]]
-fn entry() -> [[builtin(position)]] vec4<f32> {
+@stage(vertex)
+fn entry() -> @builtin(position) vec4<f32> {
   return vec4<f32>();
 }
 )";
diff --git a/src/tint/transform/multiplanar_external_texture_test.cc b/src/tint/transform/multiplanar_external_texture_test.cc
index b37ef68..0a322ef 100644
--- a/src/tint/transform/multiplanar_external_texture_test.cc
+++ b/src/tint/transform/multiplanar_external_texture_test.cc
@@ -36,7 +36,7 @@
 }
 TEST_F(MultiplanarExternalTextureTest, ShouldRunHasExternalTextureGlobal) {
   auto* src = R"(
-[[group(0), binding(0)]] var ext_tex : texture_external;
+@group(0) @binding(0) var ext_tex : texture_external;
 )";
 
   EXPECT_TRUE(ShouldRun<MultiplanarExternalTexture>(src));
@@ -1163,10 +1163,10 @@
   textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
 }
 
-[[group(0), binding(0)]] var ext_tex : ET;
-[[group(0), binding(1)]] var smp : sampler;
+@group(0) @binding(0) var ext_tex : ET;
+@group(0) @binding(1) var smp : sampler;
 
-[[stage(fragment)]]
+@stage(fragment)
 fn main() {
   f(ext_tex, smp);
 }
@@ -1226,7 +1226,7 @@
 // Tests that the the transform handles aliases to external textures
 TEST_F(MultiplanarExternalTextureTest, ExternalTextureAlias_OutOfOrder) {
   auto* src = R"(
-[[stage(fragment)]]
+@stage(fragment)
 fn main() {
   f(ext_tex, smp);
 }
@@ -1235,8 +1235,8 @@
   textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
 }
 
-[[group(0), binding(0)]] var ext_tex : ET;
-[[group(0), binding(1)]] var smp : sampler;
+@group(0) @binding(0) var ext_tex : ET;
+@group(0) @binding(1) var smp : sampler;
 
 type ET = texture_external;
 )";
diff --git a/src/tint/transform/num_workgroups_from_uniform_test.cc b/src/tint/transform/num_workgroups_from_uniform_test.cc
index 1674433..a728b53 100644
--- a/src/tint/transform/num_workgroups_from_uniform_test.cc
+++ b/src/tint/transform/num_workgroups_from_uniform_test.cc
@@ -34,8 +34,8 @@
 
 TEST_F(NumWorkgroupsFromUniformTest, ShouldRunHasNumWorkgroups) {
   auto* src = R"(
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(num_workgroups)]] num_wgs : vec3<u32>) {
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(num_workgroups) num_wgs : vec3<u32>) {
 }
 )";
 
@@ -44,8 +44,8 @@
 
 TEST_F(NumWorkgroupsFromUniformTest, Error_MissingTransformData) {
   auto* src = R"(
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(num_workgroups)]] num_wgs : vec3<u32>) {
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(num_workgroups) num_wgs : vec3<u32>) {
 }
 )";
 
diff --git a/src/tint/transform/pad_array_elements_test.cc b/src/tint/transform/pad_array_elements_test.cc
index d0e0b4b..f5e921b 100644
--- a/src/tint/transform/pad_array_elements_test.cc
+++ b/src/tint/transform/pad_array_elements_test.cc
@@ -40,7 +40,7 @@
 
 TEST_F(PadArrayElementsTest, ShouldRunHasExplicitArrayStride) {
   auto* src = R"(
-var<private> arr : [[stride(8)]] array<i32, 4>;
+var<private> arr : @stride(8) array<i32, 4>;
 )";
 
   EXPECT_TRUE(ShouldRun<PadArrayElements>(src));
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index 89164d1..fad7968 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -262,7 +262,7 @@
     // WGSL can ignore the invariant attribute on pre MSL 2.1 devices.
     // See: https://github.com/gpuweb/gpuweb/issues/893#issuecomment-745537465
     line(&helpers_) << "#if __METAL_VERSION__ >= 210";
-    line(&helpers_) << "#define " << invariant_define_name_ << " [[invariant]]";
+    line(&helpers_) << "#define " << invariant_define_name_ << " @invariant";
     line(&helpers_) << "#else";
     line(&helpers_) << "#define " << invariant_define_name_;
     line(&helpers_) << "#endif";
diff --git a/src/tint/writer/msl/generator_impl_test.cc b/src/tint/writer/msl/generator_impl_test.cc
index 01caf83..26085e3 100644
--- a/src/tint/writer/msl/generator_impl_test.cc
+++ b/src/tint/writer/msl/generator_impl_test.cc
@@ -97,7 +97,7 @@
 using namespace metal;
 
 #if __METAL_VERSION__ >= 210
-#define TINT_INVARIANT [[invariant]]
+#define TINT_INVARIANT @invariant
 #else
 #define TINT_INVARIANT
 #endif
diff --git a/test/tint/shader_io/invariant.wgsl.expected.msl b/test/tint/shader_io/invariant.wgsl.expected.msl
index 5982216..3d41ee7 100644
--- a/test/tint/shader_io/invariant.wgsl.expected.msl
+++ b/test/tint/shader_io/invariant.wgsl.expected.msl
@@ -3,7 +3,7 @@
 using namespace metal;
 
 #if __METAL_VERSION__ >= 210
-#define TINT_INVARIANT [[invariant]]
+#define TINT_INVARIANT @invariant
 #else
 #define TINT_INVARIANT
 #endif
diff --git a/test/tint/shader_io/invariant_struct_member.wgsl.expected.msl b/test/tint/shader_io/invariant_struct_member.wgsl.expected.msl
index 666b8dd..5b44709 100644
--- a/test/tint/shader_io/invariant_struct_member.wgsl.expected.msl
+++ b/test/tint/shader_io/invariant_struct_member.wgsl.expected.msl
@@ -3,7 +3,7 @@
 using namespace metal;
 
 #if __METAL_VERSION__ >= 210
-#define TINT_INVARIANT [[invariant]]
+#define TINT_INVARIANT @invariant
 #else
 #define TINT_INVARIANT
 #endif