export const description = `
Validation tests for bufferArrayView
`;

import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js';
import { WGSLLanguageFeature } from '../../../../../capability_info.js';
import { Type, elementTypeOf, kAllScalarsAndVectors } from '../../../../../util/conversion.js';
import { ShaderValidationTest } from '../../../shader_validation_test.js';

export const g = makeTestGroup(ShaderValidationTest);

g.test('must_use')
  .desc('Tests that the builtin has the @must_use attribute')
  .params(u => u.combine('must_use', [true, false] as const))
  .fn(t => {
    t.skipIfLanguageFeatureNotSupported('buffer_view');
    const wgsl = `
@group(0) @binding(0) var<storage> v : buffer;
@compute @workgroup_size(1)
fn main() {
  ${t.params.must_use ? 'let p = ' : ''}bufferArrayView<array<u32>>(&v, 0, 64);
}`;

    t.expectCompileResult(t.params.must_use, wgsl);
  });

const kTypes = objectsToRecord(kAllScalarsAndVectors);

g.test('return_store_type')
  .desc('Validates return type')
  .params(u =>
    u
      .combine('template_type', keysOf(kTypes))
      .filter(t => {
        const type = kTypes[t.template_type];
        const eleType = elementTypeOf(type);
        return (
          eleType !== Type.abstractInt && eleType !== Type.abstractFloat && eleType !== Type.bool
        );
      })
      .combine('return_type', keysOf(kTypes))
      .filter(t => {
        const type = kTypes[t.return_type];
        const eleType = elementTypeOf(type);
        return (
          eleType !== Type.abstractInt && eleType !== Type.abstractFloat && eleType !== Type.bool
        );
      })
      .beginSubcases()
      .combine('ptr', [false, true] as const)
      .combine('array', [false, true] as const)
  )
  .fn(t => {
    t.skipIfLanguageFeatureNotSupported('buffer_view');
    const template_type = kTypes[t.params.template_type];
    const return_type = kTypes[t.params.return_type];
    let ret_type = return_type.toString();
    let temp_type = template_type.toString();
    if (t.params.array) {
      ret_type = `array<${return_type.toString()}>`;
      temp_type = `array<${temp_type}>`;
    }
    if (t.params.ptr) {
      ret_type = `ptr<storage, ${ret_type}>`;
    }
    let enables = ``;
    if (template_type.requiresF16() || return_type.requiresF16()) {
      enables = `enable f16;`;
    }
    const wgsl = `
${enables}
@group(0) @binding(0) var<storage> v : buffer;
@compute @workgroup_size(1)
fn main() {
  let res : ${ret_type} = bufferArrayView<${temp_type}>(&v, 0, 128);
}`;

    t.expectCompileResult(t.params.ptr && t.params.array && template_type === return_type, wgsl);
  });

g.test('buffer_type')
  .desc('Validates the buffer parameter type')
  .params(u =>
    u
      .combine('type', [
        'unsized_ro_storage',
        'unsized_storage',
        'sized_ro_storage',
        'sized_storage',
        'sized_uniform',
        'sized_workgroup',
        'override_workgroup',
      ] as const)
      .beginSubcases()
      .combine('ptr', [false, true] as const)
  )
  .fn(t => {
    t.skipIfLanguageFeatureNotSupported('buffer_view');
    const aspace = t.params.type.substr(t.params.type.lastIndexOf('_') + 1);
    let access = '';
    switch (t.params.type) {
      case 'unsized_ro_storage':
      case 'sized_ro_storage':
        access = ', read';
        break;
      case 'unsized_storage':
      case 'sized_storage':
        access = ', read_write';
        break;
      default:
        break;
    }
    const wgsl = `
@group(0) @binding(0) var<storage> unsized_ro_storage : buffer;
@group(0) @binding(1) var<storage, read_write> unsized_storage : buffer;
@group(0) @binding(2) var<storage> sized_ro_storage : buffer<128>;
@group(0) @binding(3) var<storage, read_write> sized_storage : buffer<128>;
@group(0) @binding(4) var<uniform> sized_uniform : buffer<128>;
var<workgroup> sized_workgroup : buffer<128>;
override o : u32;
var<workgroup> override_workgroup : buffer<o>;

@compute @workgroup_size(1)
fn main() {
  let p : ptr<${aspace}, array<u32>${access}> = bufferArrayView<array<u32>>(${
    t.params.ptr ? '&' : ''
  }${t.params.type}, 0, 128);
}`;

    t.expectCompileResult(t.params.ptr, wgsl);
  });

