Skip to content

Commit

Permalink
Add more metal keywords
Browse files Browse the repository at this point in the history
Fixes #4545.
  • Loading branch information
fornwall committed Nov 17, 2023
1 parent 77f6e66 commit 7db34b1
Show file tree
Hide file tree
Showing 5 changed files with 194 additions and 91 deletions.
1 change: 1 addition & 0 deletions CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -58,6 +58,7 @@ Passing an owned value `window` to `Surface` will return a `Surface<'static>`. S
#### Naga

- Introduce a new `Scalar` struct type for use in Naga's IR, and update all frontend, middle, and backend code appropriately. By @jimblandy in [#4673](https://github.com/gfx-rs/wgpu/pull/4673).
- Add more metal keywords. By @fornwall in [#4707](https://github.com/gfx-rs/wgpu/pull/4707).

### Bug Fixes

Expand Down
274 changes: 188 additions & 86 deletions naga/src/back/msl/keywords.rs
Original file line number Diff line number Diff line change
@@ -1,114 +1,217 @@
//TODO: find a complete list
// MSLS - Metal Shading Language Specification:
// https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
//
// C++ - Standard for Programming Language C++ (N4431)
// https://www.open-std.org/jtc1/sc22/wg21/docs/papers/2015/n4431.pdf
pub const RESERVED: &[&str] = &[
// control flow
// Standard for Programming Language C++ (N4431): 2.5 Alternative tokens
"and",
"bitor",
"or",
"xor",
"compl",
"bitand",
"and_eq",
"or_eq",
"xor_eq",
"not",
"not_eq",
// Standard for Programming Language C++ (N4431): 2.11 Keywords
"alignas",
"alignof",
"asm",
"auto",
"bool",
"break",
"if",
"else",
"continue",
"goto",
"do",
"while",
"for",
"switch",
"case",
// types and values
"void",
"unsigned",
"signed",
"bool",
"catch",
"char",
"int",
"uint",
"long",
"float",
"double",
"char8_t",
"wchar_t",
"true",
"false",
"nullptr",
"union",
"char16_t",
"char32_t",
"class",
"struct",
"enum",
// other
"main",
"using",
"const",
"constexpr",
"const_cast",
"continue",
"decltype",
"sizeof",
"typeof",
"typedef",
"default",
"delete",
"do",
"double",
"dynamic_cast",
"else",
"enum",
"explicit",
"export",
"extern",
"false",
"float",
"for",
"friend",
"goto",
"if",
"inline",
"int",
"long",
"mutable",
"namespace",
"new",
"noexcept",
"nullptr",
"operator",
"private",
"protected",
"public",
"register",
"reinterpret_cast",
"return",
"short",
"signed",
"sizeof",
"static",
"static_assert",
"static_cast",
"struct",
"switch",
"template",
"typename",
"typeid",
"co_await",
"co_return",
"co_yield",
"module",
"import",
"ray_data",
"vec_step",
"visible",
"as_type",
"this",
// qualifiers
"mutable",
"static",
"volatile",
"restrict",
"const",
"non-temporal",
"dereferenceable",
"invariant",
// exceptions
"thread_local",
"throw",
"true",
"try",
"catch",
// operators
"const_cast",
"dynamic_cast",
"reinterpret_cast",
"static_cast",
"new",
"delete",
"and",
"and_eq",
"bitand",
"bitor",
"compl",
"not",
"not_eq",
"or",
"or_eq",
"xor",
"xor_eq",
"compl",
// Metal-specific
"constant",
"typedef",
"typeid",
"typename",
"union",
"unsigned",
"using",
"virtual",
"void",
"volatile",
"wchar_t",
"while",
// Metal Shading Language Specification: 1.4.4 Restrictions
"main",
// Metal Shading Language Specification: 2.1 Scalar Data Types
"int8_t",
"uchar",
"uint8_t",
"int16_t",
"ushort",
"uint16_t",
"int32_t",
"uint",
"uint32_t",
"int64_t",
"uint64_t",
"half",
"bfloat",
"size_t",
"ptrdiff_t",
// Metal Shading Language Specification: 2.2 Vector Data Types
"bool2",
"bool3",
"bool4",
"char2",
"char3",
"char4",
"short2",
"short3",
"short4",
"int2",
"int3",
"int4",
"long2",
"long3",
"long4",
"uchar2",
"uchar3",
"uchar4",
"ushort2",
"ushort3",
"ushort4",
"uint2",
"uint3",
"uint4",
"ulong2",
"ulong3",
"ulong4",
"half2",
"half3",
"half4",
"bfloat2",
"bfloat3",
"bfloat4",
"float2",
"float3",
"float4",
"vec",
// Metal Shading Language Specification: 2.2.3 Packed Vector Types
"packed_char2",
"packed_char3",
"packed_char4",
"packed_short2",
"packed_short3",
"packed_short4",
"packed_int2",
"packed_int3",
"packed_int4",
"packed_uchar2",
"packed_uchar3",
"packed_uchar4",
"packed_ushort2",
"packed_ushort3",
"packed_ushort4",
"packed_uint2",
"packed_uint3",
"packed_uint4",
"packed_half2",
"packed_half3",
"packed_half4",
"packed_bfloat2",
"packed_bfloat3",
"packed_bfloat4",
"packed_float2",
"packed_float3",
"packed_float4",
"packed_long2",
"packed_long3",
"packed_long4",
"packed_vec",
// Metal Shading Language Specification: 2.3 Matrix Data Types
"matrix",
// Metal Shading Language Specification: 2.6 Atomic Data Types
"atomic",
"atomic_int",
"atomic_uint",
"atomic_bool",
"atomic_ulong",
"atomic_float",
// Metal Shading Language Specification: 2.20 Type Conversions and Re-interpreting Data
"as_type",
// Metal Shading Language Specification: 4 Address Spaces
"device",
"constant",
"thread",
"threadgroup",
"threadgroup_imageblock",
"kernel",
"compute",
"ray_data",
"object_data",
// Metal Shading Language Specification: 5.1 Functions
"vertex",
"fragment",
"read_only",
"write_only",
"read_write",
"auto",
// Metal reserved types
"kernel",
// Metal Shading Language Specification: 6.1 Namespace and Header Files
"metal",
// C99 / C++ extension:
"restrict",
// Metal reserved types in <metal_types>:
"llong",
"ullong",
"quad",
"complex",
"imaginary",
// Metal constants
// Constants in <metal_types>:
"CHAR_BIT",
"SCHAR_MAX",
"SCHAR_MIN",
Expand Down Expand Up @@ -213,7 +316,6 @@ pub const RESERVED: &[&str] = &[
"M_SQRT1_2",
// Naga utilities
"DefaultConstructible",
"clamped_lod_e",
super::writer::FREXP_FUNCTION,
super::writer::MODF_FUNCTION,
];
2 changes: 1 addition & 1 deletion naga/src/back/msl/writer.rs
Original file line number Diff line number Diff line change
Expand Up @@ -3118,7 +3118,7 @@ impl<W: Write> Writer<W> {
super::keywords::RESERVED,
&[],
&[],
&[],
&[CLAMPED_LOD_LOAD_PREFIX],
&mut self.names,
);
self.struct_member_pads.clear();
Expand Down
4 changes: 2 additions & 2 deletions naga/tests/out/msl/interface.msl
Original file line number Diff line number Diff line change
Expand Up @@ -66,9 +66,9 @@ fragment fragment_Output fragment_(
}


struct compute_Input {
struct computeInput {
};
kernel void compute_(
kernel void compute(
metal::uint3 global_id [[thread_position_in_grid]]
, metal::uint3 local_id [[thread_position_in_threadgroup]]
, uint local_index [[thread_index_in_threadgroup]]
Expand Down
4 changes: 2 additions & 2 deletions naga/tests/out/msl/ray-query.msl
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@ constexpr metal::uint _map_intersection_type(const metal::raytracing::intersecti
}

struct Output {
uint visible_;
uint visible;
char _pad1[12];
metal::float3 normal;
};
Expand Down Expand Up @@ -72,7 +72,7 @@ kernel void main_(
}
}
RayIntersection intersection_1 = RayIntersection {_map_intersection_type(rq.intersection.type), rq.intersection.distance, rq.intersection.user_instance_id, rq.intersection.instance_id, {}, rq.intersection.geometry_id, rq.intersection.primitive_id, rq.intersection.triangle_barycentric_coord, rq.intersection.triangle_front_facing, {}, rq.intersection.object_to_world_transform, rq.intersection.world_to_object_transform};
output.visible_ = static_cast<uint>(intersection_1.kind == 0u);
output.visible = static_cast<uint>(intersection_1.kind == 0u);
metal::float3 _e25 = get_torus_normal(dir * intersection_1.t, intersection_1);
output.normal = _e25;
return;
Expand Down

0 comments on commit 7db34b1

Please sign in to comment.