diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json
index e615ee327a3c..68537bd691af 100644
--- a/src/webgpu/listing_meta.json
+++ b/src/webgpu/listing_meta.json
@@ -2475,20 +2475,13 @@
   "webgpu:shader,validation,parse,blankspace:blankspace:*": { "subcaseMS": 1.391 },
   "webgpu:shader,validation,parse,blankspace:bom:*": { "subcaseMS": 1.101 },
   "webgpu:shader,validation,parse,blankspace:null_characters:*": { "subcaseMS": 3.217 },
-  "webgpu:shader,validation,parse,break:placement:*": { "subcaseMS": 1.254 },
-  "webgpu:shader,validation,parse,break_if:placement:*": { "subcaseMS": 0.000 },
   "webgpu:shader,validation,parse,builtin:parse:*": { "subcaseMS": 3.277 },
   "webgpu:shader,validation,parse,builtin:placement:*": { "subcaseMS": 1.267 },
   "webgpu:shader,validation,parse,comments:comments:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,parse,comments:line_comment_eof:*": { "subcaseMS": 4.500 },
   "webgpu:shader,validation,parse,comments:line_comment_terminators:*": { "subcaseMS": 1.021 },
   "webgpu:shader,validation,parse,comments:unterminated_block_comment:*": { "subcaseMS": 8.950 },
-  "webgpu:shader,validation,parse,compound:parse:*": { "subcaseMS": 4.315 },
   "webgpu:shader,validation,parse,const:placement:*": { "subcaseMS": 1.167 },
-  "webgpu:shader,validation,parse,const_assert:parse:*": { "subcaseMS": 1.400 },
-  "webgpu:shader,validation,parse,continue:module_scope:*": { "subcaseMS": 0.278 },
-  "webgpu:shader,validation,parse,continue:placement:*": { "subcaseMS": 53.998 },
-  "webgpu:shader,validation,parse,continuing:placement:*": { "subcaseMS": 0.000 },
   "webgpu:shader,validation,parse,diagnostic:after_other_directives:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,parse,diagnostic:conflicting_attribute_different_location:*": { "subcaseMS": 2.257 },
   "webgpu:shader,validation,parse,diagnostic:conflicting_directive:*": { "subcaseMS": 1.244 },
@@ -2499,9 +2492,7 @@
   "webgpu:shader,validation,parse,diagnostic:valid_locations:*": { "subcaseMS": 1.368 },
   "webgpu:shader,validation,parse,diagnostic:valid_params:*": { "subcaseMS": 1.475 },
   "webgpu:shader,validation,parse,diagnostic:warning_unknown_rule:*": { "subcaseMS": 1.100 },
-  "webgpu:shader,validation,parse,discard:placement:*": { "subcaseMS": 3.357 },
   "webgpu:shader,validation,parse,enable:enable:*": { "subcaseMS": 2.303 },
-  "webgpu:shader,validation,parse,for:parse:*": { "subcaseMS": 17.898 },
   "webgpu:shader,validation,parse,identifiers:alias_name:*": { "subcaseMS": 1.262 },
   "webgpu:shader,validation,parse,identifiers:function_const_name:*": { "subcaseMS": 1.298 },
   "webgpu:shader,validation,parse,identifiers:function_let_name:*": { "subcaseMS": 1.299 },
@@ -2513,8 +2504,6 @@
   "webgpu:shader,validation,parse,identifiers:non_normalized:*": { "subcaseMS": 1.101 },
   "webgpu:shader,validation,parse,identifiers:override_name:*": { "subcaseMS": 1.228 },
   "webgpu:shader,validation,parse,identifiers:struct_name:*": { "subcaseMS": 1.230 },
-  "webgpu:shader,validation,parse,if:parse:*": { "subcaseMS": 39.663 },
-  "webgpu:shader,validation,parse,increment_decrement:parse:*": { "subcaseMS": 261.823 },
   "webgpu:shader,validation,parse,literal:abstract_float:*": { "subcaseMS": 1.411 },
   "webgpu:shader,validation,parse,literal:abstract_int:*": { "subcaseMS": 1.296 },
   "webgpu:shader,validation,parse,literal:bools:*": { "subcaseMS": 2.901 },
@@ -2522,13 +2511,11 @@
   "webgpu:shader,validation,parse,literal:f32:*": { "subcaseMS": 1.393 },
   "webgpu:shader,validation,parse,literal:i32:*": { "subcaseMS": 1.541 },
   "webgpu:shader,validation,parse,literal:u32:*": { "subcaseMS": 1.379 },
-  "webgpu:shader,validation,parse,loop:parse:*": { "subcaseMS": 16.168 },
   "webgpu:shader,validation,parse,must_use:builtin_must_use:*": { "subcaseMS": 1.400 },
   "webgpu:shader,validation,parse,must_use:builtin_no_must_use:*": { "subcaseMS": 1.206 },
   "webgpu:shader,validation,parse,must_use:call:*": { "subcaseMS": 1.275 },
   "webgpu:shader,validation,parse,must_use:declaration:*": { "subcaseMS": 1.523 },
   "webgpu:shader,validation,parse,must_use:ignore_result_of_non_must_use_that_returns_call_of_must_use:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,validation,parse,phony:parse:*": { "subcaseMS": 149.639 },
   "webgpu:shader,validation,parse,pipeline_stage:compute_parsing:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,parse,pipeline_stage:extra_on_compute_function:*": { "subcaseMS": 2.651 },
   "webgpu:shader,validation,parse,pipeline_stage:extra_on_fragment_function:*": { "subcaseMS": 1.001 },
@@ -2539,7 +2526,6 @@
   "webgpu:shader,validation,parse,pipeline_stage:vertex_parsing:*": { "subcaseMS": 1.500 },
   "webgpu:shader,validation,parse,requires:requires:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,parse,requires:wgsl_matches_api:*": { "subcaseMS": 1.000 },
-  "webgpu:shader,validation,parse,return:parse:*": { "subcaseMS": 16.893 },
   "webgpu:shader,validation,parse,semicolon:after_assignment:*": { "subcaseMS": 1.400 },
   "webgpu:shader,validation,parse,semicolon:after_call:*": { "subcaseMS": 1.301 },
   "webgpu:shader,validation,parse,semicolon:after_case:*": { "subcaseMS": 1.301 },
@@ -2593,13 +2579,7 @@
   "webgpu:shader,validation,parse,source:empty:*": { "subcaseMS": 1.101 },
   "webgpu:shader,validation,parse,source:invalid_source:*": { "subcaseMS": 1.100 },
   "webgpu:shader,validation,parse,source:valid_source:*": { "subcaseMS": 1.101 },
-  "webgpu:shader,validation,parse,statement_behavior:invalid_functions:*": { "subcaseMS": 1.107 },
-  "webgpu:shader,validation,parse,statement_behavior:invalid_statements:*": { "subcaseMS": 80.699 },
-  "webgpu:shader,validation,parse,statement_behavior:valid_functions:*": { "subcaseMS": 0.978 },
-  "webgpu:shader,validation,parse,statement_behavior:valid_statements:*": { "subcaseMS": 15.781 },
-  "webgpu:shader,validation,parse,switch:parse:*": { "subcaseMS": 16.068 },
   "webgpu:shader,validation,parse,unary_ops:all:*": { "subcaseMS": 1.000 },
-  "webgpu:shader,validation,parse,while:parse:*": { "subcaseMS": 8.641 },
   "webgpu:shader,validation,shader_io,align:multi_align:*": { "subcaseMS": 3.232 },
   "webgpu:shader,validation,shader_io,align:parsing:*": { "subcaseMS": 368.840 },
   "webgpu:shader,validation,shader_io,align:placement:*": { "subcaseMS": 20.963 },
@@ -2658,22 +2638,42 @@
   "webgpu:shader,validation,shader_io,workgroup_size:workgroup_size_function:*": { "subcaseMS": 0.800 },
   "webgpu:shader,validation,shader_io,workgroup_size:workgroup_size_var:*": { "subcaseMS": 2.101 },
   "webgpu:shader,validation,shader_io,workgroup_size:workgroup_size_vertex_shader:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,statement,break:placement:*": { "subcaseMS": 1.254 },
   "webgpu:shader,validation,statement,break_if:condition_type:*": { "subcaseMS": 542.975 },
+  "webgpu:shader,validation,statement,break_if:placement:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,statement,compound:parse:*": { "subcaseMS": 4.315 },
+  "webgpu:shader,validation,statement,const_assert:parse:*": { "subcaseMS": 1.400 },
+  "webgpu:shader,validation,statement,continue:module_scope:*": { "subcaseMS": 0.278 },
+  "webgpu:shader,validation,statement,continue:placement:*": { "subcaseMS": 53.998 },
+  "webgpu:shader,validation,statement,continuing:placement:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,statement,discard:placement:*": { "subcaseMS": 3.357 },
   "webgpu:shader,validation,statement,for:condition_type:*": { "subcaseMS": 5.714 },
+  "webgpu:shader,validation,statement,for:parse:*": { "subcaseMS": 17.898 },
   "webgpu:shader,validation,statement,if:condition_type:*": { "subcaseMS": 5.951 },
   "webgpu:shader,validation,statement,if:else_condition_type:*": { "subcaseMS": 5.717 },
+  "webgpu:shader,validation,statement,if:parse:*": { "subcaseMS": 39.663 },
   "webgpu:shader,validation,statement,increment_decrement:component:*": { "subcaseMS": 29.246 },
+  "webgpu:shader,validation,statement,increment_decrement:parse:*": { "subcaseMS": 261.823 },
   "webgpu:shader,validation,statement,increment_decrement:var_init_type:*": { "subcaseMS": 270.034 },
   "webgpu:shader,validation,statement,loop:break_if_type:*": { "subcaseMS": 5.488 },
+  "webgpu:shader,validation,statement,loop:parse:*": { "subcaseMS": 16.168 },
+  "webgpu:shader,validation,statement,phony:parse:*": { "subcaseMS": 149.639 },
   "webgpu:shader,validation,statement,phony:rhs_constructible:*": { "subcaseMS": 244.513 },
   "webgpu:shader,validation,statement,phony:rhs_with_decl:*": { "subcaseMS": 15.874 },
+  "webgpu:shader,validation,statement,return:parse:*": { "subcaseMS": 16.893 },
   "webgpu:shader,validation,statement,return:return_missing_value:*": { "subcaseMS": 3.009 },
   "webgpu:shader,validation,statement,return:return_type_match:*": { "subcaseMS": 74.912 },
   "webgpu:shader,validation,statement,return:return_unexpected_value:*": { "subcaseMS": 4.498 },
