// Copyright 2023 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
//    list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
//    this list of conditions and the following disclaimer in the documentation
//    and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
//    contributors may be used to endorse or promote products derived from
//    this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

////////////////////////////////////////////////////////////////////////////////
// MSL builtin definition file                                                //
//
// After modifying this file, run:                                            //
//    tools/run gen                                                           //
// from the Dawn source directory.                                            //
//                                                                            //
// See docs/tint/intrinsic_definition_files.md for syntax                     //
////////////////////////////////////////////////////////////////////////////////

import "src/tint/lang/core/access.def"
import "src/tint/lang/core/address_space.def"
import "src/tint/lang/core/texel_format.def"

////////////////////////////////////////////////////////////////////////////////
// Enum matchers                                                              //
////////////////////////////////////////////////////////////////////////////////

match read_write: access.read_write
match readable
  : access.read
  | access.read_write
match writable
  : access.write
  | access.read_write

match function: address_space.function
match workgroup_or_storage
  : address_space.workgroup
  | address_space.storage

match f32_texel_format
  : texel_format.r8unorm
  | texel_format.bgra8unorm
  | texel_format.rgba8unorm
  | texel_format.rgba8snorm
  | texel_format.rgba16float
  | texel_format.r32float
  | texel_format.rg32float
  | texel_format.rgba32float
match i32_texel_format
  : texel_format.rgba8sint
  | texel_format.rgba16sint
  | texel_format.r32sint
  | texel_format.rg32sint
  | texel_format.rgba32sint
match u32_texel_format
  : texel_format.rgba8uint
  | texel_format.rgba16uint
  | texel_format.r32uint
  | texel_format.rg32uint
  | texel_format.rgba32uint

////////////////////////////////////////////////////////////////////////////////
// Types                                                                      //
////////////////////////////////////////////////////////////////////////////////

type bool
type i32
type u32
type f32
type f16
type vec2<T>
type vec3<T>
type vec4<T>
@display("vec{N}<{T}>")     type vec<N: num, T>
type atomic<T>
type ptr<S: address_space, T, A: access>
type sampler
type sampler_comparison
type texture_1d<T>
type texture_2d<T>
type texture_2d_array<T>
type texture_3d<T>
type texture_cube<T>
type texture_cube_array<T>
type texture_depth_2d
type texture_depth_2d_array
type texture_depth_cube
type texture_depth_cube_array
type texture_depth_multisampled_2d
type texture_multisampled_2d<T>
type texture_storage_1d<F: texel_format, A: access>
type texture_storage_2d<F: texel_format, A: access>
type texture_storage_2d_array<F: texel_format, A: access>
type texture_storage_3d<F: texel_format, A: access>

// sampling options types
type bias
type gradient2d
type gradient3d
type gradientcube
type level

////////////////////////////////////////////////////////////////////////////////
// Type matchers                                                              //
////////////////////////////////////////////////////////////////////////////////

match iu32: i32 | u32
match fiu32: f32 | i32 | u32
match f32_f16: f32 | f16

////////////////////////////////////////////////////////////////////////////////
// Builtin Functions                                                          //
////////////////////////////////////////////////////////////////////////////////
@stage("fragment", "compute") fn atomic_compare_exchange_weak_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, ptr<function, T, read_write>, T, u32, u32) -> bool
@stage("fragment", "compute") fn atomic_exchange_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
@stage("fragment", "compute") fn atomic_fetch_add_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
@stage("fragment", "compute") fn atomic_fetch_and_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
@stage("fragment", "compute") fn atomic_fetch_max_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
@stage("fragment", "compute") fn atomic_fetch_min_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
@stage("fragment", "compute") fn atomic_fetch_or_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
@stage("fragment", "compute") fn atomic_fetch_sub_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
@stage("fragment", "compute") fn atomic_fetch_xor_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
@stage("fragment", "compute") fn atomic_load_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T
@stage("fragment", "compute") fn atomic_store_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32)

@member_function fn fence[F: texel_format, A: access](texture: texture_storage_1d<F, A>)
@member_function fn fence[F: texel_format, A: access](texture: texture_storage_2d<F, A>)
@member_function fn fence[F: texel_format, A: access](texture: texture_storage_2d_array<F, A>)
@member_function fn fence[F: texel_format, A: access](texture: texture_storage_3d<F, A>)

