webgpu,api,operation,texture_view,write:format:*

src/webgpu/api/operation/texture_view/write.spec.ts

@@ -1,6 +1,9 @@
 export const description = `
 Test the result of writing textures through texture views with various options.
 
+Reads value from a shader array, writes the value via various write methods.
+Check the texture result with the expected texel view.
+
 All x= every possible view write method: {
   - storage write {fragment, compute}
   - render pass store
@@ -13,20 +16,358 @@ TODO: Write helper for this if not already available (see resource_init, buffer_
 `;
 
 import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { GPUTest } from '../../../gpu_test.js';
+import { unreachable } from '../../../../common/util/util.js';
+import {
+  kRegularTextureFormats,
+  kTextureFormatInfo,
+  RegularTextureFormat,
+} from '../../../format_info.js';
+import { GPUTest, TextureTestMixin } from '../../../gpu_test.js';
+import { kFullscreenQuadVertexShaderCode } from '../../../util/shader.js';
+import { TexelView } from '../../../util/texture/texel_view.js';
+
+export const g = makeTestGroup(TextureTestMixin(GPUTest));
+
+const kTextureViewWriteMethods = [
+  'storage-write-fragment',
+  'storage-write-compute',
+  'render-pass-store',
+  'render-pass-resolve',
+] as const;
+type TextureViewWriteMethod = (typeof kTextureViewWriteMethods)[number];
+
+// Src color values to read from a shader array.
+const kColorsFloat = [
+  { R: 1.0, G: 0.0, B: 0.0, A: 0.8 },
+  { R: 0.0, G: 1.0, B: 0.0, A: 0.7 },
+  { R: 0.0, G: 0.0, B: 0.0, A: 0.6 },
+  { R: 0.0, G: 0.0, B: 0.0, A: 0.5 },
+  { R: 1.0, G: 1.0, B: 1.0, A: 0.4 },
+  { R: 0.7, G: 0.0, B: 0.0, A: 0.3 },
+  { R: 0.0, G: 0.8, B: 0.0, A: 0.2 },
+  { R: 0.0, G: 0.0, B: 0.9, A: 0.1 },
+  { R: 0.1, G: 0.2, B: 0.0, A: 0.3 },
+  { R: 0.4, G: 0.3, B: 0.6, A: 0.8 },
+];
+
+function FloatToIntColor(c: number) {
+  return Math.floor(c * 100);
+}
+
+const kColorsInt = kColorsFloat.map(c => {
+  return {
+    R: FloatToIntColor(c.R),
+    G: FloatToIntColor(c.G),
+    B: FloatToIntColor(c.B),
+    A: FloatToIntColor(c.A),
+  };
+});
 
