Add more metal keywords (#4707)

gfx-rs · Nov 17, 2023 · a5c93ca · a5c93ca
1 parent 77f6e66
commit a5c93ca
Show file tree

Hide file tree

Showing 5 changed files with 215 additions and 91 deletions.
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -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
 

diff --git a/naga/src/back/msl/keywords.rs b/naga/src/back/msl/keywords.rs
@@ -1,114 +1,238 @@
-//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_bool2",
+    "packed_bool3",
+    "packed_bool4",
+    "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
+    "half2x2",
+    "half2x3",
+    "half2x4",
+    "half3x2",
+    "half3x3",
+    "half3x4",
+    "half4x2",
+    "half4x3",
+    "half4x4",
+    "float2x2",
+    "float2x3",
+    "float2x4",
+    "float3x2",
+    "float3x3",
+    "float3x4",
+    "float4x2",
+    "float4x3",
+    "float4x4",
+    "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",
@@ -213,7 +337,6 @@ pub const RESERVED: &[&str] = &[
     "M_SQRT1_2",
     // Naga utilities
     "DefaultConstructible",
-    "clamped_lod_e",
     super::writer::FREXP_FUNCTION,
     super::writer::MODF_FUNCTION,
 ];
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
@@ -3118,7 +3118,7 @@ impl<W: Write> Writer<W> {
             super::keywords::RESERVED,
             &[],
             &[],
-            &[],
+            &[CLAMPED_LOD_LOAD_PREFIX],
             &mut self.names,
         );
         self.struct_member_pads.clear();

diff --git a/naga/tests/out/msl/interface.msl b/naga/tests/out/msl/interface.msl
@@ -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]]

diff --git a/naga/tests/out/msl/ray-query.msl b/naga/tests/out/msl/ray-query.msl
@@ -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;
 };
@@ -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;