263481 2023-10-20 16:11:21 -0700 [WGSL] texture1d.read(coord, level): level must be zero 2023-10-30 04:04:08 -0700 1 1 1 Unclassified WebKit WebGPU WebKit Nightly Build Unspecified Unspecified RESOLVED FIXED InRadar P2 Normal --- 1 mwyrzykowski tzagallo tzagallo webkit-bot-watchers-bugzilla webkit-bug-importer oldest_to_newest 1986583 0 mwyrzykowski 2023-10-20 16:11:21 -0700 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); 1986584 1 mwyrzykowski 2023-10-20 16:13:48 -0700 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 1986585 2 webkit-bug-importer 2023-10-20 16:13:50 -0700 <rdar://problem/117283953> 1987978 3 tzagallo 2023-10-27 06:17:42 -0700 Pull request: https://github.com/WebKit/WebKit/pull/19638 1988443 4 ews-feeder 2023-10-30 04:04:06 -0700 Committed 269931@main (f6aecdc2aadb): <https://commits.webkit.org/269931@main> Reviewed commits have been landed. Closing PR #19638 and removing active labels.