diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs
index 8a8dea664c..eb2c2ed025 100644
--- a/src/back/glsl/mod.rs
+++ b/src/back/glsl/mod.rs
@@ -1129,7 +1129,8 @@ impl<'a, W: Write> Writer<'a, W> {
         let ty_name = &self.names[&NameKey::Type(global.ty)];
         let block_name = format!(
             "{}_block_{}{:?}",
-            ty_name,
+            // avoid double underscores as they are reserved in GLSL
+            ty_name.trim_end_matches('_'),
             self.block_id.generate(),
             self.entry_point.stage,
         );
diff --git a/src/proc/namer.rs b/src/proc/namer.rs
index 2f262d7d5f..8afacb593d 100644
--- a/src/proc/namer.rs
+++ b/src/proc/namer.rs
@@ -36,6 +36,7 @@ impl Namer {
     /// - Drop leading digits.
     /// - Retain only alphanumeric and `_` characters.
     /// - Avoid prefixes in [`Namer::reserved_prefixes`].
+    /// - Replace consecutive `_` characters with a single `_` character.
     ///
     /// The return value is a valid identifier prefix in all of Naga's output languages,
     /// and it never ends with a `SEPARATOR` character.
@@ -46,6 +47,7 @@ impl Namer {
             .trim_end_matches(SEPARATOR);
 
         let base = if !string.is_empty()
+            && !string.contains("__")
             && string
                 .chars()
                 .all(|c: char| c.is_ascii_alphanumeric() || c == '_')
@@ -55,7 +57,13 @@ impl Namer {
             let mut filtered = string
                 .chars()
                 .filter(|&c| c.is_ascii_alphanumeric() || c == '_')
-                .collect::<String>();
+                .fold(String::new(), |mut s, c| {
+                    if s.ends_with('_') && c == '_' {
+                        return s;
+                    }
+                    s.push(c);
+                    s
+                });
             let stripped_len = filtered.trim_end_matches(SEPARATOR).len();
             filtered.truncate(stripped_len);
             if filtered.is_empty() {
@@ -268,4 +276,6 @@ fn test() {
     assert_eq!(namer.call("x"), "x");
     assert_eq!(namer.call("x"), "x_1");
     assert_eq!(namer.call("x1"), "x1_");
+    assert_eq!(namer.call("__x"), "_x");
+    assert_eq!(namer.call("1___x"), "_x_1");
 }
diff --git a/tests/out/glsl/math-functions.main.Fragment.glsl b/tests/out/glsl/math-functions.main.Fragment.glsl
index 37c072a6fa..ed81535ab5 100644
--- a/tests/out/glsl/math-functions.main.Fragment.glsl
+++ b/tests/out/glsl/math-functions.main.Fragment.glsl
@@ -3,55 +3,55 @@
 precision highp float;
 precision highp int;
 
-struct __modf_result_f32_ {
+struct _modf_result_f32_ {
     float fract_;
     float whole;
 };
-struct __modf_result_vec2_f32_ {
+struct _modf_result_vec2_f32_ {
     vec2 fract_;
     vec2 whole;
 };
-struct __modf_result_vec4_f32_ {
+struct _modf_result_vec4_f32_ {
     vec4 fract_;
     vec4 whole;
 };
-struct __frexp_result_f32_ {
+struct _frexp_result_f32_ {
     float fract_;
     int exp_;
 };
-struct __frexp_result_vec4_f32_ {
+struct _frexp_result_vec4_f32_ {
     vec4 fract_;
     ivec4 exp_;
 };
 
-__modf_result_f32_ naga_modf(float arg) {
+_modf_result_f32_ naga_modf(float arg) {
     float other;
     float fract = modf(arg, other);
-    return __modf_result_f32_(fract, other);
+    return _modf_result_f32_(fract, other);
 }
 
-__modf_result_vec2_f32_ naga_modf(vec2 arg) {
+_modf_result_vec2_f32_ naga_modf(vec2 arg) {
     vec2 other;
     vec2 fract = modf(arg, other);
-    return __modf_result_vec2_f32_(fract, other);
+    return _modf_result_vec2_f32_(fract, other);
 }
 
-__modf_result_vec4_f32_ naga_modf(vec4 arg) {
+_modf_result_vec4_f32_ naga_modf(vec4 arg) {
     vec4 other;
     vec4 fract = modf(arg, other);
-    return __modf_result_vec4_f32_(fract, other);
+    return _modf_result_vec4_f32_(fract, other);
 }
 
-__frexp_result_f32_ naga_frexp(float arg) {
+_frexp_result_f32_ naga_frexp(float arg) {
     int other;
     float fract = frexp(arg, other);
-    return __frexp_result_f32_(fract, other);
+    return _frexp_result_f32_(fract, other);
 }
 
-__frexp_result_vec4_f32_ naga_frexp(vec4 arg) {
+_frexp_result_vec4_f32_ naga_frexp(vec4 arg) {
     ivec4 other;
     vec4 fract = frexp(arg, other);
-    return __frexp_result_vec4_f32_(fract, other);
+    return _frexp_result_vec4_f32_(fract, other);
 }
 
 void main() {
@@ -90,13 +90,13 @@ void main() {
     uvec2 clz_d = uvec2(ivec2(31) - findMSB(uvec2(1u)));
     float lde_a = ldexp(1.0, 2);
     vec2 lde_b = ldexp(vec2(1.0, 2.0), ivec2(3, 4));
-    __modf_result_f32_ modf_a = naga_modf(1.5);
+    _modf_result_f32_ modf_a = naga_modf(1.5);
     float modf_b = naga_modf(1.5).fract_;
     float modf_c = naga_modf(1.5).whole;
-    __modf_result_vec2_f32_ modf_d = naga_modf(vec2(1.5, 1.5));
+    _modf_result_vec2_f32_ modf_d = naga_modf(vec2(1.5, 1.5));
     float modf_e = naga_modf(vec4(1.5, 1.5, 1.5, 1.5)).whole.x;
     float modf_f = naga_modf(vec2(1.5, 1.5)).fract_.y;
-    __frexp_result_f32_ frexp_a = naga_frexp(1.5);
+    _frexp_result_f32_ frexp_a = naga_frexp(1.5);
     float frexp_b = naga_frexp(1.5).fract_;
     int frexp_c = naga_frexp(1.5).exp_;
     int frexp_d = naga_frexp(vec4(1.5, 1.5, 1.5, 1.5)).exp_.x;
diff --git a/tests/out/glsl/padding.vertex.Vertex.glsl b/tests/out/glsl/padding.vertex.Vertex.glsl
index 1810cffc69..82d8d31fa9 100644
--- a/tests/out/glsl/padding.vertex.Vertex.glsl
+++ b/tests/out/glsl/padding.vertex.Vertex.glsl
@@ -20,9 +20,9 @@ struct Test3_ {
 };
 uniform Test_block_0Vertex { Test _group_0_binding_0_vs; };
 
-uniform Test2__block_1Vertex { Test2_ _group_0_binding_1_vs; };
+uniform Test2_block_1Vertex { Test2_ _group_0_binding_1_vs; };
 
-uniform Test3__block_2Vertex { Test3_ _group_0_binding_2_vs; };
+uniform Test3_block_2Vertex { Test3_ _group_0_binding_2_vs; };
 
 
 void main() {
diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl
index fc5cadb65e..53d3acf0c1 100644
--- a/tests/out/hlsl/math-functions.hlsl
+++ b/tests/out/hlsl/math-functions.hlsl
@@ -1,63 +1,63 @@
-struct __modf_result_f32_ {
+struct _modf_result_f32_ {
     float fract;
     float whole;
 };
 
-struct __modf_result_vec2_f32_ {
+struct _modf_result_vec2_f32_ {
     float2 fract;
     float2 whole;
 };
 
-struct __modf_result_vec4_f32_ {
+struct _modf_result_vec4_f32_ {
     float4 fract;
     float4 whole;
 };
 
-struct __frexp_result_f32_ {
+struct _frexp_result_f32_ {
     float fract;
     int exp_;
 };
 
-struct __frexp_result_vec4_f32_ {
+struct _frexp_result_vec4_f32_ {
     float4 fract;
     int4 exp_;
 };
 
-__modf_result_f32_ naga_modf(float arg) {
+_modf_result_f32_ naga_modf(float arg) {
     float other;
-    __modf_result_f32_ result;
+    _modf_result_f32_ result;
     result.fract = modf(arg, other);
     result.whole = other;
     return result;
 }
 
-__modf_result_vec2_f32_ naga_modf(float2 arg) {
+_modf_result_vec2_f32_ naga_modf(float2 arg) {
     float2 other;
-    __modf_result_vec2_f32_ result;
+    _modf_result_vec2_f32_ result;
     result.fract = modf(arg, other);
     result.whole = other;
     return result;
 }
 
-__modf_result_vec4_f32_ naga_modf(float4 arg) {
+_modf_result_vec4_f32_ naga_modf(float4 arg) {
     float4 other;
-    __modf_result_vec4_f32_ result;
+    _modf_result_vec4_f32_ result;
     result.fract = modf(arg, other);
     result.whole = other;
     return result;
 }
 
-__frexp_result_f32_ naga_frexp(float arg) {
+_frexp_result_f32_ naga_frexp(float arg) {
     float other;
-    __frexp_result_f32_ result;
+    _frexp_result_f32_ result;
     result.fract = sign(arg) * frexp(arg, other);
     result.exp_ = other;
     return result;
 }
 
-__frexp_result_vec4_f32_ naga_frexp(float4 arg) {
+_frexp_result_vec4_f32_ naga_frexp(float4 arg) {
     float4 other;
-    __frexp_result_vec4_f32_ result;
+    _frexp_result_vec4_f32_ result;
     result.fract = sign(arg) * frexp(arg, other);
     result.exp_ = other;
     return result;
@@ -100,13 +100,13 @@ void main()
     uint2 clz_d = ((31u).xx - firstbithigh((1u).xx));
     float lde_a = ldexp(1.0, 2);
     float2 lde_b = ldexp(float2(1.0, 2.0), int2(3, 4));
-    __modf_result_f32_ modf_a = naga_modf(1.5);
+    _modf_result_f32_ modf_a = naga_modf(1.5);
     float modf_b = naga_modf(1.5).fract;
     float modf_c = naga_modf(1.5).whole;
-    __modf_result_vec2_f32_ modf_d = naga_modf(float2(1.5, 1.5));
+    _modf_result_vec2_f32_ modf_d = naga_modf(float2(1.5, 1.5));
     float modf_e = naga_modf(float4(1.5, 1.5, 1.5, 1.5)).whole.x;
     float modf_f = naga_modf(float2(1.5, 1.5)).fract.y;
-    __frexp_result_f32_ frexp_a = naga_frexp(1.5);
+    _frexp_result_f32_ frexp_a = naga_frexp(1.5);
     float frexp_b = naga_frexp(1.5).fract;
     int frexp_c = naga_frexp(1.5).exp_;
     int frexp_d = naga_frexp(float4(1.5, 1.5, 1.5, 1.5)).exp_.x;
diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl
index dccb90ad6c..d0ae8c5047 100644
--- a/tests/out/msl/math-functions.msl
+++ b/tests/out/msl/math-functions.msl
@@ -4,55 +4,55 @@
 
 using metal::uint;
 
-struct __modf_result_f32_ {
+struct _modf_result_f32_ {
     float fract;
     float whole;
 };
-struct __modf_result_vec2_f32_ {
+struct _modf_result_vec2_f32_ {
     metal::float2 fract;
     metal::float2 whole;
 };
-struct __modf_result_vec4_f32_ {
+struct _modf_result_vec4_f32_ {
     metal::float4 fract;
     metal::float4 whole;
 };
-struct __frexp_result_f32_ {
+struct _frexp_result_f32_ {
     float fract;
     int exp;
 };
-struct __frexp_result_vec4_f32_ {
+struct _frexp_result_vec4_f32_ {
     metal::float4 fract;
     metal::int4 exp;
 };
 
-__modf_result_f32_ naga_modf(float arg) {
+_modf_result_f32_ naga_modf(float arg) {
     float other;
     float fract = metal::modf(arg, other);
-    return __modf_result_f32_{ fract, other };
+    return _modf_result_f32_{ fract, other };
 }
 
-__modf_result_vec2_f32_ naga_modf(metal::float2 arg) {
+_modf_result_vec2_f32_ naga_modf(metal::float2 arg) {
     metal::float2 other;
     metal::float2 fract = metal::modf(arg, other);
-    return __modf_result_vec2_f32_{ fract, other };
+    return _modf_result_vec2_f32_{ fract, other };
 }
 
-__modf_result_vec4_f32_ naga_modf(metal::float4 arg) {
+_modf_result_vec4_f32_ naga_modf(metal::float4 arg) {
     metal::float4 other;
     metal::float4 fract = metal::modf(arg, other);
-    return __modf_result_vec4_f32_{ fract, other };
+    return _modf_result_vec4_f32_{ fract, other };
 }
 
-__frexp_result_f32_ naga_frexp(float arg) {
+_frexp_result_f32_ naga_frexp(float arg) {
     int other;
     float fract = metal::frexp(arg, other);
-    return __frexp_result_f32_{ fract, other };
+    return _frexp_result_f32_{ fract, other };
 }
 
-__frexp_result_vec4_f32_ naga_frexp(metal::float4 arg) {
+_frexp_result_vec4_f32_ naga_frexp(metal::float4 arg) {
     int4 other;
     metal::float4 fract = metal::frexp(arg, other);
-    return __frexp_result_vec4_f32_{ fract, other };
+    return _frexp_result_vec4_f32_{ fract, other };
 }
 
 fragment void main_(
@@ -95,13 +95,13 @@ fragment void main_(
     metal::uint2 clz_d = metal::clz(metal::uint2(1u));
     float lde_a = metal::ldexp(1.0, 2);
     metal::float2 lde_b = metal::ldexp(metal::float2(1.0, 2.0), metal::int2(3, 4));
-    __modf_result_f32_ modf_a = naga_modf(1.5);
+    _modf_result_f32_ modf_a = naga_modf(1.5);
     float modf_b = naga_modf(1.5).fract;
     float modf_c = naga_modf(1.5).whole;
-    __modf_result_vec2_f32_ modf_d = naga_modf(metal::float2(1.5, 1.5));
+    _modf_result_vec2_f32_ modf_d = naga_modf(metal::float2(1.5, 1.5));
     float modf_e = naga_modf(metal::float4(1.5, 1.5, 1.5, 1.5)).whole.x;
     float modf_f = naga_modf(metal::float2(1.5, 1.5)).fract.y;
-    __frexp_result_f32_ frexp_a = naga_frexp(1.5);
+    _frexp_result_f32_ frexp_a = naga_frexp(1.5);
     float frexp_b = naga_frexp(1.5).fract;
     int frexp_c = naga_frexp(1.5).exp;
     int frexp_d = naga_frexp(metal::float4(1.5, 1.5, 1.5, 1.5)).exp.x;