From aa5b1d8f0cd98dc64bd93499aaef8a07e57eadc2 Mon Sep 17 00:00:00 2001
From: Gregg Tavares <github@greggman.com>
Date: Thu, 15 Aug 2024 10:27:50 -0700
Subject: [PATCH] Fix for texture_utils.ts

The issue is on Chrome and Firefox on Intel Mac,
when sampling between 2 mip levels using `textureSampleLevel`
in a compute shader, the weights used for mixing are
very unexpected. The same issue doesn't happen in Safari TP
so there is probably a fix.

For now though, the same issue doesn't happen when using
a fragment shader. So, switched to using a fragment
shader to look up these weights. This is more appropriate
for the current tests because the tests are running in
fragment shaders.

Will add an issue to test all stages.
---
 .../expression/call/builtin/texture_utils.ts  | 40 +++++++++++++++----
 1 file changed, 32 insertions(+), 8 deletions(-)

diff --git a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts
index 38e94675f3e0..7e8b1e168482 100644
--- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts
+++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts
@@ -134,16 +134,32 @@ async function initMipGradientValuesForDevice(t: GPUTest) {
         @group(0) @binding(1) var smp: sampler;
         @group(0) @binding(2) var<storage, read_write> result: array<f32>;
 
-        @compute @workgroup_size(1) fn cs(@builtin(global_invocation_id) id: vec3u) {
-          let mipLevel = f32(id.x) / ${kMipGradientSteps};
-          result[id.x] = textureSampleLevel(tex, smp, vec2f(0.5), mipLevel).r;
+        @vertex fn vs(@builtin(vertex_index) vNdx: u32) -> @builtin(position) vec4f {
+          let pos = array(
+            vec2f(-1,  3),
+            vec2f( 3, -1),
+            vec2f(-1, -1),
+          );
+          return vec4f(pos[vNdx], 0, 1);
+        }
+        @fragment fn fs(@builtin(position) pos: vec4f) -> @location(0) vec4f {
+          let mipLevel = floor(pos.x) / ${kMipGradientSteps};
+          result[u32(pos.x)] = textureSampleLevel(tex, smp, vec2f(0.5), mipLevel).r;
+          return vec4f(0);
         }
       `,
     });
 
-    const pipeline = device.createComputePipeline({
+    const pipeline = device.createRenderPipeline({
       layout: 'auto',
-      compute: { module },
+      vertex: { module },
+      fragment: { module, targets: [{ format: 'rgba8unorm' }] },
+    });
+
+    const target = t.createTextureTracked({
+      size: [kMipGradientSteps + 1, 1, 1],
+      format: 'rgba8unorm',
+      usage: GPUTextureUsage.RENDER_ATTACHMENT,
     });
 
     const texture = t.createTextureTracked({
@@ -186,10 +202,18 @@ async function initMipGradientValuesForDevice(t: GPUTest) {
     });
 
     const encoder = device.createCommandEncoder();
-    const pass = encoder.beginComputePass();
+    const pass = encoder.beginRenderPass({
+      colorAttachments: [
+        {
+          view: target.createView(),
+          loadOp: 'clear',
+          storeOp: 'store',
+        },
+      ],
+    });
     pass.setPipeline(pipeline);
     pass.setBindGroup(0, bindGroup);
-    pass.dispatchWorkgroups(kMipGradientSteps + 1);
+    pass.draw(3);
     pass.end();
     encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, resultBuffer.size);
     device.queue.submit([encoder.finish()]);
@@ -215,7 +239,7 @@ async function initMipGradientValuesForDevice(t: GPUTest) {
     //                 standard
     // step  mipLevel    gpu        AMD
     // ----  --------  --------  ----------
-    //  0:   0         0           1
+    //  0:   0         0           0
     //  1:   0.0625    0.0625      0
     //  2:   0.125     0.125       0.03125
     //  3:   0.1875    0.1875      0.109375