+  "webgpu:shader,validation,statement,statement_behavior:invalid_functions:*": { "subcaseMS": 1.107 },
+  "webgpu:shader,validation,statement,statement_behavior:invalid_statements:*": { "subcaseMS": 80.699 },
+  "webgpu:shader,validation,statement,statement_behavior:valid_functions:*": { "subcaseMS": 0.978 },
+  "webgpu:shader,validation,statement,statement_behavior:valid_statements:*": { "subcaseMS": 15.781 },
   "webgpu:shader,validation,statement,switch:case_types_match:*": { "subcaseMS": 4.140 },
   "webgpu:shader,validation,statement,switch:condition_type:*": { "subcaseMS": 5.402 },
   "webgpu:shader,validation,statement,switch:condition_type_match_case_type:*": { "subcaseMS": 3.750 },
+  "webgpu:shader,validation,statement,switch:parse:*": { "subcaseMS": 16.068 },
   "webgpu:shader,validation,statement,while:condition_type:*": { "subcaseMS": 5.328 },
+  "webgpu:shader,validation,statement,while:parse:*": { "subcaseMS": 8.641 },
   "webgpu:shader,validation,types,alias:any_type:*": { "subcaseMS": 42.329 },
   "webgpu:shader,validation,types,alias:match_non_alias:*": { "subcaseMS": 3.767 },
   "webgpu:shader,validation,types,alias:no_direct_recursion:*": { "subcaseMS": 1.450 },
