export const description = `
createComputePipeline and createComputePipelineAsync validation tests.

Note: entry point matching tests are in shader_module/entry_point.spec.ts
`;

import { AllFeaturesMaxLimitsGPUTest } from '../.././gpu_test.js';
import { makeTestGroup } from '../../../common/framework/test_group.js';
import { keysOf } from '../../../common/util/data_tables.js';
import {
  isTextureFormatUsableWithStorageAccessMode,
  kPossibleStorageTextureFormats,
} from '../../format_info.js';
import { kValue } from '../../util/constants.js';
import { TShaderStage, getShaderWithEntryPoint } from '../../util/shader.js';

import {
  kAPIResources,
  getWGSLShaderForResource,
  getAPIBindGroupLayoutForResource,
  doResourcesMatch,
} from './utils.js';
import * as vtu from './validation_test_utils.js';

class F extends AllFeaturesMaxLimitsGPUTest {
  getShaderModule(
    shaderStage: TShaderStage = 'compute',
    entryPoint: string = 'main'
  ): GPUShaderModule {
    return this.device.createShaderModule({
      code: getShaderWithEntryPoint(shaderStage, entryPoint),
    });
  }
}

export const g = makeTestGroup(F);

g.test('basic')
  .desc(
    `
Control case for createComputePipeline and createComputePipelineAsync.
Call the API with valid compute shader and matching valid entryPoint, making sure that the test function working well.
`
  )
  .params(u => u.combine('isAsync', [true, false]))
  .fn(t => {
    const { isAsync } = t.params;
    vtu.doCreateComputePipelineTest(t, isAsync, true, {
      layout: 'auto',
      compute: { module: t.getShaderModule('compute', 'main'), entryPoint: 'main' },
    });
  });

g.test('shader_module,invalid')
  .desc(
    `
Tests calling createComputePipeline(Async) with a invalid compute shader, and check that the APIs catch this error.
`
  )
  .params(u => u.combine('isAsync', [true, false]))
  .fn(t => {
    const { isAsync } = t.params;
    vtu.doCreateComputePipelineTest(t, isAsync, false, {
      layout: 'auto',
      compute: {
        module: vtu.createInvalidShaderModule(t),
        entryPoint: 'main',
      },
    });
  });

