[Webkit-unassigned] [Bug 263481] New: [WGSL] texture1d.read(coord, level): level is not casted to uint in generated metal

[bugzilla-daemon at webkit.org]

https://bugs.webkit.org/show_bug.cgi?id=263481

            Bug ID: 263481
           Summary: [WGSL] texture1d.read(coord, level): level is not
                    casted to uint in generated metal
           Product: WebKit
           Version: WebKit Nightly Build
          Hardware: Unspecified
                OS: Unspecified
            Status: NEW
          Severity: Normal
          Priority: P2
         Component: WebGPU
          Assignee: webkit-unassigned at lists.webkit.org
          Reporter: mwyrzykowski at apple.com

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);

-- 
You are receiving this mail because:
You are the assignee for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.webkit.org/pipermail/webkit-unassigned/attachments/20231020/8463c96b/attachment-0001.htm>