diff --git a/src/webgpu/shader/validation/parse/break_if.spec.ts b/src/webgpu/shader/validation/parse/break_if.spec.ts
deleted file mode 100644
index c365b9185c65..000000000000
--- a/src/webgpu/shader/validation/parse/break_if.spec.ts
+++ /dev/null
@@ -1,92 +0,0 @@
-export const description = `Validation tests for break if`;
-
-import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../common/util/data_tables.js';
-import { ShaderValidationTest } from '../shader_validation_test.js';
-
-export const g = makeTestGroup(ShaderValidationTest);
-
-const kTests = {
-  compound_break: {
-    src: '{ break if true; }',
-    pass: false,
-  },
-  loop_break: {
-    src: 'loop { break if true; }',
-    pass: false,
-  },
-  loop_if_break: {
-    src: 'loop { if true { break if false; } }',
-    pass: false,
-  },
-  continuing_break_if: {
-    src: 'loop { continuing { break if true; } }',
-    pass: true,
-  },
-  continuing_break_if_parens: {
-    src: 'loop { continuing { break if (true); } }',
-    pass: true,
-  },
-  continuing_break_if_not_last: {
-    src: 'loop { continuing { break if (true); let a = 4;} }',
-    pass: false,
-  },
-  while_break: {
-    src: 'while true { break if true; }',
-    pass: false,
-  },
-  while_if_break: {
-    src: 'while true { if true { break if true; } }',
-    pass: false,
-  },
-  for_break: {
-    src: 'for (;;) { break if true; }',
-    pass: false,
-  },
-  for_if_break: {
-    src: 'for (;;) { if true { break if true; } }',
-    pass: false,
-  },
-  switch_case_break: {
-    src: 'switch(1) { default: { break if true; } }',
-    pass: false,
-  },
-  switch_case_if_break: {
-    src: 'switch(1) { default: { if true { break if true; } } }',
-    pass: false,
-  },
-  break: {
-    src: 'break if true;',
-    pass: false,
-  },
-  return_break: {
-    src: 'return break if true;',
-    pass: false,
-  },
-  if_break: {
-    src: 'if true { break if true; }',
-    pass: false,
-  },
-  continuing_if_break: {
-    src: 'loop { continuing { if (true) { break if true; } } }',
-    pass: false,
-  },
-  switch_break: {
-    src: 'switch(1) { break if true; }',
-    pass: false,
-  },
-};
-
-g.test('placement')
-  .desc('Test that break if placement is validated correctly')
-  .params(u => u.combine('stmt', keysOf(kTests)))
-  .fn(t => {
-    const code = `
-@vertex
-fn vtx() -> @builtin(position) vec4f {
-  ${kTests[t.params.stmt].src}
-  return vec4f(1);
-}
-    `;
-    t.expectCompileResult(kTests[t.params.stmt].pass, code);
-  });
diff --git a/src/webgpu/shader/validation/parse/for.spec.ts b/src/webgpu/shader/validation/parse/for.spec.ts
deleted file mode 100644
index 0e75b2e4a452..000000000000
--- a/src/webgpu/shader/validation/parse/for.spec.ts
+++ /dev/null
@@ -1,65 +0,0 @@
-export const description = `Validation parser tests for 'for' statements`;
-
-import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../common/util/data_tables.js';
-import { ShaderValidationTest } from '../shader_validation_test.js';
-
-export const g = makeTestGroup(ShaderValidationTest);
-
-const kTests = {
-  break: { wgsl: `for (;;) { break; }`, pass: true },
-
-  init_var: { wgsl: `for (var i = 1;;) { break; }`, pass: true },
-  init_var_type: { wgsl: `for (var i : i32 = 1;;) { break; }`, pass: true },
-  init_var_function: { wgsl: `for (var<function> i = 1;;) { break; }`, pass: true },
-  init_var_function_type: { wgsl: `for (var<function> i : i32 = 1;;) { break; }`, pass: true },
-  init_let: { wgsl: `for (let i = 1;;) { break; }`, pass: true },
-  init_let_type: { wgsl: `for (let i : u32 = 1;;) { break; }`, pass: true },
-  init_const: { wgsl: `for (const i = 1;;) { break; }`, pass: true },
-  init_const_type: { wgsl: `for (const i : f32 = 1;;) { break; }`, pass: true },
-  init_call: { wgsl: `for (x();;) { break; }`, pass: true },
-  init_phony: { wgsl: `for (_ = v;;) { break; }`, pass: true },
-  init_increment: { wgsl: `for (v++;;) { break; }`, pass: true },
-  init_compound_assign: { wgsl: `for (v += 3;;) { break; }`, pass: true },
-
-  cond_true: { wgsl: `for (;true;) { break; }`, pass: true },
-
-  cont_call: { wgsl: `for (;;x()) { break; }`, pass: true },
-  cont_phony: { wgsl: `for (;;_ = v) { break; }`, pass: true },
-  cont_increment: { wgsl: `for (;;v++) { break; }`, pass: true },
-  cont_compound_assign: { wgsl: `for (;;v += 3) { break; }`, pass: true },
-
-  init_cond: { wgsl: `for (var i = 1; i < 5;) {}`, pass: true },
-  cond_cont: { wgsl: `for (;v < 5; v++) {}`, pass: true },
-  init_cond_cont: { wgsl: `for (var i = 0; i < 5; i++) {}`, pass: true },
-  init_shadow: { wgsl: `for (var f = 0; f < 5; f++) {}`, pass: true },
-
-  no_semicolon: { wgsl: `for () { break; }`, pass: false },
-  one_semicolon: { wgsl: `for (;) { break; }`, pass: false },
-  no_paren: { wgsl: `for ;; { break; }`, pass: false },
-  empty: { wgsl: `for (;;) {}`, pass: false }, // note: fails due to behavior-analysis
-  init_expr: { wgsl: `for (true;;) { break; }`, pass: false },
-  cond_stmt: { wgsl: `for (;var i = 0;) { break; }`, pass: false },
-  cont_expr: { wgsl: `for (;;true) { break; }`, pass: false },
-  cont_var: { wgsl: `for (;;var i = 1) { break; }`, pass: false },
-  cont_var_type: { wgsl: `for (;;var i : i32 = 1) { break; }`, pass: false },
-  cont_var_function: { wgsl: `for (;;var<function> i = 1) { break; }`, pass: false },
-  cont_var_function_type: { wgsl: `for (;;var<function> i : i32 = 1) { break; }`, pass: false },
-  cont_let: { wgsl: `for (;;let i = 1) { break; }`, pass: false },
-  cont_let_type: { wgsl: `for (;;let i : u32 = 1) { break; }`, pass: false },
-};
-
-g.test('parse')
-  .desc(`Test that 'for' statements are parsed correctly.`)
-  .params(u => u.combine('test', keysOf(kTests)))
-  .fn(t => {
-    const code = `
-fn f() {
-  var v = 1;
-  ${kTests[t.params.test].wgsl}
-}
-
-fn x() {}
-`;
-    t.expectCompileResult(kTests[t.params.test].pass, code);
-  });
diff --git a/src/webgpu/shader/validation/parse/if.spec.ts b/src/webgpu/shader/validation/parse/if.spec.ts
deleted file mode 100644
index 4f0a10f9a90f..000000000000
--- a/src/webgpu/shader/validation/parse/if.spec.ts
+++ /dev/null
@@ -1,70 +0,0 @@
-export const description = `Validation parser tests for 'if' statements`;
-
-import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../common/util/data_tables.js';
-import { ShaderValidationTest } from '../shader_validation_test.js';
-
-export const g = makeTestGroup(ShaderValidationTest);
-
-const kTests = {
-  true: { wgsl: `if true {}`, pass: true },
-  paren_true: { wgsl: `if (true) {}`, pass: true },
-  expr: { wgsl: `if expr {}`, pass: true },
-  paren_expr: { wgsl: `if (expr) {}`, pass: true },
-
-  true_else: { wgsl: `if true {} else {}`, pass: true },
-  paren_true_else: { wgsl: `if (true) {} else {}`, pass: true },
-  expr_else: { wgsl: `if expr {} else {}`, pass: true },
-  paren_expr_else: { wgsl: `if (expr) {} else {}`, pass: true },
-
-  true_else_if_true: { wgsl: `if true {} else if true {}`, pass: true },
-  paren_true_else_if_paren_true: { wgsl: `if (true) {} else if (true) {}`, pass: true },
-  true_else_if_paren_true: { wgsl: `if true {} else if (true) {}`, pass: true },
-  paren_true_else_if_true: { wgsl: `if (true) {} else if true {}`, pass: true },
-
-  expr_else_if_expr: { wgsl: `if expr {} else if expr {}`, pass: true },
-  paren_expr_else_if_paren_expr: { wgsl: `if (expr) {} else if (expr) {}`, pass: true },
-  expr_else_if_paren_expr: { wgsl: `if expr {} else if (expr) {}`, pass: true },
-  paren_expr_else_if_expr: { wgsl: `if (expr) {} else if expr {}`, pass: true },
-
-  if: { wgsl: `if`, pass: false },
-  block: { wgsl: `if{}`, pass: false },
-  semicolon: { wgsl: `if;`, pass: false },
-  true_lbrace: { wgsl: `if true {`, pass: false },
-  true_rbrace: { wgsl: `if true }`, pass: false },
-
-  lparen_true: { wgsl: `if (true {}`, pass: false },
-  rparen_true: { wgsl: `if )true {}`, pass: false },
-  true_lparen: { wgsl: `if true( {}`, pass: false },
-  true_rparen: { wgsl: `if true) {}`, pass: false },
-
-  true_else_if_no_block: { wgsl: `if true {} else if `, pass: false },
-  true_else_if: { wgsl: `if true {} else if {}`, pass: false },
-  true_else_if_semicolon: { wgsl: `if true {} else if ;`, pass: false },
-  true_else_if_true_lbrace: { wgsl: `if true {} else if true {`, pass: false },
-  true_else_if_true_rbrace: { wgsl: `if true {} else if true }`, pass: false },
-
-  true_else_if_lparen_true: { wgsl: `if true {} else if (true {}`, pass: false },
-  true_else_if_rparen_true: { wgsl: `if true {} else if )true {}`, pass: false },
-  true_else_if_true_lparen: { wgsl: `if true {} else if true( {}`, pass: false },
-  true_else_if_true_rparen: { wgsl: `if true {} else if true) {}`, pass: false },
-
-  else: { wgsl: `else { }`, pass: false },
-  else_if: { wgsl: `else if true { }`, pass: false },
-  true_elif: { wgsl: `if (true) { } elif (true) {}`, pass: false },
-  true_elsif: { wgsl: `if (true) { } elsif (true) {}`, pass: false },
-  elif: { wgsl: `elif (true) {}`, pass: false },
-  elsif: { wgsl: `elsif (true) {}`, pass: false },
-};
-
-g.test('parse')
-  .desc(`Test that 'if' statements are parsed correctly.`)
-  .params(u => u.combine('test', keysOf(kTests)))
-  .fn(t => {
-    const code = `
-fn f() {
-  let expr = true;
-  ${kTests[t.params.test].wgsl}
-}`;
-    t.expectCompileResult(kTests[t.params.test].pass, code);
-  });
diff --git a/src/webgpu/shader/validation/parse/increment_decrement.spec.ts b/src/webgpu/shader/validation/parse/increment_decrement.spec.ts
deleted file mode 100644
index 8d97f1f2d48c..000000000000
--- a/src/webgpu/shader/validation/parse/increment_decrement.spec.ts
+++ /dev/null
@@ -1,143 +0,0 @@
-export const description = `Validation parser tests for increment and decrement statements`;
-
-import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../common/util/data_tables.js';
-import { ShaderValidationTest } from '../shader_validation_test.js';
-
-export const g = makeTestGroup(ShaderValidationTest);
-
-interface Case {
-  // Statement in a function to execute.
-  wgsl: string;
-  // True if the test should pass
-  pass: boolean;
-  // Module-scope program fragment, if any. Usually a declaration
-  gdecl?: string;
-}
-
-const kTests = {
-  var: { wgsl: 'a++;', pass: true },
-  vector: { wgsl: 'v++;', pass: false },
-  paren_var_paren: { wgsl: '(a)++;', pass: true },
-  star_and_var: { wgsl: '*&a++;', pass: true },
-  paren_star_and_var_paren: { wgsl: '(*&a)++;', pass: true },
-  many_star_and_var: { wgsl: '*&*&*&a++;', pass: true },
-
-  space: { wgsl: 'a ++;', pass: true },
-  tab: { wgsl: 'a\t++;', pass: true },
-  newline: { wgsl: 'a\n++;', pass: true },
-  cr: { wgsl: 'a\r++;', pass: true },
-  space_space: { wgsl: 'a ++ ;', pass: true },
-  plus_space_plus: { wgsl: 'a+ +;', pass: false },
-  minux_space_minus: { wgsl: 'a- -;', pass: false },
-
-  no_var: { wgsl: '++;', pass: false },
-  no_semi: { wgsl: 'a++', pass: false },
-  prefix: { wgsl: '++a;', pass: false },
-
-  postfix_x: { wgsl: 'v++.x;', pass: false },
-  postfix_r: { wgsl: 'v++.r;', pass: false },
-  postfix_index: { wgsl: 'v++[0];', pass: false },
-  postfix_field: { wgsl: 'a++.foo;', pass: false },
-
-  literal_i32: { wgsl: '12i++;', pass: false },
-  literal_u32: { wgsl: '12u++;', pass: false },
-  literal_abstract_int: { wgsl: '12++;', pass: false },
-  literal_abstract_float: { wgsl: '12.0++;', pass: false },
-  literal_f32: { wgsl: '12.0f++;', pass: false },
-
-  assign_to: { wgsl: 'a++ = 1;', pass: false },
-
-  at_global: { wgsl: '', pass: false, gdecl: 'var<private> g:i32; g++;' },
-  private: { wgsl: 'g++;', pass: true, gdecl: 'var<private> g:i32;' },
-  workgroup: { wgsl: 'g++;', pass: true, gdecl: 'var<workgroup> g:i32;' },
-  storage_rw: {
-    wgsl: 'g++;',
-    pass: true,
-    gdecl: '@group(0) @binding(0) var<storage,read_write> g: i32;',
-  },
-  storage_r: {
-    wgsl: 'g++;',
-    pass: false,
-    gdecl: '@group(0) @binding(0) var<storage,read> g: i32;',
-  },
-  storage: { wgsl: 'g++;', pass: false, gdecl: '@group(0) @binding(0) var<storage,read> g: i32;' },
-  uniform: { wgsl: 'g++;', pass: false, gdecl: '@group(0) @binding(0) var<uniform> g: i32;' },
-  texture: { wgsl: 'g++;', pass: false, gdecl: '@group(0) @binding(0) var g: texture_2d<u32>;' },
-  texture_x: {
-    wgsl: 'g.x++;',
-    pass: false,
-    gdecl: '@group(0) @binding(0) var g: texture_2d<u32>;',
-  },
-  texture_storage: {
-    wgsl: 'g++;',
-    pass: false,
-    gdecl: '@group(0) @binding(0) var g: texture_storage_2d<r32uint>;',
-  },
-  texture_storage_x: {
-    wgsl: 'g.x++;',
-    pass: false,
-    gdecl: '@group(0) @binding(0) var g: texture_storage_2d<r32uint>;',
-  },
-  sampler: { wgsl: 'g++;', pass: false, gdecl: '@group(0) @binding(0) var g: sampler;' },
-  sampler_comparison: {
-    wgsl: 'g++;',
-    pass: false,
-    gdecl: '@group(0) @binding(0) var g: sampler_comparison;',
-  },
-  override: { wgsl: 'g++;', pass: false, gdecl: 'override g:i32;' },
-  global_const: { wgsl: 'g++;', pass: false, gdecl: 'const g:i32 = 0;' },
-  workgroup_atomic: { wgsl: 'g++;', pass: false, gdecl: 'var<workgroup> g:atomic<i32>;' },
-  storage_atomic: {
-    wgsl: 'g++;',
-    pass: false,
-    gdecl: '@group(0) @binding(0) var<storage,read_write> g:atomic<u32>;',
-  },
-
-  subexpr: { wgsl: 'a = b++;', pass: false },
-  expr_paren: { wgsl: '(a++);', pass: false },
-  expr_add: { wgsl: '0 + a++;', pass: false },
-  expr_negate: { wgsl: '-a++;', pass: false },
-  inc_inc: { wgsl: 'a++++;', pass: false },
-  inc_space_inc: { wgsl: 'a++ ++;', pass: false },
-  inc_dec: { wgsl: 'a++--;', pass: false },
-  inc_space_dec: { wgsl: 'a++ --;', pass: false },
-  paren_inc: { wgsl: '(a++)++;', pass: false },
-  paren_dec: { wgsl: '(a++)--;', pass: false },
-
-  in_block: { wgsl: '{ a++; }', pass: true },
-  in_for_init: { wgsl: 'for (a++;false;) {}', pass: true },
-  in_for_cond: { wgsl: 'for (;a++;) {}', pass: false },
-  in_for_update: { wgsl: 'for (;false;a++) {}', pass: true },
-  in_for_update_semi: { wgsl: 'for (;false;a++;) {}', pass: false },
-  in_continuing: { wgsl: 'loop { continuing { a++; break if true;}}', pass: true },
-
-  let: { wgsl: 'let c = a; c++;', pass: false },
-  const: { wgsl: 'const c = 1; c++;', pass: false },
-  builtin: { wgsl: 'max++', pass: false },
-  enum: { wgsl: 'r32uint++', pass: false },
-  param: { wgsl: '', pass: false, gdecl: 'fn bump(p: i32) { p++;}' },
-};
-
-g.test('parse')
-  .desc(`Test that increment and decrement statements are parsed correctly.`)
-  .params(u => u.combine('test', keysOf(kTests)).combine('direction', ['up', 'down']))
-  .fn(t => {
-    const c = kTests[t.params.test] as Case;
-    let { wgsl, gdecl } = c;
-    gdecl = gdecl ?? '';
-    if (t.params.direction === 'down') {
-      wgsl = wgsl.replace('++', '--');
-      gdecl = gdecl.replace('++', '--');
-    }
-
-    const code = `
-${gdecl}
-fn f() {
-  var a: u32;
-  var b: u32;
-  var v: vec4u;
-  ${wgsl}
-}`;
-    t.expectCompileResult(c.pass, code);
-  });
diff --git a/src/webgpu/shader/validation/parse/loop.spec.ts b/src/webgpu/shader/validation/parse/loop.spec.ts
deleted file mode 100644
index fc930974b4d2..000000000000
--- a/src/webgpu/shader/validation/parse/loop.spec.ts
+++ /dev/null
@@ -1,65 +0,0 @@
-export const description = `Validation parser tests for 'loop' statements`;
-
-import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../common/util/data_tables.js';
-import { ShaderValidationTest } from '../shader_validation_test.js';
-
-export const g = makeTestGroup(ShaderValidationTest);
-
-const kTests = {
-  break: { wgsl: `loop { break; }`, pass: true },
-  return: { wgsl: `loop { return; }`, pass: true },
-  break_continuing: { wgsl: `loop { break; continuing {} }`, pass: true },
-  var_break: { wgsl: `loop { var a = 1; break; }`, pass: true },
-  var_break_continuing_inc: {
-    wgsl: `loop { var a = 1; break; continuing { a += 1; }}`,
-    pass: true,
-  },
-  var_break_continuing_discard: {
-    wgsl: `loop { var a = 1; break; continuing { discard; }}`,
-    pass: true,
-  },
-  continuing_break_if: {
-    wgsl: `loop { continuing { break if true; } }`,
-    pass: true,
-  },
-
-  expr_break: { wgsl: `loop expr { break; }`, pass: false },
-  loop: { wgsl: `loop`, pass: false },
-  continuing_break: { wgsl: `loop { continuing {} break; }`, pass: false },
-  break_continuing_continue: { wgsl: `loop { break; continuing { continue; } }`, pass: false },
-  break_continuing_return: { wgsl: `loop { break; continuing { return; } }`, pass: false },
-  break_continuing_if_break: {
-    wgsl: `loop { break; continuing { if true { break; } }`,
-    pass: false,
-  },
-  break_continuing_if_return: {
-    wgsl: `loop { break; continuing { if true { return; } }`,
-    pass: false,
-  },
-  break_continuing_lbrace: { wgsl: `loop { break; continuing { }`, pass: false },
-  break_continuing_rbrace: { wgsl: `loop { break; continuing } }`, pass: false },
-  continuing: { wgsl: `loop { continuing {} }`, pass: false },
-  semicolon: { wgsl: `loop;`, pass: false },
-  lbrace: { wgsl: `loop {`, pass: false },
-  rbrace: { wgsl: `loop }`, pass: false },
-  lparen: { wgsl: `loop ({}`, pass: false },
-  rparen: { wgsl: `loop ){}`, pass: false },
-
-  // note: these parse, but fails due to behavior-analysis
-  continue: { wgsl: `loop { continue; }`, pass: false },
-  discard: { wgsl: `loop { discard; }`, pass: false },
-  empty: { wgsl: `loop{}`, pass: false },
-};
-
-g.test('parse')
-  .desc(`Test that 'loop' statements are parsed correctly.`)
-  .params(u => u.combine('test', keysOf(kTests)))
-  .fn(t => {
-    const code = `
-fn f() {
-  let expr = true;
-  ${kTests[t.params.test].wgsl}
-}`;
-    t.expectCompileResult(kTests[t.params.test].pass, code);
-  });
diff --git a/src/webgpu/shader/validation/parse/phony.spec.ts b/src/webgpu/shader/validation/parse/phony.spec.ts
deleted file mode 100644
index 9ba100361229..000000000000
--- a/src/webgpu/shader/validation/parse/phony.spec.ts
+++ /dev/null
@@ -1,48 +0,0 @@
-export const description = `Validation parser tests for phony assignment statements`;
-
-import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../common/util/data_tables.js';
-import { ShaderValidationTest } from '../shader_validation_test.js';
-
-export const g = makeTestGroup(ShaderValidationTest);
-
-const kTests = {
-  literal: { wgsl: `_ = 1;`, pass: true },
-  expr: { wgsl: `_ = (1+v);`, pass: true },
-  var: { wgsl: `_ = v;`, pass: true },
-
-  in_for_init: { wgsl: `for (_ = v;false;) {}`, pass: true },
-  in_for_init_semi: { wgsl: `for (_ = v;;false;) {}`, pass: false },
-  in_for_update: { wgsl: `for (;false; _ = v) {}`, pass: true },
-  in_for_update_semi: { wgsl: `for (;false; _ = v;) {}`, pass: false },
-
-  in_block: { wgsl: `{_ = v;}`, pass: true },
-  in_continuing: { wgsl: `loop { continuing { _ = v; break if true;}}`, pass: true },
-
-  in_paren: { wgsl: `(_ = v;)`, pass: false },
-
-  underscore: { wgsl: `_`, pass: false },
-  underscore_semi: { wgsl: `_;`, pass: false },
-  underscore_equal: { wgsl: `_=`, pass: false },
-  underscore_equal_semi: { wgsl: `_=;`, pass: false },
-  underscore_equal_underscore_semi: { wgsl: `_=_;`, pass: false },
-  paren_underscore_paren: { wgsl: `(_) = 1;`, pass: false },
-  // LHS is not a reference type
-  star_ampersand_undsscore: { wgsl: `*&_ = 1;`, pass: false },
-  compound: { wgsl: `_ += 1;`, pass: false },
-  equality: { wgsl: `_ == 1;`, pass: false },
-  block: { wgsl: `_ = {};`, pass: false },
-  return: { wgsl: `_ = return;`, pass: false },
-};
-
-g.test('parse')
-  .desc(`Test that 'phony assignment' statements are parsed correctly.`)
-  .params(u => u.combine('test', keysOf(kTests)))
-  .fn(t => {
-    const code = `
-fn f() {
-  var v: u32;
-  ${kTests[t.params.test].wgsl}
-}`;
-    t.expectCompileResult(kTests[t.params.test].pass, code);
-  });
diff --git a/src/webgpu/shader/validation/parse/return.spec.ts b/src/webgpu/shader/validation/parse/return.spec.ts
deleted file mode 100644
index 12a29b2f8744..000000000000
--- a/src/webgpu/shader/validation/parse/return.spec.ts
+++ /dev/null
@@ -1,51 +0,0 @@
-export const description = `Validation parser tests for 'return' statements`;
-
-import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../common/util/data_tables.js';
-import { ShaderValidationTest } from '../shader_validation_test.js';
-
-export const g = makeTestGroup(ShaderValidationTest);
-
-const kTests = {
-  no_expr: { wgsl: `return;`, pass_value: false, pass_no_value: true },
-  v: { wgsl: `return v;`, pass_value: true, pass_no_value: false },
-  literal: { wgsl: `return 10;`, pass_value: true, pass_no_value: false },
-  expr: { wgsl: `return 1 + 2;`, pass_value: true, pass_no_value: false },
-  paren_expr: { wgsl: `return (1 + 2);`, pass_value: true, pass_no_value: false },
-  call: { wgsl: `return x();`, pass_value: true, pass_no_value: false },
-
-  v_no_semicolon: { wgsl: `return v`, pass_value: false, pass_no_value: false },
-  expr_no_semicolon: { wgsl: `return 1 + 2`, pass_value: false, pass_no_value: false },
-  phony_assign: { wgsl: `return _ = 1;`, pass_value: false, pass_no_value: false },
-  increment: { wgsl: `return v++;`, pass_value: false, pass_no_value: false },
-  compound_assign: { wgsl: `return v += 4;`, pass_value: false, pass_no_value: false },
-  lparen_literal: { wgsl: `return (4;`, pass_value: false, pass_no_value: false },
-  literal_lparen: { wgsl: `return 4(;`, pass_value: false, pass_no_value: false },
-  rparen_literal: { wgsl: `return )4;`, pass_value: false, pass_no_value: false },
-  literal_rparen: { wgsl: `return 4);`, pass_value: false, pass_no_value: false },
-  lparen_literal_lparen: { wgsl: `return (4(;`, pass_value: false, pass_no_value: false },
-  rparen_literal_rparen: { wgsl: `return )4);`, pass_value: false, pass_no_value: false },
-};
-
-g.test('parse')
-  .desc(`Test that 'return' statements are parsed correctly.`)
-  .params(u =>
-    u.combine('test', keysOf(kTests)).combine('fn_returns_value', [false, true] as const)
-  )
-  .fn(t => {
-    const code = `
-fn f() ${t.params.fn_returns_value ? '-> i32' : ''} {
-  let v = 42;
-  ${kTests[t.params.test].wgsl}
-}
-fn x() -> i32 {
-  return 1;
-}
-`;
-    t.expectCompileResult(
-      t.params.fn_returns_value
-        ? kTests[t.params.test].pass_value
-        : kTests[t.params.test].pass_no_value,
-      code
-    );
-  });
diff --git a/src/webgpu/shader/validation/parse/switch.spec.ts b/src/webgpu/shader/validation/parse/switch.spec.ts
deleted file mode 100644
index 9671c07ff50c..000000000000
--- a/src/webgpu/shader/validation/parse/switch.spec.ts
+++ /dev/null
@@ -1,78 +0,0 @@
-export const description = `Validation parser tests for 'switch' statements`;
-
-import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../common/util/data_tables.js';
-import { ShaderValidationTest } from '../shader_validation_test.js';
-
-export const g = makeTestGroup(ShaderValidationTest);
-
-const kTests = {
-  L_default: { wgsl: `switch L { default {} }`, pass: true },
-  L_paren_default: { wgsl: `switch (L) { default {} }`, pass: true },
-  L_case_1_2_default: { wgsl: `switch L { case 1, 2 {} default {} }`, pass: true },
-  L_case_1_case_2_default: { wgsl: `switch L { case 1 {} case 2 {} default {} }`, pass: true },
-  L_case_1_colon_case_2_colon_default_colon: {
-    wgsl: `switch L { case 1: {} case 2: {} default: {} }`,
-    pass: true,
-  },
-  L_case_1_colon_default_colon: { wgsl: `switch L { case 1: {} default: {} }`, pass: true },
-  L_case_1_colon_default: { wgsl: `switch L { case 1: {} default {} }`, pass: true },
-  L_case_1_default_2: { wgsl: `switch L { case 1, default, 2 {} }`, pass: true },
-  L_case_1_default_case_2: { wgsl: `switch L { case 1 {} default {} case 2 {} }`, pass: true },
-  L_case_1_default_colon: { wgsl: `switch L { case 1 {} default: {} }`, pass: true },
-  L_case_1_default: { wgsl: `switch L { case 1 {} default {} }`, pass: true },
-  L_case_2_1_default: { wgsl: `switch L { case 2, 1 {} default {} }`, pass: true },
-  L_case_2_case_1_default: { wgsl: `switch L { case 2 {} case 1 {} default {} }`, pass: true },
-  L_case_2_default_case_1: { wgsl: `switch L { case 2 {} default {} case 1 {} }`, pass: true },
-  L_case_builtin_default: { wgsl: `switch L { case max(1,2) {} default {} }`, pass: true },
-  L_case_C1_case_C2_default: { wgsl: `switch L { case C1 {} case C2 {} default {} }`, pass: true },
-  L_case_C1_default: { wgsl: `switch L { case C1 {} default {} }`, pass: true },
-  L_case_default_1: { wgsl: `switch L { case default, 1 {} }`, pass: true },
-  L_case_default_2_1: { wgsl: `switch L { case default, 2, 1 {} }`, pass: true },
-  L_case_default_2_case_1: { wgsl: `switch L { case default, 2 {} case 1 {} }`, pass: true },
-  L_case_default: { wgsl: `switch L { case default {} }`, pass: true },
-  L_case_expr_default: { wgsl: `switch L { case 1+1 {} default {} }`, pass: true },
-  L_default_break: { wgsl: `switch L { default { break; } }`, pass: true },
-  L_default_case_1_2: { wgsl: `switch L { default {} case 1, 2 {} }`, pass: true },
-  L_default_case_1_break: { wgsl: `switch L { default {} case 1 { break; } }`, pass: true },
-  L_default_case_1_case_2: { wgsl: `switch L { default {} case 1 {} case 2 {} }`, pass: true },
-  L_default_case_1_colon_break: { wgsl: `switch L { default {} case 1: { break; } }`, pass: true },
-  L_default_case_2_case_1: { wgsl: `switch L { default {} case 2 {} case 1 {} }`, pass: true },
-  L_default_colon_break: { wgsl: `switch L { default: { break; } }`, pass: true },
-  L_default_colon: { wgsl: `switch L { default: {} }`, pass: true },
-
-  L_no_block: { wgsl: `switch L`, pass: false },
-  L_empty_block: { wgsl: `switch L {}`, pass: false },
-  L_no_default: { wgsl: `switch L { case 1 {} }`, pass: false },
-  L_default_default: { wgsl: `switch L { default, default {} }`, pass: false },
-  L_default_block_default_block: { wgsl: `switch L { default {} default {} }`, pass: false },
-  L_case_1_case_1_default: { wgsl: `switch L { case 1 {} case 1 {} default {} }`, pass: false },
-  L_case_C1_case_C1_default: { wgsl: `switch L { case C1 {} case C1 {} default {} }`, pass: false },
-  L_case_C2_case_expr_default: {
-    wgsl: `switch L { case C2 {} case 1+1 {} default {} }`,
-    pass: false,
-  },
-  L_default_1: { wgsl: `switch L { default, 1 {} }`, pass: false },
-  L_default_2_case_1: { wgsl: `switch L { default, 2 {} case 1 {} }`, pass: false },
-
-  no_cond: { wgsl: `switch { default{} }`, pass: false },
-  no_cond_no_block: { wgsl: `switch;`, pass: false },
-  lparen_L: { wgsl: `switch (L { default {}}`, pass: false },
-  L_lparen: { wgsl: `switch L) { default {}}`, pass: false },
-  lparen_L_lparen: { wgsl: `switch )L) { default {}}`, pass: false },
-  rparen_L_rparen: { wgsl: `switch (L( { default {}}`, pass: false },
-};
-
-g.test('parse')
-  .desc(`Test that 'switch' statements are parsed correctly.`)
-  .params(u => u.combine('test', keysOf(kTests)))
-  .fn(t => {
-    const code = `
-fn f() {
-  let L = 1;
-  const C1 = 1;
-  const C2 = 2;
-  ${kTests[t.params.test].wgsl}
-}`;
-    t.expectCompileResult(kTests[t.params.test].pass, code);
-  });
diff --git a/src/webgpu/shader/validation/parse/while.spec.ts b/src/webgpu/shader/validation/parse/while.spec.ts
deleted file mode 100644
index 4e005f9ebfdc..000000000000
--- a/src/webgpu/shader/validation/parse/while.spec.ts
+++ /dev/null
@@ -1,42 +0,0 @@
-export const description = `Validation parser tests for 'while' statements`;
-
-import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../common/util/data_tables.js';
-import { ShaderValidationTest } from '../shader_validation_test.js';
-
-export const g = makeTestGroup(ShaderValidationTest);
-
-const kTests = {
-  true: { wgsl: `while true {}`, pass: true },
-  paren_true: { wgsl: `while (true) {}`, pass: true },
-  true_break: { wgsl: `while true { break; }`, pass: true },
-  true_discard: { wgsl: `while true { discard; }`, pass: true },
-  true_return: { wgsl: `while true { return; }`, pass: true },
-  expr: { wgsl: `while expr {}`, pass: true },
-  paren_expr: { wgsl: `while (expr) {}`, pass: true },
-
-  while: { wgsl: `while`, pass: false },
-  block: { wgsl: `while{}`, pass: false },
-  semicolon: { wgsl: `while;`, pass: false },
-  true_lbrace: { wgsl: `while true {`, pass: false },
-  true_rbrace: { wgsl: `while true }`, pass: false },
-
-  lparen_true: { wgsl: `while (true {}`, pass: false },
-  rparen_true: { wgsl: `while )true {}`, pass: false },
-  true_lparen: { wgsl: `while true( {}`, pass: false },
-  true_rparen: { wgsl: `while true) {}`, pass: false },
-  lparen_true_lparen: { wgsl: `while (true( {}`, pass: false },
-  rparen_true_rparen: { wgsl: `while )true) {}`, pass: false },
-};
-
-g.test('parse')
-  .desc(`Test that 'while' statements are parsed correctly.`)
-  .params(u => u.combine('test', keysOf(kTests)))
-  .fn(t => {
-    const code = `
-fn f() {
-  let expr = true;
-  ${kTests[t.params.test].wgsl}
-}`;
-    t.expectCompileResult(kTests[t.params.test].pass, code);
-  });
diff --git a/src/webgpu/shader/validation/parse/break.spec.ts b/src/webgpu/shader/validation/statement/break.spec.ts
similarity index 100%
rename from src/webgpu/shader/validation/parse/break.spec.ts
rename to src/webgpu/shader/validation/statement/break.spec.ts
diff --git a/src/webgpu/shader/validation/statement/break_if.spec.ts b/src/webgpu/shader/validation/statement/break_if.spec.ts
index 1cc291fd21dd..2841c94530a4 100644
--- a/src/webgpu/shader/validation/statement/break_if.spec.ts
+++ b/src/webgpu/shader/validation/statement/break_if.spec.ts
@@ -35,3 +35,88 @@ fn f() {
     const pass = t.params.type === 'bool';
     t.expectCompileResult(pass, code);
   });