g.test('shader_module,compute')
  .desc(
    `
Tests calling createComputePipeline(Async) with valid but different stage shader and matching entryPoint,
and check that the APIs only accept compute shader.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combine('shaderModuleStage', ['compute', 'vertex', 'fragment'] as TShaderStage[])
  )
  .fn(t => {
    const { isAsync, shaderModuleStage } = t.params;
    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.getShaderModule(shaderModuleStage, 'main'),
        entryPoint: 'main',
      },
    };
    vtu.doCreateComputePipelineTest(t, isAsync, shaderModuleStage === 'compute', descriptor);
  });

g.test('shader_module,device_mismatch')
  .desc(
    'Tests createComputePipeline(Async) cannot be called with a shader module created from another device'
  )
  .paramsSubcasesOnly(u => u.combine('isAsync', [true, false]).combine('mismatched', [true, false]))
  .beforeAllSubcases(t => t.usesMismatchedDevice())
  .fn(t => {
    const { isAsync, mismatched } = t.params;

    const sourceDevice = mismatched ? t.mismatchedDevice : t.device;

    const module = sourceDevice.createShaderModule({
      code: '@compute @workgroup_size(1) fn main() {}',
    });

    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module,
        entryPoint: 'main',
      },
    };

    vtu.doCreateComputePipelineTest(t, isAsync, !mismatched, descriptor);
  });

g.test('pipeline_layout,device_mismatch')
  .desc(
    'Tests createComputePipeline(Async) cannot be called with a pipeline layout created from another device'
  )
  .paramsSubcasesOnly(u => u.combine('isAsync', [true, false]).combine('mismatched', [true, false]))
  .beforeAllSubcases(t => t.usesMismatchedDevice())
  .fn(t => {
    const { isAsync, mismatched } = t.params;
    const sourceDevice = mismatched ? t.mismatchedDevice : t.device;

    const layout = sourceDevice.createPipelineLayout({ bindGroupLayouts: [] });

    const descriptor = {
      layout,
      compute: {
        module: t.getShaderModule('compute', 'main'),
        entryPoint: 'main',
      },
    };

    vtu.doCreateComputePipelineTest(t, isAsync, !mismatched, descriptor);
  });

g.test('limits,workgroup_storage_size')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for compute using <= device.limits.maxComputeWorkgroupStorageSize bytes of workgroup storage.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combineWithParams([
        { type: 'vec4<f32>', _typeSize: 16 },
        { type: 'mat4x4<f32>', _typeSize: 64 },
      ])
      .beginSubcases()
      .combine('countDeltaFromLimit', [0, 1])
  )
  .fn(t => {
    const { isAsync, type, _typeSize, countDeltaFromLimit } = t.params;
    const countAtLimit = Math.floor(t.device.limits.maxComputeWorkgroupStorageSize / _typeSize);
    const count = countAtLimit + countDeltaFromLimit;

    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.device.createShaderModule({
          code: `
          var<workgroup> data: array<${type}, ${count}>;
          @compute @workgroup_size(64) fn main () {
            _ = data;
          }
          `,
        }),
        entryPoint: 'main',
      },
    };
    vtu.doCreateComputePipelineTest(t, isAsync, count <= countAtLimit, descriptor);
  });

g.test('limits,invocations_per_workgroup')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for compute using <= device.limits.maxComputeInvocationsPerWorkgroup per workgroup.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combine('size', [
        // Assume maxComputeWorkgroupSizeX/Y >= 129, maxComputeWorkgroupSizeZ >= 33
        [128, 1, 2],
        [129, 1, 2],
        [2, 128, 1],
        [2, 129, 1],
        [1, 8, 32],
        [1, 8, 33],
      ])
  )
  .fn(t => {
    const { isAsync, size } = t.params;

    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.device.createShaderModule({
          code: `
          @compute @workgroup_size(${size.join(',')}) fn main () {
          }
          `,
        }),
        entryPoint: 'main',
      },
    };

    vtu.doCreateComputePipelineTest(
      t,
      isAsync,
      size[0] * size[1] * size[2] <= t.device.limits.maxComputeInvocationsPerWorkgroup,
      descriptor
    );
  });

g.test('limits,invocations_per_workgroup,each_component')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for compute workgroup_size attribute has each component no more than their limits.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combine('size', [
        // Assume maxComputeInvocationsPerWorkgroup >= 256
        [64],
        [256, 1, 1],
        [257, 1, 1],
        [1, 256, 1],
        [1, 257, 1],
        [1, 1, 63],
        [1, 1, 64],
        [1, 1, 65],
      ] as const)
  )
  .fn(t => {
    const { isAsync, size } = t.params;

    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.device.createShaderModule({
          code: `
          @compute @workgroup_size(${size.join(',')}) fn main () {
          }
          `,
        }),
        entryPoint: 'main',
      },
    };

    const workgroupX = size[0];
    const workgroupY = size[1] ?? 1;
    const workgroupZ = size[2] ?? 1;

    const _success =
      workgroupX <= t.device.limits.maxComputeWorkgroupSizeX &&
      workgroupY <= t.device.limits.maxComputeWorkgroupSizeY &&
      workgroupZ <= t.device.limits.maxComputeWorkgroupSizeZ;
    vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
  });

g.test('overrides,identifier')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for overridable constants identifiers.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combineWithParams([
        { constants: {}, _success: true },
        { constants: { c0: 0 }, _success: true },
        { constants: { c0: 0, c1: 1 }, _success: true },
        { constants: { 'c0\0': 0 }, _success: false },
        { constants: { c9: 0 }, _success: false },
        { constants: { 1: 0 }, _success: true },
        { constants: { c3: 0 }, _success: false }, // pipeline constant id is specified for c3
        { constants: { 2: 0 }, _success: false },
        { constants: { 1000: 0 }, _success: true },
        { constants: { 9999: 0 }, _success: false },
        { constants: { 1000: 0, c2: 0 }, _success: false },
        { constants: { 数: 0 }, _success: true },
        { constants: { séquençage: 0 }, _success: false }, // test unicode is not normalized
      ] as { constants: Record<string, GPUPipelineConstantValue>; _success: boolean }[])
  )
  .fn(t => {
    const { isAsync, constants, _success } = t.params;

    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.device.createShaderModule({
          code: `
            override c0: bool = true;      // type: bool
            override c1: u32 = 0u;          // default override
            override 数: u32 = 0u;          // non-ASCII
            override séquençage: u32 = 0u;  // normalizable unicode (WGSL does not normalize)
            @id(1000) override c2: u32 = 10u;  // default
            @id(1) override c3: u32 = 11u;     // default
            @compute @workgroup_size(1) fn main () {
              // make sure the overridable constants are not optimized out
              _ = u32(c0);
              _ = u32(c1);
              _ = u32(c2 + séquençage);
              _ = u32(c3 + 数);
            }`,
        }),
        entryPoint: 'main',
        constants,
      },
    };

    vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
  });

g.test('overrides,uninitialized')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for uninitialized overridable constants.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combineWithParams([
        { constants: {}, _success: false },
        { constants: { c0: 0, c2: 0, c8: 0 }, _success: false }, // c5 is missing
        { constants: { c0: 0, c2: 0, c5: 0, c8: 0 }, _success: true },
        { constants: { c0: 0, c2: 0, c5: 0, c8: 0, c1: 0 }, _success: true },
      ] as { constants: Record<string, GPUPipelineConstantValue>; _success: boolean }[])
  )
  .fn(t => {
    const { isAsync, constants, _success } = t.params;

    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.device.createShaderModule({
          code: `
            override c0: bool;              // type: bool
            override c1: bool = false;      // default override
            override c2: f32;               // type: float32
            override c3: f32 = 0.0;         // default override
            override c4: f32 = 4.0;         // default
            override c5: i32;               // type: int32
            override c6: i32 = 0;           // default override
            override c7: i32 = 7;           // default
            override c8: u32;               // type: uint32
            override c9: u32 = 0u;          // default override
            @id(1000) override c10: u32 = 10u;  // default
            @compute @workgroup_size(1) fn main () {
              // make sure the overridable constants are not optimized out
              _ = u32(c0);
              _ = u32(c1);
              _ = u32(c2);
              _ = u32(c3);
              _ = u32(c4);
              _ = u32(c5);
              _ = u32(c6);
              _ = u32(c7);
              _ = u32(c8);
              _ = u32(c9);
              _ = u32(c10);
            }`,
        }),
        entryPoint: 'main',
        constants,
      },
    };

    vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
  });

g.test('overrides,value,type_error')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for constant values like inf, NaN will results in TypeError.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combineWithParams([
        { constants: { cf: 1 }, _success: true }, // control
        { constants: { cf: NaN }, _success: false },
        { constants: { cf: Number.POSITIVE_INFINITY }, _success: false },
        { constants: { cf: Number.NEGATIVE_INFINITY }, _success: false },
      ] as const)
  )
  .fn(t => {
    const { isAsync, constants, _success } = t.params;

    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.device.createShaderModule({
          code: `
            override cf: f32 = 0.0;
            @compute @workgroup_size(1) fn main () {
              _ = cf;
            }`,
        }),
        entryPoint: 'main',
        constants,
      },
    };

    vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor, 'TypeError');
  });

g.test('overrides,value,validation_error')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for unrepresentable constant values in compute stage.

TODO(#2060): test with last_castable_pipeline_override.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combineWithParams([
        { constants: { cu: kValue.u32.min }, _success: true },
        { constants: { cu: kValue.u32.min - 1 }, _success: false },
        { constants: { cu: kValue.u32.max }, _success: true },
        { constants: { cu: kValue.u32.max + 1 }, _success: false },
        { constants: { ci: kValue.i32.negative.min }, _success: true },
        { constants: { ci: kValue.i32.negative.min - 1 }, _success: false },
        { constants: { ci: kValue.i32.positive.max }, _success: true },
        { constants: { ci: kValue.i32.positive.max + 1 }, _success: false },
        { constants: { cf: kValue.f32.negative.min }, _success: true },
        {
          constants: { cf: kValue.f32.negative.first_non_castable_pipeline_override },
          _success: false,
        },
        { constants: { cf: kValue.f32.positive.max }, _success: true },
        {
          constants: { cf: kValue.f32.positive.first_non_castable_pipeline_override },
          _success: false,
        },
        // Conversion to boolean can't fail
        { constants: { cb: Number.MAX_VALUE }, _success: true },
        { constants: { cb: kValue.i32.negative.min - 1 }, _success: true },
      ] as { constants: Record<string, GPUPipelineConstantValue>; _success: boolean }[])
  )
  .fn(t => {
    const { isAsync, constants, _success } = t.params;
    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.device.createShaderModule({
          code: `
          override cb: bool = false;
          override cu: u32 = 0u;
          override ci: i32 = 0;
          override cf: f32 = 0.0;
          @compute @workgroup_size(1) fn main () {
            _ = cb;
            _ = cu;
            _ = ci;
            _ = cf;
          }`,
        }),
        entryPoint: 'main',
        constants,
      },
    };

    vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
  });

g.test('overrides,entry_point,validation_error')
  .desc(
    `
Tests that pipeline constant (override) errors only trigger on entry point usage.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combine('pipeEntryPoint', ['main_success', 'main_pipe_error'])
  )
  .fn(t => {
    const { isAsync, pipeEntryPoint } = t.params;
    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.device.createShaderModule({
          code: `
          override cu: u32 = 0u;
          override cx: u32 = 1u/cu;
          @compute @workgroup_size(1) fn main_success () {
            _ = cu;
          }
          @compute @workgroup_size(1) fn main_pipe_error () {
            _ = cx;
          }`,
        }),
        entryPoint: pipeEntryPoint,
      },
    };

    vtu.doCreateComputePipelineTest(t, isAsync, pipeEntryPoint === 'main_success', descriptor);
  });

