blob: 47f6c3c2c064b5e035817beae53b60362dd5cf9d [file] [log] [blame]
dan sinclair0ebbc5c2023-11-29 17:53:24 +00001// Copyright 2023 The Dawn & Tint Authors
2//
3// Redistribution and use in source and binary forms, with or without
4// modification, are permitted provided that the following conditions are met:
5//
6// 1. Redistributions of source code must retain the above copyright notice, this
7// list of conditions and the following disclaimer.
8//
9// 2. Redistributions in binary form must reproduce the above copyright notice,
10// this list of conditions and the following disclaimer in the documentation
11// and/or other materials provided with the distribution.
12//
13// 3. Neither the name of the copyright holder nor the names of its
14// contributors may be used to endorse or promote products derived from
15// this software without specific prior written permission.
16//
17// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
18// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
19// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
20// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
21// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
22// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
23// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
24// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
25// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
26// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27
28////////////////////////////////////////////////////////////////////////////////
29// MSL builtin definition file //
Ben Claytone4351242024-02-21 08:11:33 +000030//
31// After modifying this file, run: //
32// tools/run gen //
33// from the Dawn source directory. //
34// //
35// See docs/tint/intrinsic_definition_files.md for syntax //
dan sinclair0ebbc5c2023-11-29 17:53:24 +000036////////////////////////////////////////////////////////////////////////////////
37
James Price0272fd22024-06-05 12:24:04 +000038import "src/tint/lang/core/access.def"
39import "src/tint/lang/core/address_space.def"
James Price456cf282024-06-11 18:24:06 +000040import "src/tint/lang/core/texel_format.def"
James Price0272fd22024-06-05 12:24:04 +000041
42////////////////////////////////////////////////////////////////////////////////
43// Enum matchers //
44////////////////////////////////////////////////////////////////////////////////
45
46match read_write: access.read_write
47
James Pricefae957f2024-06-05 12:24:04 +000048match function: address_space.function
James Price0272fd22024-06-05 12:24:04 +000049match workgroup_or_storage
50 : address_space.workgroup
51 | address_space.storage
52
53////////////////////////////////////////////////////////////////////////////////
54// Types //
55////////////////////////////////////////////////////////////////////////////////
56
James Pricefae957f2024-06-05 12:24:04 +000057type bool
James Price0272fd22024-06-05 12:24:04 +000058type i32
dan sinclair0ebbc5c2023-11-29 17:53:24 +000059type u32
James Price54883c22024-06-11 18:24:06 +000060type f32
61type vec2<T>
62type vec3<T>
63type vec4<T>
James Price0272fd22024-06-05 12:24:04 +000064type atomic<T>
65type ptr<S: address_space, T, A: access>
James Price54883c22024-06-11 18:24:06 +000066type sampler
67type texture_1d<T>
68type texture_2d<T>
69type texture_2d_array<T>
70type texture_3d<T>
71type texture_cube<T>
72type texture_cube_array<T>
73type texture_depth_2d
74type texture_depth_2d_array
75type texture_depth_cube
76type texture_depth_cube_array
James Price456cf282024-06-11 18:24:06 +000077type texture_depth_multisampled_2d
78type texture_multisampled_2d<T>
79type texture_storage_1d<F: texel_format, A: access>
80type texture_storage_2d<F: texel_format, A: access>
81type texture_storage_2d_array<F: texel_format, A: access>
82type texture_storage_3d<F: texel_format, A: access>
James Price0272fd22024-06-05 12:24:04 +000083
84////////////////////////////////////////////////////////////////////////////////
85// Type matchers //
86////////////////////////////////////////////////////////////////////////////////
87
88match iu32: i32 | u32
James Price456cf282024-06-11 18:24:06 +000089match fiu32: f32 | i32 | u32
dan sinclair0ebbc5c2023-11-29 17:53:24 +000090
91////////////////////////////////////////////////////////////////////////////////
92// Builtin Functions //
93////////////////////////////////////////////////////////////////////////////////
James Pricefae957f2024-06-05 12:24:04 +000094@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
James Price0272fd22024-06-05 12:24:04 +000095@stage("fragment", "compute") fn atomic_exchange_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
96@stage("fragment", "compute") fn atomic_fetch_add_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
97@stage("fragment", "compute") fn atomic_fetch_and_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
98@stage("fragment", "compute") fn atomic_fetch_max_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
99@stage("fragment", "compute") fn atomic_fetch_min_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
100@stage("fragment", "compute") fn atomic_fetch_or_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
101@stage("fragment", "compute") fn atomic_fetch_sub_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
102@stage("fragment", "compute") fn atomic_fetch_xor_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
103@stage("fragment", "compute") fn atomic_load_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T
104@stage("fragment", "compute") fn atomic_store_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32)
James Price54883c22024-06-11 18:24:06 +0000105
James Price456cf282024-06-11 18:24:06 +0000106@member_function fn get_width[T: fiu32](texture: texture_1d<T>) -> u32
107@member_function fn get_width[T: fiu32](texture: texture_2d<T>, u32) -> u32
108@member_function fn get_width[T: fiu32](texture: texture_2d_array<T>, u32) -> u32
109@member_function fn get_width[T: fiu32](texture: texture_3d<T>, u32) -> u32
110@member_function fn get_width[T: fiu32](texture: texture_cube<T>, u32) -> u32
111@member_function fn get_width[T: fiu32](texture: texture_cube_array<T>, u32) -> u32
112@member_function fn get_width[T: fiu32](texture: texture_multisampled_2d<T>) -> u32
113@member_function fn get_width(texture: texture_depth_2d, u32) -> u32
114@member_function fn get_width(texture: texture_depth_2d_array, u32) -> u32
115@member_function fn get_width(texture: texture_depth_cube, u32) -> u32
116@member_function fn get_width(texture: texture_depth_cube_array, u32) -> u32
117@member_function fn get_width(texture: texture_depth_multisampled_2d) -> u32
118@member_function fn get_width[F: texel_format, A: access](texture: texture_storage_1d<F, A>) -> u32
119@member_function fn get_width[F: texel_format, A: access](texture: texture_storage_2d<F, A>, u32) -> u32
120@member_function fn get_width[F: texel_format, A: access](texture: texture_storage_2d_array<F, A>, u32) -> u32
121@member_function fn get_width[F: texel_format, A: access](texture: texture_storage_3d<F, A>, u32) -> u32
122
123@member_function fn get_height[T: fiu32](texture: texture_2d<T>, u32) -> u32
124@member_function fn get_height[T: fiu32](texture: texture_2d_array<T>, u32) -> u32
125@member_function fn get_height[T: fiu32](texture: texture_3d<T>, u32) -> u32
126@member_function fn get_height[T: fiu32](texture: texture_cube<T>, u32) -> u32
127@member_function fn get_height[T: fiu32](texture: texture_cube_array<T>, u32) -> u32
128@member_function fn get_height[T: fiu32](texture: texture_multisampled_2d<T>) -> u32
129@member_function fn get_height(texture: texture_depth_2d, u32) -> u32
130@member_function fn get_height(texture: texture_depth_2d_array, u32) -> u32
131@member_function fn get_height(texture: texture_depth_cube, u32) -> u32
132@member_function fn get_height(texture: texture_depth_cube_array, u32) -> u32
133@member_function fn get_height(texture: texture_depth_multisampled_2d) -> u32
134@member_function fn get_height[F: texel_format, A: access](texture: texture_storage_2d<F, A>, u32) -> u32
135@member_function fn get_height[F: texel_format, A: access](texture: texture_storage_2d_array<F, A>, u32) -> u32
136@member_function fn get_height[F: texel_format, A: access](texture: texture_storage_3d<F, A>, u32) -> u32
137
138@member_function fn get_depth[T: fiu32](texture: texture_3d<T>, u32) -> u32
139@member_function fn get_depth[F: texel_format, A: access](texture: texture_storage_3d<F, A>, u32) -> u32
140
James Price54883c22024-06-11 18:24:06 +0000141@member_function @stage("fragment") fn sample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32>
142@member_function @stage("fragment") fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
143@member_function @stage("fragment") fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32>
144@member_function @stage("fragment") fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A) -> vec4<f32>
145@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>
146@member_function @stage("fragment") fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
147@member_function @stage("fragment") fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32>
148@member_function @stage("fragment") fn sample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
149@member_function @stage("fragment") fn sample[A: iu32](texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: A) -> vec4<f32>
150@member_function @stage("fragment") fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> f32
151@member_function @stage("fragment") fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> f32
152@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A) -> f32
153@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
154@member_function @stage("fragment") fn sample(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> f32
155@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: A) -> f32
156
dan sinclair0ebbc5c2023-11-29 17:53:24 +0000157@stage("compute") fn threadgroup_barrier(u32)
158