+
+const kTests = {
+  compound_break: {
+    src: '{ break if true; }',
+    pass: false,
+  },
+  loop_break: {
+    src: 'loop { break if true; }',
+    pass: false,
+  },
+  loop_if_break: {
+    src: 'loop { if true { break if false; } }',
+    pass: false,
+  },
+  continuing_break_if: {
+    src: 'loop { continuing { break if true; } }',
+    pass: true,
+  },
+  continuing_break_if_parens: {
+    src: 'loop { continuing { break if (true); } }',
+    pass: true,
+  },
+  continuing_break_if_not_last: {
+    src: 'loop { continuing { break if (true); let a = 4;} }',
+    pass: false,
+  },
+  while_break: {
+    src: 'while true { break if true; }',
+    pass: false,
+  },
+  while_if_break: {
+    src: 'while true { if true { break if true; } }',
+    pass: false,
+  },
+  for_break: {
+    src: 'for (;;) { break if true; }',
+    pass: false,
+  },
+  for_if_break: {
+    src: 'for (;;) { if true { break if true; } }',
+    pass: false,
+  },
+  switch_case_break: {
+    src: 'switch(1) { default: { break if true; } }',
+    pass: false,
+  },
+  switch_case_if_break: {
+    src: 'switch(1) { default: { if true { break if true; } } }',
+    pass: false,
+  },
+  break: {
+    src: 'break if true;',
+    pass: false,
+  },
+  return_break: {
+    src: 'return break if true;',
+    pass: false,
+  },
+  if_break: {
+    src: 'if true { break if true; }',
+    pass: false,
+  },
+  continuing_if_break: {
+    src: 'loop { continuing { if (true) { break if true; } } }',
+    pass: false,
+  },
+  switch_break: {
+    src: 'switch(1) { break if true; }',
+    pass: false,
+  },
+};
+
+g.test('placement')
+  .desc('Test that break if placement is validated correctly')
+  .params(u => u.combine('stmt', keysOf(kTests)))
+  .fn(t => {
+    const code = `
+@vertex
+fn vtx() -> @builtin(position) vec4f {
+  ${kTests[t.params.stmt].src}
+  return vec4f(1);
+}
+    `;
+    t.expectCompileResult(kTests[t.params.stmt].pass, code);
+  });
diff --git a/src/webgpu/shader/validation/parse/compound.spec.ts b/src/webgpu/shader/validation/statement/compound.spec.ts
similarity index 100%
rename from src/webgpu/shader/validation/parse/compound.spec.ts
rename to src/webgpu/shader/validation/statement/compound.spec.ts
diff --git a/src/webgpu/shader/validation/parse/const_assert.spec.ts b/src/webgpu/shader/validation/statement/const_assert.spec.ts
similarity index 100%
rename from src/webgpu/shader/validation/parse/const_assert.spec.ts
rename to src/webgpu/shader/validation/statement/const_assert.spec.ts
diff --git a/src/webgpu/shader/validation/parse/continue.spec.ts b/src/webgpu/shader/validation/statement/continue.spec.ts
similarity index 100%
rename from src/webgpu/shader/validation/parse/continue.spec.ts
rename to src/webgpu/shader/validation/statement/continue.spec.ts
diff --git a/src/webgpu/shader/validation/parse/continuing.spec.ts b/src/webgpu/shader/validation/statement/continuing.spec.ts
similarity index 100%
rename from src/webgpu/shader/validation/parse/continuing.spec.ts
rename to src/webgpu/shader/validation/statement/continuing.spec.ts
diff --git a/src/webgpu/shader/validation/parse/discard.spec.ts b/src/webgpu/shader/validation/statement/discard.spec.ts
similarity index 100%
rename from src/webgpu/shader/validation/parse/discard.spec.ts
rename to src/webgpu/shader/validation/statement/discard.spec.ts
diff --git a/src/webgpu/shader/validation/statement/for.spec.ts b/src/webgpu/shader/validation/statement/for.spec.ts
index 647e1897bc63..0b4795be1edb 100644
--- a/src/webgpu/shader/validation/statement/for.spec.ts
+++ b/src/webgpu/shader/validation/statement/for.spec.ts
@@ -34,3 +34,61 @@ fn f() -> bool {
     const pass = t.params.type === 'bool';
     t.expectCompileResult(pass, code);
   });