g.test('overrides,value,validation_error,f16')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for unrepresentable f16 constant values in compute stage.

TODO(#2060): Tighten the cases around the valid/invalid boundary once we have WGSL spec
clarity on whether values like f16.positive.last_castable_pipeline_override would be valid. See issue.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combineWithParams([
        { constants: { cf16: kValue.f16.negative.min }, _success: true },
        {
          constants: { cf16: kValue.f16.negative.first_non_castable_pipeline_override },
          _success: false,
        },
        { constants: { cf16: kValue.f16.positive.max }, _success: true },
        {
          constants: { cf16: kValue.f16.positive.first_non_castable_pipeline_override },
          _success: false,
        },
        { constants: { cf16: kValue.f32.negative.min }, _success: false },
        { constants: { cf16: kValue.f32.positive.max }, _success: false },
        {
          constants: { cf16: kValue.f32.negative.first_non_castable_pipeline_override },
          _success: false,
        },
        {
          constants: { cf16: kValue.f32.positive.first_non_castable_pipeline_override },
          _success: false,
        },
      ] as const)
  )
  .fn(t => {
    t.skipIfDeviceDoesNotHaveFeature('shader-f16');
    const { isAsync, constants, _success } = t.params;

    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.device.createShaderModule({
          code: `
          enable f16;

          override cf16: f16 = 0.0h;
          @compute @workgroup_size(1) fn main () {
            _ = cf16;
          }`,
        }),
        entryPoint: 'main',
        constants,
      },
    };

    vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
  });

