Naga generates bad MSL for textureLoad access on 1d textures

[jimblandy]

Naga compiles the following WGSL input to Metal Shading Language that does not validate:

// Load a texel from a 1d texture.
[[group(0), binding(0)]]
var tex: texture_1d<u32>;

[[stage(compute), workgroup_size(16)]]
fn main(
    [[builtin(local_invocation_id)]] local_id: vec3<u32>,
) {
    var x = textureLoad(tex, i32(local_id.x), i32(local_id.y));
}

The generated MSL is:

// language: metal1.1
#include <metal_stdlib>
#include <simd/simd.h>


struct main_Input {
};
kernel void main_(
  metal::uint3 local_id [[thread_position_in_threadgroup]]
, metal::texture1d<uint, metal::access::sample> tex [[user(fake0)]]
) {
    metal::uint4 x;
    metal::uint4 _e6 = tex.read(metal::uint(static_cast<int>(local_id.x)), static_cast<int>(local_id.y));
    x = _e6;
    return;
}

The error messages are a bunch of complaints explaining why each overload of read doesn't apply, but the gist of it is, we would like this overload:

  METAL_FUNC vec<T, 4> read(uint coord, uint lod = 0) const thread METAL_CONST_ARG(lod) METAL_ZERO_ARG(lod)

but we don't get it because the level-of-detail argument is not a compile-time constant.

1D textures can't have mipmaps. At the WGSL level, this means that textureNumLevels always returns 1 on 1D textures, and any value of the textureLoad's level argument other than zero is out of bounds. But at the MSL level, the level argument to texture1d::read must be a constexpr equal to zero. So when Naga emits the expression static_cast<int>(local_id.y) for that argument, Metal validation fails.

[jimblandy]

The same problem arises with WGSL's textureDimensions: it accepts a level argument, but MSL requires that it be a constexpr equal to zero.