Mistranslation of `loop` with `continuing` block
Author: Patryk27Created Jun 6, 2023Updated Sep 18, 2026
Labelstype: bugarea: naga back-endnagalang: GLSLlang: Metallang: HLSL
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:
... which I'm not sure what is supposed to do, but just commenting-out self.named_expressions.remove(&handle); generates correct code :eyes:
Source: gfx-rs/wgpu