const kOverridesWorkgroupSizeShaders = {
  u32: `
override x: u32 = 1u;
override y: u32 = 1u;
override z: u32 = 1u;
@compute @workgroup_size(x, y, z) fn main () {
  _ = 0u;
}
`,
  i32: `
override x: i32 = 1;
override y: i32 = 1;
override z: i32 = 1;
@compute @workgroup_size(x, y, z) fn main () {
  _ = 0u;
}
`,
};

g.test('overrides,workgroup_size')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for overridable constants used for workgroup size.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combine('type', ['u32', 'i32'] as const)
      .combineWithParams([
        { constants: {}, _success: true },
        { constants: { x: 0, y: 0, z: 0 }, _success: false },
        { constants: { x: 1, y: -1, z: 1 }, _success: false },
        { constants: { x: 1, y: 0, z: 0 }, _success: false },
        { constants: { x: 16, y: 1, z: 1 }, _success: true },
      ] as { constants: Record<string, GPUPipelineConstantValue>; _success: boolean }[])
  )
  .fn(t => {
    const { isAsync, type, constants, _success } = t.params;

    const descriptor = {
      layout: 'auto' as const,
      compute: {
        module: t.device.createShaderModule({
          code: kOverridesWorkgroupSizeShaders[type],
        }),
        entryPoint: 'main',
        constants,
      },
    };

    vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
  });

g.test('overrides,workgroup_size,limits')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for overridable constants for workgroupSize exceeds device limits.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
      .combine('type', ['u32', 'i32'] as const)
  )
  .fn(t => {
    const { isAsync, type } = t.params;

    const limits = t.device.limits;

    const testFn = (x: number, y: number, z: number, _success: boolean) => {
      const descriptor = {
        layout: 'auto' as const,
        compute: {
          module: t.device.createShaderModule({
            code: kOverridesWorkgroupSizeShaders[type],
          }),
          entryPoint: 'main',
          constants: {
            x,
            y,
            z,
          },
        },
      };

      vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
    };

    testFn(limits.maxComputeWorkgroupSizeX, 1, 1, true);
    testFn(limits.maxComputeWorkgroupSizeX + 1, 1, 1, false);
    testFn(1, limits.maxComputeWorkgroupSizeY, 1, true);
    testFn(1, limits.maxComputeWorkgroupSizeY + 1, 1, false);
    testFn(1, 1, limits.maxComputeWorkgroupSizeZ, true);
    testFn(1, 1, limits.maxComputeWorkgroupSizeZ + 1, false);
    testFn(
      limits.maxComputeWorkgroupSizeX,
      limits.maxComputeWorkgroupSizeY,
      limits.maxComputeWorkgroupSizeZ,
      limits.maxComputeWorkgroupSizeX *
        limits.maxComputeWorkgroupSizeY *
        limits.maxComputeWorkgroupSizeZ <=
        limits.maxComputeInvocationsPerWorkgroup
    );
  });