+
+const kTests = {
+  break: { wgsl: `for (;;) { break; }`, pass: true },
+
+  init_var: { wgsl: `for (var i = 1;;) { break; }`, pass: true },
+  init_var_type: { wgsl: `for (var i : i32 = 1;;) { break; }`, pass: true },
+  init_var_function: { wgsl: `for (var<function> i = 1;;) { break; }`, pass: true },
+  init_var_function_type: { wgsl: `for (var<function> i : i32 = 1;;) { break; }`, pass: true },
+  init_let: { wgsl: `for (let i = 1;;) { break; }`, pass: true },
+  init_let_type: { wgsl: `for (let i : u32 = 1;;) { break; }`, pass: true },
+  init_const: { wgsl: `for (const i = 1;;) { break; }`, pass: true },
+  init_const_type: { wgsl: `for (const i : f32 = 1;;) { break; }`, pass: true },
+  init_call: { wgsl: `for (x();;) { break; }`, pass: true },
+  init_phony: { wgsl: `for (_ = v;;) { break; }`, pass: true },
+  init_increment: { wgsl: `for (v++;;) { break; }`, pass: true },
+  init_compound_assign: { wgsl: `for (v += 3;;) { break; }`, pass: true },
+
+  cond_true: { wgsl: `for (;true;) { break; }`, pass: true },
+
+  cont_call: { wgsl: `for (;;x()) { break; }`, pass: true },
+  cont_phony: { wgsl: `for (;;_ = v) { break; }`, pass: true },
+  cont_increment: { wgsl: `for (;;v++) { break; }`, pass: true },
+  cont_compound_assign: { wgsl: `for (;;v += 3) { break; }`, pass: true },
+
+  init_cond: { wgsl: `for (var i = 1; i < 5;) {}`, pass: true },
+  cond_cont: { wgsl: `for (;v < 5; v++) {}`, pass: true },
+  init_cond_cont: { wgsl: `for (var i = 0; i < 5; i++) {}`, pass: true },
+  init_shadow: { wgsl: `for (var f = 0; f < 5; f++) {}`, pass: true },
+
+  no_semicolon: { wgsl: `for () { break; }`, pass: false },
+  one_semicolon: { wgsl: `for (;) { break; }`, pass: false },
+  no_paren: { wgsl: `for ;; { break; }`, pass: false },
+  empty: { wgsl: `for (;;) {}`, pass: false }, // note: fails due to behavior-analysis
+  init_expr: { wgsl: `for (true;;) { break; }`, pass: false },
+  cond_stmt: { wgsl: `for (;var i = 0;) { break; }`, pass: false },
+  cont_expr: { wgsl: `for (;;true) { break; }`, pass: false },
+  cont_var: { wgsl: `for (;;var i = 1) { break; }`, pass: false },
+  cont_var_type: { wgsl: `for (;;var i : i32 = 1) { break; }`, pass: false },
+  cont_var_function: { wgsl: `for (;;var<function> i = 1) { break; }`, pass: false },
+  cont_var_function_type: { wgsl: `for (;;var<function> i : i32 = 1) { break; }`, pass: false },
+  cont_let: { wgsl: `for (;;let i = 1) { break; }`, pass: false },
+  cont_let_type: { wgsl: `for (;;let i : u32 = 1) { break; }`, pass: false },
+};
+
+g.test('parse')
+  .desc(`Test that 'for' statements are parsed correctly.`)
+  .params(u => u.combine('test', keysOf(kTests)))
+  .fn(t => {
+    const code = `
+fn f() {
+  var v = 1;
+  ${kTests[t.params.test].wgsl}
+}
+
+fn x() {}
+`;
+    t.expectCompileResult(kTests[t.params.test].pass, code);
+  });
diff --git a/src/webgpu/shader/validation/statement/if.spec.ts b/src/webgpu/shader/validation/statement/if.spec.ts
index 91f19674aee4..bb05eace7a14 100644
--- a/src/webgpu/shader/validation/statement/if.spec.ts
+++ b/src/webgpu/shader/validation/statement/if.spec.ts
@@ -63,3 +63,66 @@ fn f(c : bool) -> bool {
     const pass = t.params.type === 'bool';
     t.expectCompileResult(pass, code);
   });