-export const g = makeTestGroup(GPUTest);
+const kTextureSize = 16;
+
+function writeTextureAndGetExpectedTexelView(
+  t: GPUTest,
+  method: TextureViewWriteMethod,
+  view: GPUTextureView,
+  format: RegularTextureFormat,
+  sampleCount: number
+) {
+  const info = kTextureFormatInfo[format];
+  const isFloatType = info.color.type === 'float' || info.color.type === 'unfilterable-float';
+  const kColors = isFloatType ? kColorsFloat : kColorsInt;
+  const expectedTexelView = TexelView.fromTexelsAsColors(
+    format,
+    coords => {
+      const pixelPos = coords.y * kTextureSize + coords.x;
+      return kColors[pixelPos % kColors.length];
+    },
+    { clampToFormatRange: true }
+  );
+  const vecType = isFloatType ? 'vec4f' : info.color.type === 'sint' ? 'vec4i' : 'vec4u';
+  const kColorArrayShaderString = `array<${vecType}, ${kColors.length}>(
+      ${kColors.map(t => `${vecType}(${t.R}, ${t.G}, ${t.B}, ${t.A}) `).join(',')}
+    )`;
+
+  switch (method) {
+    case 'storage-write-compute':
+      {
+        const pipeline = t.device.createComputePipeline({
+          layout: 'auto',
+          compute: {
+            module: t.device.createShaderModule({
+              code: `
+                @group(0) @binding(0) var dst: texture_storage_2d<${format}, write>;
+                @compute @workgroup_size(1, 1) fn main(
+                  @builtin(global_invocation_id) global_id: vec3<u32>,
+                ) {
+                  const src = ${kColorArrayShaderString};
+                  let coord = vec2u(global_id.xy);
+                  let idx = coord.x + coord.y * ${kTextureSize};
+                  textureStore(dst, coord, src[idx % ${kColors.length}]);
+                }`,
+            }),
+            entryPoint: 'main',
+          },
+        });
+        const commandEncoder = t.device.createCommandEncoder();
+        const pass = commandEncoder.beginComputePass();
+        pass.setPipeline(pipeline);
+        pass.setBindGroup(
+          0,
+          t.device.createBindGroup({
+            layout: pipeline.getBindGroupLayout(0),
+            entries: [
+              {
+                binding: 0,
+                resource: view,
+              },
+            ],
+          })
+        );
+        pass.dispatchWorkgroups(kTextureSize, kTextureSize);
+        pass.end();
+        t.device.queue.submit([commandEncoder.finish()]);
+      }
+      break;
+
+    case 'storage-write-fragment':
+      {
+        // Create a placeholder color attachment texture,
+        // The size of which equals that of format texture we are testing,
+        // so that we have the same number of fragments and texels.
+        const kPlaceholderTextureFormat = 'rgba8unorm';
+        const placeholderTexture = t.trackForCleanup(
+          t.device.createTexture({
+            format: kPlaceholderTextureFormat,
+            size: [kTextureSize, kTextureSize],
+            usage: GPUTextureUsage.RENDER_ATTACHMENT,
+          })
+        );
+
+        const pipeline = t.device.createRenderPipeline({
+          layout: 'auto',
+          vertex: {
+            module: t.device.createShaderModule({
+              code: kFullscreenQuadVertexShaderCode,
+            }),
+          },
+          fragment: {
+            module: t.device.createShaderModule({
+              code: `
+                @group(0) @binding(0) var dst: texture_storage_2d<${format}, write>;
+                @fragment fn main(
+                  @builtin(position) fragCoord: vec4<f32>,
+                ) {
+                  const src = ${kColorArrayShaderString};
+                  let coord = vec2u(fragCoord.xy);
+                  let idx = coord.x + coord.y * ${kTextureSize};
+                  textureStore(dst, coord, src[idx % ${kColors.length}]);
+                }`,
+            }),
+            // Set writeMask to 0 as the fragment shader has no output.
+            targets: [
+              {
+                format: kPlaceholderTextureFormat,
+                writeMask: 0,
+              },
+            ],
+          },
+        });
+        const commandEncoder = t.device.createCommandEncoder();
+        const pass = commandEncoder.beginRenderPass({
+          colorAttachments: [
+            {
+              view: placeholderTexture.createView(),
+              loadOp: 'clear',
+              storeOp: 'discard',
+            },
+          ],
+        });
+        pass.setPipeline(pipeline);
+        pass.setBindGroup(
+          0,
+          t.device.createBindGroup({
+            layout: pipeline.getBindGroupLayout(0),
+            entries: [
+              {
+                binding: 0,
+                resource: view,
+              },
+            ],
+          })
+        );
+        pass.draw(6);
+        pass.end();
+        t.device.queue.submit([commandEncoder.finish()]);
+      }
+      break;
+
+    case 'render-pass-store':
+    case 'render-pass-resolve':
+      {
+        // Create a placeholder color attachment texture for the store target when tesing texture is used as resolve target.
+        const targetView =
+          method === 'render-pass-store'
+            ? view
+            : t
+                .trackForCleanup(
+                  t.device.createTexture({
+                    format,
+                    size: [kTextureSize, kTextureSize],
+                    usage: GPUTextureUsage.RENDER_ATTACHMENT,
+                    sampleCount: 4,
+                  })
+                )
+                .createView();
+        const resolveView = method === 'render-pass-store' ? undefined : view;
+        const multisampleCount = method === 'render-pass-store' ? sampleCount : 4;
+
+        const pipeline = t.device.createRenderPipeline({
+          layout: 'auto',
+          vertex: {
+            module: t.device.createShaderModule({
+              code: kFullscreenQuadVertexShaderCode,
+            }),
+          },
+          fragment: {
+            module: t.device.createShaderModule({
+              code: `
+                @fragment fn main(
+                  @builtin(position) fragCoord: vec4<f32>,
+                ) -> @location(0) ${vecType} {
+                  const src = ${kColorArrayShaderString};
+                  let coord = vec2u(fragCoord.xy);
+                  let idx = coord.x + coord.y * ${kTextureSize};
+                  return src[idx % ${kColors.length}];
+                }`,
+            }),
+            targets: [
+              {
+                format,
+              },
+            ],
+          },
+          multisample: {
+            count: multisampleCount,
+          },
+        });
+        const commandEncoder = t.device.createCommandEncoder();
+        const pass = commandEncoder.beginRenderPass({
+          colorAttachments: [
+            {
+              view: targetView,
+              resolveTarget: resolveView,
+              loadOp: 'clear',
+              storeOp: 'store',
+            },
+          ],
+        });
+        pass.setPipeline(pipeline);
+        pass.draw(6);
+        pass.end();
+        t.device.queue.submit([commandEncoder.finish()]);
+      }
+      break;
+    default:
+      unreachable();
+  }
+
+  return expectedTexelView;
+}
 
 g.test('format')
   .desc(
     `Views of every allowed format.
 
+Read values from color array in the shader, and write it to the texture view via different write methods.
+
 - x= every texture format
 - x= sampleCount {1, 4} if valid
 - x= every possible view write method (see above)
+
+TODO: Test sampleCount > 1 for 'render-pass-store' after extending copySinglePixelTextureToBufferUsingComputePass
+      to read multiple pixels from multisampled textures. [1]
+TODO: Test rgb10a2uint when TexelRepresentation.numericRange is made per-component. [2]
 `
   )
