fix(namer): escape, rather than strip, non-ASCII ident. characters

CHANGELOG.md

@@ -231,6 +231,7 @@ By @cwfitzgerald in [#8162](https://github.com/gfx-rs/wgpu/pull/8162).
 
 - [wgsl-in] Allow a trailing comma in `@blend_src(…)` attributes. By @ErichDonGubler in [#8137](https://github.com/gfx-rs/wgpu/pull/8137).
 - [wgsl-in] Allow a trailing comma in the list of `case` values inside a `switch`. By @reima in [#8165](https://github.com/gfx-rs/wgpu/pull/8165).
+- Escape, rather than strip, identifiers with Unicode. By @ErichDonGubler in [7995](https://github.com/gfx-rs/wgpu/pull/7995).
 
 ### Documentation
 

naga/src/proc/namer.rs

@@ -118,20 +118,37 @@ 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 c = match c {
+                    // Make several common characters in C++-ish types become snake case
+                    // separators.
+                    ':' | '<' | '>' | ',' => '_',
+                    c => 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)
         };

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();
+}

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;
 }
 

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;
+}
+

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_Sint_4_ {
     int old_value;
     bool exchanged;
 };
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_Uint_4_ {
     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_Sint_4_ _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;

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_Sint_4_ {
     int old_value;
     bool exchanged;
 };
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_Uint_4_ {
     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_Uint_4_ _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;

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_Uint_4_ {
     uint old_value;
     bool exchanged;
 };
-struct _atomic_compare_exchange_resultSint4_ {
+struct _atomic_compare_exchange_result_Sint_4_ {
     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_Uint_4_ _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_Sint_4_ _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_Uint_4_ _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_Sint_4_ _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_Uint_4_ _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_Sint_4_ _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_Uint_4_ _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_Sint_4_ _e344; _e344.old_value = atomicCompSwap(workgroup_struct.atomic_arr[1], 1, 2);
     _e344.exchanged = (_e344.old_value == 1);
     return;
 }

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;
 }
 

naga/tests/out/hlsl/spv-fetch_depth.hlsl

@@ -19,7 +19,7 @@ void function()
 }
 
 [numthreads(32, 1, 1)]
-void cullfetch_depth()
+void cull_fetch_depth()
 {
     function();
 }

naga/tests/out/hlsl/spv-fetch_depth.ron

@@ -5,7 +5,7 @@
     ],
     compute:[
         (
-            entry_point:"cullfetch_depth",
+            entry_point:"cull_fetch_depth",
             target_profile:"cs_5_1",
         ),
     ],

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;
+}

naga/tests/out/hlsl/wgsl-7995-unicode-idents.ron

@@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"main",
+            target_profile:"cs_5_1",
+        ),
+    ],
+)

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_Sint_8_ {
     int64_t old_value;
     bool exchanged;
     int _end_pad_0;
 };
 
-struct _atomic_compare_exchange_resultUint8_ {
+struct _atomic_compare_exchange_result_Uint_8_ {
     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_Sint_8_ _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_Uint_8_ _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;

naga/tests/out/hlsl/wgsl-atomicCompareExchange.hlsl

@@ -1,9 +1,9 @@
-struct _atomic_compare_exchange_resultSint4_ {
+struct _atomic_compare_exchange_result_Sint_4_ {
     int old_value;
     bool exchanged;
 };
 
-struct _atomic_compare_exchange_resultUint4_ {
+struct _atomic_compare_exchange_result_Uint_4_ {
     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_Sint_4_ _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_Uint_4_ _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;

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_Uint_8_ {
     uint64_t old_value;
     bool exchanged;
     int _end_pad_0;
 };
 
-struct _atomic_compare_exchange_resultSint8_ {
+struct _atomic_compare_exchange_result_Sint_8_ {
     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_Uint_8_ _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_Sint_8_ _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_Uint_8_ _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_Sint_8_ _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_Uint_8_ _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_Sint_8_ _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_Uint_8_ _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_Sint_8_ _e328; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], 1L, 2L, _e328.old_value);
     _e328.exchanged = (_e328.old_value == 1L);
     return;
 }