g.test('offset_type')
  .desc('Validates the offset parameter type')
  .params(u => u.combine('type', keysOf(kTypes)))
  .fn(t => {
    t.skipIfLanguageFeatureNotSupported('buffer_view');
    const type = kTypes[t.params.type];
    const enables = type.requiresF16() ? 'enable f16;' : '';
    const wgsl = `
${enables}
@group(0) @binding(0) var<storage> v : buffer;
@compute @workgroup_size(1)
fn main() {
  let p = bufferArrayView<array<u32>>(&v, ${type.create(0).wgsl()}, 128);
}`;

    t.expectCompileResult(
      type === Type.abstractInt || type === Type.u32 || type === Type.i32,
      wgsl
    );
  });

g.test('size_type')
  .desc('Validates the offset parameter type')
  .params(u => u.combine('type', keysOf(kTypes)))
  .fn(t => {
    t.skipIfLanguageFeatureNotSupported('buffer_view');
    const type = kTypes[t.params.type];
    const enables = type.requiresF16() ? 'enable f16;' : '';
    const wgsl = `
${enables}
@group(0) @binding(0) var<storage> v : buffer;
@compute @workgroup_size(1)
fn main() {
  let p = bufferArrayView<array<u32>>(&v, 0, ${type.create(128).wgsl()});
}`;

    t.expectCompileResult(
      type === Type.abstractInt || type === Type.u32 || type === Type.i32,
      wgsl
    );
  });

interface EarlyEvalCase {
  code: string;
  valid: boolean | 'pipeline';
  constants?: Record<string, GPUPipelineConstantValue>;
  ptr_param?: boolean;
}