-  .unimplemented();
+  .params(u =>
+    u //
+      .combine('method', kTextureViewWriteMethods)
+      .combine('format', kRegularTextureFormats)
+      .combine('sampleCount', [1, 4])
+      .filter(({ format, method, sampleCount }) => {
+        const info = kTextureFormatInfo[format];
+
+        if (sampleCount > 1 && !info.multisample) {
+          return false;
+        }
+
+        // [2]
+        if (format === 'rgb10a2uint') {
+          return false;
+        }
+
+        switch (method) {
+          case 'storage-write-compute':
+          case 'storage-write-fragment':
+            return info.color?.storage && sampleCount === 1;
+          case 'render-pass-store':
+            // [1]
+            if (sampleCount > 1) {
+              return false;
+            }
+            return !!info.colorRender;
+          case 'render-pass-resolve':
+            return !!info.colorRender?.resolve && sampleCount === 1;
+        }
+        return true;
+      })
+  )
+  .beforeAllSubcases(t => {
+    const { format, method } = t.params;
+    t.skipIfTextureFormatNotSupported(format);
+
+    switch (method) {
+      case 'storage-write-compute':
+      case 'storage-write-fragment':
+        // Still need to filter again for compat mode.
+        t.skipIfTextureFormatNotUsableAsStorageTexture(format);
+        break;
+    }
+  })
+  .fn(t => {
+    const { format, method, sampleCount } = t.params;
+
+    const usage =
+      GPUTextureUsage.COPY_SRC |
+      (method.includes('storage')
+        ? GPUTextureUsage.STORAGE_BINDING
+        : GPUTextureUsage.RENDER_ATTACHMENT);
+
+    const texture = t.trackForCleanup(
+      t.device.createTexture({
+        format,
+        usage,
+        size: [kTextureSize, kTextureSize],
+        sampleCount,
+      })
+    );
+
+    const view = texture.createView();
+    const expectedTexelView = writeTextureAndGetExpectedTexelView(
+      t,
+      method,
+      view,
+      format,
+      sampleCount
+    );
+
+    // [1] Use copySinglePixelTextureToBufferUsingComputePass to check multisampled texture.
+    t.expectTexelViewComparisonIsOkInTexture({ texture }, expectedTexelView, [
+      kTextureSize,
+      kTextureSize,
+    ]);
+  });
 
 g.test('dimension')
   .desc(

src/webgpu/util/shader.ts

@@ -11,6 +11,27 @@ export const kDefaultFragmentShaderCode = `
   return vec4<f32>(1.0, 1.0, 1.0, 1.0);
 }`;
 
+// MAINTENANCE_TODO(#3344): deduplicate fullscreen quad shader code.
+export const kFullscreenQuadVertexShaderCode = `
+  struct VertexOutput {
+    @builtin(position) Position : vec4<f32>
+  };
+
+  @vertex fn main(@builtin(vertex_index) VertexIndex : u32) -> VertexOutput {
+    var pos = array<vec2<f32>, 6>(
+        vec2<f32>( 1.0,  1.0),
+        vec2<f32>( 1.0, -1.0),
+        vec2<f32>(-1.0, -1.0),
+        vec2<f32>( 1.0,  1.0),
+        vec2<f32>(-1.0, -1.0),
+        vec2<f32>(-1.0,  1.0));
+
+    var output : VertexOutput;
+    output.Position = vec4<f32>(pos[VertexIndex], 0.0, 1.0);
+    return output;
+  }
+`;
+
 const kPlainTypeInfo = {
   i32: {
     suffix: '',