[msl-out] Generated metal uses packed_float3 in an expression

[scoopr]

With this wgsl

struct ObjectOutputData {
    model_view_proj: mat4x4<f32>;
    inv_squared_scale: vec3<f32>;
    material_idx: u32;
};

var<private> global : ObjectOutputData;

fn foo()  {
    let zero3 = vec3<f32>(0.0, 0.0, 0.0);
    let zero4 = vec4<f32>(0.0, 0.0, 0.0, 0.0);
    let data : ObjectOutputData = global;
    let zero_mat3: mat3x3<f32> = mat3x3<f32>(zero3, zero3, zero3);
    let foo = zero_mat3 * data.inv_squared_scale;
}

the converted struct ends up as

struct ObjectOutputData {
    metal::float4x4 model_view_proj;
    packed_float3 inv_squared_scale;
    metal::uint material_idx;
};

and the expression uses it directly

metal::float3 foo_1 = zero_mat3_ * data.inv_squared_scale;

But I believe it should be cast first, like float3(data.inv_squared_scale)

A workaround can be to not have something with align(4) (like the u32 here) next to vec3. Just moving the u32 to first member works, or having some vec4 or mat4x4 in between also works. Then it won't be generated as packed_float3.

Originally found with rend3 trunk version of opaque.vert.cpu.wgsl, cc @cwfitzgerald

[teoxoy]

I'm not sure when this got fixed but it's no longer an issue as can be seen here: a335188

[scoopr]

Given the example at the top (fixed for the struct semicolon change), with naga master as of a3d968e

$ cargo run naga -- test.wgsl test.metal
$ xcrun metal test.metal

still yields exactly the same error, the struct field is of a packed type, but it is used in an expression that expects non-packed.

I believe in the example you gave, it is merely doing an assignment with the field, which I suppose is fine, but the arithmetic is not. In fact in the example above, assigning the field to a local variable first is a valid workaround.

But thanks for looking in to it!

[teoxoy]

Ah, I see what's going on. The test I added takes advantage of baking the load expression.

naga/src/back/mod.rs

Lines 139 to 142 in a3d968e

	// TODO: We need a better fix for named `Load` expressions
	// More info - https://github.com/gfx-rs/naga/pull/914
	// And https://github.com/gfx-rs/naga/issues/910
	crate::Expression::Load { .. } => 1,

However the issue with the example shared in this issue is that the load happens when the global is assigned to data.

struct Test {
    v: vec3<f32>,
    s: u32,
};

var<private> global: Test;

fn foo() {
    let data = global; // <- if you change this to a `var` you can see it will work (since the load will happen later)
    let foo = mat3x3<f32>() * data.v;
}

I'll mark the PR as a draft and see if there is a better solution tomorrow.