const kEarlyEvalCases: Record<string, EarlyEvalCase> = {
  ro_storage_buffer: {
    code: `
@group(0) @binding(0) var<storage> v : buffer<16>;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, 0, 16);
}`,
    valid: true,
  },
  storage_buffer: {
    code: `
@group(0) @binding(0) var<storage, read_write> v : buffer<16>;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, 0, 16);
}`,
    valid: true,
  },
  uniform_buffer: {
    code: `
@group(0) @binding(0) var<uniform> v : buffer<16>;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, 0, 16);
}`,
    valid: true,
  },
  workgroup_buffer: {
    code: `
var<workgroup> v : buffer<16>;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, 0, 16);
}`,
    valid: true,
  },
  const_size_too_small_for_type: {
    code: `
@group(0) @binding(0) var<storage> v : buffer;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, 0, 12);
}`,
    valid: false,
  },
  buffer_too_small_for_const_size: {
    code: `
@group(0) @binding(0) var<storage> v : buffer<12>;
fn foo() {
  let p = bufferArrayView<array<u32>>(&v, 0, 16);
}`,
    valid: false,
  },
  buffer_too_small_for_type: {
    code: `
@group(0) @binding(0) var<storage> v : buffer<12>;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, 0, 16);
}`,
    valid: false,
  },
  const_size_too_small_for_const_offset_and_type: {
    code: `
var<workgroup> v : buffer<128>;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, 16, 30);
}`,
    valid: false,
  },
  buffer_size_too_small_for_const_offset_and_type: {
    code: `
var<workgroup> v : buffer<30>;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, 16, 128);
}`,
    valid: false,
  },
  override_size_too_small_for_type: {
    code: `
override size : u32;
@group(0) @binding(0) var<storage> v : buffer;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, 0, size);
}`,
    valid: 'pipeline',
    constants: { size: 12 },
  },
  buffer_too_small_for_override_size: {
    code: `
override buffer_size : u32;
override size : u32;
var<workgroup> v : buffer<buffer_size>;
fn foo() {
  let p = bufferArrayView<array<u32>>(&v, 0, size);
}`,
    valid: 'pipeline',
    constants: { buffer_size: 12, size: 16 },
  },
  override_buffer_too_small_for_type: {
    code: `
override buffer_size : u32;
override size : u32;
var<workgroup> v : buffer<buffer_size>;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, 0, size);
}`,
    valid: 'pipeline',
    constants: { buffer_size: 12, size: 16 },
  },
  override_size_too_small_for_override_offset_and_type: {
    code: `
override offset : u32;
override size : u32;
@group(0) @binding(0) var<uniform> v : buffer<128>;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, offset, size);
}`,
    valid: 'pipeline',
    constants: { offset: 16, size: 30 },
  },
  buffer_size_too_small_for_override_offset_and_type: {
    code: `
override buffer_size : u32;
override offset : u32;
var<workgroup> v : buffer<buffer_size>;
fn foo() {
  let p = bufferArrayView<array<vec4u>>(&v, offset, 128);
}`,
    valid: 'pipeline',
    constants: { buffer_size: 30, offset: 16 },
  },
  const_offset_not_aligned: {
    code: `
@group(0) @binding(0) var<storage, read_write> v : buffer;
fn foo() {
  let p = bufferarrayview<array<vec4u>>(&v, 12, 128);
}`,
    valid: false,
  },
  override_offset_not_aligned: {
    code: `
override offset : u32;
@group(0) @binding(0) var<storage, read_write> v : buffer;
fn foo() {
  let p = bufferarrayview<array<vec4u>>(&v, offset, 128);
}`,
    valid: false,
    constants: { offset: 12 },
  },
  const_offset_plus_type_size_out_of_range: {
    code: `
@group(0) @binding(0) var<storage, read_write> v : buffer;
struct S {
  a: vec4u,
  b: vec4f,
}
fn foo() {
  let p = buffer<array<S>>(&v, 4294967279, 128);
}`,
    valid: false,
  },
  override_offset_plus_type_size_out_of_range: {
    code: `
override offset : u32;
@group(0) @binding(0) var<storage, read_write> v : buffer;
struct S {
  a: vec4u,
  b: vec4f,
}
fn foo() {
  let p = buffer<array<S>>(&v, offset, 128);
}`,
    valid: false,
    constants: { offset: 4294967279 },
  },
  const_multiple_of_stride: {
    code: `
struct S {
  a : vec4u,
  b : array<vec2u>,
}
var<workgroup> v : buffer<256>;
fn foo() {
  let p = bufferArrayView<S>(&v, 0, 24);
}`,
    valid: true,
  },
  const_not_multiple_of_stride: {
    code: `
struct S {
  a : vec4u,
  b : array<vec2u>,
}
var<workgroup> v : buffer<256>;
fn foo() {
  let p = bufferArrayView<S>(&v, 0, 26);
}`,
    valid: false,
  },
  override_not_multiple_of_stride: {
    code: `
override size : u32;
struct S {
  a : vec4u,
  b : array<vec2u>,
}
var<workgroup> v : buffer<256>;
fn foo() {
  let p = bufferArrayView<S>(&v, 0, size);
}`,
    valid: 'pipeline',
    constants: { size: 26 },
  },
  const_buffer_too_small_through_unsized_function: {
    code: `
@group(0) @binding(0) var<storage> buffer<16>;
fn bar(p : ptr<storage, buffer>) {
  let q = bufferArrayView<array<u32>>(p, 0, 32);
}
fn foo() {
  bar(&v);
}`,
    valid: false,
    ptr_param: true,
  },
  override_buffer_too_small_through_unsized_function: {
    code: `
override size : u32;
@group(0) @binding(0) var<storage> v : buffer<16>;
fn bar(p : ptr<storage, buffer>) {
  let q = bufferArrayView<array<u32>>(p, 0, size);
}
fn foo() {
  bar(&v);
}`,
    valid: 'pipeline',
    constants: { size: 32 },
    ptr_param: true,
  },
  sized_buffer_to_smaller_buffer_too_small: {
    code: `
var<workgroup> v : buffer<256>;
fn bar(p : ptr<workgroup, buffer<16>>) {
  let q = bufferArrayView<array<u32>>(q, 16, 4);
}
fn foo() {
  bar(&v);
}`,
    valid: false,
    ptr_param: true,
  },
};