g.test('overrides,workgroup_size,limits,workgroup_storage_size')
  .desc(
    `
Tests calling createComputePipeline(Async) validation for overridable constants for workgroupStorageSize exceeds device limits.
`
  )
  .params(u =>
    u //
      .combine('isAsync', [true, false])
  )
  .fn(t => {
    const { isAsync } = t.params;

    const limits = t.device.limits;

    const kVec4Size = 16;
    const maxVec4Count = limits.maxComputeWorkgroupStorageSize / kVec4Size;
    const kMat4Size = 64;
    const maxMat4Count = limits.maxComputeWorkgroupStorageSize / kMat4Size;

    const testFn = (vec4Count: number, mat4Count: number, _success: boolean) => {
      const descriptor = {
        layout: 'auto' as const,
        compute: {
          module: t.device.createShaderModule({
            code: `
              override a: u32;
              override b: u32;
              ${vec4Count <= 0 ? '' : 'var<workgroup> vec4_data: array<vec4<f32>, a>;'}
              ${mat4Count <= 0 ? '' : 'var<workgroup> mat4_data: array<mat4x4<f32>, b>;'}
              @compute @workgroup_size(1) fn main() {
                ${vec4Count <= 0 ? '' : '_ = vec4_data[0];'}
                ${mat4Count <= 0 ? '' : '_ = mat4_data[0];'}
              }`,
          }),
          entryPoint: 'main',
          constants: {
            a: vec4Count,
            b: mat4Count,
          },
        },
      };

      vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
    };

    testFn(1, 1, true);
    testFn(maxVec4Count + 1, 0, false);
    testFn(0, maxMat4Count + 1, false);
  });

g.test('resource_compatibility')
  .desc(
    'Tests validation of resource (bind group) compatibility between pipeline layout and WGSL shader'
  )
  .params(u =>
    u //
      .combine('apiResource', keysOf(kAPIResources))
      .beginSubcases()
      .combine('isAsync', [true, false] as const)
      .combine('wgslResource', keysOf(kAPIResources))
  )
  .fn(t => {
    const apiResource = kAPIResources[t.params.apiResource];
    const wgslResource = kAPIResources[t.params.wgslResource];
    t.skipIf(
      wgslResource.storageTexture !== undefined &&
        wgslResource.storageTexture.access !== 'write-only' &&
        !t.hasLanguageFeature('readonly_and_readwrite_storage_textures'),
      'Storage textures require language feature'
    );
    t.skipIfTextureViewDimensionNotSupported(wgslResource.texture?.viewDimension);

    const layout = t.device.createPipelineLayout({
      bindGroupLayouts: [
        getAPIBindGroupLayoutForResource(t.device, GPUShaderStage.COMPUTE, apiResource),
      ],
    });

    const descriptor = {
      layout,
      compute: {
        module: t.device.createShaderModule({
          code: getWGSLShaderForResource('compute', wgslResource),
        }),
        entryPoint: 'main',
      },
    };
    vtu.doCreateComputePipelineTest(
      t,
      t.params.isAsync,
      doResourcesMatch(apiResource, wgslResource),
      descriptor
    );
  });

g.test('storage_texture,format')
  .desc(
    `
Test that a pipeline with auto layout and storage texture access combo that is not supported
generates a validation error at createComputePipeline(Async)
  `
  )
  .params(u =>
    u //
      .combine('format', kPossibleStorageTextureFormats)
      .beginSubcases()
      .combine('isAsync', [true, false] as const)
      .combine('access', ['read', 'write', 'read_write'] as const)
      .combine('dimension', ['1d', '2d', '3d'] as const)
  )
  .fn(t => {
    const { format, isAsync, access, dimension } = t.params;
    t.skipIfTextureFormatNotSupported(format);

    const code = `
      @group(0) @binding(0) var tex: texture_storage_${dimension}<${format}, ${access}>;
      @compute @workgroup_size(1) fn main() {
        _ = tex;
      }
    `;
    const module = t.device.createShaderModule({ code });

    const success = isTextureFormatUsableWithStorageAccessMode(t.device.features, format, access);
    const descriptor: GPUComputePipelineDescriptor = {
      layout: 'auto',
      compute: { module },
    };
    vtu.doCreateComputePipelineTest(t, isAsync, success, descriptor);
  });