Bug 263481 – [WGSL] texture1d.read(coord, level): level must be zero

• Status: RESOLVED
• Resolution: FIXED
• Component: WebGPU

[Mike Wyrzykowski]

The following WGSL:
            struct Constants {
              level : i32
            };

            @group(0) @binding(0) var<uniform> constants : Constants;
            @group(0) @binding(1) var myTexture : texture_1d<f32>;

            struct Result {
              values : array<f32>
            };
            @group(0) @binding(3) var<storage, read_write> result : Result;

            @compute @workgroup_size(1)
            fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
              let flatIndex : u32 = 1u * (
                2u * 1u * GlobalInvocationID.z +
                2u * GlobalInvocationID.y +
                GlobalInvocationID.x
              );
              let texel : vec4<f32> = textureLoad(
                myTexture, i32(GlobalInvocationID.x), constants.level);

              for (var i : u32 = 0u; i < 1u; i = i + 1u) {
                result.values[flatIndex + i] = texel.r;
              }
            }

produces this Metal code, but doesn't compile because "global0.level" in the call to read is expected to be a uint, but it is an int.

#include <metal_stdlib>

using namespace metal;

template<typename T>
T __pack(T unpacked)
{
    return unpacked;
}

template<typename T>
T __unpack(T packed)
{
    return packed;
}

struct type0 {
    int level;
};

struct type1 {
    array<float, 1> values;
};

struct __ArgumentBufferT_0 {
    const constant type0& global0 [[id(0)]];
    texture1d<float, access::sample> global1 [[id(1)]];
    device type1& global2 [[id(3)]];
};

[[kernel]]  void function0(vec<unsigned, 3> parameter0 [[thread_position_in_grid]], constant __ArgumentBufferT_0& __ArgumentBuffer_0 [[buffer(0)]])
{
    device type1& global2 = __ArgumentBuffer_0.global2;
    const constant type0& global0 = __ArgumentBuffer_0.global0;
    texture1d<float, access::sample> global1 = __ArgumentBuffer_0.global1;
    unsigned local0 = (1u * (((2u * parameter0.z) + (2u * parameter0.y)) + parameter0.x));
    vec<float, 4> local1 = global1.read(uint(int(parameter0.x)), global0.level);
    for (unsigned local2 = 0u; (local2 < 1u); local2 = (local2 + 1u)) {
        global2.values[(local0 + local2)] = local1.r;
    }
}


MSL compilation error: Error Domain=MTLLibraryErrorDomain Code=3 "program_source:37:36: error: no matching member function for call to 'read'
    vec<float, 4> local1 = global1.read(uint(int(parameter0.x)), global0.level);

[Mike Wyrzykowski]

This can be reproduced by running https://gpuweb.github.io/cts/standalone/?q=webgpu:api,operation,resource_init,texture_zero:uninitialized_texture_is_zero:dimension=%221d%22;readMethod=%22Sample%22;format=%22r8unorm%22

[Radar WebKit Bug Importer]

<rdar://problem/117283953>

[Tadeu Zagallo]

Pull request: https://github.com/WebKit/WebKit/pull/19638

[EWS]

Committed 269931@main (f6aecdc2aadb): <https://commits.webkit.org/269931@main>

Reviewed commits have been landed. Closing PR #19638 and removing active labels.