+
+const kTests = {
+  true: { wgsl: `if true {}`, pass: true },
+  paren_true: { wgsl: `if (true) {}`, pass: true },
+  expr: { wgsl: `if expr {}`, pass: true },
+  paren_expr: { wgsl: `if (expr) {}`, pass: true },
+
+  true_else: { wgsl: `if true {} else {}`, pass: true },
+  paren_true_else: { wgsl: `if (true) {} else {}`, pass: true },
+  expr_else: { wgsl: `if expr {} else {}`, pass: true },
+  paren_expr_else: { wgsl: `if (expr) {} else {}`, pass: true },
+
+  true_else_if_true: { wgsl: `if true {} else if true {}`, pass: true },
+  paren_true_else_if_paren_true: { wgsl: `if (true) {} else if (true) {}`, pass: true },
+  true_else_if_paren_true: { wgsl: `if true {} else if (true) {}`, pass: true },
+  paren_true_else_if_true: { wgsl: `if (true) {} else if true {}`, pass: true },
+
+  expr_else_if_expr: { wgsl: `if expr {} else if expr {}`, pass: true },
+  paren_expr_else_if_paren_expr: { wgsl: `if (expr) {} else if (expr) {}`, pass: true },
+  expr_else_if_paren_expr: { wgsl: `if expr {} else if (expr) {}`, pass: true },
+  paren_expr_else_if_expr: { wgsl: `if (expr) {} else if expr {}`, pass: true },
+
+  if: { wgsl: `if`, pass: false },
+  block: { wgsl: `if{}`, pass: false },
+  semicolon: { wgsl: `if;`, pass: false },
+  true_lbrace: { wgsl: `if true {`, pass: false },
+  true_rbrace: { wgsl: `if true }`, pass: false },
+
+  lparen_true: { wgsl: `if (true {}`, pass: false },
+  rparen_true: { wgsl: `if )true {}`, pass: false },
+  true_lparen: { wgsl: `if true( {}`, pass: false },
+  true_rparen: { wgsl: `if true) {}`, pass: false },
+
+  true_else_if_no_block: { wgsl: `if true {} else if `, pass: false },
+  true_else_if: { wgsl: `if true {} else if {}`, pass: false },
+  true_else_if_semicolon: { wgsl: `if true {} else if ;`, pass: false },
+  true_else_if_true_lbrace: { wgsl: `if true {} else if true {`, pass: false },
+  true_else_if_true_rbrace: { wgsl: `if true {} else if true }`, pass: false },
+
+  true_else_if_lparen_true: { wgsl: `if true {} else if (true {}`, pass: false },
+  true_else_if_rparen_true: { wgsl: `if true {} else if )true {}`, pass: false },
+  true_else_if_true_lparen: { wgsl: `if true {} else if true( {}`, pass: false },
+  true_else_if_true_rparen: { wgsl: `if true {} else if true) {}`, pass: false },
+
+  else: { wgsl: `else { }`, pass: false },
+  else_if: { wgsl: `else if true { }`, pass: false },
+  true_elif: { wgsl: `if (true) { } elif (true) {}`, pass: false },
+  true_elsif: { wgsl: `if (true) { } elsif (true) {}`, pass: false },
+  elif: { wgsl: `elif (true) {}`, pass: false },
+  elsif: { wgsl: `elsif (true) {}`, pass: false },
+};
+
+g.test('parse')
+  .desc(`Test that 'if' statements are parsed correctly.`)
+  .params(u => u.combine('test', keysOf(kTests)))
+  .fn(t => {
+    const code = `
+fn f() {
+  let expr = true;
+  ${kTests[t.params.test].wgsl}
+}`;
+    t.expectCompileResult(kTests[t.params.test].pass, code);
+  });
diff --git a/src/webgpu/shader/validation/statement/increment_decrement.spec.ts b/src/webgpu/shader/validation/statement/increment_decrement.spec.ts
index 452f1cc7efbc..56cadc1c032a 100644
--- a/src/webgpu/shader/validation/statement/increment_decrement.spec.ts
+++ b/src/webgpu/shader/validation/statement/increment_decrement.spec.ts
@@ -132,3 +132,139 @@ fn f() {
 }`;
     t.expectCompileResult(pass, code);
   });
+
+interface ParseCase {
+  // Statement in a function to execute.
+  wgsl: string;
+  // True if the test should pass
+  pass: boolean;
+  // Module-scope program fragment, if any. Usually a declaration
+  gdecl?: string;
+}
+
+const kTests = {
+  var: { wgsl: 'a++;', pass: true },
+  vector: { wgsl: 'v++;', pass: false },
+  paren_var_paren: { wgsl: '(a)++;', pass: true },
+  star_and_var: { wgsl: '*&a++;', pass: true },
+  paren_star_and_var_paren: { wgsl: '(*&a)++;', pass: true },
+  many_star_and_var: { wgsl: '*&*&*&a++;', pass: true },
+
+  space: { wgsl: 'a ++;', pass: true },
+  tab: { wgsl: 'a\t++;', pass: true },
+  newline: { wgsl: 'a\n++;', pass: true },
+  cr: { wgsl: 'a\r++;', pass: true },
+  space_space: { wgsl: 'a ++ ;', pass: true },
+  plus_space_plus: { wgsl: 'a+ +;', pass: false },
+  minux_space_minus: { wgsl: 'a- -;', pass: false },
+
+  no_var: { wgsl: '++;', pass: false },
+  no_semi: { wgsl: 'a++', pass: false },
+  prefix: { wgsl: '++a;', pass: false },
+
+  postfix_x: { wgsl: 'v++.x;', pass: false },
+  postfix_r: { wgsl: 'v++.r;', pass: false },
+  postfix_index: { wgsl: 'v++[0];', pass: false },
+  postfix_field: { wgsl: 'a++.foo;', pass: false },
+
+  literal_i32: { wgsl: '12i++;', pass: false },
+  literal_u32: { wgsl: '12u++;', pass: false },
+  literal_abstract_int: { wgsl: '12++;', pass: false },
+  literal_abstract_float: { wgsl: '12.0++;', pass: false },
+  literal_f32: { wgsl: '12.0f++;', pass: false },
+
+  assign_to: { wgsl: 'a++ = 1;', pass: false },
+
+  at_global: { wgsl: '', pass: false, gdecl: 'var<private> g:i32; g++;' },
+  private: { wgsl: 'g++;', pass: true, gdecl: 'var<private> g:i32;' },
+  workgroup: { wgsl: 'g++;', pass: true, gdecl: 'var<workgroup> g:i32;' },
+  storage_rw: {
+    wgsl: 'g++;',
+    pass: true,
+    gdecl: '@group(0) @binding(0) var<storage,read_write> g: i32;',
+  },
+  storage_r: {
+    wgsl: 'g++;',
+    pass: false,
+    gdecl: '@group(0) @binding(0) var<storage,read> g: i32;',
+  },
+  storage: { wgsl: 'g++;', pass: false, gdecl: '@group(0) @binding(0) var<storage,read> g: i32;' },
+  uniform: { wgsl: 'g++;', pass: false, gdecl: '@group(0) @binding(0) var<uniform> g: i32;' },
+  texture: { wgsl: 'g++;', pass: false, gdecl: '@group(0) @binding(0) var g: texture_2d<u32>;' },
+  texture_x: {
+    wgsl: 'g.x++;',
+    pass: false,
+    gdecl: '@group(0) @binding(0) var g: texture_2d<u32>;',
+  },
+  texture_storage: {
+    wgsl: 'g++;',
+    pass: false,
+    gdecl: '@group(0) @binding(0) var g: texture_storage_2d<r32uint>;',
+  },
+  texture_storage_x: {
+    wgsl: 'g.x++;',
+    pass: false,
+    gdecl: '@group(0) @binding(0) var g: texture_storage_2d<r32uint>;',
+  },
+  sampler: { wgsl: 'g++;', pass: false, gdecl: '@group(0) @binding(0) var g: sampler;' },
+  sampler_comparison: {
+    wgsl: 'g++;',
+    pass: false,
+    gdecl: '@group(0) @binding(0) var g: sampler_comparison;',
+  },
+  override: { wgsl: 'g++;', pass: false, gdecl: 'override g:i32;' },
+  global_const: { wgsl: 'g++;', pass: false, gdecl: 'const g:i32 = 0;' },
+  workgroup_atomic: { wgsl: 'g++;', pass: false, gdecl: 'var<workgroup> g:atomic<i32>;' },
+  storage_atomic: {
+    wgsl: 'g++;',
+    pass: false,
+    gdecl: '@group(0) @binding(0) var<storage,read_write> g:atomic<u32>;',
+  },
+
+  subexpr: { wgsl: 'a = b++;', pass: false },
+  expr_paren: { wgsl: '(a++);', pass: false },
+  expr_add: { wgsl: '0 + a++;', pass: false },
+  expr_negate: { wgsl: '-a++;', pass: false },
+  inc_inc: { wgsl: 'a++++;', pass: false },
+  inc_space_inc: { wgsl: 'a++ ++;', pass: false },
+  inc_dec: { wgsl: 'a++--;', pass: false },
+  inc_space_dec: { wgsl: 'a++ --;', pass: false },
+  paren_inc: { wgsl: '(a++)++;', pass: false },
+  paren_dec: { wgsl: '(a++)--;', pass: false },
+
+  in_block: { wgsl: '{ a++; }', pass: true },
+  in_for_init: { wgsl: 'for (a++;false;) {}', pass: true },
+  in_for_cond: { wgsl: 'for (;a++;) {}', pass: false },
+  in_for_update: { wgsl: 'for (;false;a++) {}', pass: true },
+  in_for_update_semi: { wgsl: 'for (;false;a++;) {}', pass: false },
+  in_continuing: { wgsl: 'loop { continuing { a++; break if true;}}', pass: true },
+
+  let: { wgsl: 'let c = a; c++;', pass: false },
+  const: { wgsl: 'const c = 1; c++;', pass: false },
+  builtin: { wgsl: 'max++', pass: false },
+  enum: { wgsl: 'r32uint++', pass: false },
+  param: { wgsl: '', pass: false, gdecl: 'fn bump(p: i32) { p++;}' },
+};
+
+g.test('parse')
+  .desc(`Test that increment and decrement statements are parsed correctly.`)
+  .params(u => u.combine('test', keysOf(kTests)).combine('direction', ['up', 'down']))
+  .fn(t => {
+    const c = kTests[t.params.test] as ParseCase;
+    let { wgsl, gdecl } = c;
+    gdecl = gdecl ?? '';
+    if (t.params.direction === 'down') {
+      wgsl = wgsl.replace('++', '--');
+      gdecl = gdecl.replace('++', '--');
+    }
+
+    const code = `
+${gdecl}
+fn f() {
+  var a: u32;
+  var b: u32;
+  var v: vec4u;
+  ${wgsl}
+}`;
+    t.expectCompileResult(c.pass, code);
+  });
diff --git a/src/webgpu/shader/validation/statement/loop.spec.ts b/src/webgpu/shader/validation/statement/loop.spec.ts
index f0dad5e5c7a9..bb759a4181a5 100644
--- a/src/webgpu/shader/validation/statement/loop.spec.ts
+++ b/src/webgpu/shader/validation/statement/loop.spec.ts
@@ -35,3 +35,61 @@ fn f() {
     const pass = t.params.type === 'bool';
     t.expectCompileResult(pass, code);
   });
+
+const kTests = {
+  break: { wgsl: `loop { break; }`, pass: true },
+  return: { wgsl: `loop { return; }`, pass: true },
+  break_continuing: { wgsl: `loop { break; continuing {} }`, pass: true },
+  var_break: { wgsl: `loop { var a = 1; break; }`, pass: true },
+  var_break_continuing_inc: {
+    wgsl: `loop { var a = 1; break; continuing { a += 1; }}`,
+    pass: true,
+  },
+  var_break_continuing_discard: {
+    wgsl: `loop { var a = 1; break; continuing { discard; }}`,
+    pass: true,
+  },
+  continuing_break_if: {
+    wgsl: `loop { continuing { break if true; } }`,
+    pass: true,
+  },
+
+  expr_break: { wgsl: `loop expr { break; }`, pass: false },
+  loop: { wgsl: `loop`, pass: false },
+  continuing_break: { wgsl: `loop { continuing {} break; }`, pass: false },
+  break_continuing_continue: { wgsl: `loop { break; continuing { continue; } }`, pass: false },
+  break_continuing_return: { wgsl: `loop { break; continuing { return; } }`, pass: false },
+  break_continuing_if_break: {
+    wgsl: `loop { break; continuing { if true { break; } }`,
+    pass: false,
+  },
+  break_continuing_if_return: {
+    wgsl: `loop { break; continuing { if true { return; } }`,
+    pass: false,
+  },
+  break_continuing_lbrace: { wgsl: `loop { break; continuing { }`, pass: false },
+  break_continuing_rbrace: { wgsl: `loop { break; continuing } }`, pass: false },
+  continuing: { wgsl: `loop { continuing {} }`, pass: false },
+  semicolon: { wgsl: `loop;`, pass: false },
+  lbrace: { wgsl: `loop {`, pass: false },
+  rbrace: { wgsl: `loop }`, pass: false },
+  lparen: { wgsl: `loop ({}`, pass: false },
+  rparen: { wgsl: `loop ){}`, pass: false },
+
+  // note: these parse, but fails due to behavior-analysis
+  continue: { wgsl: `loop { continue; }`, pass: false },
+  discard: { wgsl: `loop { discard; }`, pass: false },
+  empty: { wgsl: `loop{}`, pass: false },
+};
+
+g.test('parse')
+  .desc(`Test that 'loop' statements are parsed correctly.`)
+  .params(u => u.combine('test', keysOf(kTests)))
+  .fn(t => {
+    const code = `
+fn f() {
+  let expr = true;
+  ${kTests[t.params.test].wgsl}
+}`;
+    t.expectCompileResult(kTests[t.params.test].pass, code);
+  });
diff --git a/src/webgpu/shader/validation/statement/phony.spec.ts b/src/webgpu/shader/validation/statement/phony.spec.ts
index 8701f41c3f3e..dfd8edc4c639 100644
--- a/src/webgpu/shader/validation/statement/phony.spec.ts
+++ b/src/webgpu/shader/validation/statement/phony.spec.ts
@@ -129,3 +129,44 @@ fn f() {
 }`;
     t.expectCompileResult(pass, code);
   });