@member_function fn gather[T: fiu32](texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>, component: u32) -> vec4<T>
@member_function fn gather[T: fiu32, A: iu32](texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>, component: u32) -> vec4<T>
@member_function fn gather[T: fiu32](texture: texture_cube<T>, sampler: sampler, coords: vec3<f32>, component: u32) -> vec4<T>
@member_function fn gather[T: fiu32, A: iu32](texture: texture_cube_array<T>, sampler: sampler, coords: vec3<f32>, array_index: A, component: u32) -> vec4<T>
@member_function fn gather(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32>
@member_function fn gather[A: iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> vec4<f32>
@member_function fn gather(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
@member_function fn gather[A: iu32](texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: A) -> vec4<f32>

@member_function fn gather_compare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> vec4<f32>
@member_function fn gather_compare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> vec4<f32>
@member_function fn gather_compare[A: iu32](texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: A, depth_ref: f32) -> vec4<f32>
@member_function fn gather_compare[A: iu32](texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: A, depth_ref: f32, @const offset: vec2<i32>) -> vec4<f32>
@member_function fn gather_compare(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> vec4<f32>
@member_function fn gather_compare[A: iu32](texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: A, depth_ref: f32) -> vec4<f32>

@member_function fn get_width[T: fiu32](texture: texture_1d<T>) -> u32
@member_function fn get_width[T: fiu32](texture: texture_2d<T>, u32) -> u32
@member_function fn get_width[T: fiu32](texture: texture_2d_array<T>, u32) -> u32
@member_function fn get_width[T: fiu32](texture: texture_3d<T>, u32) -> u32
@member_function fn get_width[T: fiu32](texture: texture_cube<T>, u32) -> u32
@member_function fn get_width[T: fiu32](texture: texture_cube_array<T>, u32) -> u32
@member_function fn get_width[T: fiu32](texture: texture_multisampled_2d<T>) -> u32
@member_function fn get_width(texture: texture_depth_2d, u32) -> u32
@member_function fn get_width(texture: texture_depth_2d_array, u32) -> u32
@member_function fn get_width(texture: texture_depth_cube, u32) -> u32
@member_function fn get_width(texture: texture_depth_cube_array, u32) -> u32
@member_function fn get_width(texture: texture_depth_multisampled_2d) -> u32
@member_function fn get_width[F: texel_format, A: access](texture: texture_storage_1d<F, A>) -> u32
@member_function fn get_width[F: texel_format, A: access](texture: texture_storage_2d<F, A>, u32) -> u32
@member_function fn get_width[F: texel_format, A: access](texture: texture_storage_2d_array<F, A>, u32) -> u32
@member_function fn get_width[F: texel_format, A: access](texture: texture_storage_3d<F, A>, u32) -> u32

@member_function fn get_height[T: fiu32](texture: texture_2d<T>, u32) -> u32
@member_function fn get_height[T: fiu32](texture: texture_2d_array<T>, u32) -> u32
@member_function fn get_height[T: fiu32](texture: texture_3d<T>, u32) -> u32
@member_function fn get_height[T: fiu32](texture: texture_cube<T>, u32) -> u32
@member_function fn get_height[T: fiu32](texture: texture_cube_array<T>, u32) -> u32
@member_function fn get_height[T: fiu32](texture: texture_multisampled_2d<T>) -> u32
@member_function fn get_height(texture: texture_depth_2d, u32) -> u32
@member_function fn get_height(texture: texture_depth_2d_array, u32) -> u32
@member_function fn get_height(texture: texture_depth_cube, u32) -> u32
@member_function fn get_height(texture: texture_depth_cube_array, u32) -> u32
@member_function fn get_height(texture: texture_depth_multisampled_2d) -> u32
@member_function fn get_height[F: texel_format, A: access](texture: texture_storage_2d<F, A>, u32) -> u32
@member_function fn get_height[F: texel_format, A: access](texture: texture_storage_2d_array<F, A>, u32) -> u32
@member_function fn get_height[F: texel_format, A: access](texture: texture_storage_3d<F, A>, u32) -> u32

@member_function fn get_depth[T: fiu32](texture: texture_3d<T>, u32) -> u32
@member_function fn get_depth[F: texel_format, A: access](texture: texture_storage_3d<F, A>, u32) -> u32

@member_function fn get_array_size[T: fiu32](texture: texture_2d_array<T>) -> u32
@member_function fn get_array_size[T: fiu32](texture: texture_cube_array<T>) -> u32
@member_function fn get_array_size(texture: texture_depth_2d_array) -> u32
@member_function fn get_array_size(texture: texture_depth_cube_array) -> u32
@member_function fn get_array_size[F: texel_format, A: access](texture: texture_storage_2d_array<F, A>) -> u32

@member_function fn get_num_mip_levels[T: fiu32](texture: texture_1d<T>) -> u32
@member_function fn get_num_mip_levels[T: fiu32](texture: texture_2d<T>) -> u32
@member_function fn get_num_mip_levels[T: fiu32](texture: texture_2d_array<T>) -> u32
@member_function fn get_num_mip_levels[T: fiu32](texture: texture_3d<T>) -> u32
@member_function fn get_num_mip_levels[T: fiu32](texture: texture_cube<T>) -> u32
@member_function fn get_num_mip_levels[T: fiu32](texture: texture_cube_array<T>) -> u32
@member_function fn get_num_mip_levels(texture: texture_depth_2d) -> u32
@member_function fn get_num_mip_levels(texture: texture_depth_2d_array) -> u32
@member_function fn get_num_mip_levels(texture: texture_depth_cube) -> u32
@member_function fn get_num_mip_levels(texture: texture_depth_cube_array) -> u32

@member_function fn get_num_samples[T: fiu32](texture: texture_multisampled_2d<T>) -> u32
@member_function fn get_num_samples(texture: texture_depth_multisampled_2d) -> u32

@member_function fn read[T: fiu32](texture: texture_1d<T>, coords: u32) -> vec4<T>
@member_function fn read[T: fiu32, L: iu32](texture: texture_2d<T>, coords: vec2<u32>, level: L) -> vec4<T>
@member_function fn read[T: fiu32, A: iu32, L: iu32](texture: texture_2d_array<T>, coords: vec2<u32>, array_index: A, level: L) -> vec4<T>
@member_function fn read[T: fiu32, L: iu32](texture: texture_3d<T>, coords: vec3<u32>, level: L) -> vec4<T>
@member_function fn read[T: fiu32, S: iu32](texture: texture_multisampled_2d<T>, coords: vec2<u32>, sample_index: S) -> vec4<T>
@member_function fn read[L: iu32](texture: texture_depth_2d, coords: vec2<u32>, level: L) -> f32
@member_function fn read[A: iu32, L: iu32](texture: texture_depth_2d_array, coords: vec2<u32>, array_index: A, level: L) -> f32
@member_function fn read[S: iu32](texture: texture_depth_multisampled_2d, coords: vec2<u32>, sample_index: S) -> f32
@member_function fn read(texture: texture_storage_1d<f32_texel_format, readable>, coords: u32) -> vec4<f32>
@member_function fn read(texture: texture_storage_2d<f32_texel_format, readable>, coords: vec2<u32>) -> vec4<f32>
@member_function fn read[A: iu32](texture: texture_storage_2d_array<f32_texel_format, readable>, coords: vec2<u32>, array_index: A) -> vec4<f32>
@member_function fn read(texture: texture_storage_3d<f32_texel_format, readable>, coords: vec3<u32>) -> vec4<f32>
@member_function fn read(texture: texture_storage_1d<i32_texel_format, readable>, coords: u32) -> vec4<i32>
@member_function fn read(texture: texture_storage_2d<i32_texel_format, readable>, coords: vec2<u32>) -> vec4<i32>
@member_function fn read[A: iu32](texture: texture_storage_2d_array<i32_texel_format, readable>, coords: vec2<u32>, array_index: A) -> vec4<i32>
@member_function fn read(texture: texture_storage_3d<i32_texel_format, readable>, coords: vec3<u32>) -> vec4<i32>
@member_function fn read(texture: texture_storage_1d<u32_texel_format, readable>, coords: u32) -> vec4<u32>
@member_function fn read(texture: texture_storage_2d<u32_texel_format, readable>, coords: vec2<u32>) -> vec4<u32>
@member_function fn read[A: iu32](texture: texture_storage_2d_array<u32_texel_format, readable>, coords: vec2<u32>, array_index: A) -> vec4<u32>
@member_function fn read(texture: texture_storage_3d<u32_texel_format, readable>, coords: vec3<u32>) -> vec4<u32>

// Sample with implicit LOD.
@member_function @stage("fragment") fn sample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32>
@member_function @stage("fragment") fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
@member_function @stage("fragment") fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32>
@member_function @stage("fragment") fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A) -> vec4<f32>
@member_function @stage("fragment") fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> vec4<f32>
@member_function @stage("fragment") fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
@member_function @stage("fragment") fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32>
@member_function @stage("fragment") fn sample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
@member_function @stage("fragment") fn sample[A: iu32](texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: A) -> vec4<f32>
@member_function @stage("fragment") fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> f32
@member_function @stage("fragment") fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> f32
@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A) -> f32
@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> f32
@member_function @stage("fragment") fn sample(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> f32
@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: A) -> f32

// Sample with explicit LOD.
@member_function fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, level: level) -> vec4<f32>
@member_function fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, level: level, @const offset: vec2<i32>) -> vec4<f32>
@member_function fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, level: level) -> vec4<f32>
@member_function fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, level: level, @const offset: vec2<i32>) -> vec4<f32>
@member_function fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, level: level) -> vec4<f32>
@member_function fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, level: level, @const offset: vec3<i32>) -> vec4<f32>
@member_function fn sample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>, level: level) -> vec4<f32>
@member_function fn sample[A: iu32](texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: A, level: level) -> vec4<f32>
@member_function fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, level: level) -> f32
@member_function fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, level: level, @const offset: vec2<i32>) -> f32
@member_function fn sample[A: iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A, level: level) -> f32
@member_function fn sample[A: iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A, level: level, @const offset: vec2<i32>) -> f32
@member_function fn sample(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>, level: level) -> f32
@member_function fn sample[A: iu32](texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: A, level: level) -> f32

