src/tint/lang/msl/msl.def - dawn - Git at Google

 // 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 vec2<T>
 type vec3<T>
 type vec4<T>
 type atomic<T>
 type ptr<S: address_space, T, A: access>
 type sampler
 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 level

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

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

 ////////////////////////////////////////////////////////////////////////////////
 // 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 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_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 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>

 @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>)

 @stage("compute") fn threadgroup_barrier(u32)
	// 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 vec2<T>
	type vec3<T>
	type vec4<T>
	type atomic<T>
	type ptr<S: address_space, T, A: access>
	type sampler
	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 level

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

	match iu32: i32 \| u32
	match fiu32: f32 \| i32 \| u32

	////////////////////////////////////////////////////////////////////////////////
	// 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 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_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 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>

	@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>)

	@stage("compute") fn threadgroup_barrier(u32)