+
+const kTests = {
+  literal: { wgsl: `_ = 1;`, pass: true },
+  expr: { wgsl: `_ = (1+v);`, pass: true },
+  var: { wgsl: `_ = v;`, pass: true },
+
+  in_for_init: { wgsl: `for (_ = v;false;) {}`, pass: true },
+  in_for_init_semi: { wgsl: `for (_ = v;;false;) {}`, pass: false },
+  in_for_update: { wgsl: `for (;false; _ = v) {}`, pass: true },
+  in_for_update_semi: { wgsl: `for (;false; _ = v;) {}`, pass: false },
+
+  in_block: { wgsl: `{_ = v;}`, pass: true },
+  in_continuing: { wgsl: `loop { continuing { _ = v; break if true;}}`, pass: true },
+
+  in_paren: { wgsl: `(_ = v;)`, pass: false },
+
+  underscore: { wgsl: `_`, pass: false },
+  underscore_semi: { wgsl: `_;`, pass: false },
+  underscore_equal: { wgsl: `_=`, pass: false },
+  underscore_equal_semi: { wgsl: `_=;`, pass: false },
+  underscore_equal_underscore_semi: { wgsl: `_=_;`, pass: false },
+  paren_underscore_paren: { wgsl: `(_) = 1;`, pass: false },
+  // LHS is not a reference type
+  star_ampersand_undsscore: { wgsl: `*&_ = 1;`, pass: false },
+  compound: { wgsl: `_ += 1;`, pass: false },
+  equality: { wgsl: `_ == 1;`, pass: false },
+  block: { wgsl: `_ = {};`, pass: false },
+  return: { wgsl: `_ = return;`, pass: false },
+};
+
+g.test('parse')
+  .desc(`Test that 'phony assignment' statements are parsed correctly.`)
+  .params(u => u.combine('test', keysOf(kTests)))
+  .fn(t => {
+    const code = `
+fn f() {
+  var v: u32;
+  ${kTests[t.params.test].wgsl}
+}`;
+    t.expectCompileResult(kTests[t.params.test].pass, code);
+  });
diff --git a/src/webgpu/shader/validation/statement/return.spec.ts b/src/webgpu/shader/validation/statement/return.spec.ts
index 7475dfc5bdc2..81294b676c65 100644
--- a/src/webgpu/shader/validation/statement/return.spec.ts
+++ b/src/webgpu/shader/validation/statement/return.spec.ts
@@ -1,6 +1,7 @@
 export const description = `Validation tests for 'return' statements'`;
 
 import { makeTestGroup } from '../../../../common/framework/test_group.js';
+import { keysOf } from '../../../../common/util/data_tables.js';
 import { isConvertible, scalarTypeOf, Type } from '../../../util/conversion.js';
 import { ShaderValidationTest } from '../shader_validation_test.js';
 
@@ -109,3 +110,47 @@ fn f() -> ${fnReturnType} {
     const pass = isConvertible(returnValueType, fnReturnType);
     t.expectCompileResult(pass, code);
   });