g.test('early_eval_errors')
  .desc('Test shader-creation and pipeline-creation errors')
  .params(u => u.combine('case', keysOf(kEarlyEvalCases)))
  .fn(t => {
    const testcase = kEarlyEvalCases[t.params.case];
    t.skipIfLanguageFeatureNotSupported('buffer_view');
    if (testcase.ptr_param === true) {
      t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
    }

    t.expectCompileResult(testcase.valid !== false, testcase.code);
    if (testcase.valid !== false) {
      t.expectPipelineResult({
        expectedResult: testcase.valid === true,
        code: testcase.code,
        constants: testcase.constants,
        statements: ['foo();'],
      });
    }
  });

interface LHSCallCase {
  code: string;
  requires?: WGSLLanguageFeature[];
}

const kLHSCallCases: Record<string, LHSCallCase> = {
  u32_array: {
    code: `
@group(0) @binding(0) var<storage, read_write> v : buffer;
@compute @workgroup_size(1)
fn main() {
  (*bufferArrayView<array<u32>>(&v, 0, 128))[0] = 123u;
}`,
  },
  f32_vector_letter: {
    code: `
@group(0) @binding(0) var<storage, read_write> v : buffer;
@compute @workgroup_size(1)
fn main() {
  (*bufferArrayView<array<vec4f>>(&v, 0, 128))[0].x = 13.3f;
}`,
  },
  f32_struct_access: {
    code: `
struct S {
  member_a: f32,
}
@group(0) @binding(0) var<storage, read_write> v : buffer;
@compute @workgroup_size(1)
fn main() {
  (*bufferArrayView<array<S>>(&v, 0, 128))[0].member_a = 13.3f;
}`,
  },
  u32_array_ptr_composite: {
    code: `
@group(0) @binding(0) var<storage, read_write> v : buffer;
@compute @workgroup_size(1)
fn main() {
  bufferArrayView<array<u32>>(&v, 0, 128)[0] = 123u;
}`,
    requires: ['pointer_composite_access'],
  },
  vec4u_swizzle_assign: {
    code: `
@group(0) @binding(0) var<storage, read_write> v : buffer;
@compute @workgroup_size(1)
fn main() {
  (*bufferArrayView<array<vec4u>>(&v, 0, 128))[0].zxy = vec3u(42);
}`,
    requires: ['swizzle_assignment'],
  },
  vec4u_ptr_composite_swizzle_assign: {
    code: `
@group(0) @binding(0) var<storage, read_write> v : buffer;
@compute @workgroup_size(1)
fn main() {
  bufferArrayView<array<vec4u>>(&v, 0, 128)[0].zxy = vec3u(42);
}`,
    requires: ['pointer_composite_access', 'swizzle_assignment'],
  },
  compound_assign: {
    code: `
var<workgroup> v : buffer<256>;
@compute @workgroup_size(1)
fn main() {
  (*bufferArrayView<array<u32>>(&v, 0, 128))[0] += 123u;
}`,
  },
};

g.test('lhs_call')
  .desc('Validate that bufferView can be on the LHS of an assignment')
  .params(u => u.combine('case', keysOf(kLHSCallCases)))
  .fn(t => {
    t.skipIfLanguageFeatureNotSupported('buffer_view');

    const testcase = kLHSCallCases[t.params.case];
    let valid = true;
    const features = testcase.requires ?? [];
    features.forEach(f => {
      valid &&= t.hasLanguageFeature(f);
    });
    t.expectCompileResult(valid, testcase.code);
  });
