export const description = `Test memory layout requirements`;

import { makeTestGroup } from '../../../common/framework/test_group.js';
import { keysOf } from '../../../common/util/data_tables.js';
import { iterRange } from '../../../common/util/util.js';
import { AllFeaturesMaxLimitsGPUTest } from '../../gpu_test.js';

export const g = makeTestGroup(AllFeaturesMaxLimitsGPUTest);

interface LayoutCase {
  type: string;
  decl?: string;
  read_assign: string;
  write_assign: string;
  offset: number;
  f16?: boolean;
  f32?: boolean;
  skip_uniform?: boolean;
}

const kLayoutCases: Record<string, LayoutCase> = {
  vec2u_align8: {
    type: `S_vec2u_align`,
    decl: `struct S_vec2u_align {
      x : u32,
      y : vec2u,
    }`,
    read_assign: `out = in.y[1]`,
    write_assign: `out.y[1] = in`,
    offset: 12,
  },
  vec3u_align16: {
    type: `S_vec3u_align`,
    decl: `struct S_vec3u_align {
      x : u32,
      y : vec3u,
    }`,
    read_assign: `out = in.y[2]`,
    write_assign: `out.y[2] = in`,
    offset: 24,
  },
  vec4u_align16: {
    type: `S_vec4u_align`,
    decl: `struct S_vec4u_align {
      x : u32,
      y : vec4u,
    }`,
    read_assign: `out = in.y[0]`,
    write_assign: `out.y[0] = in`,
    offset: 16,
  },
  struct_align32: {
    type: `S_align32`,
    decl: `struct S_align32 {
      x : u32,
      @align(32) y : u32,
    }`,
    read_assign: `out = in.y;`,
    write_assign: `out.y = in`,
    offset: 32,
  },
  vec2h_align4: {
    type: `S_vec2h_align`,
    decl: `struct S_vec2h_align {
      x : f16,
      y : vec2h,
    }`,
    read_assign: `out = u32(in.y[0])`,
    write_assign: `out.y[0] = f16(in)`,
    offset: 4,
    f16: true,
  },
  vec3h_align8: {
    type: `S_vec3h_align`,
    decl: `struct S_vec3h_align {
      x : f16,
      y : vec3h,
    }`,
    read_assign: `out = u32(in.y[2])`,
    write_assign: `out.y[2] = f16(in)`,
    offset: 12,
    f16: true,
  },
  vec4h_align8: {
    type: `S_vec4h_align`,
    decl: `struct S_vec4h_align {
      x : f16,
      y : vec4h,
    }`,
    read_assign: `out = u32(in.y[2])`,
    write_assign: `out.y[2] = f16(in)`,
    offset: 12,
    f16: true,
  },
  vec2f_align8: {
    type: `S_vec2f_align`,
    decl: `struct S_vec2f_align {
      x : u32,
      y : vec2f,
    }`,
    read_assign: `out = u32(in.y[1])`,
    write_assign: `out.y[1] = f32(in)`,
    offset: 12,
    f32: true,
  },
  vec3f_align16: {
    type: `S_vec3f_align`,
    decl: `struct S_vec3f_align {
      x : u32,
      y : vec3f,
    }`,
    read_assign: `out = u32(in.y[2])`,
    write_assign: `out.y[2] = f32(in)`,
    offset: 24,
    f32: true,
  },
  vec4f_align16: {
    type: `S_vec4f_align`,
    decl: `struct S_vec4f_align {
      x : u32,
      y : vec4f,
    }`,
    read_assign: `out = u32(in.y[0])`,
    write_assign: `out.y[0] = f32(in)`,
    offset: 16,
    f32: true,
  },
  vec3i_size12: {
    type: `S_vec3i_size`,
    decl: `struct S_vec3i_size {
      x : vec3i,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 12,
  },
  vec3h_size6: {
    type: `S_vec3h_size`,
    decl: `struct S_vec3h_size {
      x : vec3h,
      y : f16,
      z : f16,
    }`,
    read_assign: `out = u32(in.z)`,
    write_assign: `out.z = f16(in)`,
    offset: 8,
    f16: true,
  },
  size80: {
    type: `S_size80`,
    decl: `struct S_size80 {
      @size(80) x : u32,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 80,
  },
  atomic_align4: {
    type: `S_atomic_align`,
    decl: `struct S_atomic_align {
      x : u32,
      y : atomic<u32>,
    }`,
    read_assign: `out = atomicLoad(&in.y)`,
    write_assign: `atomicStore(&out.y, in)`,
    offset: 4,
  },
  atomic_size4: {
    type: `S_atomic_size`,
    decl: `struct S_atomic_size {
      x : atomic<u32>,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 4,
  },
  mat2x2f_align8: {
    type: `S_mat2x2f_align`,
    decl: `struct S_mat2x2f_align {
      x : u32,
      y : mat2x2f,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f32(in)`,
    offset: 8,
    f32: true,
  },
  mat3x2f_align8: {
    type: `S_mat3x2f_align`,
    decl: `struct S_mat3x2f_align {
      x : u32,
      y : mat3x2f,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f32(in)`,
    offset: 8,
    f32: true,
  },
  mat4x2f_align8: {
    type: `S_mat4x2f_align`,
    decl: `struct S_mat4x2f_align {
      x : u32,
      y : mat4x2f,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f32(in)`,
    offset: 8,
    f32: true,
  },
  mat2x3f_align16: {
    type: `S_mat2x3f_align`,
    decl: `struct S_mat2x3f_align {
      x : u32,
      y : mat2x3f,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f32(in)`,
    offset: 16,
    f32: true,
  },
  mat3x3f_align16: {
    type: `S_mat3x3f_align`,
    decl: `struct S_mat3x3f_align {
      x : u32,
      y : mat3x3f,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f32(in)`,
    offset: 16,
    f32: true,
  },
  mat4x3f_align16: {
    type: `S_mat4x3f_align`,
    decl: `struct S_mat4x3f_align {
      x : u32,
      y : mat4x3f,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f32(in)`,
    offset: 16,
    f32: true,
  },
  mat2x4f_align16: {
    type: `S_mat2x4f_align`,
    decl: `struct S_mat2x4f_align {
      x : u32,
      y : mat2x4f,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f32(in)`,
    offset: 16,
    f32: true,
  },
  mat3x4f_align16: {
    type: `S_mat3x4f_align`,
    decl: `struct S_mat3x4f_align {
      x : u32,
      y : mat3x4f,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f32(in)`,
    offset: 16,
    f32: true,
  },
  mat4x4f_align16: {
    type: `S_mat4x4f_align`,
    decl: `struct S_mat4x4f_align {
      x : u32,
      y : mat4x4f,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f32(in)`,
    offset: 16,
    f32: true,
  },
  mat2x2h_align4: {
    type: `S_mat2x2h_align`,
    decl: `struct S_mat2x2h_align {
      x : u32,
      y : mat2x2h,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f16(in)`,
    offset: 4,
    f16: true,
  },
  mat3x2h_align4: {
    type: `S_mat3x2h_align`,
    decl: `struct S_mat3x2h_align {
      x : u32,
      y : mat3x2h,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f16(in)`,
    offset: 4,
    f16: true,
  },
  mat4x2h_align4: {
    type: `S_mat4x2h_align`,
    decl: `struct S_mat4x2h_align {
      x : u32,
      y : mat4x2h,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f16(in)`,
    offset: 4,
    f16: true,
  },
  mat2x3h_align8: {
    type: `S_mat2x3h_align`,
    decl: `struct S_mat2x3h_align {
      x : u32,
      y : mat2x3h,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f16(in)`,
    offset: 8,
    f16: true,
  },
  mat3x3h_align8: {
    type: `S_mat3x3h_align`,
    decl: `struct S_mat3x3h_align {
      x : u32,
      y : mat2x3h,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f16(in)`,
    offset: 8,
    f16: true,
  },
  mat4x3h_align8: {
    type: `S_mat4x3h_align`,
    decl: `struct S_mat4x3h_align {
      x : u32,
      y : mat4x3h,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f16(in)`,
    offset: 8,
    f16: true,
  },
  mat2x4h_align8: {
    type: `S_mat2x4h_align`,
    decl: `struct S_mat2x4h_align {
      x : u32,
      y : mat2x4h,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f16(in)`,
    offset: 8,
    f16: true,
  },
  mat3x4h_align8: {
    type: `S_mat3x4h_align`,
    decl: `struct S_mat3x4h_align {
      x : u32,
      y : mat3x4h,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f16(in)`,
    offset: 8,
    f16: true,
  },
  mat4x4h_align8: {
    type: `S_mat4x4h_align`,
    decl: `struct S_mat4x4h_align {
      x : u32,
      y : mat4x4h,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f16(in)`,
    offset: 8,
    f16: true,
  },
  mat2x2f_size: {
    type: `S_mat2x2f_size`,
    decl: `struct S_mat2x2f_size {
      x : mat2x2f,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 16,
  },
  mat3x2f_size: {
    type: `S_mat3x2f_size`,
    decl: `struct S_mat3x2f_size {
      x : mat3x2f,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 24,
  },
  mat4x2f_size: {
    type: `S_mat4x2f_size`,
    decl: `struct S_mat4x2f_size {
      x : mat4x2f,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 32,
  },
  mat2x3f_size: {
    type: `S_mat2x3f_size`,
    decl: `struct S_mat2x3f_size {
      x : mat2x3f,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 32,
  },
  mat3x3f_size: {
    type: `S_mat3x3f_size`,
    decl: `struct S_mat3x3f_size {
      x : mat3x3f,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 48,
  },
  mat4x3f_size: {
    type: `S_mat4x3f_size`,
    decl: `struct S_mat4x3f_size {
      x : mat4x3f,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 64,
  },
  mat2x4f_size: {
    type: `S_mat2x4f_size`,
    decl: `struct S_mat2x4f_size {
      x : mat2x4f,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 32,
  },
  mat3x4f_size: {
    type: `S_mat3x4f_size`,
    decl: `struct S_mat3x4f_size {
      x : mat3x4f,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 48,
  },
  mat4x4f_size: {
    type: `S_mat4x4f_size`,
    decl: `struct S_mat4x4f_size {
      x : mat4x4f,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 64,
  },
  mat2x2h_size: {
    type: `S_mat2x2h_size`,
    decl: `struct S_mat2x2h_size {
      x : mat2x2h,
      y : f16,
    }`,
    read_assign: `out = u32(in.y)`,
    write_assign: `out.y = f16(in)`,
    offset: 8,
    f16: true,
  },
  mat3x2h_size: {
    type: `S_mat3x2h_size`,
    decl: `struct S_mat3x2h_size {
      x : mat3x2h,
      y : f16,
    }`,
    read_assign: `out = u32(in.y)`,
    write_assign: `out.y = f16(in)`,
    offset: 12,
    f16: true,
  },
  mat4x2h_size: {
    type: `S_mat4x2h_size`,
    decl: `struct S_mat4x2h_size {
      x : mat4x2h,
      y : f16,
    }`,
    read_assign: `out = u32(in.y)`,
    write_assign: `out.y = f16(in)`,
    offset: 16,
    f16: true,
  },
  mat2x3h_size: {
    type: `S_mat2x3h_size`,
    decl: `struct S_mat2x3h_size {
      x : mat2x3h,
      y : f16,
    }`,
    read_assign: `out = u32(in.y)`,
    write_assign: `out.y = f16(in)`,
    offset: 16,
    f16: true,
  },
  mat3x3h_size: {
    type: `S_mat3x3h_size`,
    decl: `struct S_mat3x3h_size {
      x : mat3x3h,
      y : f16,
    }`,
    read_assign: `out = u32(in.y)`,
    write_assign: `out.y = f16(in)`,
    offset: 24,
    f16: true,
  },
  mat4x3h_size: {
    type: `S_mat4x3h_size`,
    decl: `struct S_mat4x3h_size {
      x : mat4x3h,
      y : f16,
    }`,
    read_assign: `out = u32(in.y)`,
    write_assign: `out.y = f16(in)`,
    offset: 32,
    f16: true,
  },
  mat2x4h_size: {
    type: `S_mat2x4h_size`,
    decl: `struct S_mat2x4h_size {
      x : mat2x4h,
      y : f16,
    }`,
    read_assign: `out = u32(in.y)`,
    write_assign: `out.y = f16(in)`,
    offset: 16,
    f16: true,
  },
  mat3x4h_size: {
    type: `S_mat3x4h_size`,
    decl: `struct S_mat3x4h_size {
      x : mat3x4h,
      y : f16,
    }`,
    read_assign: `out = u32(in.y)`,
    write_assign: `out.y = f16(in)`,
    offset: 24,
    f16: true,
  },
  mat4x4h_size: {
    type: `S_mat4x4h_size`,
    decl: `struct S_mat4x4h_size {
      x : mat4x4h,
      y : f16,
    }`,
    read_assign: `out = u32(in.y)`,
    write_assign: `out.y = f16(in)`,
    offset: 32,
    f16: true,
  },
  struct_align_vec2i: {
    type: `S_struct_align_vec2i`,
    decl: `struct Inner {
      x : u32,
      y : vec2i,
    }
    struct S_struct_align_vec2i {
      x : u32,
      y : Inner,
    }`,
    read_assign: `out = in.y.x`,
    write_assign: `out.y.x = in`,
    offset: 8,
    skip_uniform: true,
  },
  struct_align_vec3i: {
    type: `S_struct_align_vec3i`,
    decl: `struct Inner {
      x : u32,
      y : vec3i,
    }
    struct S_struct_align_vec3i {
      x : u32,
      y : Inner,
    }`,
    read_assign: `out = in.y.x`,
    write_assign: `out.y.x = in`,
    offset: 16,
  },
  struct_align_vec4i: {
    type: `S_struct_align_vec4i`,
    decl: `struct Inner {
      x : u32,
      y : vec4i,
    }
    struct S_struct_align_vec4i {
      x : u32,
      y : Inner,
    }`,
    read_assign: `out = in.y.x`,
    write_assign: `out.y.x = in`,
    offset: 16,
  },
  struct_align_vec2h: {
    type: `S_struct_align_vec2h`,
    decl: `struct Inner {
      x : f16,
      y : vec2h,
    }
    struct S_struct_align_vec2h {
      x : f16,
      y : Inner,
    }`,
    read_assign: `out = u32(in.y.x)`,
    write_assign: `out.y.x = f16(in)`,
    offset: 4,
    f16: true,
    skip_uniform: true,
  },
  struct_align_vec3h: {
    type: `S_struct_align_vec3h`,
    decl: `struct Inner {
      x : f16,
      y : vec3h,
    }
    struct S_struct_align_vec3h {
      x : f16,
      y : Inner,
    }`,
    read_assign: `out = u32(in.y.x)`,
    write_assign: `out.y.x = f16(in)`,
    offset: 8,
    f16: true,
    skip_uniform: true,
  },
  struct_align_vec4h: {
    type: `S_struct_align_vec4h`,
    decl: `struct Inner {
      x : f16,
      y : vec4h,
    }
    struct S_struct_align_vec4h {
      x : f16,
      y : Inner,
    }`,
    read_assign: `out = u32(in.y.x)`,
    write_assign: `out.y.x = f16(in)`,
    offset: 8,
    f16: true,
    skip_uniform: true,
  },
  struct_size_roundup: {
    type: `S_struct_size_roundup`,
    decl: `struct Inner {
      x : vec3u,
    }
    struct S_struct_size_roundup {
      x : Inner,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 16,
  },
  struct_inner_size: {
    type: `S_struct_inner_size`,
    decl: `struct Inner {
      @size(112) x : u32,
    }
    struct S_struct_inner_size {
      x : Inner,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 112,
  },
  struct_inner_align: {
    type: `S_struct_inner_align`,
    decl: `struct Inner {
      @align(64) x : u32,
    }
    struct S_struct_inner_align {
      x : u32,
      y : Inner,
    }`,
    read_assign: `out = in.y.x`,
    write_assign: `out.y.x = in`,
    offset: 64,
  },
  struct_inner_size_and_align: {
    type: `S_struct_inner_size_and_align`,
    decl: `struct Inner {
      @align(32) @size(33) x : u32,
    }
    struct S_struct_inner_size_and_align {
      x : Inner,
      y : Inner,
    }`,
    read_assign: `out = in.y.x`,
    write_assign: `out.y.x = in`,
    offset: 64,
  },
  struct_override_size: {
    type: `S_struct_override_size`,
    decl: `struct Inner {
      @size(32) x : u32,
    }
    struct S_struct_override_size {
      @size(64) x : Inner,
      y : u32,
    }`,
    read_assign: `out = in.y`,
    write_assign: `out.y = in`,
    offset: 64,
  },
  struct_double_align: {
    type: `S_struct_double_align`,
    decl: `struct Inner {
      x : u32,
      @align(32) y : u32,
    }
    struct S_struct_double_align {
      x : u32,
      @align(64) y : Inner,
    }`,
    read_assign: `out = in.y.y`,
    write_assign: `out.y.y = in`,
    offset: 96,
  },
  array_vec3u_align: {
    type: `S_array_vec3u_align`,
    decl: `struct S_array_vec3u_align {
      x : u32,
      y : array<vec3u, 2>,
    }`,
    read_assign: `out = in.y[0][0]`,
    write_assign: `out.y[0][0] = in`,
    offset: 16,
  },
  array_vec3h_align: {
    type: `S_array_vec3h_align`,
    decl: `struct S_array_vec3h_align {
      x : f16,
      y : array<vec3h, 2>,
    }`,
    read_assign: `out = u32(in.y[0][0])`,
    write_assign: `out.y[0][0] = f16(in)`,
    offset: 8,
    f16: true,
    skip_uniform: true,
  },
  array_vec3u_stride: {
    type: `S_array_vec3u_stride`,
    decl: `struct S_array_vec3u_stride {
      x : array<vec3u, 4>,
    }`,
    read_assign: `out = in.x[1][0]`,
    write_assign: `out.x[1][0] = in`,
    offset: 16,
  },
  array_vec3h_stride: {
    type: `S_array_vec3h_stride`,
    decl: `struct S_array_vec3h_stride {
      x : array<vec3h, 4>,
    }`,
    read_assign: `out = u32(in.x[1][0])`,
    write_assign: `out.x[1][0] = f16(in)`,
    offset: 8,
    f16: true,
    skip_uniform: true,
  },
  array_stride_size: {
    type: `array<S_stride, 4>`,
    decl: `struct S_stride {
      @size(16) x : u32,
    }`,
    read_assign: `out = in[2].x`,
    write_assign: `out[2].x = in`,
    offset: 32,
  },
  array_mat2x2f_stride: {
    type: `array<mat2x2f, 4>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f32(in)`,
    offset: 16,
    f32: true,
  },
  array_mat2x2h_stride: {
    type: `array<mat2x2h, 2>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f16(in)`,
    offset: 8,
    f16: true,
    skip_uniform: true,
  },
  array_mat3x2f_stride: {
    type: `array<mat3x2f, 3>`,
    read_assign: `out = u32(in[2][0][0])`,
    write_assign: `out[2][0][0] = f32(in)`,
    offset: 48,
    f32: true,
    skip_uniform: true,
  },
  array_mat3x2h_stride: {
    type: `array<mat3x2h, 2>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f16(in)`,
    offset: 12,
    f16: true,
    skip_uniform: true,
  },
  array_mat4x2f_stride: {
    type: `array<mat4x2f, 4>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f32(in)`,
    offset: 32,
    f32: true,
  },
  array_mat4x2h_stride: {
    type: `array<mat4x2h, 2>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f16(in)`,
    offset: 16,
    f16: true,
  },
  array_mat2x3f_stride: {
    type: `array<mat2x3f, 4>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f32(in)`,
    offset: 32,
    f32: true,
  },
  array_mat2x3h_stride: {
    type: `array<mat2x3h, 2>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f16(in)`,
    offset: 16,
    f16: true,
  },
  array_mat3x3f_stride: {
    type: `array<mat3x3f, 3>`,
    read_assign: `out = u32(in[2][0][0])`,
    write_assign: `out[2][0][0] = f32(in)`,
    offset: 96,
    f32: true,
  },
  array_mat3x3h_stride: {
    type: `array<mat3x3h, 2>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f16(in)`,
    offset: 24,
    f16: true,
    skip_uniform: true,
  },
  array_mat4x3f_stride: {
    type: `array<mat4x3f, 4>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f32(in)`,
    offset: 64,
    f32: true,
  },
  array_mat4x3h_stride: {
    type: `array<mat4x3h, 2>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f16(in)`,
    offset: 32,
    f16: true,
  },
  array_mat2x4f_stride: {
    type: `array<mat2x4f, 4>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f32(in)`,
    offset: 32,
    f32: true,
  },
  array_mat2x4h_stride: {
    type: `array<mat2x4h, 2>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f16(in)`,
    offset: 16,
    f16: true,
  },
  array_mat3x4f_stride: {
    type: `array<mat3x4f, 3>`,
    read_assign: `out = u32(in[2][0][0])`,
    write_assign: `out[2][0][0] = f32(in)`,
    offset: 96,
    f32: true,
  },
  array_mat3x4h_stride: {
    type: `array<mat3x4h, 2>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f16(in)`,
    offset: 24,
    f16: true,
    skip_uniform: true,
  },
  array_mat4x4f_stride: {
    type: `array<mat4x4f, 4>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f32(in)`,
    offset: 64,
    f32: true,
  },
  array_mat4x4h_stride: {
    type: `array<mat4x4h, 2>`,
    read_assign: `out = u32(in[1][0][0])`,
    write_assign: `out[1][0][0] = f16(in)`,
    offset: 32,
    f16: true,
  },
};

g.test('read_layout')
  .desc('Test reading memory layouts')
  .params(u =>
    u
      .combine('case', keysOf(kLayoutCases))
      .combine('aspace', ['storage', 'uniform', 'workgroup', 'function', 'private'] as const)
      .beginSubcases()
  )
  .beforeAllSubcases(t => {
    const testcase = kLayoutCases[t.params.case];
    // Don't test atomics in workgroup due to initialization boilerplate.
    t.skipIf(
      testcase.type.includes('atomic') && t.params.aspace !== 'storage',
      `Skipping atomic test for non-storage address space`
    );

    // If the `uniform_buffer_standard_layout` feature is supported, the `uniform` address space has
    // the same layout constraints as `storage`.
    const ubo_std_layout = t.hasLanguageFeature('uniform_buffer_standard_layout');

    t.skipIf(
      !ubo_std_layout && testcase.skip_uniform === true && t.params.aspace === 'uniform',
      `Uniform requires 16 byte alignment`
    );
  })
  .fn(t => {
    const testcase = kLayoutCases[t.params.case];
    if (testcase.f16) {
      t.skipIfDeviceDoesNotHaveFeature('shader-f16');
    }

    let code = `
${testcase.f16 ? 'enable f16;' : ''}
${testcase.decl ?? ''}

@group(0) @binding(1)
var<storage, read_write> out : u32;
`;

    if (t.params.aspace === 'uniform') {
      code += `@group(0) @binding(0)
      var<${t.params.aspace}> in : ${testcase.type};`;
    } else if (t.params.aspace === 'storage') {
      // Use read_write for input data to support atomics.
      code += `@group(0) @binding(0)
      var<${t.params.aspace}, read_write> in : ${testcase.type};`;
    } else {
      code += `@group(0) @binding(0)
      var<storage> pre_in : ${testcase.type};`;
      if (t.params.aspace === 'workgroup') {
        code += `
        var<workgroup> in : ${testcase.type};`;
      } else if (t.params.aspace === 'private') {
        code += `
        var<private> in : ${testcase.type};`;
      }
    }

    code += `
@compute @workgroup_size(1,1,1)
fn main() {
`;

    if (
      t.params.aspace === 'workgroup' ||
      t.params.aspace === 'function' ||
      t.params.aspace === 'private'
    ) {
      if (t.params.aspace === 'function') {
        code += `var in : ${testcase.type};\n`;
      }
      code += `in = pre_in;`;
      if (t.params.aspace === 'workgroup') {
        code += `workgroupBarrier();\n`;
      }
    }

    code += `\n${testcase.read_assign};\n}`;

    let usage = GPUBufferUsage.COPY_SRC;
    if (t.params.aspace === 'uniform') {
      usage |= GPUBufferUsage.UNIFORM;
    } else {
      usage |= GPUBufferUsage.STORAGE;
    }

    // Magic number is 42 in various representations.
    const inMagicNumber = testcase.f16 ? 0x5140 : testcase.f32 ? 0x42280000 : 42;
    const in_buffer = t.makeBufferWithContents(
      new Uint32Array([
        ...iterRange(128, x => {
          if (x * 4 === testcase.offset) {
            return inMagicNumber;
          } else {
            return 0;
          }
        }),
      ]),
      usage
    );

    const out_buffer = t.makeBufferWithContents(
      new Uint32Array([...iterRange(1, x => 0)]),
      GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST
    );

    const pipeline = t.device.createComputePipeline({
      layout: 'auto',
      compute: {
        module: t.device.createShaderModule({
          code,
        }),
        entryPoint: 'main',
      },
    });

    const bg = t.device.createBindGroup({
      layout: pipeline.getBindGroupLayout(0),
      entries: [
        {
          binding: 0,
          resource: {
            buffer: in_buffer,
          },
        },
        {
          binding: 1,
          resource: {
            buffer: out_buffer,
          },
        },
      ],
    });

    const encoder = t.device.createCommandEncoder();
    const pass = encoder.beginComputePass();
    pass.setPipeline(pipeline);
    pass.setBindGroup(0, bg);
    pass.dispatchWorkgroups(1, 1, 1);
    pass.end();
    t.queue.submit([encoder.finish()]);

    t.expectGPUBufferValuesEqual(out_buffer, new Uint32Array([42]));
  });

g.test('write_layout')
  .desc('Test writing memory layouts')
  .params(u =>
    u
      .combine('case', keysOf(kLayoutCases))
      .combine('aspace', ['storage', 'workgroup', 'function', 'private'] as const)
      .beginSubcases()
  )
  .beforeAllSubcases(t => {
    const testcase = kLayoutCases[t.params.case];
    // Don't test atomics in workgroup due to initialization boilerplate.
    t.skipIf(
      testcase.type.includes('atomic') && t.params.aspace !== 'storage',
      `Skipping atomic test for non-storage address space`
    );
  })
  .fn(t => {
    const testcase = kLayoutCases[t.params.case];
    if (testcase.f16) {
      t.skipIfDeviceDoesNotHaveFeature('shader-f16');
    }

    let code = `
${testcase.f16 ? 'enable f16;' : ''}
${testcase.decl ?? ''}

@group(0) @binding(0)
var<storage> in : u32;
`;

    if (t.params.aspace === 'storage') {
      code += `@group(0) @binding(1)
      var<storage, read_write> out : ${testcase.type};\n`;
    } else {
      code += `@group(0) @binding(1)
      var<storage, read_write> post_out : ${testcase.type};\n`;

      if (t.params.aspace === 'workgroup') {
        code += `var<workgroup> out : ${testcase.type};\n`;
      } else if (t.params.aspace === 'private') {
        code += `var<private> out : ${testcase.type};\n`;
      }
    }

    code += `
@compute @workgroup_size(1,1,1)
fn main() {
`;

    if (t.params.aspace === 'function') {
      code += `var out : ${testcase.type};\n`;
    }

    code += `${testcase.write_assign};\n`;
    if (
      t.params.aspace === 'workgroup' ||
      t.params.aspace === 'function' ||
      t.params.aspace === 'private'
    ) {
      if (t.params.aspace === 'workgroup') {
        code += `workgroupBarrier();\n`;
      }
      code += `post_out = out;`;
    }

    code += `\n}`;

    const in_buffer = t.makeBufferWithContents(
      new Uint32Array([42]),
      GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE
    );

    const out_buffer = t.makeBufferWithContents(
      new Uint32Array([...iterRange(128, x => 0)]),
      GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST
    );

    const pipeline = t.device.createComputePipeline({
      layout: 'auto',
      compute: {
        module: t.device.createShaderModule({
          code,
        }),
        entryPoint: 'main',
      },
    });

    const bg = t.device.createBindGroup({
      layout: pipeline.getBindGroupLayout(0),
      entries: [
        {
          binding: 0,
          resource: {
            buffer: in_buffer,
          },
        },
        {
          binding: 1,
          resource: {
            buffer: out_buffer,
          },
        },
      ],
    });

    const encoder = t.device.createCommandEncoder();
    const pass = encoder.beginComputePass();
    pass.setPipeline(pipeline);
    pass.setBindGroup(0, bg);
    pass.dispatchWorkgroups(1, 1, 1);
    pass.end();
    t.queue.submit([encoder.finish()]);

    // Magic number is 42 in various representations.
    const outMagicNumber = testcase.f16 ? 0x5140 : testcase.f32 ? 0x42280000 : 42;
    const expect = new Uint32Array([
      ...iterRange(128, x => {
        if (x * 4 === testcase.offset) {
          return outMagicNumber;
        } else {
          return 0;
        }
      }),
    ]);

    t.expectGPUBufferValuesEqual(out_buffer, expect);
  });