// Sample with mip level bias.
@member_function fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, bias: bias) -> vec4<f32>
@member_function fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, bias: bias, @const offset: vec2<i32>) -> vec4<f32>
@member_function fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, bias: bias) -> vec4<f32>
@member_function fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, bias: bias, @const offset: vec2<i32>) -> vec4<f32>
@member_function fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, bias: bias) -> vec4<f32>
@member_function fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, bias: bias, @const offset: vec3<i32>) -> vec4<f32>
@member_function fn sample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>, bias: bias) -> vec4<f32>
@member_function fn sample[A: iu32](texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: A, bias: bias) -> vec4<f32>

// Sample with explicit derivatives.
@member_function fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, gradient2d) -> vec4<f32>
@member_function fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, gradient2d, @const offset: vec2<i32>) -> vec4<f32>
@member_function fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, gradient2d) -> vec4<f32>
@member_function fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, gradient2d, @const offset: vec2<i32>) -> vec4<f32>
@member_function fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, gradient3d) -> vec4<f32>
@member_function fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, gradient3d, @const offset: vec3<i32>) -> vec4<f32>
@member_function fn sample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>, gradientcube) -> vec4<f32>
@member_function fn sample[A: iu32](texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: A, gradientcube) -> vec4<f32>

// Sample compare with implicit LOD.
@member_function @stage("fragment") fn sample_compare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> f32
@member_function @stage("fragment") fn sample_compare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> f32
@member_function @stage("fragment") fn sample_compare[A: iu32](texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: A, depth_ref: f32) -> f32
@member_function @stage("fragment") fn sample_compare[A: iu32](texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: A, depth_ref: f32, @const offset: vec2<i32>) -> f32
@member_function @stage("fragment") fn sample_compare(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> f32
@member_function @stage("fragment") fn sample_compare[A: iu32](texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: A, depth_ref: f32) -> f32

// Sample compare with explicit LOD.
@member_function fn sample_compare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, level: level) -> f32
@member_function fn sample_compare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, level: level, @const offset: vec2<i32>) -> f32
@member_function fn sample_compare[A: iu32](texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: A, depth_ref: f32, level: level) -> f32
@member_function fn sample_compare[A: iu32](texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: A, depth_ref: f32, level: level, @const offset: vec2<i32>) -> f32
@member_function fn sample_compare(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32, level: level) -> f32
@member_function fn sample_compare[A: iu32](texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: A, depth_ref: f32, level: level) -> f32

