Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Add more metal keywords #4707

Merged
merged 2 commits into from
Nov 17, 2023
Merged

Add more metal keywords #4707

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
7db34b1
Add more metal keywords
fornwall Nov 17, 2023
52ee978
Add packed_booln, halfnxm and floatnxm
fornwall Nov 17, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions CHANGELOG.md
View file
Open in desktop
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
295 changes: 209 additions & 86 deletions naga/src/back/msl/keywords.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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",
Expand Down Expand Up @@ -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,
];
2 changes: 1 addition & 1 deletion naga/src/back/msl/writer.rs
View file
Open in desktop
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
View file
Open in desktop
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
View file
Open in desktop
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
Loading