blob: 8225c604e0bdf831084b27a23f7f1b2575dbab03 [file] [log] [blame]
// 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.
syntax = "proto3";
package tint.core.ir.binary.pb;
message Module {
repeated Type types = 1;
repeated Value values = 2;
repeated ConstantValue constant_values = 3;
repeated Function functions = 4;
repeated Block blocks = 5;
uint32 root_block = 6; // Module.blocks
}
////////////////////////////////////////////////////////////////////////////////
// Types
////////////////////////////////////////////////////////////////////////////////
message Type {
oneof kind {
TypeBasic basic = 1;
TypeVector vector = 2;
TypeMatrix matrix = 3;
TypeArray array = 4;
TypePointer pointer = 5;
TypeStruct struct = 6;
TypeAtomic atomic = 7;
TypeDepthTexture depth_texture = 8;
TypeSampledTexture sampled_texture = 9;
TypeMultisampledTexture multisampled_texture = 10;
TypeDepthMultisampledTexture depth_multisampled_texture = 11;
TypeStorageTexture storage_texture = 12;
TypeExternalTexture external_texture = 13;
TypeSampler sampler = 14;
TypeInputAttachment input_attachment = 15;
TypeSubgroupMatrix subgroup_matrix_left = 16;
TypeSubgroupMatrix subgroup_matrix_right = 17;
TypeSubgroupMatrix subgroup_matrix_result = 18;
}
}
// Non-compound types
enum TypeBasic {
void = 0;
bool = 1;
i32 = 2;
u32 = 3;
f32 = 4;
f16 = 5;
}
message TypeVector {
uint32 width = 1;
uint32 element_type = 2; // Module.types
}
message TypeMatrix {
uint32 num_columns = 1;
uint32 num_rows = 2;
uint32 element_type = 3; // Module.types
}
message TypeArray {
uint32 element = 1; // Module.types
uint32 stride = 2;
uint32 count = 3;
}
message TypePointer {
AddressSpace address_space = 1;
uint32 store_type = 2; // Module.types
AccessControl access = 3;
}
message TypeStruct {
string name = 1;
repeated TypeStructMember member = 2;
}
message TypeStructMember {
string name = 1;
uint32 type = 2;
uint32 size = 3;
uint32 align = 4;
optional AttributesStructMember attributes = 5;
}
message TypeAtomic {
uint32 type = 1; // Module.types
}
message TypeDepthTexture {
TextureDimension dimension = 1;
}
message TypeSampledTexture {
TextureDimension dimension = 1;
uint32 sub_type = 2; // Module.types
}
message TypeMultisampledTexture {
TextureDimension dimension = 1;
uint32 sub_type = 2; // Module.types
}
message TypeDepthMultisampledTexture {
TextureDimension dimension = 1;
}
message TypeStorageTexture {
TextureDimension dimension = 1;
TexelFormat texel_format = 2;
AccessControl access = 3;
}
message TypeExternalTexture {}
message TypeSampler {
SamplerKind kind = 1;
}
message TypeInputAttachment {
uint32 sub_type = 1; // Module.types
}
message TypeSubgroupMatrix {
uint32 sub_type = 1; // Module.types
uint32 rows = 2;
uint32 columns = 3;
}
////////////////////////////////////////////////////////////////////////////////
// Values
////////////////////////////////////////////////////////////////////////////////
message Value {
oneof kind {
uint32 function = 1; // Module.functions
InstructionResult instruction_result = 2;
FunctionParameter function_parameter = 3;
BlockParameter block_parameter = 4;
uint32 constant = 5; // Module.constant_values
}
}
message InstructionResult {
uint32 type = 1; // Module.types
optional string name = 2;
}
message FunctionParameter {
uint32 type = 1; // Module.types
optional string name = 2;
optional AttributesFunctionParameter attributes = 3;
}
message BlockParameter {
uint32 type = 1; // Module.types
optional string name = 2;
}
////////////////////////////////////////////////////////////////////////////////
// ConstantValues
////////////////////////////////////////////////////////////////////////////////
message ConstantValue {
oneof kind {
ConstantValueScalar scalar = 1;
ConstantValueComposite composite = 2;
ConstantValueSplat splat = 3;
}
}
message ConstantValueScalar {
oneof kind {
bool bool = 1;
int32 i32 = 2;
uint32 u32 = 3;
float f32 = 4;
float f16 = 5;
}
}
message ConstantValueComposite {
uint32 type = 1; // Module.types
repeated uint32 elements = 2; // Module.constant_values
}
message ConstantValueSplat {
uint32 type = 1; // Module.types
uint32 elements = 2; // Module.constant_values
uint32 count = 3; // splat count
}
////////////////////////////////////////////////////////////////////////////////
// Functions
////////////////////////////////////////////////////////////////////////////////
message Function {
uint32 return_type = 1; // Module.types
uint32 block = 2; // Module.blocks
optional string name = 3;
optional PipelineStage pipeline_stage = 4;
optional WorkgroupSize workgroup_size = 5;
repeated uint32 parameters = 6; // Module.values
optional uint32 return_location = 7;
optional Interpolation return_interpolation = 8;
optional BuiltinValue return_builtin = 9;
bool return_invariant = 10;
}
enum PipelineStage {
Compute = 0;
Fragment = 1;
Vertex = 2;
}
message WorkgroupSize {
uint32 x = 1;
uint32 y = 2;
uint32 z = 3;
}
////////////////////////////////////////////////////////////////////////////////
// Blocks
////////////////////////////////////////////////////////////////////////////////
message Block {
repeated uint32 parameters = 1; // Module.values
repeated Instruction instructions = 2;
bool is_multi_in = 3;
}
////////////////////////////////////////////////////////////////////////////////
// Instructions
////////////////////////////////////////////////////////////////////////////////
message Instruction {
repeated uint32 operands = 1; // Module.values
repeated uint32 results = 2; // Module.values
oneof kind {
InstructionReturn return = 3;
InstructionUnary unary = 4;
InstructionBinary binary = 5;
InstructionDiscard discard = 6;
InstructionLet let = 7;
InstructionVar var = 8;
InstructionBitcast bitcast = 9;
InstructionConstruct construct = 10;
InstructionConvert convert = 11;
InstructionAccess access = 12;
InstructionUserCall user_call = 13;
InstructionBuiltinCall builtin_call = 14;
InstructionLoad load = 15;
InstructionStore store = 16;
InstructionLoadVectorElement load_vector_element = 17;
InstructionStoreVectorElement store_vector_element = 18;
InstructionSwizzle swizzle = 19;
InstructionIf if = 20;
InstructionSwitch switch = 21;
InstructionLoop loop = 22;
InstructionExitIf exit_if = 23;
InstructionExitSwitch exit_switch = 24;
InstructionExitLoop exit_loop = 25;
InstructionNextIteration next_iteration = 26;
InstructionContinue continue = 27;
InstructionBreakIf break_if = 28;
InstructionUnreachable unreachable = 29;
}
}
message InstructionReturn {}
message InstructionUnary {
UnaryOp op = 1;
}
message InstructionBinary {
BinaryOp op = 1;
}
message InstructionBitcast {}
message InstructionConstruct {}
message InstructionDiscard {}
message InstructionLet {}
message InstructionVar {
optional BindingPoint binding_point = 1;
optional uint32 input_attachment_index = 2;
}
message InstructionConvert {}
message InstructionAccess {}
message InstructionUserCall {}
message InstructionBuiltinCall {
BuiltinFn builtin = 1;
}
message InstructionLoad {}
message InstructionStore {}
message InstructionLoadVectorElement {}
message InstructionStoreVectorElement {}
message InstructionSwizzle {
repeated uint32 indices = 1;
}
message InstructionIf {
optional uint32 true = 1; // Module.blocks
optional uint32 false = 2; // Module.blocks
}
message InstructionSwitch {
repeated SwitchCase cases = 1; // Module.blocks
}
message InstructionLoop {
optional uint32 initializer = 1; // Module.blocks
optional uint32 body = 2; // Module.blocks
optional uint32 continuing = 3; // Module.blocks
}
message InstructionExitIf {}
message InstructionExitSwitch {}
message InstructionExitLoop {}
message SwitchCase {
uint32 block = 1; // Module.blocks
repeated uint32 selectors = 2; // Module.constant_values
bool is_default = 3;
}
message BindingPoint {
uint32 group = 1;
uint32 binding = 2;
}
message InstructionNextIteration {}
message InstructionContinue {}
message InstructionBreakIf {
uint32 num_next_iter_values = 1;
}
message InstructionUnreachable {}
////////////////////////////////////////////////////////////////////////////////
// Attributes
////////////////////////////////////////////////////////////////////////////////
message AttributesStructMember {
optional uint32 location = 1;
optional uint32 blend_src = 2;
optional uint32 color = 3;
optional BuiltinValue builtin = 4;
optional Interpolation interpolation = 5;
bool invariant = 6;
}
message AttributesFunctionParameter {
optional BuiltinValue builtin = 1;
optional uint32 location = 2;
optional uint32 color = 3;
optional Interpolation interpolation = 4;
optional BindingPoint binding_point = 5;
bool invariant = 6;
}
message Interpolation {
InterpolationType type = 1;
optional InterpolationSampling sampling = 2;
}
////////////////////////////////////////////////////////////////////////////////
// Enums
////////////////////////////////////////////////////////////////////////////////
enum AddressSpace {
function = 0;
handle = 1;
pixel_local = 2;
private = 3;
push_constant = 4;
storage = 5;
uniform = 6;
workgroup = 7;
}
enum AccessControl {
read = 0;
write = 1;
read_write = 2;
}
enum UnaryOp {
complement = 0;
negation = 1;
address_of = 2;
indirection = 3;
not = 4;
}
enum BinaryOp {
add_ = 0;
subtract = 1;
multiply = 2;
divide = 3;
modulo = 4;
and = 5;
or_ = 6;
xor_ = 7;
equal = 8;
not_equal = 9;
less_than = 10;
greater_than = 11;
less_than_equal = 12;
greater_than_equal = 13;
shift_left = 14;
shift_right = 15;
logical_and = 16;
logical_or = 17;
}
enum TextureDimension {
_1d = 0;
_2d = 1;
_2d_array = 2;
_3d = 3;
cube = 4;
cube_array = 5;
}
enum TexelFormat {
bgra8_unorm = 0;
r8_unorm = 1;
r32_float = 2;
r32_sint = 3;
r32_uint = 4;
rg32_float = 5;
rg32_sint = 6;
rg32_uint = 7;
rgba16_float = 8;
rgba16_sint = 9;
rgba16_uint = 10;
rgba32_float = 11;
rgba32_sint = 12;
rgba32_uint = 13;
rgba8_sint = 14;
rgba8_snorm = 15;
rgba8_uint = 16;
rgba8_unorm = 17;
}
enum SamplerKind {
sampler = 0;
comparison = 1;
}
enum InterpolationType {
flat = 0;
linear = 1;
perspective = 2;
}
enum InterpolationSampling {
center = 0;
centroid = 1;
sample = 2;
first = 3;
either = 4;
}
enum BuiltinValue {
point_size = 0;
frag_depth = 1;
front_facing = 2;
global_invocation_id = 3;
instance_index = 4;
local_invocation_id = 5;
local_invocation_index = 6;
num_workgroups = 7;
position = 8;
sample_index = 9;
sample_mask = 10;
subgroup_invocation_id = 11;
subgroup_size = 12;
vertex_index = 13;
workgroup_id = 14;
clip_distances = 15;
}
enum BuiltinFn {
abs = 0;
acos = 1;
acosh = 2;
all = 3;
any = 4;
array_length = 5;
asin = 6;
asinh = 7;
atan = 8;
atan2 = 9;
atanh = 10;
ceil = 11;
clamp = 12;
cos = 13;
cosh = 14;
count_leading_zeros = 15;
count_one_bits = 16;
count_trailing_zeros = 17;
cross = 18;
degrees = 19;
determinant = 20;
distance = 21;
dot = 22;
dot4i8_packed = 23;
dot4u8_packed = 24;
dpdx = 25;
dpdx_coarse = 26;
dpdx_fine = 27;
dpdy = 28;
dpdy_coarse = 29;
dpdy_fine = 30;
exp = 31;
exp2 = 32;
extract_bits = 33;
face_forward = 34;
first_leading_bit = 35;
first_trailing_bit = 36;
floor = 37;
fma = 38;
fract = 39;
frexp = 40;
fwidth = 41;
fwidth_coarse = 42;
fwidth_fine = 43;
insert_bits = 44;
inverse_sqrt = 45;
ldexp = 46;
length = 47;
log = 48;
log2 = 49;
max = 50;
min = 51;
mix = 52;
modf = 53;
normalize = 54;
pack2x16_float = 55;
pack2x16_snorm = 56;
pack2x16_unorm = 57;
pack4x8_snorm = 58;
pack4x8_unorm = 59;
pack4xi8 = 60;
pack4xu8 = 61;
pack4xi8_clamp = 62;
pack4xu8_clamp = 63;
pow = 64;
quantize_to_f16 = 65;
radians = 66;
reflect = 67;
refract = 68;
reverse_bits = 69;
round = 70;
saturate = 71;
select = 72;
sign = 73;
sin = 74;
sinh = 75;
smoothstep = 76;
sqrt = 77;
step = 78;
storage_barrier = 79;
tan = 80;
tanh = 81;
transpose = 82;
trunc = 83;
unpack2x16_float = 84;
unpack2x16_snorm = 85;
unpack2x16_unorm = 86;
unpack4x8_snorm = 87;
unpack4x8_unorm = 88;
unpack4xi8 = 89;
unpack4xu8 = 90;
workgroup_barrier = 91;
texture_barrier = 92;
texture_dimensions = 93;
texture_gather = 94;
texture_gather_compare = 95;
texture_num_layers = 96;
texture_num_levels = 97;
texture_num_samples = 98;
texture_sample = 99;
texture_sample_bias = 100;
texture_sample_compare = 101;
texture_sample_compare_level = 102;
texture_sample_grad = 103;
texture_sample_level = 104;
texture_sample_base_clamp_to_edge = 105;
texture_store = 106;
texture_load = 107;
atomic_load = 108;
atomic_store = 109;
atomic_add = 110;
atomic_sub = 111;
atomic_max = 112;
atomic_min = 113;
atomic_and = 114;
atomic_or = 115;
atomic_xor = 116;
atomic_exchange = 117;
atomic_compare_exchange_weak = 118;
subgroup_ballot = 119;
subgroup_broadcast = 120;
input_attachment_load = 121;
subgroup_add = 122;
subgroup_exclusive_add = 123;
subgroup_mul = 124;
subgroup_exclusive_mul = 125;
subgroup_and = 126;
subgroup_or = 127;
subgroup_xor = 128;
subgroup_min = 129;
subgroup_max = 130;
subgroup_any = 131;
subgroup_all = 132;
subgroup_elect = 133;
subgroup_broadcast_first = 134;
subgroup_shuffle = 135;
subgroup_shuffle_xor = 136;
subgroup_shuffle_up = 137;
subgroup_shuffle_down = 138;
quad_broadcast = 139;
quad_swap_x = 140;
quad_swap_y = 141;
quad_swap_diagonal = 142;
}