@member_function fn write(texture: texture_storage_1d<f32_texel_format, writable>, value: vec4<f32>, coords: u32)
@member_function fn write(texture: texture_storage_2d<f32_texel_format, writable>, value: vec4<f32>, coords: vec2<u32>)
@member_function fn write[A: iu32](texture: texture_storage_2d_array<f32_texel_format, writable>, value: vec4<f32>, coords: vec2<u32>, array_index: A)
@member_function fn write(texture: texture_storage_3d<f32_texel_format, writable>, value: vec4<f32>, coords: vec3<u32>)
@member_function fn write(texture: texture_storage_1d<i32_texel_format, writable>, value: vec4<i32>, coords: u32)
@member_function fn write(texture: texture_storage_2d<i32_texel_format, writable>, value: vec4<i32>, coords: vec2<u32>)
@member_function fn write[A: iu32](texture: texture_storage_2d_array<i32_texel_format, writable>, value: vec4<i32>, coords: vec2<u32>, array_index: A)
@member_function fn write(texture: texture_storage_3d<i32_texel_format, writable>, value: vec4<i32>, coords: vec3<u32>)
@member_function fn write(texture: texture_storage_1d<u32_texel_format, writable>, value: vec4<u32>, coords: u32)
@member_function fn write(texture: texture_storage_2d<u32_texel_format, writable>, value: vec4<u32>, coords: vec2<u32>)
@member_function fn write[A: iu32](texture: texture_storage_2d_array<u32_texel_format, writable>, value: vec4<u32>, coords: vec2<u32>, array_index: A)
@member_function fn write(texture: texture_storage_3d<u32_texel_format, writable>, value: vec4<u32>, coords: vec3<u32>)

fn distance[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> T
fn length[N: num, T: f32_f16](vec<N, T>) -> T
@stage("compute") fn threadgroup_barrier(u32)