+
+const kTests = {
+  no_expr: { wgsl: `return;`, pass_value: false, pass_no_value: true },
+  v: { wgsl: `return v;`, pass_value: true, pass_no_value: false },
+  literal: { wgsl: `return 10;`, pass_value: true, pass_no_value: false },
+  expr: { wgsl: `return 1 + 2;`, pass_value: true, pass_no_value: false },
+  paren_expr: { wgsl: `return (1 + 2);`, pass_value: true, pass_no_value: false },
+  call: { wgsl: `return x();`, pass_value: true, pass_no_value: false },
+
+  v_no_semicolon: { wgsl: `return v`, pass_value: false, pass_no_value: false },
+  expr_no_semicolon: { wgsl: `return 1 + 2`, pass_value: false, pass_no_value: false },
+  phony_assign: { wgsl: `return _ = 1;`, pass_value: false, pass_no_value: false },
+  increment: { wgsl: `return v++;`, pass_value: false, pass_no_value: false },
+  compound_assign: { wgsl: `return v += 4;`, pass_value: false, pass_no_value: false },
+  lparen_literal: { wgsl: `return (4;`, pass_value: false, pass_no_value: false },
+  literal_lparen: { wgsl: `return 4(;`, pass_value: false, pass_no_value: false },
+  rparen_literal: { wgsl: `return )4;`, pass_value: false, pass_no_value: false },
+  literal_rparen: { wgsl: `return 4);`, pass_value: false, pass_no_value: false },
+  lparen_literal_lparen: { wgsl: `return (4(;`, pass_value: false, pass_no_value: false },
+  rparen_literal_rparen: { wgsl: `return )4);`, pass_value: false, pass_no_value: false },
+};
+
+g.test('parse')
+  .desc(`Test that 'return' statements are parsed correctly.`)
+  .params(u =>
+    u.combine('test', keysOf(kTests)).combine('fn_returns_value', [false, true] as const)
+  )
+  .fn(t => {
+    const code = `
+fn f() ${t.params.fn_returns_value ? '-> i32' : ''} {
+  let v = 42;
+  ${kTests[t.params.test].wgsl}
+}
+fn x() -> i32 {
+  return 1;
+}
+`;
+    t.expectCompileResult(
+      t.params.fn_returns_value
+        ? kTests[t.params.test].pass_value
+        : kTests[t.params.test].pass_no_value,
+      code
+    );
+  });
diff --git a/src/webgpu/shader/validation/parse/statement_behavior.spec.ts b/src/webgpu/shader/validation/statement/statement_behavior.spec.ts
similarity index 100%
rename from src/webgpu/shader/validation/parse/statement_behavior.spec.ts
rename to src/webgpu/shader/validation/statement/statement_behavior.spec.ts
diff --git a/src/webgpu/shader/validation/statement/switch.spec.ts b/src/webgpu/shader/validation/statement/switch.spec.ts
index 30760427c1d7..c685378c8042 100644
--- a/src/webgpu/shader/validation/statement/switch.spec.ts
+++ b/src/webgpu/shader/validation/statement/switch.spec.ts
@@ -99,3 +99,74 @@ switch 1 {
       t.params.case_b_type === 'abstract-int';
     t.expectCompileResult(pass, code);
   });
+
+const kTests = {
+  L_default: { wgsl: `switch L { default {} }`, pass: true },
+  L_paren_default: { wgsl: `switch (L) { default {} }`, pass: true },
+  L_case_1_2_default: { wgsl: `switch L { case 1, 2 {} default {} }`, pass: true },
+  L_case_1_case_2_default: { wgsl: `switch L { case 1 {} case 2 {} default {} }`, pass: true },
+  L_case_1_colon_case_2_colon_default_colon: {
+    wgsl: `switch L { case 1: {} case 2: {} default: {} }`,
+    pass: true,
+  },
+  L_case_1_colon_default_colon: { wgsl: `switch L { case 1: {} default: {} }`, pass: true },
+  L_case_1_colon_default: { wgsl: `switch L { case 1: {} default {} }`, pass: true },
+  L_case_1_default_2: { wgsl: `switch L { case 1, default, 2 {} }`, pass: true },
+  L_case_1_default_case_2: { wgsl: `switch L { case 1 {} default {} case 2 {} }`, pass: true },
+  L_case_1_default_colon: { wgsl: `switch L { case 1 {} default: {} }`, pass: true },
+  L_case_1_default: { wgsl: `switch L { case 1 {} default {} }`, pass: true },
+  L_case_2_1_default: { wgsl: `switch L { case 2, 1 {} default {} }`, pass: true },
+  L_case_2_case_1_default: { wgsl: `switch L { case 2 {} case 1 {} default {} }`, pass: true },
+  L_case_2_default_case_1: { wgsl: `switch L { case 2 {} default {} case 1 {} }`, pass: true },
+  L_case_builtin_default: { wgsl: `switch L { case max(1,2) {} default {} }`, pass: true },
+  L_case_C1_case_C2_default: { wgsl: `switch L { case C1 {} case C2 {} default {} }`, pass: true },
+  L_case_C1_default: { wgsl: `switch L { case C1 {} default {} }`, pass: true },
+  L_case_default_1: { wgsl: `switch L { case default, 1 {} }`, pass: true },
+  L_case_default_2_1: { wgsl: `switch L { case default, 2, 1 {} }`, pass: true },
+  L_case_default_2_case_1: { wgsl: `switch L { case default, 2 {} case 1 {} }`, pass: true },
+  L_case_default: { wgsl: `switch L { case default {} }`, pass: true },
+  L_case_expr_default: { wgsl: `switch L { case 1+1 {} default {} }`, pass: true },
+  L_default_break: { wgsl: `switch L { default { break; } }`, pass: true },
+  L_default_case_1_2: { wgsl: `switch L { default {} case 1, 2 {} }`, pass: true },
+  L_default_case_1_break: { wgsl: `switch L { default {} case 1 { break; } }`, pass: true },
+  L_default_case_1_case_2: { wgsl: `switch L { default {} case 1 {} case 2 {} }`, pass: true },
+  L_default_case_1_colon_break: { wgsl: `switch L { default {} case 1: { break; } }`, pass: true },
+  L_default_case_2_case_1: { wgsl: `switch L { default {} case 2 {} case 1 {} }`, pass: true },
+  L_default_colon_break: { wgsl: `switch L { default: { break; } }`, pass: true },
+  L_default_colon: { wgsl: `switch L { default: {} }`, pass: true },
+
+  L_no_block: { wgsl: `switch L`, pass: false },
+  L_empty_block: { wgsl: `switch L {}`, pass: false },
+  L_no_default: { wgsl: `switch L { case 1 {} }`, pass: false },
+  L_default_default: { wgsl: `switch L { default, default {} }`, pass: false },
+  L_default_block_default_block: { wgsl: `switch L { default {} default {} }`, pass: false },
+  L_case_1_case_1_default: { wgsl: `switch L { case 1 {} case 1 {} default {} }`, pass: false },
+  L_case_C1_case_C1_default: { wgsl: `switch L { case C1 {} case C1 {} default {} }`, pass: false },
+  L_case_C2_case_expr_default: {
+    wgsl: `switch L { case C2 {} case 1+1 {} default {} }`,
+    pass: false,
+  },
+  L_default_1: { wgsl: `switch L { default, 1 {} }`, pass: false },
+  L_default_2_case_1: { wgsl: `switch L { default, 2 {} case 1 {} }`, pass: false },
+
+  no_cond: { wgsl: `switch { default{} }`, pass: false },
+  no_cond_no_block: { wgsl: `switch;`, pass: false },
+  lparen_L: { wgsl: `switch (L { default {}}`, pass: false },
+  L_lparen: { wgsl: `switch L) { default {}}`, pass: false },
+  lparen_L_lparen: { wgsl: `switch )L) { default {}}`, pass: false },
+  rparen_L_rparen: { wgsl: `switch (L( { default {}}`, pass: false },
+};
+
+g.test('parse')
+  .desc(`Test that 'switch' statements are parsed correctly.`)
+  .params(u => u.combine('test', keysOf(kTests)))
+  .fn(t => {
+    const code = `
+fn f() {
+  let L = 1;
+  const C1 = 1;
+  const C2 = 2;
+  ${kTests[t.params.test].wgsl}
+}`;
+    t.expectCompileResult(kTests[t.params.test].pass, code);
+  });
diff --git a/src/webgpu/shader/validation/statement/while.spec.ts b/src/webgpu/shader/validation/statement/while.spec.ts
index da57512f5433..34832cbda49c 100644
--- a/src/webgpu/shader/validation/statement/while.spec.ts
+++ b/src/webgpu/shader/validation/statement/while.spec.ts
@@ -34,3 +34,38 @@ fn f() -> bool {
     const pass = t.params.type === 'bool';
     t.expectCompileResult(pass, code);
   });
+
+const kTests = {
+  true: { wgsl: `while true {}`, pass: true },
+  paren_true: { wgsl: `while (true) {}`, pass: true },
+  true_break: { wgsl: `while true { break; }`, pass: true },
+  true_discard: { wgsl: `while true { discard; }`, pass: true },
+  true_return: { wgsl: `while true { return; }`, pass: true },
+  expr: { wgsl: `while expr {}`, pass: true },
+  paren_expr: { wgsl: `while (expr) {}`, pass: true },
+
+  while: { wgsl: `while`, pass: false },
+  block: { wgsl: `while{}`, pass: false },
+  semicolon: { wgsl: `while;`, pass: false },
+  true_lbrace: { wgsl: `while true {`, pass: false },
+  true_rbrace: { wgsl: `while true }`, pass: false },
+
+  lparen_true: { wgsl: `while (true {}`, pass: false },
+  rparen_true: { wgsl: `while )true {}`, pass: false },
+  true_lparen: { wgsl: `while true( {}`, pass: false },
+  true_rparen: { wgsl: `while true) {}`, pass: false },
+  lparen_true_lparen: { wgsl: `while (true( {}`, pass: false },
+  rparen_true_rparen: { wgsl: `while )true) {}`, pass: false },
+};
+
+g.test('parse')
+  .desc(`Test that 'while' statements are parsed correctly.`)
+  .params(u => u.combine('test', keysOf(kTests)))
+  .fn(t => {
+    const code = `
+fn f() {
+  let expr = true;
+  ${kTests[t.params.test].wgsl}
+}`;
+    t.expectCompileResult(kTests[t.params.test].pass, code);
+  });