wgpu: Mistranslation of `loop` with `continuing` block

Hi,

Converting this wgsl:

var<private> global: vec3<u32>;
@group(0) @binding(2) 
var global_1: texture_storage_2d<rgba16float,read_write>;

fn function() {
    var phi_79_: i32;
    var phi_89_: i32;

    let _e17 = global;
    phi_79_ = 0;
    loop {
        let _e22 = phi_79_;
        let _e23 = (_e22 < 8);
        if _e23 {
            phi_89_ = (_e22 + 1);
        } else {
            textureStore(global_1, vec2<u32>(_e17.x, _e17.y), vec4<f32>(0.10000000149011612, 0.20000000298023224, 0.30000001192092896, 1.0));
            phi_89_ = 0;
        }
        let _e26 = phi_89_;
        continue;
        continuing {
            _ = !(select(false, true, _e23));
            phi_79_ = _e26;
            break if !(select(false, true, _e23));
        }
    }
    return;
}

@compute @workgroup_size(8, 8, 1) 
fn main(@builtin(global_invocation_id) param: vec3<u32>) {
    global = param;
    function();
}

… yields an invalid msl:

// language: metal2.0
#include <metal_stdlib>
#include <simd/simd.h>

using metal::uint;

constant metal::float4 const_type_9_ = {0.10000000149011612, 0.20000000298023224, 0.30000001192092896, 1.0};

void function(
    thread metal::uint3& global,
    metal::texture2d<float, metal::access::read_write> global_1
) {
    int phi_79_ = {};
    int phi_89_ = {};
    metal::uint3 _e17 = global;
    phi_79_ = 0;
    bool loop_init = true;
    while(true) {
        if (!loop_init) {
            bool unnamed = !(((phi_79_ < 8) ? true : false));
            phi_79_ = phi_89_;
            if (!(((phi_79_ < 8) ? true : false))) {
                break;
            }
        }
        loop_init = false;
        int _e22 = phi_79_;
        bool _e23 = _e22 < 8;
        if (_e23) {
            phi_89_ = _e22 + 1;
        } else {
            global_1.write(const_type_9_, metal::uint2(metal::uint2(_e17.x, _e17.y)));
            phi_89_ = 0;
        }
        int _e26 = phi_89_;
        continue;
    }
    return;
}

struct main_Input {
};
kernel void main_(
  metal::uint3 param [[thread_position_in_grid]]
, metal::texture2d<float, metal::access::read_write> global_1 [[user(fake0)]]
) {
    metal::uint3 global = {};
    global = param;
    function(global, global_1);
}

… with a pretty subtle mistake related to the if (!loop_init) part:

if (!(((phi_79_ < 8) ? true : false))) {

Note that it uses phi_79_ with value after the assignment, instead of using the value from-before the assignment, which causes this shader to be effectively an infinite loop (later optimized out by Metal Shader Compiler into a no-op, apparently).

I’ve narrowed down the issue to this place:

https://github.com/gfx-rs/naga/blob/b7f4006e46313da063f8f2f930230767c9740239/src/back/msl/writer.rs#L2960

… which I’m not sure what is supposed to do, but just commenting-out self.named_expressions.remove(&handle); generates correct code 👀