diff --git a/CHANGELOG.md b/CHANGELOG.md
index 4360ce16004..83328636d15 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -88,6 +88,10 @@ Fix `STATUS_HEAP_CORRUPTION` crash when concurrently calling `create_sampler`. B
 - Fixed build error occurring when the `profiling` dependency is configured to have profiling active. By @kpreid in [#7916](https://github.com/gfx-rs/wgpu/pull/7916).
 - Emit a validation error instead of panicking when a query set index is OOB. By @ErichDonGubler in [#7908](https://github.com/gfx-rs/wgpu/pull/7908).
 
+#### Naga
+
+- Escape, rather than strip, identifiers with Unicode. By @ErichDonGubler in [7995](https://github.com/gfx-rs/wgpu/pull/7995).
+
 ## v26.0.0 (2025-07-09)
 
 ### Major Features
@@ -145,9 +149,9 @@ let (device, queue) = adapter
     .unwrap();
 ```
 
-More examples of this 
+More examples of this
 
-By @Vecvec in [#7829](https://github.com/gfx-rs/wgpu/pull/7829). 
+By @Vecvec in [#7829](https://github.com/gfx-rs/wgpu/pull/7829).
 
 ### Naga
 
@@ -231,7 +235,6 @@ By @Vecvec in [#7829](https://github.com/gfx-rs/wgpu/pull/7829).
   - The definition of `enum CommandEncoderError` has changed significantly, to reflect which errors can be raised by `CommandEncoder.finish()`. There are also some errors that no longer appear directly in `CommandEncoderError`, and instead appear nested within the `RenderPass` or `ComputePass` variants.
   - `CopyError` has been removed. Errors that were previously a `CopyError` are now a `CommandEncoderError` returned by `finish()`. (The detailed reasons for copies to fail were and still are described by `TransferError`, which was previously a variant of `CopyError`, and is now a variant of `CommandEncoderError`).
 
-
 #### Naga
 
 - Mark `readonly_and_readwrite_storage_textures` & `packed_4x8_integer_dot_product` language extensions as implemented. By @teoxoy in [#7543](https://github.com/gfx-rs/wgpu/pull/7543)
diff --git a/naga/src/proc/namer.rs b/naga/src/proc/namer.rs
index 0b812ff036a..4911a6831e6 100644
--- a/naga/src/proc/namer.rs
+++ b/naga/src/proc/namer.rs
@@ -131,20 +131,31 @@ impl Namer {
         {
             Cow::Borrowed(string)
         } else {
-            let mut filtered = string
-                .chars()
-                .filter(|&c| c.is_ascii_alphanumeric() || c == '_')
-                .fold(String::new(), |mut s, c| {
-                    if s.ends_with('_') && c == '_' {
-                        return s;
-                    }
+            let mut filtered = string.chars().fold(String::new(), |mut s, c| {
+                let had_underscore_at_end = s.ends_with('_');
+                if had_underscore_at_end && c == '_' {
+                    return s;
+                }
+                if c.is_ascii_alphanumeric() || c == '_' {
                     s.push(c);
-                    s
-                });
+                } else {
+                    use core::fmt::Write as _;
+                    if !s.is_empty() && !had_underscore_at_end {
+                        s.push('_');
+                    }
+                    write!(s, "u{:04x}_", c as u32).unwrap();
+                }
+                s
+            });
             let stripped_len = filtered.trim_end_matches(SEPARATOR).len();
             filtered.truncate(stripped_len);
             if filtered.is_empty() {
                 filtered.push_str("unnamed");
+            } else if filtered.starts_with(|c: char| c.is_ascii_digit()) {
+                unreachable!(
+                    "internal error: invalid identifier starting with ASCII digit {:?}",
+                    filtered.chars().nth(0)
+                )
             }
             Cow::Owned(filtered)
         };
diff --git a/naga/tests/in/wgsl/7995-unicode-idents.wgsl b/naga/tests/in/wgsl/7995-unicode-idents.wgsl
new file mode 100644
index 00000000000..64e932ed09d
--- /dev/null
+++ b/naga/tests/in/wgsl/7995-unicode-idents.wgsl
@@ -0,0 +1,14 @@
+// NOTE: This allows us to suppress compaction below, to force the handling of identifiers
+// containing Unicode.
+@group(0) @binding(0)
+var<storage> asdf: f32;
+
+fn compute() -> f32 {
+  let θ2 = asdf + 9001.0;
+  return θ2;
+}
+
+@compute @workgroup_size(1, 1)
+fn main() {
+  compute();
+}
diff --git a/naga/tests/out/glsl/spv-do-while.main.Fragment.glsl b/naga/tests/out/glsl/spv-do-while.main.Fragment.glsl
index fb8f6e1efa7..a79915b8881 100644
--- a/naga/tests/out/glsl/spv-do-while.main.Fragment.glsl
+++ b/naga/tests/out/glsl/spv-do-while.main.Fragment.glsl
@@ -4,7 +4,7 @@ precision highp float;
 precision highp int;
 
 
-void fb1_(inout bool cond) {
+void f_u0028_b1_u003b(inout bool cond) {
     bool loop_init = true;
     while(true) {
         if (!loop_init) {
@@ -22,7 +22,7 @@ void fb1_(inout bool cond) {
 void main_1() {
     bool param = false;
     param = false;
-    fb1_(param);
+    f_u0028_b1_u003b(param);
     return;
 }
 
diff --git a/naga/tests/out/glsl/wgsl-7995-unicode-idents.main.Compute.glsl b/naga/tests/out/glsl/wgsl-7995-unicode-idents.main.Compute.glsl
new file mode 100644
index 00000000000..25fb2d7548f
--- /dev/null
+++ b/naga/tests/out/glsl/wgsl-7995-unicode-idents.main.Compute.glsl
@@ -0,0 +1,21 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(std430) readonly buffer type_block_0Compute { float _group_0_binding_0_cs; };
+
+
+float compute() {
+    float _e1 = _group_0_binding_0_cs;
+    float u03b8_2_ = (_e1 + 9001.0);
+    return u03b8_2_;
+}
+
+void main() {
+    float _e0 = compute();
+    return;
+}
+
diff --git a/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_i32.Compute.glsl b/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_i32.Compute.glsl
index 92d9c4355ff..af3f9a13c99 100644
--- a/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_i32.Compute.glsl
+++ b/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_i32.Compute.glsl
@@ -5,11 +5,11 @@ precision highp int;
 
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
-struct _atomic_compare_exchange_resultSint4_ {
+struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
     int old_value;
     bool exchanged;
 };
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
     uint old_value;
     bool exchanged;
 };
@@ -50,7 +50,7 @@ void main() {
                     int new = floatBitsToInt((intBitsToFloat(_e14) + 1.0));
                     uint _e20 = i;
                     int _e22 = old;
-                    _atomic_compare_exchange_resultSint4_ _e23; _e23.old_value = atomicCompSwap(_group_0_binding_0_cs[_e20], _e22, new);
+                    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e23; _e23.old_value = atomicCompSwap(_group_0_binding_0_cs[_e20], _e22, new);
                     _e23.exchanged = (_e23.old_value == _e22);
                     old = _e23.old_value;
                     exchanged = _e23.exchanged;
diff --git a/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_u32.Compute.glsl b/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_u32.Compute.glsl
index bbb4dea843b..bd32bb95781 100644
--- a/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_u32.Compute.glsl
+++ b/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_u32.Compute.glsl
@@ -5,11 +5,11 @@ precision highp int;
 
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
-struct _atomic_compare_exchange_resultSint4_ {
+struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
     int old_value;
     bool exchanged;
 };
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
     uint old_value;
     bool exchanged;
 };
@@ -50,7 +50,7 @@ void main() {
                     uint new = floatBitsToUint((uintBitsToFloat(_e14) + 1.0));
                     uint _e20 = i_1;
                     uint _e22 = old_1;
-                    _atomic_compare_exchange_resultUint4_ _e23; _e23.old_value = atomicCompSwap(_group_0_binding_1_cs[_e20], _e22, new);
+                    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e23; _e23.old_value = atomicCompSwap(_group_0_binding_1_cs[_e20], _e22, new);
                     _e23.exchanged = (_e23.old_value == _e22);
                     old_1 = _e23.old_value;
                     exchanged_1 = _e23.exchanged;
diff --git a/naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl b/naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl
index 7e1b5061ba1..5713a76340f 100644
--- a/naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl
+++ b/naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl
@@ -9,11 +9,11 @@ struct Struct {
     uint atomic_scalar;
     int atomic_arr[2];
 };
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
     uint old_value;
     bool exchanged;
 };
-struct _atomic_compare_exchange_resultSint4_ {
+struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
     int old_value;
     bool exchanged;
 };
@@ -135,21 +135,21 @@ void main() {
     int _e295 = atomicExchange(workgroup_atomic_arr[1], 1);
     uint _e299 = atomicExchange(workgroup_struct.atomic_scalar, 1u);
     int _e304 = atomicExchange(workgroup_struct.atomic_arr[1], 1);
-    _atomic_compare_exchange_resultUint4_ _e308; _e308.old_value = atomicCompSwap(_group_0_binding_0_cs, 1u, 2u);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e308; _e308.old_value = atomicCompSwap(_group_0_binding_0_cs, 1u, 2u);
     _e308.exchanged = (_e308.old_value == 1u);
-    _atomic_compare_exchange_resultSint4_ _e313; _e313.old_value = atomicCompSwap(_group_0_binding_1_cs[1], 1, 2);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e313; _e313.old_value = atomicCompSwap(_group_0_binding_1_cs[1], 1, 2);
     _e313.exchanged = (_e313.old_value == 1);
-    _atomic_compare_exchange_resultUint4_ _e318; _e318.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e318; _e318.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_scalar, 1u, 2u);
     _e318.exchanged = (_e318.old_value == 1u);
-    _atomic_compare_exchange_resultSint4_ _e324; _e324.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_arr[1], 1, 2);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e324; _e324.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_arr[1], 1, 2);
     _e324.exchanged = (_e324.old_value == 1);
-    _atomic_compare_exchange_resultUint4_ _e328; _e328.old_value = atomicCompSwap(workgroup_atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e328; _e328.old_value = atomicCompSwap(workgroup_atomic_scalar, 1u, 2u);
     _e328.exchanged = (_e328.old_value == 1u);
-    _atomic_compare_exchange_resultSint4_ _e333; _e333.old_value = atomicCompSwap(workgroup_atomic_arr[1], 1, 2);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e333; _e333.old_value = atomicCompSwap(workgroup_atomic_arr[1], 1, 2);
     _e333.exchanged = (_e333.old_value == 1);
-    _atomic_compare_exchange_resultUint4_ _e338; _e338.old_value = atomicCompSwap(workgroup_struct.atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e338; _e338.old_value = atomicCompSwap(workgroup_struct.atomic_scalar, 1u, 2u);
     _e338.exchanged = (_e338.old_value == 1u);
-    _atomic_compare_exchange_resultSint4_ _e344; _e344.old_value = atomicCompSwap(workgroup_struct.atomic_arr[1], 1, 2);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e344; _e344.old_value = atomicCompSwap(workgroup_struct.atomic_arr[1], 1, 2);
     _e344.exchanged = (_e344.old_value == 1);
     return;
 }
diff --git a/naga/tests/out/hlsl/spv-do-while.hlsl b/naga/tests/out/hlsl/spv-do-while.hlsl
index 5724831de7b..cdd5273a6f3 100644
--- a/naga/tests/out/hlsl/spv-do-while.hlsl
+++ b/naga/tests/out/hlsl/spv-do-while.hlsl
@@ -1,4 +1,4 @@
-void fb1_(inout bool cond)
+void f_u0028_b1_u003b(inout bool cond)
 {
     uint2 loop_bound = uint2(4294967295u, 4294967295u);
     bool loop_init = true;
@@ -22,7 +22,7 @@ void main_1()
     bool param = (bool)0;
 
     param = false;
-    fb1_(param);
+    f_u0028_b1_u003b(param);
     return;
 }
 
diff --git a/naga/tests/out/hlsl/spv-fetch_depth.hlsl b/naga/tests/out/hlsl/spv-fetch_depth.hlsl
index 71788135f1a..909ac96ad3a 100644
--- a/naga/tests/out/hlsl/spv-fetch_depth.hlsl
+++ b/naga/tests/out/hlsl/spv-fetch_depth.hlsl
@@ -19,7 +19,7 @@ void function()
 }
 
 [numthreads(32, 1, 1)]
-void cullfetch_depth()
+void cull_u003a_u003a_fetch_depth()
 {
     function();
 }
diff --git a/naga/tests/out/hlsl/spv-fetch_depth.ron b/naga/tests/out/hlsl/spv-fetch_depth.ron
index cb8a85c5f8c..2651131aa2c 100644
--- a/naga/tests/out/hlsl/spv-fetch_depth.ron
+++ b/naga/tests/out/hlsl/spv-fetch_depth.ron
@@ -5,7 +5,7 @@
     ],
     compute:[
         (
-            entry_point:"cullfetch_depth",
+            entry_point:"cull_u003a_u003a_fetch_depth",
             target_profile:"cs_5_1",
         ),
     ],
diff --git a/naga/tests/out/hlsl/wgsl-7995-unicode-idents.hlsl b/naga/tests/out/hlsl/wgsl-7995-unicode-idents.hlsl
new file mode 100644
index 00000000000..3bc8261789f
--- /dev/null
+++ b/naga/tests/out/hlsl/wgsl-7995-unicode-idents.hlsl
@@ -0,0 +1,15 @@
+ByteAddressBuffer asdf : register(t0);
+
+float compute()
+{
+    float _e1 = asfloat(asdf.Load(0));
+    float u03b8_2_ = (_e1 + 9001.0);
+    return u03b8_2_;
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    const float _e0 = compute();
+    return;
+}
diff --git a/naga/tests/out/hlsl/wgsl-7995-unicode-idents.ron b/naga/tests/out/hlsl/wgsl-7995-unicode-idents.ron
new file mode 100644
index 00000000000..a07b03300b1
--- /dev/null
+++ b/naga/tests/out/hlsl/wgsl-7995-unicode-idents.ron
@@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"main",
+            target_profile:"cs_5_1",
+        ),
+    ],
+)
diff --git a/naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.hlsl b/naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.hlsl
index e0e4ca4021c..b9086e4a797 100644
--- a/naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.hlsl
+++ b/naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.hlsl
@@ -5,13 +5,13 @@ struct NagaConstants {
 };
 ConstantBuffer<NagaConstants> _NagaConstants: register(b0, space1);
 
-struct _atomic_compare_exchange_resultSint8_ {
+struct _atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e {
     int64_t old_value;
     bool exchanged;
     int _end_pad_0;
 };
 
-struct _atomic_compare_exchange_resultUint8_ {
+struct _atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e {
     uint64_t old_value;
     bool exchanged;
     int _end_pad_0;
@@ -63,7 +63,7 @@ void test_atomic_compare_exchange_i64_()
                     int64_t new_ = (_e14 + 10L);
                     uint _e19 = i;
                     int64_t _e21 = old;
-                    _atomic_compare_exchange_resultSint8_ _e22; arr_i64_.InterlockedCompareExchange64(_e19*8, _e21, new_, _e22.old_value);
+                    _atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e _e22; arr_i64_.InterlockedCompareExchange64(_e19*8, _e21, new_, _e22.old_value);
                     _e22.exchanged = (_e22.old_value == _e21);
                     old = _e22.old_value;
                     exchanged = _e22.exchanged;
@@ -115,7 +115,7 @@ void test_atomic_compare_exchange_u64_()
                     uint64_t new_1 = (_e14 + 10uL);
                     uint _e19 = i_1;
                     uint64_t _e21 = old_1;
-                    _atomic_compare_exchange_resultUint8_ _e22; arr_u64_.InterlockedCompareExchange64(_e19*8, _e21, new_1, _e22.old_value);
+                    _atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e _e22; arr_u64_.InterlockedCompareExchange64(_e19*8, _e21, new_1, _e22.old_value);
                     _e22.exchanged = (_e22.old_value == _e21);
                     old_1 = _e22.old_value;
                     exchanged_1 = _e22.exchanged;
diff --git a/naga/tests/out/hlsl/wgsl-atomicCompareExchange.hlsl b/naga/tests/out/hlsl/wgsl-atomicCompareExchange.hlsl
index d523f7a9d2f..907b8303731 100644
--- a/naga/tests/out/hlsl/wgsl-atomicCompareExchange.hlsl
+++ b/naga/tests/out/hlsl/wgsl-atomicCompareExchange.hlsl
@@ -1,9 +1,9 @@
-struct _atomic_compare_exchange_resultSint4_ {
+struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
     int old_value;
     bool exchanged;
 };
 
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
     uint old_value;
     bool exchanged;
 };
@@ -54,7 +54,7 @@ void test_atomic_compare_exchange_i32_()
                     int new_ = asint((asfloat(_e14) + 1.0));
                     uint _e20 = i;
                     int _e22 = old;
-                    _atomic_compare_exchange_resultSint4_ _e23; arr_i32_.InterlockedCompareExchange(_e20*4, _e22, new_, _e23.old_value);
+                    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e23; arr_i32_.InterlockedCompareExchange(_e20*4, _e22, new_, _e23.old_value);
                     _e23.exchanged = (_e23.old_value == _e22);
                     old = _e23.old_value;
                     exchanged = _e23.exchanged;
@@ -106,7 +106,7 @@ void test_atomic_compare_exchange_u32_()
                     uint new_1 = asuint((asfloat(_e14) + 1.0));
                     uint _e20 = i_1;
                     uint _e22 = old_1;
-                    _atomic_compare_exchange_resultUint4_ _e23; arr_u32_.InterlockedCompareExchange(_e20*4, _e22, new_1, _e23.old_value);
+                    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e23; arr_u32_.InterlockedCompareExchange(_e20*4, _e22, new_1, _e23.old_value);
                     _e23.exchanged = (_e23.old_value == _e22);
                     old_1 = _e23.old_value;
                     exchanged_1 = _e23.exchanged;
diff --git a/naga/tests/out/hlsl/wgsl-atomicOps-int64.hlsl b/naga/tests/out/hlsl/wgsl-atomicOps-int64.hlsl
index 5c05b4ea9f5..93a806afdd9 100644
--- a/naga/tests/out/hlsl/wgsl-atomicOps-int64.hlsl
+++ b/naga/tests/out/hlsl/wgsl-atomicOps-int64.hlsl
@@ -10,13 +10,13 @@ struct Struct {
     int64_t atomic_arr[2];
 };
 
-struct _atomic_compare_exchange_resultUint8_ {
+struct _atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e {
     uint64_t old_value;
     bool exchanged;
     int _end_pad_0;
 };
 
-struct _atomic_compare_exchange_resultSint8_ {
+struct _atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e {
     int64_t old_value;
     bool exchanged;
     int _end_pad_0;
@@ -126,21 +126,21 @@ void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_Group
     int64_t _e279; InterlockedExchange(workgroup_atomic_arr[1], 1L, _e279);
     uint64_t _e283; InterlockedExchange(workgroup_struct.atomic_scalar, 1uL, _e283);
     int64_t _e288; InterlockedExchange(workgroup_struct.atomic_arr[1], 1L, _e288);
-    _atomic_compare_exchange_resultUint8_ _e292; storage_atomic_scalar.InterlockedCompareExchange64(0, 1uL, 2uL, _e292.old_value);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e _e292; storage_atomic_scalar.InterlockedCompareExchange64(0, 1uL, 2uL, _e292.old_value);
     _e292.exchanged = (_e292.old_value == 1uL);
-    _atomic_compare_exchange_resultSint8_ _e297; storage_atomic_arr.InterlockedCompareExchange64(8, 1L, 2L, _e297.old_value);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e _e297; storage_atomic_arr.InterlockedCompareExchange64(8, 1L, 2L, _e297.old_value);
     _e297.exchanged = (_e297.old_value == 1L);
-    _atomic_compare_exchange_resultUint8_ _e302; storage_struct.InterlockedCompareExchange64(0, 1uL, 2uL, _e302.old_value);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e _e302; storage_struct.InterlockedCompareExchange64(0, 1uL, 2uL, _e302.old_value);
     _e302.exchanged = (_e302.old_value == 1uL);
-    _atomic_compare_exchange_resultSint8_ _e308; storage_struct.InterlockedCompareExchange64(8+8, 1L, 2L, _e308.old_value);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e _e308; storage_struct.InterlockedCompareExchange64(8+8, 1L, 2L, _e308.old_value);
     _e308.exchanged = (_e308.old_value == 1L);
-    _atomic_compare_exchange_resultUint8_ _e312; InterlockedCompareExchange(workgroup_atomic_scalar, 1uL, 2uL, _e312.old_value);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e _e312; InterlockedCompareExchange(workgroup_atomic_scalar, 1uL, 2uL, _e312.old_value);
     _e312.exchanged = (_e312.old_value == 1uL);
-    _atomic_compare_exchange_resultSint8_ _e317; InterlockedCompareExchange(workgroup_atomic_arr[1], 1L, 2L, _e317.old_value);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e _e317; InterlockedCompareExchange(workgroup_atomic_arr[1], 1L, 2L, _e317.old_value);
     _e317.exchanged = (_e317.old_value == 1L);
-    _atomic_compare_exchange_resultUint8_ _e322; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1uL, 2uL, _e322.old_value);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_8_u003e _e322; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1uL, 2uL, _e322.old_value);
     _e322.exchanged = (_e322.old_value == 1uL);
-    _atomic_compare_exchange_resultSint8_ _e328; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], 1L, 2L, _e328.old_value);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_8_u003e _e328; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], 1L, 2L, _e328.old_value);
     _e328.exchanged = (_e328.old_value == 1L);
     return;
 }
diff --git a/naga/tests/out/hlsl/wgsl-atomicOps.hlsl b/naga/tests/out/hlsl/wgsl-atomicOps.hlsl
index 5771c898d94..47b586d7dba 100644
--- a/naga/tests/out/hlsl/wgsl-atomicOps.hlsl
+++ b/naga/tests/out/hlsl/wgsl-atomicOps.hlsl
@@ -3,12 +3,12 @@ struct Struct {
     int atomic_arr[2];
 };
 
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
     uint old_value;
     bool exchanged;
 };
 
-struct _atomic_compare_exchange_resultSint4_ {
+struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
     int old_value;
     bool exchanged;
 };
@@ -117,21 +117,21 @@ void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_Group
     int _e295; InterlockedExchange(workgroup_atomic_arr[1], int(1), _e295);
     uint _e299; InterlockedExchange(workgroup_struct.atomic_scalar, 1u, _e299);
     int _e304; InterlockedExchange(workgroup_struct.atomic_arr[1], int(1), _e304);
-    _atomic_compare_exchange_resultUint4_ _e308; storage_atomic_scalar.InterlockedCompareExchange(0, 1u, 2u, _e308.old_value);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e308; storage_atomic_scalar.InterlockedCompareExchange(0, 1u, 2u, _e308.old_value);
     _e308.exchanged = (_e308.old_value == 1u);
-    _atomic_compare_exchange_resultSint4_ _e313; storage_atomic_arr.InterlockedCompareExchange(4, int(1), int(2), _e313.old_value);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e313; storage_atomic_arr.InterlockedCompareExchange(4, int(1), int(2), _e313.old_value);
     _e313.exchanged = (_e313.old_value == int(1));
-    _atomic_compare_exchange_resultUint4_ _e318; storage_struct.InterlockedCompareExchange(0, 1u, 2u, _e318.old_value);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e318; storage_struct.InterlockedCompareExchange(0, 1u, 2u, _e318.old_value);
     _e318.exchanged = (_e318.old_value == 1u);
-    _atomic_compare_exchange_resultSint4_ _e324; storage_struct.InterlockedCompareExchange(4+4, int(1), int(2), _e324.old_value);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e324; storage_struct.InterlockedCompareExchange(4+4, int(1), int(2), _e324.old_value);
     _e324.exchanged = (_e324.old_value == int(1));
-    _atomic_compare_exchange_resultUint4_ _e328; InterlockedCompareExchange(workgroup_atomic_scalar, 1u, 2u, _e328.old_value);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e328; InterlockedCompareExchange(workgroup_atomic_scalar, 1u, 2u, _e328.old_value);
     _e328.exchanged = (_e328.old_value == 1u);
-    _atomic_compare_exchange_resultSint4_ _e333; InterlockedCompareExchange(workgroup_atomic_arr[1], int(1), int(2), _e333.old_value);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e333; InterlockedCompareExchange(workgroup_atomic_arr[1], int(1), int(2), _e333.old_value);
     _e333.exchanged = (_e333.old_value == int(1));
-    _atomic_compare_exchange_resultUint4_ _e338; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1u, 2u, _e338.old_value);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e338; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1u, 2u, _e338.old_value);
     _e338.exchanged = (_e338.old_value == 1u);
-    _atomic_compare_exchange_resultSint4_ _e344; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], int(1), int(2), _e344.old_value);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e344; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], int(1), int(2), _e344.old_value);
     _e344.exchanged = (_e344.old_value == int(1));
     return;
 }
diff --git a/naga/tests/out/msl/spv-do-while.msl b/naga/tests/out/msl/spv-do-while.msl
index 7e5d111b505..d58be5ef508 100644
--- a/naga/tests/out/msl/spv-do-while.msl
+++ b/naga/tests/out/msl/spv-do-while.msl
@@ -5,7 +5,7 @@
 using metal::uint;
 
 
-void fb1_(
+void f_u0028_b1_u003b(
     thread bool& cond
 ) {
     uint2 loop_bound = uint2(4294967295u);
@@ -29,7 +29,7 @@ void main_1(
 ) {
     bool param = {};
     param = false;
-    fb1_(param);
+    f_u0028_b1_u003b(param);
     return;
 }
 
diff --git a/naga/tests/out/msl/spv-fetch_depth.msl b/naga/tests/out/msl/spv-fetch_depth.msl
index 867a284372c..642cf305a5c 100644
--- a/naga/tests/out/msl/spv-fetch_depth.msl
+++ b/naga/tests/out/msl/spv-fetch_depth.msl
@@ -22,7 +22,7 @@ void function(
     return;
 }
 
-kernel void cullfetch_depth(
+kernel void cull_u003a_u003a_fetch_depth(
   device type_2& global [[user(fake0)]]
 , device type_4 const& global_1 [[user(fake0)]]
 , metal::depth2d<float, metal::access::sample> global_2 [[user(fake0)]]
diff --git a/naga/tests/out/msl/wgsl-7995-unicode-idents.msl b/naga/tests/out/msl/wgsl-7995-unicode-idents.msl
new file mode 100644
index 00000000000..9f912706d1c
--- /dev/null
+++ b/naga/tests/out/msl/wgsl-7995-unicode-idents.msl
@@ -0,0 +1,21 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+
+float compute(
+    device float const& asdf
+) {
+    float _e1 = asdf;
+    float u03b8_2_ = _e1 + 9001.0;
+    return u03b8_2_;
+}
+
+kernel void main_(
+  device float const& asdf [[user(fake0)]]
+) {
+    float _e0 = compute(asdf);
+    return;
+}
diff --git a/naga/tests/out/msl/wgsl-atomicCompareExchange.msl b/naga/tests/out/msl/wgsl-atomicCompareExchange.msl
index 72c418b9ed3..3c372b55e63 100644
--- a/naga/tests/out/msl/wgsl-atomicCompareExchange.msl
+++ b/naga/tests/out/msl/wgsl-atomicCompareExchange.msl
@@ -10,19 +10,19 @@ struct type_2 {
 struct type_4 {
     metal::atomic_uint inner[128];
 };
-struct _atomic_compare_exchange_resultSint4_ {
+struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
     int old_value;
     bool exchanged;
     char _pad2[3];
 };
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
     uint old_value;
     bool exchanged;
     char _pad2[3];
 };
 
 template <typename A>
-_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
+_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
     device A *atomic_ptr,
     int cmp,
     int v
@@ -31,10 +31,10 @@ _atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit
         atomic_ptr, &cmp, v,
         metal::memory_order_relaxed, metal::memory_order_relaxed
     );
-    return _atomic_compare_exchange_resultSint4_{cmp, swapped};
+    return _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e{cmp, swapped};
 }
 template <typename A>
-_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
+_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
     threadgroup A *atomic_ptr,
     int cmp,
     int v
@@ -43,11 +43,11 @@ _atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit
         atomic_ptr, &cmp, v,
         metal::memory_order_relaxed, metal::memory_order_relaxed
     );
-    return _atomic_compare_exchange_resultSint4_{cmp, swapped};
+    return _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e{cmp, swapped};
 }
 
 template <typename A>
-_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
+_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
     device A *atomic_ptr,
     uint cmp,
     uint v
@@ -56,10 +56,10 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
         atomic_ptr, &cmp, v,
         metal::memory_order_relaxed, metal::memory_order_relaxed
     );
-    return _atomic_compare_exchange_resultUint4_{cmp, swapped};
+    return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
 }
 template <typename A>
-_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
+_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
     threadgroup A *atomic_ptr,
     uint cmp,
     uint v
@@ -68,7 +68,7 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
         atomic_ptr, &cmp, v,
         metal::memory_order_relaxed, metal::memory_order_relaxed
     );
-    return _atomic_compare_exchange_resultUint4_{cmp, swapped};
+    return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
 }
 constant uint SIZE = 128u;
 
@@ -112,7 +112,7 @@ kernel void test_atomic_compare_exchange_i32_(
                     int new_ = as_type<int>(as_type<float>(_e14) + 1.0);
                     uint _e20 = i;
                     int _e22 = old;
-                    _atomic_compare_exchange_resultSint4_ _e23 = naga_atomic_compare_exchange_weak_explicit(&arr_i32_.inner[_e20], _e22, new_);
+                    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e23 = naga_atomic_compare_exchange_weak_explicit(&arr_i32_.inner[_e20], _e22, new_);
                     old = _e23.old_value;
                     exchanged = _e23.exchanged;
                 }
@@ -163,7 +163,7 @@ kernel void test_atomic_compare_exchange_u32_(
                     uint new_1 = as_type<uint>(as_type<float>(_e14) + 1.0);
                     uint _e20 = i_1;
                     uint _e22 = old_1;
-                    _atomic_compare_exchange_resultUint4_ _e23 = naga_atomic_compare_exchange_weak_explicit(&arr_u32_.inner[_e20], _e22, new_1);
+                    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e23 = naga_atomic_compare_exchange_weak_explicit(&arr_u32_.inner[_e20], _e22, new_1);
                     old_1 = _e23.old_value;
                     exchanged_1 = _e23.exchanged;
                 }
diff --git a/naga/tests/out/msl/wgsl-atomicOps.msl b/naga/tests/out/msl/wgsl-atomicOps.msl
index 63cab2a80f5..8266c1bd764 100644
--- a/naga/tests/out/msl/wgsl-atomicOps.msl
+++ b/naga/tests/out/msl/wgsl-atomicOps.msl
@@ -11,19 +11,19 @@ struct Struct {
     metal::atomic_uint atomic_scalar;
     type_2 atomic_arr;
 };
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
     uint old_value;
     bool exchanged;
     char _pad2[3];
 };
-struct _atomic_compare_exchange_resultSint4_ {
+struct _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e {
     int old_value;
     bool exchanged;
     char _pad2[3];
 };
 
 template <typename A>
-_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
+_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
     device A *atomic_ptr,
     uint cmp,
     uint v
@@ -32,10 +32,10 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
         atomic_ptr, &cmp, v,
         metal::memory_order_relaxed, metal::memory_order_relaxed
     );
-    return _atomic_compare_exchange_resultUint4_{cmp, swapped};
+    return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
 }
 template <typename A>
-_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
+_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
     threadgroup A *atomic_ptr,
     uint cmp,
     uint v
@@ -44,11 +44,11 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
         atomic_ptr, &cmp, v,
         metal::memory_order_relaxed, metal::memory_order_relaxed
     );
-    return _atomic_compare_exchange_resultUint4_{cmp, swapped};
+    return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
 }
 
 template <typename A>
-_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
+_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
     device A *atomic_ptr,
     int cmp,
     int v
@@ -57,10 +57,10 @@ _atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit
         atomic_ptr, &cmp, v,
         metal::memory_order_relaxed, metal::memory_order_relaxed
     );
-    return _atomic_compare_exchange_resultSint4_{cmp, swapped};
+    return _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e{cmp, swapped};
 }
 template <typename A>
-_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
+_atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
     threadgroup A *atomic_ptr,
     int cmp,
     int v
@@ -69,7 +69,7 @@ _atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit
         atomic_ptr, &cmp, v,
         metal::memory_order_relaxed, metal::memory_order_relaxed
     );
-    return _atomic_compare_exchange_resultSint4_{cmp, swapped};
+    return _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e{cmp, swapped};
 }
 
 struct cs_mainInput {
@@ -182,13 +182,13 @@ kernel void cs_main(
     int _e295 = metal::atomic_exchange_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
     uint _e299 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
     int _e304 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
-    _atomic_compare_exchange_resultUint4_ _e308 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_scalar, 1u, 2u);
-    _atomic_compare_exchange_resultSint4_ _e313 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_arr.inner[1], 1, 2);
-    _atomic_compare_exchange_resultUint4_ _e318 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_scalar, 1u, 2u);
-    _atomic_compare_exchange_resultSint4_ _e324 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_arr.inner[1], 1, 2);
-    _atomic_compare_exchange_resultUint4_ _e328 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_scalar, 1u, 2u);
-    _atomic_compare_exchange_resultSint4_ _e333 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_arr.inner[1], 1, 2);
-    _atomic_compare_exchange_resultUint4_ _e338 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_scalar, 1u, 2u);
-    _atomic_compare_exchange_resultSint4_ _e344 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_arr.inner[1], 1, 2);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e308 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e313 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_arr.inner[1], 1, 2);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e318 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e324 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_arr.inner[1], 1, 2);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e328 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e333 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_arr.inner[1], 1, 2);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e338 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_result_u003c_Sint_u002c_4_u003e _e344 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_arr.inner[1], 1, 2);
     return;
 }
diff --git a/naga/tests/out/msl/wgsl-overrides-atomicCompareExchangeWeak.msl b/naga/tests/out/msl/wgsl-overrides-atomicCompareExchangeWeak.msl
index 48f09ee6de2..a15984d9c16 100644
--- a/naga/tests/out/msl/wgsl-overrides-atomicCompareExchangeWeak.msl
+++ b/naga/tests/out/msl/wgsl-overrides-atomicCompareExchangeWeak.msl
@@ -4,14 +4,14 @@
 
 using metal::uint;
 
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e {
     uint old_value;
     bool exchanged;
     char _pad2[3];
 };
 
 template <typename A>
-_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
+_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
     device A *atomic_ptr,
     uint cmp,
     uint v
@@ -20,10 +20,10 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
         atomic_ptr, &cmp, v,
         metal::memory_order_relaxed, metal::memory_order_relaxed
     );
-    return _atomic_compare_exchange_resultUint4_{cmp, swapped};
+    return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
 }
 template <typename A>
-_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
+_atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e naga_atomic_compare_exchange_weak_explicit(
     threadgroup A *atomic_ptr,
     uint cmp,
     uint v
@@ -32,7 +32,7 @@ _atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit
         atomic_ptr, &cmp, v,
         metal::memory_order_relaxed, metal::memory_order_relaxed
     );
-    return _atomic_compare_exchange_resultUint4_{cmp, swapped};
+    return _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e{cmp, swapped};
 }
 constant int o = 2;
 
@@ -44,6 +44,6 @@ kernel void f(
         metal::atomic_store_explicit(&a, 0, metal::memory_order_relaxed);
     }
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
-    _atomic_compare_exchange_resultUint4_ _e5 = naga_atomic_compare_exchange_weak_explicit(&a, 2u, 1u);
+    _atomic_compare_exchange_result_u003c_Uint_u002c_4_u003e _e5 = naga_atomic_compare_exchange_weak_explicit(&a, 2u, 1u);
     return;
 }
diff --git a/naga/tests/out/spv/wgsl-7995-unicode-idents.spvasm b/naga/tests/out/spv/wgsl-7995-unicode-idents.spvasm
new file mode 100644
index 00000000000..a051e436ab5
--- /dev/null
+++ b/naga/tests/out/spv/wgsl-7995-unicode-idents.spvasm
@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 24
+OpCapability Shader
+OpExtension "SPV_KHR_storage_buffer_storage_class"
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %19 "main"
+OpExecutionMode %19 LocalSize 1 1 1
+OpDecorate %4 NonWritable
+OpDecorate %4 DescriptorSet 0
+OpDecorate %4 Binding 0
+OpDecorate %5 Block
+OpMemberDecorate %5 0 Offset 0
+%2 = OpTypeVoid
+%3 = OpTypeFloat 32
+%5 = OpTypeStruct %3
+%6 = OpTypePointer StorageBuffer %5
+%4 = OpVariable  %6  StorageBuffer
+%9 = OpTypeFunction %3
+%10 = OpTypePointer StorageBuffer %3
+%12 = OpTypeInt 32 0
+%11 = OpConstant  %12  0
+%14 = OpConstant  %3  9001
+%20 = OpTypeFunction %2
+%8 = OpFunction  %3  None %9
+%7 = OpLabel
+%13 = OpAccessChain  %10  %4 %11
+OpBranch %15
+%15 = OpLabel
+%16 = OpLoad  %3  %13
+%17 = OpFAdd  %3  %16 %14
+OpReturnValue %17
+OpFunctionEnd
+%19 = OpFunction  %2  None %20
+%18 = OpLabel
+%21 = OpAccessChain  %10  %4 %11
+OpBranch %22
+%22 = OpLabel
+%23 = OpFunctionCall  %3  %8
+OpReturn
+OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/wgsl/spv-atomic_compare_exchange.wgsl b/naga/tests/out/wgsl/spv-atomic_compare_exchange.wgsl
index 21456b064c9..88476c5057d 100644
--- a/naga/tests/out/wgsl/spv-atomic_compare_exchange.wgsl
+++ b/naga/tests/out/wgsl/spv-atomic_compare_exchange.wgsl
@@ -61,6 +61,6 @@ fn function() {
 }
 
 @compute @workgroup_size(32, 1, 1) 
-fn stagetest_atomic_compare_exchange() {
+fn stage_u003a_u003a_test_atomic_compare_exchange() {
     function();
 }
diff --git a/naga/tests/out/wgsl/spv-atomic_exchange.wgsl b/naga/tests/out/wgsl/spv-atomic_exchange.wgsl
index f954a850795..e17ac9e5879 100644
--- a/naga/tests/out/wgsl/spv-atomic_exchange.wgsl
+++ b/naga/tests/out/wgsl/spv-atomic_exchange.wgsl
@@ -75,6 +75,6 @@ fn function() {
 }
 
 @compute @workgroup_size(32, 1, 1) 
-fn stagetest_atomic_exchange() {
+fn stage_u003a_u003a_test_atomic_exchange() {
     function();
 }
diff --git a/naga/tests/out/wgsl/spv-atomic_i_add_sub.wgsl b/naga/tests/out/wgsl/spv-atomic_i_add_sub.wgsl
index d7798bd411d..b783267ad37 100644
--- a/naga/tests/out/wgsl/spv-atomic_i_add_sub.wgsl
+++ b/naga/tests/out/wgsl/spv-atomic_i_add_sub.wgsl
@@ -21,6 +21,6 @@ fn function() {
 }
 
 @compute @workgroup_size(32, 1, 1) 
-fn stagetest_atomic_i_add_sub() {
+fn stage_u003a_u003a_test_atomic_i_add_sub() {
     function();
 }
diff --git a/naga/tests/out/wgsl/spv-atomic_i_decrement.wgsl b/naga/tests/out/wgsl/spv-atomic_i_decrement.wgsl
index 190256e1c97..e0bf0a2488a 100644
--- a/naga/tests/out/wgsl/spv-atomic_i_decrement.wgsl
+++ b/naga/tests/out/wgsl/spv-atomic_i_decrement.wgsl
@@ -32,6 +32,6 @@ fn function() {
 }
 
 @compute @workgroup_size(32, 1, 1) 
-fn stagetest_atomic_i_decrement() {
+fn stage_u003a_u003a_test_atomic_i_decrement() {
     function();
 }
diff --git a/naga/tests/out/wgsl/spv-atomic_i_increment.wgsl b/naga/tests/out/wgsl/spv-atomic_i_increment.wgsl
index f7282aa3e64..830bf98d971 100644
--- a/naga/tests/out/wgsl/spv-atomic_i_increment.wgsl
+++ b/naga/tests/out/wgsl/spv-atomic_i_increment.wgsl
@@ -37,6 +37,6 @@ fn function() {
 }
 
 @compute @workgroup_size(32, 1, 1) 
-fn stagetest_atomic_i_increment() {
+fn stage_u003a_u003a_test_atomic_i_increment() {
     function();
 }
diff --git a/naga/tests/out/wgsl/spv-atomic_load_and_store.wgsl b/naga/tests/out/wgsl/spv-atomic_load_and_store.wgsl
index 9bfaa4ef55e..d33f92171f8 100644
--- a/naga/tests/out/wgsl/spv-atomic_load_and_store.wgsl
+++ b/naga/tests/out/wgsl/spv-atomic_load_and_store.wgsl
@@ -67,6 +67,6 @@ fn function() {
 }
 
 @compute @workgroup_size(32, 1, 1) 
-fn stagetest_atomic_load_and_store() {
+fn stage_u003a_u003a_test_atomic_load_and_store() {
     function();
 }
diff --git a/naga/tests/out/wgsl/spv-builtin-accessed-outside-entrypoint.wgsl b/naga/tests/out/wgsl/spv-builtin-accessed-outside-entrypoint.wgsl
index 8e1b885bf72..2941b766365 100644
--- a/naga/tests/out/wgsl/spv-builtin-accessed-outside-entrypoint.wgsl
+++ b/naga/tests/out/wgsl/spv-builtin-accessed-outside-entrypoint.wgsl
@@ -8,7 +8,7 @@ struct gl_PerVertex {
 var<private> unnamed: gl_PerVertex = gl_PerVertex(vec4<f32>(0f, 0f, 0f, 1f), 1f, array<f32, 1>(), array<f32, 1>());
 var<private> gl_VertexIndex_1: i32;
 
-fn builtin_usage() {
+fn builtin_usage_u0028_() {
     let _e9 = gl_VertexIndex_1;
     let _e12 = gl_VertexIndex_1;
     unnamed.gl_Position = vec4<f32>(select(1f, -4f, (_e9 == 0i)), select(-1f, 4f, (_e12 == 2i)), 0f, 1f);
@@ -16,7 +16,7 @@ fn builtin_usage() {
 }
 
 fn main_1() {
-    builtin_usage();
+    builtin_usage_u0028_();
     return;
 }
 
diff --git a/naga/tests/out/wgsl/spv-do-while.wgsl b/naga/tests/out/wgsl/spv-do-while.wgsl
index 0b7e3afbcce..3a6d43d95dc 100644
--- a/naga/tests/out/wgsl/spv-do-while.wgsl
+++ b/naga/tests/out/wgsl/spv-do-while.wgsl
@@ -1,4 +1,4 @@
-fn fb1_(cond: ptr<function, bool>) {
+fn f_u0028_b1_u003b(cond: ptr<function, bool>) {
     loop {
         continue;
         continuing {
@@ -13,7 +13,7 @@ fn main_1() {
     var param: bool;
 
     param = false;
-    fb1_((&param));
+    f_u0028_b1_u003b((&param));
     return;
 }
 
diff --git a/naga/tests/out/wgsl/spv-fetch_depth.wgsl b/naga/tests/out/wgsl/spv-fetch_depth.wgsl
index 3e206bc2676..d6a7d5173e8 100644
--- a/naga/tests/out/wgsl/spv-fetch_depth.wgsl
+++ b/naga/tests/out/wgsl/spv-fetch_depth.wgsl
@@ -21,6 +21,6 @@ fn function() {
 }
 
 @compute @workgroup_size(32, 1, 1) 
-fn cullfetch_depth() {
+fn cull_u003a_u003a_fetch_depth() {
     function();
 }
diff --git a/naga/tests/out/wgsl/wgsl-7995-unicode-idents.wgsl b/naga/tests/out/wgsl/wgsl-7995-unicode-idents.wgsl
new file mode 100644
index 00000000000..9a93d697454
--- /dev/null
+++ b/naga/tests/out/wgsl/wgsl-7995-unicode-idents.wgsl
@@ -0,0 +1,14 @@
+@group(0) @binding(0) 
+var<storage> asdf: f32;
+
+fn compute() -> f32 {
+    let _e1 = asdf;
+    let u03b8_2_ = (_e1 + 9001f);
+    return u03b8_2_;
+}
+
+@compute @workgroup_size(1, 1, 1) 
+fn main() {
+    let _e0 = compute();
+    return;
+}