diff --git a/crates/rustc_codegen_spirv/src/attr.rs b/crates/rustc_codegen_spirv/src/attr.rs index b35dbb6c3c..cc7f7b52fd 100644 --- a/crates/rustc_codegen_spirv/src/attr.rs +++ b/crates/rustc_codegen_spirv/src/attr.rs @@ -5,6 +5,7 @@ use crate::codegen_cx::CodegenCx; use crate::symbols::Symbols; use rspirv::spirv::{BuiltIn, ExecutionMode, ExecutionModel, StorageClass}; +use rustc_ast::{LitKind, MetaItemInner, MetaItemLit}; use rustc_hir as hir; use rustc_hir::def_id::LocalModDefId; use rustc_hir::intravisit::{self, Visitor}; @@ -12,7 +13,8 @@ use rustc_hir::{Attribute, CRATE_HIR_ID, HirId, MethodKind, Target}; use rustc_middle::hir::nested_filter; use rustc_middle::query::Providers; use rustc_middle::ty::TyCtxt; -use rustc_span::{Span, Symbol}; +use rustc_span::{Ident, Span, Symbol}; +use smallvec::SmallVec; use std::rc::Rc; // FIXME(eddyb) replace with `ArrayVec<[Word; 3]>`. @@ -152,7 +154,7 @@ impl AggregatedSpirvAttributes { // NOTE(eddyb) `span_delayed_bug` ensures that if attribute checking fails // to see an attribute error, it will cause an ICE instead. - for parse_attr_result in crate::symbols::parse_attrs_for_checking(&cx.sym, attrs) { + for parse_attr_result in parse_attrs_for_checking(&cx.sym, attrs) { let (span, parsed_attr) = match parse_attr_result { Ok(span_and_parsed_attr) => span_and_parsed_attr, Err((span, msg)) => { @@ -278,7 +280,7 @@ impl CheckSpirvAttrVisitor<'_> { fn check_spirv_attributes(&self, hir_id: HirId, target: Target) { let mut aggregated_attrs = AggregatedSpirvAttributes::default(); - let parse_attrs = |attrs| crate::symbols::parse_attrs_for_checking(&self.sym, attrs); + let parse_attrs = |attrs| parse_attrs_for_checking(&self.sym, attrs); let attrs = self.tcx.hir_attrs(hir_id); for parse_attr_result in parse_attrs(attrs) { @@ -512,3 +514,334 @@ pub(crate) fn provide(providers: &mut Providers) { ..*providers }; } + +// FIXME(eddyb) find something nicer for the error type. +type ParseAttrError = (Span, String); + +#[allow(clippy::get_first)] +fn parse_attrs_for_checking<'a>( + sym: &'a Symbols, + attrs: &'a [Attribute], +) -> impl Iterator> + 'a { + attrs + .iter() + .map(move |attr| { + // parse the #[rust_gpu::spirv(...)] attr and return the inner list + match attr { + Attribute::Unparsed(item) => { + // #[...] + let s = &item.path.segments; + if let Some(rust_gpu) = s.get(0) && rust_gpu.name == sym.rust_gpu { + // #[rust_gpu ...] + match s.get(1) { + Some(command) if command.name == sym.spirv_attr_with_version => { + // #[rust_gpu::spirv ...] + if let Some(args) = attr.meta_item_list() { + // #[rust_gpu::spirv(...)] + Ok(parse_spirv_attr(sym, args.iter())) + } else { + // #[rust_gpu::spirv] + Err(( + attr.span(), + "#[spirv(..)] attribute must have at least one argument" + .to_string(), + )) + } + } + _ => { + // #[rust_gpu::...] but not a know version + let spirv = sym.spirv_attr_with_version.as_str(); + Err(( + attr.span(), + format!("unknown `rust_gpu` attribute, expected `rust_gpu::{spirv}`. \ + Do the versions of `spirv-std` and `rustc_codegen_spirv` match?"), + )) + } + } + } else { + // #[...] but not #[rust_gpu ...] + Ok(Default::default()) + } + } + Attribute::Parsed(_) => Ok(Default::default()), + } + }) + .flat_map(|result| { + result + .unwrap_or_else(|err| SmallVec::from_iter([Err(err)])) + .into_iter() + }) +} + +fn parse_spirv_attr<'a>( + sym: &Symbols, + iter: impl Iterator, +) -> SmallVec<[Result<(Span, SpirvAttribute), ParseAttrError>; 4]> { + iter.map(|arg| { + let span = arg.span(); + let parsed_attr = + if arg.has_name(sym.descriptor_set) { + SpirvAttribute::DescriptorSet(parse_attr_int_value(arg)?) + } else if arg.has_name(sym.binding) { + SpirvAttribute::Binding(parse_attr_int_value(arg)?) + } else if arg.has_name(sym.input_attachment_index) { + SpirvAttribute::InputAttachmentIndex(parse_attr_int_value(arg)?) + } else if arg.has_name(sym.spec_constant) { + SpirvAttribute::SpecConstant(parse_spec_constant_attr(sym, arg)?) + } else { + let name = match arg.ident() { + Some(i) => i, + None => { + return Err(( + span, + "#[spirv(..)] attribute argument must be single identifier".to_string(), + )); + } + }; + sym.attributes.get(&name.name).map_or_else( + || Err((name.span, "unknown argument to spirv attribute".to_string())), + |a| { + Ok(match a { + SpirvAttribute::Entry(entry) => SpirvAttribute::Entry( + parse_entry_attrs(sym, arg, &name, entry.execution_model)?, + ), + _ => a.clone(), + }) + }, + )? + }; + Ok((span, parsed_attr)) + }) + .collect() +} + +fn parse_spec_constant_attr( + sym: &Symbols, + arg: &MetaItemInner, +) -> Result { + let mut id = None; + let mut default = None; + + if let Some(attrs) = arg.meta_item_list() { + for attr in attrs { + if attr.has_name(sym.id) { + if id.is_none() { + id = Some(parse_attr_int_value(attr)?); + } else { + return Err((attr.span(), "`id` may only be specified once".into())); + } + } else if attr.has_name(sym.default) { + if default.is_none() { + default = Some(parse_attr_int_value(attr)?); + } else { + return Err((attr.span(), "`default` may only be specified once".into())); + } + } else { + return Err((attr.span(), "expected `id = ...` or `default = ...`".into())); + } + } + } + Ok(SpecConstant { + id: id.ok_or_else(|| (arg.span(), "expected `spec_constant(id = ...)`".into()))?, + default, + }) +} + +fn parse_attr_int_value(arg: &MetaItemInner) -> Result { + let arg = match arg.meta_item() { + Some(arg) => arg, + None => return Err((arg.span(), "attribute must have value".to_string())), + }; + match arg.name_value_literal() { + Some(&MetaItemLit { + kind: LitKind::Int(x, ..), + .. + }) if x <= u32::MAX as u128 => Ok(x.get() as u32), + _ => Err((arg.span, "attribute value must be integer".to_string())), + } +} + +fn parse_local_size_attr(arg: &MetaItemInner) -> Result<[u32; 3], ParseAttrError> { + let arg = match arg.meta_item() { + Some(arg) => arg, + None => return Err((arg.span(), "attribute must have value".to_string())), + }; + match arg.meta_item_list() { + Some(tuple) if !tuple.is_empty() && tuple.len() < 4 => { + let mut local_size = [1; 3]; + for (idx, lit) in tuple.iter().enumerate() { + match lit { + MetaItemInner::Lit(MetaItemLit { + kind: LitKind::Int(x, ..), + .. + }) if *x <= u32::MAX as u128 => local_size[idx] = x.get() as u32, + _ => return Err((lit.span(), "must be a u32 literal".to_string())), + } + } + Ok(local_size) + } + Some([]) => Err(( + arg.span, + "#[spirv(compute(threads(x, y, z)))] must have the x dimension specified, trailing ones may be elided".to_string(), + )), + Some(tuple) if tuple.len() > 3 => Err(( + arg.span, + "#[spirv(compute(threads(x, y, z)))] is three dimensional".to_string(), + )), + _ => Err(( + arg.span, + "#[spirv(compute(threads(x, y, z)))] must have 1 to 3 parameters, trailing ones may be elided".to_string(), + )), + } +} + +// for a given entry, gather up the additional attributes +// in this case ExecutionMode's, some have extra arguments +// others are specified with x, y, or z components +// ie #[spirv(fragment(origin_lower_left))] or #[spirv(gl_compute(local_size_x=64, local_size_y=8))] +fn parse_entry_attrs( + sym: &Symbols, + arg: &MetaItemInner, + name: &Ident, + execution_model: ExecutionModel, +) -> Result { + use ExecutionMode::*; + use ExecutionModel::*; + let mut entry = Entry::from(execution_model); + let mut origin_mode: Option = None; + let mut local_size: Option<[u32; 3]> = None; + let mut local_size_hint: Option<[u32; 3]> = None; + // Reserved + //let mut max_workgroup_size_intel: Option<[u32; 3]> = None; + if let Some(attrs) = arg.meta_item_list() { + for attr in attrs { + if let Some(attr_name) = attr.ident() { + if let Some((execution_mode, extra_dim)) = sym.execution_modes.get(&attr_name.name) + { + use crate::symbols::ExecutionModeExtraDim::*; + let val = match extra_dim { + None | Tuple => Option::None, + _ => Some(parse_attr_int_value(attr)?), + }; + match execution_mode { + OriginUpperLeft | OriginLowerLeft => { + origin_mode.replace(*execution_mode); + } + LocalSize => { + if local_size.is_none() { + local_size.replace(parse_local_size_attr(attr)?); + } else { + return Err(( + attr_name.span, + String::from( + "`#[spirv(compute(threads))]` may only be specified once", + ), + )); + } + } + LocalSizeHint => { + let val = val.unwrap(); + if local_size_hint.is_none() { + local_size_hint.replace([1, 1, 1]); + } + let local_size_hint = local_size_hint.as_mut().unwrap(); + match extra_dim { + X => { + local_size_hint[0] = val; + } + Y => { + local_size_hint[1] = val; + } + Z => { + local_size_hint[2] = val; + } + _ => unreachable!(), + } + } + // Reserved + /*MaxWorkgroupSizeINTEL => { + let val = val.unwrap(); + if max_workgroup_size_intel.is_none() { + max_workgroup_size_intel.replace([1, 1, 1]); + } + let max_workgroup_size_intel = max_workgroup_size_intel.as_mut() + .unwrap(); + match extra_dim { + X => { + max_workgroup_size_intel[0] = val; + }, + Y => { + max_workgroup_size_intel[1] = val; + }, + Z => { + max_workgroup_size_intel[2] = val; + }, + _ => unreachable!(), + } + },*/ + _ => { + if let Some(val) = val { + entry + .execution_modes + .push((*execution_mode, ExecutionModeExtra::new([val]))); + } else { + entry + .execution_modes + .push((*execution_mode, ExecutionModeExtra::new([]))); + } + } + } + } else if attr_name.name == sym.entry_point_name { + match attr.value_str() { + Some(sym) => { + entry.name = Some(sym); + } + None => { + return Err(( + attr_name.span, + format!( + "#[spirv({name}(..))] unknown attribute argument {attr_name}" + ), + )); + } + } + } else { + return Err(( + attr_name.span, + format!("#[spirv({name}(..))] unknown attribute argument {attr_name}",), + )); + } + } else { + return Err(( + arg.span(), + format!("#[spirv({name}(..))] attribute argument must be single identifier"), + )); + } + } + } + match entry.execution_model { + Fragment => { + let origin_mode = origin_mode.unwrap_or(OriginUpperLeft); + entry + .execution_modes + .push((origin_mode, ExecutionModeExtra::new([]))); + } + GLCompute | MeshNV | TaskNV | TaskEXT | MeshEXT => { + if let Some(local_size) = local_size { + entry + .execution_modes + .push((LocalSize, ExecutionModeExtra::new(local_size))); + } else { + return Err(( + arg.span(), + String::from( + "The `threads` argument must be specified when using `#[spirv(compute)]`, `#[spirv(mesh_nv)]`, `#[spirv(task_nv)]`, `#[spirv(task_ext)]` or `#[spirv(mesh_ext)]`", + ), + )); + } + } + //TODO: Cover more defaults + _ => {} + } + Ok(entry) +} diff --git a/crates/rustc_codegen_spirv/src/lib.rs b/crates/rustc_codegen_spirv/src/lib.rs index 3e57d8648d..d565c6074c 100644 --- a/crates/rustc_codegen_spirv/src/lib.rs +++ b/crates/rustc_codegen_spirv/src/lib.rs @@ -133,6 +133,8 @@ mod custom_decorations; mod custom_insts; mod link; mod linker; +#[path = "../../spirv_attr_version.rs"] +mod spirv_attr_version; mod spirv_type; mod spirv_type_constraints; mod symbols; diff --git a/crates/rustc_codegen_spirv/src/symbols.rs b/crates/rustc_codegen_spirv/src/symbols.rs index 4dfef5b4d3..475e100f10 100644 --- a/crates/rustc_codegen_spirv/src/symbols.rs +++ b/crates/rustc_codegen_spirv/src/symbols.rs @@ -1,11 +1,9 @@ -use crate::attr::{Entry, ExecutionModeExtra, IntrinsicType, SpecConstant, SpirvAttribute}; +use crate::attr::{IntrinsicType, SpirvAttribute}; use crate::builder::libm_intrinsics; +use crate::spirv_attr_version::spirv_attr_with_version; use rspirv::spirv::{BuiltIn, ExecutionMode, ExecutionModel, StorageClass}; -use rustc_ast::ast::{LitKind, MetaItemInner, MetaItemLit}; use rustc_data_structures::fx::FxHashMap; -use rustc_hir::Attribute; -use rustc_span::Span; -use rustc_span::symbol::{Ident, Symbol}; +use rustc_span::symbol::Symbol; use std::rc::Rc; /// Various places in the codebase (mostly attribute parsing) need to compare rustc Symbols to particular keywords. @@ -16,21 +14,21 @@ use std::rc::Rc; pub struct Symbols { pub discriminant: Symbol, pub rust_gpu: Symbol, - pub spirv: Symbol, + pub spirv_attr_with_version: Symbol, pub libm: Symbol, pub entry_point_name: Symbol, pub spv_khr_vulkan_memory_model: Symbol, - descriptor_set: Symbol, - binding: Symbol, - input_attachment_index: Symbol, + pub descriptor_set: Symbol, + pub binding: Symbol, + pub input_attachment_index: Symbol, - spec_constant: Symbol, - id: Symbol, - default: Symbol, + pub spec_constant: Symbol, + pub id: Symbol, + pub default: Symbol, - attributes: FxHashMap, - execution_modes: FxHashMap, + pub attributes: FxHashMap, + pub execution_modes: FxHashMap, pub libm_intrinsics: FxHashMap, } @@ -204,7 +202,7 @@ const EXECUTION_MODELS: &[(&str, ExecutionModel)] = { }; #[derive(Copy, Clone, Debug)] -enum ExecutionModeExtraDim { +pub enum ExecutionModeExtraDim { None, Value, X, @@ -407,7 +405,7 @@ impl Symbols { Self { discriminant: Symbol::intern("discriminant"), rust_gpu: Symbol::intern("rust_gpu"), - spirv: Symbol::intern("spirv"), + spirv_attr_with_version: Symbol::intern(&spirv_attr_with_version()), libm: Symbol::intern("libm"), entry_point_name: Symbol::intern("entry_point_name"), spv_khr_vulkan_memory_model: Symbol::intern("SPV_KHR_vulkan_memory_model"), @@ -436,321 +434,3 @@ impl Symbols { SYMBOLS.with(Rc::clone) } } - -// FIXME(eddyb) find something nicer for the error type. -type ParseAttrError = (Span, String); - -// FIXME(eddyb) maybe move this to `attr`? -pub(crate) fn parse_attrs_for_checking<'a>( - sym: &'a Symbols, - attrs: &'a [Attribute], -) -> impl Iterator> + 'a { - attrs.iter().flat_map(move |attr| { - let (whole_attr_error, args) = match attr { - Attribute::Unparsed(item) => { - // #[...] - let s = &item.path.segments; - if s.len() > 1 && s[0].name == sym.rust_gpu { - // #[rust_gpu ...] - if s.len() != 2 || s[1].name != sym.spirv { - // #[rust_gpu::...] but not #[rust_gpu::spirv] - ( - Some(Err(( - attr.span(), - "unknown `rust_gpu` attribute, expected `rust_gpu::spirv`" - .to_string(), - ))), - Default::default(), - ) - } else if let Some(args) = attr.meta_item_list() { - // #[rust_gpu::spirv(...)] - (None, args) - } else { - // #[rust_gpu::spirv] - ( - Some(Err(( - attr.span(), - "#[rust_gpu::spirv(..)] attribute must have at least one argument" - .to_string(), - ))), - Default::default(), - ) - } - } else { - // #[...] but not #[rust_gpu ...] - (None, Default::default()) - } - } - Attribute::Parsed(_) => (None, Default::default()), - }; - - whole_attr_error - .into_iter() - .chain(args.into_iter().map(move |ref arg| { - let span = arg.span(); - let parsed_attr = if arg.has_name(sym.descriptor_set) { - SpirvAttribute::DescriptorSet(parse_attr_int_value(arg)?) - } else if arg.has_name(sym.binding) { - SpirvAttribute::Binding(parse_attr_int_value(arg)?) - } else if arg.has_name(sym.input_attachment_index) { - SpirvAttribute::InputAttachmentIndex(parse_attr_int_value(arg)?) - } else if arg.has_name(sym.spec_constant) { - SpirvAttribute::SpecConstant(parse_spec_constant_attr(sym, arg)?) - } else { - let name = match arg.ident() { - Some(i) => i, - None => { - return Err(( - span, - "#[spirv(..)] attribute argument must be single identifier" - .to_string(), - )); - } - }; - sym.attributes.get(&name.name).map_or_else( - || Err((name.span, "unknown argument to spirv attribute".to_string())), - |a| { - Ok(match a { - SpirvAttribute::Entry(entry) => SpirvAttribute::Entry( - parse_entry_attrs(sym, arg, &name, entry.execution_model)?, - ), - _ => a.clone(), - }) - }, - )? - }; - Ok((span, parsed_attr)) - })) - }) -} - -fn parse_spec_constant_attr( - sym: &Symbols, - arg: &MetaItemInner, -) -> Result { - let mut id = None; - let mut default = None; - - if let Some(attrs) = arg.meta_item_list() { - for attr in attrs { - if attr.has_name(sym.id) { - if id.is_none() { - id = Some(parse_attr_int_value(attr)?); - } else { - return Err((attr.span(), "`id` may only be specified once".into())); - } - } else if attr.has_name(sym.default) { - if default.is_none() { - default = Some(parse_attr_int_value(attr)?); - } else { - return Err((attr.span(), "`default` may only be specified once".into())); - } - } else { - return Err((attr.span(), "expected `id = ...` or `default = ...`".into())); - } - } - } - Ok(SpecConstant { - id: id.ok_or_else(|| (arg.span(), "expected `spec_constant(id = ...)`".into()))?, - default, - }) -} - -fn parse_attr_int_value(arg: &MetaItemInner) -> Result { - let arg = match arg.meta_item() { - Some(arg) => arg, - None => return Err((arg.span(), "attribute must have value".to_string())), - }; - match arg.name_value_literal() { - Some(&MetaItemLit { - kind: LitKind::Int(x, ..), - .. - }) if x <= u32::MAX as u128 => Ok(x.get() as u32), - _ => Err((arg.span, "attribute value must be integer".to_string())), - } -} - -fn parse_local_size_attr(arg: &MetaItemInner) -> Result<[u32; 3], ParseAttrError> { - let arg = match arg.meta_item() { - Some(arg) => arg, - None => return Err((arg.span(), "attribute must have value".to_string())), - }; - match arg.meta_item_list() { - Some(tuple) if !tuple.is_empty() && tuple.len() < 4 => { - let mut local_size = [1; 3]; - for (idx, lit) in tuple.iter().enumerate() { - match lit { - MetaItemInner::Lit(MetaItemLit { - kind: LitKind::Int(x, ..), - .. - }) if *x <= u32::MAX as u128 => local_size[idx] = x.get() as u32, - _ => return Err((lit.span(), "must be a u32 literal".to_string())), - } - } - Ok(local_size) - } - Some([]) => Err(( - arg.span, - "#[spirv(compute(threads(x, y, z)))] must have the x dimension specified, trailing ones may be elided".to_string(), - )), - Some(tuple) if tuple.len() > 3 => Err(( - arg.span, - "#[spirv(compute(threads(x, y, z)))] is three dimensional".to_string(), - )), - _ => Err(( - arg.span, - "#[spirv(compute(threads(x, y, z)))] must have 1 to 3 parameters, trailing ones may be elided".to_string(), - )), - } -} - -// for a given entry, gather up the additional attributes -// in this case ExecutionMode's, some have extra arguments -// others are specified with x, y, or z components -// ie #[spirv(fragment(origin_lower_left))] or #[spirv(gl_compute(local_size_x=64, local_size_y=8))] -fn parse_entry_attrs( - sym: &Symbols, - arg: &MetaItemInner, - name: &Ident, - execution_model: ExecutionModel, -) -> Result { - use ExecutionMode::*; - use ExecutionModel::*; - let mut entry = Entry::from(execution_model); - let mut origin_mode: Option = None; - let mut local_size: Option<[u32; 3]> = None; - let mut local_size_hint: Option<[u32; 3]> = None; - // Reserved - //let mut max_workgroup_size_intel: Option<[u32; 3]> = None; - if let Some(attrs) = arg.meta_item_list() { - for attr in attrs { - if let Some(attr_name) = attr.ident() { - if let Some((execution_mode, extra_dim)) = sym.execution_modes.get(&attr_name.name) - { - use ExecutionModeExtraDim::*; - let val = match extra_dim { - None | Tuple => Option::None, - _ => Some(parse_attr_int_value(attr)?), - }; - match execution_mode { - OriginUpperLeft | OriginLowerLeft => { - origin_mode.replace(*execution_mode); - } - LocalSize => { - if local_size.is_none() { - local_size.replace(parse_local_size_attr(attr)?); - } else { - return Err(( - attr_name.span, - String::from( - "`#[spirv(compute(threads))]` may only be specified once", - ), - )); - } - } - LocalSizeHint => { - let val = val.unwrap(); - if local_size_hint.is_none() { - local_size_hint.replace([1, 1, 1]); - } - let local_size_hint = local_size_hint.as_mut().unwrap(); - match extra_dim { - X => { - local_size_hint[0] = val; - } - Y => { - local_size_hint[1] = val; - } - Z => { - local_size_hint[2] = val; - } - _ => unreachable!(), - } - } - // Reserved - /*MaxWorkgroupSizeINTEL => { - let val = val.unwrap(); - if max_workgroup_size_intel.is_none() { - max_workgroup_size_intel.replace([1, 1, 1]); - } - let max_workgroup_size_intel = max_workgroup_size_intel.as_mut() - .unwrap(); - match extra_dim { - X => { - max_workgroup_size_intel[0] = val; - }, - Y => { - max_workgroup_size_intel[1] = val; - }, - Z => { - max_workgroup_size_intel[2] = val; - }, - _ => unreachable!(), - } - },*/ - _ => { - if let Some(val) = val { - entry - .execution_modes - .push((*execution_mode, ExecutionModeExtra::new([val]))); - } else { - entry - .execution_modes - .push((*execution_mode, ExecutionModeExtra::new([]))); - } - } - } - } else if attr_name.name == sym.entry_point_name { - match attr.value_str() { - Some(sym) => { - entry.name = Some(sym); - } - None => { - return Err(( - attr_name.span, - format!( - "#[spirv({name}(..))] unknown attribute argument {attr_name}" - ), - )); - } - } - } else { - return Err(( - attr_name.span, - format!("#[spirv({name}(..))] unknown attribute argument {attr_name}",), - )); - } - } else { - return Err(( - arg.span(), - format!("#[spirv({name}(..))] attribute argument must be single identifier"), - )); - } - } - } - match entry.execution_model { - Fragment => { - let origin_mode = origin_mode.unwrap_or(OriginUpperLeft); - entry - .execution_modes - .push((origin_mode, ExecutionModeExtra::new([]))); - } - GLCompute | MeshNV | TaskNV | TaskEXT | MeshEXT => { - if let Some(local_size) = local_size { - entry - .execution_modes - .push((LocalSize, ExecutionModeExtra::new(local_size))); - } else { - return Err(( - arg.span(), - String::from( - "The `threads` argument must be specified when using `#[spirv(compute)]`, `#[spirv(mesh_nv)]`, `#[spirv(task_nv)]`, `#[spirv(task_ext)]` or `#[spirv(mesh_ext)]`", - ), - )); - } - } - //TODO: Cover more defaults - _ => {} - } - Ok(entry) -} diff --git a/crates/spirv-std/macros/src/lib.rs b/crates/spirv-std/macros/src/lib.rs index c1f0d57ab9..f407e26561 100644 --- a/crates/spirv-std/macros/src/lib.rs +++ b/crates/spirv-std/macros/src/lib.rs @@ -72,13 +72,16 @@ #![doc = include_str!("../README.md")] mod image; +#[path = "../../../spirv_attr_version.rs"] +mod spirv_attr_version; use proc_macro::TokenStream; use proc_macro2::{Delimiter, Group, Span, TokenTree}; use syn::{ImplItemFn, visit_mut::VisitMut}; -use quote::{ToTokens, quote}; +use crate::spirv_attr_version::spirv_attr_with_version; +use quote::{ToTokens, TokenStreamExt, format_ident, quote}; use std::fmt::Write; /// A macro for creating SPIR-V `OpTypeImage` types. Always produces a @@ -143,36 +146,46 @@ pub fn Image(item: TokenStream) -> TokenStream { /// `#[cfg_attr(target_arch="spirv", rust_gpu::spirv(..))]`. #[proc_macro_attribute] pub fn spirv(attr: TokenStream, item: TokenStream) -> TokenStream { + let spirv = format_ident!("{}", &spirv_attr_with_version()); let mut tokens: Vec = Vec::new(); // prepend with #[rust_gpu::spirv(..)] let attr: proc_macro2::TokenStream = attr.into(); - tokens.extend(quote! { #[cfg_attr(target_arch="spirv", rust_gpu::spirv(#attr))] }); + tokens.extend(quote! { #[cfg_attr(target_arch="spirv", rust_gpu::#spirv(#attr))] }); let item: proc_macro2::TokenStream = item.into(); for tt in item { match tt { TokenTree::Group(group) if group.delimiter() == Delimiter::Parenthesis => { - let mut sub_tokens = Vec::new(); + let mut group_tokens = proc_macro2::TokenStream::new(); + let mut last_token_hashtag = false; for tt in group.stream() { + let is_token_hashtag = + matches!(&tt, TokenTree::Punct(punct) if punct.as_char() == '#'); match tt { TokenTree::Group(group) if group.delimiter() == Delimiter::Bracket - && matches!(group.stream().into_iter().next(), Some(TokenTree::Ident(ident)) if ident == "spirv") - && matches!(sub_tokens.last(), Some(TokenTree::Punct(p)) if p.as_char() == '#') => + && last_token_hashtag + && matches!(group.stream().into_iter().next(), Some(TokenTree::Ident(ident)) if ident == "spirv") => { // group matches [spirv ...] - let inner = group.stream(); // group stream doesn't include the brackets - sub_tokens.extend( - quote! { [cfg_attr(target_arch="spirv", rust_gpu::#inner)] }, + // group stream doesn't include the brackets + let inner = group + .stream() + .into_iter() + .skip(1) + .collect::(); + group_tokens.extend( + quote! { [cfg_attr(target_arch="spirv", rust_gpu::#spirv #inner)] }, ); } - _ => sub_tokens.push(tt), + _ => group_tokens.append(tt), } + last_token_hashtag = is_token_hashtag; } tokens.push(TokenTree::from(Group::new( Delimiter::Parenthesis, - sub_tokens.into_iter().collect(), + group_tokens, ))); } _ => tokens.push(tt), diff --git a/crates/spirv_attr_version.rs b/crates/spirv_attr_version.rs new file mode 100644 index 0000000000..50bd89c129 --- /dev/null +++ b/crates/spirv_attr_version.rs @@ -0,0 +1,21 @@ +//! This is placed outside any crate, and included by both `spirv_std` and `rustc_codegen_spirv`. +//! I could have made a new crate, shared between the two, but decided against having even more small crates for sharing +//! types. Instead, you get this single small file to specify the versioned spirv attribute. +//! +//! This also ensures that the macros below take the *exact* version of the two crates above, and not some dependency +//! that both of them depend on. + +/// The spirv attribute with version tag +/// +/// ```ignore +/// # we don't know the namespace of our function +/// let spirv = spirv_attr_with_version(); +/// let attr = format!("#[rust_gpu::{spirv}(vertex)]"); +/// // version here may be out-of-date +/// assert_eq!("#[rust_gpu::spirv_v0_9(vertex)]", attr); +/// ``` +pub fn spirv_attr_with_version() -> String { + let major: u32 = env!("CARGO_PKG_VERSION_MAJOR").parse().unwrap(); + let minor: u32 = env!("CARGO_PKG_VERSION_MINOR").parse().unwrap(); + format!("spirv_v{major}_{minor}") +} diff --git a/tests/compiletests/ui/spirv-attr/invalid-target.rs b/tests/compiletests/ui/spirv-attr/invalid-target.rs index 2dbd5bb106..5adce99ab3 100644 --- a/tests/compiletests/ui/spirv-attr/invalid-target.rs +++ b/tests/compiletests/ui/spirv-attr/invalid-target.rs @@ -1,4 +1,4 @@ -// Tests that `#[rust_gpu::spirv(...)]` attributes cannot be applied to the wrong "targets" +// Tests that `#[rust_gpu::spirv_v0_9(...)]` attributes cannot be applied to the wrong "targets" // (i.e. various kinds of definitions and other syntactic categories). // build-fail @@ -11,15 +11,15 @@ )] // NOTE(eddyb) in the interest of keeping this test manageable, only one of -// each of the following categories of `#[rust_gpu::spirv(...)]` attributes is used: +// each of the following categories of `#[rust_gpu::spirv_v0_9(...)]` attributes is used: // * entry: `vertex` // * storage class: `uniform` // * builtin: `position` // NOTE(eddyb) accounting for the number of errors this test actually produces: // * 437 errors, all "attribute is only valid on" (see `invalid-target.stderr`) -// * 41 uses of `#[rust_gpu::spirv(...)]` in this test -// * at most 11 attributes per `#[rust_gpu::spirv(...)]`, so an upper bound of `41*11 = 451` +// * 41 uses of `#[rust_gpu::spirv_v0_9(...)]` in this test +// * at most 11 attributes per `#[rust_gpu::spirv_v0_9(...)]`, so an upper bound of `41*11 = 451` // * the difference between 451 and 437 is 14, i.e. valid attributes, made up of: // * 4 on `_Struct` // * 4 on functions, i.e. 1 on each of: @@ -29,9 +29,9 @@ // * `_fn` // * 6 on `_entry_param` -// NOTE(shesp) Directly using `#[rust_gpu::spirv(...)]` because macro attributes are invalid in most contexts +// NOTE(shesp) Directly using `#[rust_gpu::spirv_v0_9(...)]` because macro attributes are invalid in most contexts -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only @@ -40,48 +40,48 @@ macro_rules! _macro { () => {}; } -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] extern crate spirv_std as _; -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] use spirv_std as _; -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] mod _mod {} -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] extern "C" { - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] type _ForeignTy; - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] static _FOREIGN_STATIC: (); - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only @@ -89,28 +89,28 @@ extern "C" { fn _foreign_fn(); } -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] static _STATIC: () = (); -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] const _CONST: () = (); -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] type _TyAlias = (); -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only @@ -121,19 +121,19 @@ fn _opaque_ty_definer() -> _OpaqueTy { () } -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] enum _Enum { - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] _Variant { - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only @@ -142,13 +142,13 @@ enum _Enum { }, } -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] union _Union { - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only @@ -156,12 +156,12 @@ union _Union { _field: (), } -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] struct _Struct { - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only @@ -169,113 +169,113 @@ struct _Struct { _field: (), } -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] impl _Struct { - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] const _INHERENT_ASSOC_CONST: () = (); - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] fn _inherent_method() {} } -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] trait _TraitAlias = Copy; -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] trait _Trait { - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] type _AssocTy; - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] const _TRAIT_ASSOC_CONST: (); - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] fn _trait_method(); - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] fn _trait_method_with_default() {} } -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] impl _Trait for () { - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] type _AssocTy = (); - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] const _TRAIT_ASSOC_CONST: () = (); - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] fn _trait_method() {} } -#[rust_gpu::spirv( +#[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] fn _fn( - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only )] _entry_param: (), ) { - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] let _statement = (); - let _closure = #[rust_gpu::spirv( + let _closure = #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only @@ -283,7 +283,7 @@ fn _fn( || {}; ( - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only @@ -293,7 +293,7 @@ fn _fn( ); match () { - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only @@ -303,17 +303,17 @@ fn _fn( } fn _fn_with_generics< - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] '_lifetime_param, - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only )] _TyParam, - #[rust_gpu::spirv( + #[rust_gpu::spirv_v0_9( sampler, block, sampled_image, generic_image_type, // struct-only vertex, // fn-only uniform, position, descriptor_set = 0, binding = 0, flat, invariant, // param-only