1use std::num::NonZeroU32;
2
3use crate::front::wgsl::error::{Error, ExpectedToken, InvalidAssignmentType};
4use crate::front::wgsl::index::Index;
5use crate::front::wgsl::parse::number::Number;
6use crate::front::wgsl::parse::{ast, conv};
7use crate::front::Typifier;
8use crate::proc::{
9 ensure_block_returns, Alignment, ConstantEvaluator, Emitter, Layouter, ResolveContext,
10};
11use crate::{Arena, FastHashMap, FastIndexMap, Handle, Span};
12
13mod construction;
14mod conversion;
15
16macro_rules! resolve_inner {
30 ($ctx:ident, $expr:expr) => {{
31 $ctx.grow_types($expr)?;
32 $ctx.typifier()[$expr].inner_with(&$ctx.module.types)
33 }};
34}
35pub(super) use resolve_inner;
36
37macro_rules! resolve_inner_binary {
45 ($ctx:ident, $left:expr, $right:expr) => {{
46 $ctx.grow_types($left)?;
47 $ctx.grow_types($right)?;
48 (
49 $ctx.typifier()[$left].inner_with(&$ctx.module.types),
50 $ctx.typifier()[$right].inner_with(&$ctx.module.types),
51 )
52 }};
53}
54
55macro_rules! resolve {
65 ($ctx:ident, $expr:expr) => {{
66 $ctx.grow_types($expr)?;
67 &$ctx.typifier()[$expr]
68 }};
69}
70pub(super) use resolve;
71
72pub struct GlobalContext<'source, 'temp, 'out> {
74 ast_expressions: &'temp Arena<ast::Expression<'source>>,
76
77 types: &'temp Arena<ast::Type<'source>>,
79
80 globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,
84
85 module: &'out mut crate::Module,
87
88 const_typifier: &'temp mut Typifier,
89
90 global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
91}
92
93impl<'source> GlobalContext<'source, '_, '_> {
94 fn as_const(&mut self) -> ExpressionContext<'source, '_, '_> {
95 ExpressionContext {
96 ast_expressions: self.ast_expressions,
97 globals: self.globals,
98 types: self.types,
99 module: self.module,
100 const_typifier: self.const_typifier,
101 expr_type: ExpressionContextType::Constant(None),
102 global_expression_kind_tracker: self.global_expression_kind_tracker,
103 }
104 }
105
106 fn as_override(&mut self) -> ExpressionContext<'source, '_, '_> {
107 ExpressionContext {
108 ast_expressions: self.ast_expressions,
109 globals: self.globals,
110 types: self.types,
111 module: self.module,
112 const_typifier: self.const_typifier,
113 expr_type: ExpressionContextType::Override,
114 global_expression_kind_tracker: self.global_expression_kind_tracker,
115 }
116 }
117
118 fn ensure_type_exists(
119 &mut self,
120 name: Option<String>,
121 inner: crate::TypeInner,
122 ) -> Handle<crate::Type> {
123 self.module
124 .types
125 .insert(crate::Type { inner, name }, Span::UNDEFINED)
126 }
127}
128
129pub struct StatementContext<'source, 'temp, 'out> {
131 ast_expressions: &'temp Arena<ast::Expression<'source>>,
137
138 types: &'temp Arena<ast::Type<'source>>,
143
144 globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,
148
149 local_table:
164 &'temp mut FastHashMap<Handle<ast::Local>, Declared<Typed<Handle<crate::Expression>>>>,
165
166 const_typifier: &'temp mut Typifier,
167 typifier: &'temp mut Typifier,
168 function: &'out mut crate::Function,
169 named_expressions: &'out mut FastIndexMap<Handle<crate::Expression>, (String, Span)>,
172 module: &'out mut crate::Module,
173
174 local_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
184 global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
185}
186
187impl<'a, 'temp> StatementContext<'a, 'temp, '_> {
188 fn as_const<'t>(
189 &'t mut self,
190 block: &'t mut crate::Block,
191 emitter: &'t mut Emitter,
192 ) -> ExpressionContext<'a, 't, 't>
193 where
194 'temp: 't,
195 {
196 ExpressionContext {
197 globals: self.globals,
198 types: self.types,
199 ast_expressions: self.ast_expressions,
200 const_typifier: self.const_typifier,
201 global_expression_kind_tracker: self.global_expression_kind_tracker,
202 module: self.module,
203 expr_type: ExpressionContextType::Constant(Some(LocalExpressionContext {
204 local_table: self.local_table,
205 function: self.function,
206 block,
207 emitter,
208 typifier: self.typifier,
209 local_expression_kind_tracker: self.local_expression_kind_tracker,
210 })),
211 }
212 }
213
214 fn as_expression<'t>(
215 &'t mut self,
216 block: &'t mut crate::Block,
217 emitter: &'t mut Emitter,
218 ) -> ExpressionContext<'a, 't, 't>
219 where
220 'temp: 't,
221 {
222 ExpressionContext {
223 globals: self.globals,
224 types: self.types,
225 ast_expressions: self.ast_expressions,
226 const_typifier: self.const_typifier,
227 global_expression_kind_tracker: self.global_expression_kind_tracker,
228 module: self.module,
229 expr_type: ExpressionContextType::Runtime(LocalExpressionContext {
230 local_table: self.local_table,
231 function: self.function,
232 block,
233 emitter,
234 typifier: self.typifier,
235 local_expression_kind_tracker: self.local_expression_kind_tracker,
236 }),
237 }
238 }
239
240 fn as_global(&mut self) -> GlobalContext<'a, '_, '_> {
241 GlobalContext {
242 ast_expressions: self.ast_expressions,
243 globals: self.globals,
244 types: self.types,
245 module: self.module,
246 const_typifier: self.const_typifier,
247 global_expression_kind_tracker: self.global_expression_kind_tracker,
248 }
249 }
250
251 fn invalid_assignment_type(&self, expr: Handle<crate::Expression>) -> InvalidAssignmentType {
252 if let Some(&(_, span)) = self.named_expressions.get(&expr) {
253 InvalidAssignmentType::ImmutableBinding(span)
254 } else {
255 match self.function.expressions[expr] {
256 crate::Expression::Swizzle { .. } => InvalidAssignmentType::Swizzle,
257 crate::Expression::Access { base, .. } => self.invalid_assignment_type(base),
258 crate::Expression::AccessIndex { base, .. } => self.invalid_assignment_type(base),
259 _ => InvalidAssignmentType::Other,
260 }
261 }
262 }
263}
264
265pub struct LocalExpressionContext<'temp, 'out> {
266 local_table: &'temp FastHashMap<Handle<ast::Local>, Declared<Typed<Handle<crate::Expression>>>>,
271
272 function: &'out mut crate::Function,
273 block: &'temp mut crate::Block,
274 emitter: &'temp mut Emitter,
275 typifier: &'temp mut Typifier,
276
277 local_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
282}
283
284pub enum ExpressionContextType<'temp, 'out> {
286 Runtime(LocalExpressionContext<'temp, 'out>),
293
294 Constant(Option<LocalExpressionContext<'temp, 'out>>),
301
302 Override,
309}
310
311pub struct ExpressionContext<'source, 'temp, 'out> {
349 ast_expressions: &'temp Arena<ast::Expression<'source>>,
351 types: &'temp Arena<ast::Type<'source>>,
352
353 globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,
357
358 module: &'out mut crate::Module,
362
363 const_typifier: &'temp mut Typifier,
367 global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
368
369 expr_type: ExpressionContextType<'temp, 'out>,
372}
373
374impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
375 fn as_const(&mut self) -> ExpressionContext<'source, '_, '_> {
376 ExpressionContext {
377 globals: self.globals,
378 types: self.types,
379 ast_expressions: self.ast_expressions,
380 const_typifier: self.const_typifier,
381 module: self.module,
382 expr_type: ExpressionContextType::Constant(None),
383 global_expression_kind_tracker: self.global_expression_kind_tracker,
384 }
385 }
386
387 fn as_global(&mut self) -> GlobalContext<'source, '_, '_> {
388 GlobalContext {
389 ast_expressions: self.ast_expressions,
390 globals: self.globals,
391 types: self.types,
392 module: self.module,
393 const_typifier: self.const_typifier,
394 global_expression_kind_tracker: self.global_expression_kind_tracker,
395 }
396 }
397
398 fn as_const_evaluator(&mut self) -> ConstantEvaluator {
399 match self.expr_type {
400 ExpressionContextType::Runtime(ref mut rctx) => ConstantEvaluator::for_wgsl_function(
401 self.module,
402 &mut rctx.function.expressions,
403 rctx.local_expression_kind_tracker,
404 rctx.emitter,
405 rctx.block,
406 false,
407 ),
408 ExpressionContextType::Constant(Some(ref mut rctx)) => {
409 ConstantEvaluator::for_wgsl_function(
410 self.module,
411 &mut rctx.function.expressions,
412 rctx.local_expression_kind_tracker,
413 rctx.emitter,
414 rctx.block,
415 true,
416 )
417 }
418 ExpressionContextType::Constant(None) => ConstantEvaluator::for_wgsl_module(
419 self.module,
420 self.global_expression_kind_tracker,
421 false,
422 ),
423 ExpressionContextType::Override => ConstantEvaluator::for_wgsl_module(
424 self.module,
425 self.global_expression_kind_tracker,
426 true,
427 ),
428 }
429 }
430
431 fn append_expression(
432 &mut self,
433 expr: crate::Expression,
434 span: Span,
435 ) -> Result<Handle<crate::Expression>, Error<'source>> {
436 let mut eval = self.as_const_evaluator();
437 eval.try_eval_and_append(expr, span)
438 .map_err(|e| Error::ConstantEvaluatorError(e.into(), span))
439 }
440
441 fn const_access(&self, handle: Handle<crate::Expression>) -> Option<u32> {
442 match self.expr_type {
443 ExpressionContextType::Runtime(ref ctx) => {
444 if !ctx.local_expression_kind_tracker.is_const(handle) {
445 return None;
446 }
447
448 self.module
449 .to_ctx()
450 .eval_expr_to_u32_from(handle, &ctx.function.expressions)
451 .ok()
452 }
453 ExpressionContextType::Constant(Some(ref ctx)) => {
454 assert!(ctx.local_expression_kind_tracker.is_const(handle));
455 self.module
456 .to_ctx()
457 .eval_expr_to_u32_from(handle, &ctx.function.expressions)
458 .ok()
459 }
460 ExpressionContextType::Constant(None) => {
461 self.module.to_ctx().eval_expr_to_u32(handle).ok()
462 }
463 ExpressionContextType::Override => None,
464 }
465 }
466
467 fn get_expression_span(&self, handle: Handle<crate::Expression>) -> Span {
468 match self.expr_type {
469 ExpressionContextType::Runtime(ref ctx)
470 | ExpressionContextType::Constant(Some(ref ctx)) => {
471 ctx.function.expressions.get_span(handle)
472 }
473 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
474 self.module.global_expressions.get_span(handle)
475 }
476 }
477 }
478
479 fn typifier(&self) -> &Typifier {
480 match self.expr_type {
481 ExpressionContextType::Runtime(ref ctx)
482 | ExpressionContextType::Constant(Some(ref ctx)) => ctx.typifier,
483 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
484 self.const_typifier
485 }
486 }
487 }
488
489 fn local(
490 &mut self,
491 local: &Handle<ast::Local>,
492 span: Span,
493 ) -> Result<Typed<Handle<crate::Expression>>, Error<'source>> {
494 match self.expr_type {
495 ExpressionContextType::Runtime(ref ctx) => Ok(ctx.local_table[local].runtime()),
496 ExpressionContextType::Constant(Some(ref ctx)) => ctx.local_table[local]
497 .const_time()
498 .ok_or(Error::UnexpectedOperationInConstContext(span)),
499 _ => Err(Error::UnexpectedOperationInConstContext(span)),
500 }
501 }
502
503 fn runtime_expression_ctx(
504 &mut self,
505 span: Span,
506 ) -> Result<&mut LocalExpressionContext<'temp, 'out>, Error<'source>> {
507 match self.expr_type {
508 ExpressionContextType::Runtime(ref mut ctx) => Ok(ctx),
509 ExpressionContextType::Constant(_) | ExpressionContextType::Override => {
510 Err(Error::UnexpectedOperationInConstContext(span))
511 }
512 }
513 }
514
515 fn gather_component(
516 &mut self,
517 expr: Handle<crate::Expression>,
518 component_span: Span,
519 gather_span: Span,
520 ) -> Result<crate::SwizzleComponent, Error<'source>> {
521 match self.expr_type {
522 ExpressionContextType::Runtime(ref rctx) => {
523 if !rctx.local_expression_kind_tracker.is_const(expr) {
524 return Err(Error::ExpectedConstExprConcreteIntegerScalar(
525 component_span,
526 ));
527 }
528
529 let index = self
530 .module
531 .to_ctx()
532 .eval_expr_to_u32_from(expr, &rctx.function.expressions)
533 .map_err(|err| match err {
534 crate::proc::U32EvalError::NonConst => {
535 Error::ExpectedConstExprConcreteIntegerScalar(component_span)
536 }
537 crate::proc::U32EvalError::Negative => {
538 Error::ExpectedNonNegative(component_span)
539 }
540 })?;
541 crate::SwizzleComponent::XYZW
542 .get(index as usize)
543 .copied()
544 .ok_or(Error::InvalidGatherComponent(component_span))
545 }
546 ExpressionContextType::Constant(_) | ExpressionContextType::Override => {
549 Err(Error::UnexpectedOperationInConstContext(gather_span))
550 }
551 }
552 }
553
554 fn register_type(
565 &mut self,
566 handle: Handle<crate::Expression>,
567 ) -> Result<Handle<crate::Type>, Error<'source>> {
568 self.grow_types(handle)?;
569 let typifier = match self.expr_type {
573 ExpressionContextType::Runtime(ref ctx)
574 | ExpressionContextType::Constant(Some(ref ctx)) => ctx.typifier,
575 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
576 &*self.const_typifier
577 }
578 };
579 Ok(typifier.register_type(handle, &mut self.module.types))
580 }
581
582 fn grow_types(
603 &mut self,
604 handle: Handle<crate::Expression>,
605 ) -> Result<&mut Self, Error<'source>> {
606 let empty_arena = Arena::new();
607 let resolve_ctx;
608 let typifier;
609 let expressions;
610 match self.expr_type {
611 ExpressionContextType::Runtime(ref mut ctx)
612 | ExpressionContextType::Constant(Some(ref mut ctx)) => {
613 resolve_ctx = ResolveContext::with_locals(
614 self.module,
615 &ctx.function.local_variables,
616 &ctx.function.arguments,
617 );
618 typifier = &mut *ctx.typifier;
619 expressions = &ctx.function.expressions;
620 }
621 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
622 resolve_ctx = ResolveContext::with_locals(self.module, &empty_arena, &[]);
623 typifier = self.const_typifier;
624 expressions = &self.module.global_expressions;
625 }
626 };
627 typifier
628 .grow(handle, expressions, &resolve_ctx)
629 .map_err(Error::InvalidResolve)?;
630
631 Ok(self)
632 }
633
634 fn image_data(
635 &mut self,
636 image: Handle<crate::Expression>,
637 span: Span,
638 ) -> Result<(crate::ImageClass, bool), Error<'source>> {
639 match *resolve_inner!(self, image) {
640 crate::TypeInner::Image { class, arrayed, .. } => Ok((class, arrayed)),
641 _ => Err(Error::BadTexture(span)),
642 }
643 }
644
645 fn prepare_args<'b>(
646 &mut self,
647 args: &'b [Handle<ast::Expression<'source>>],
648 min_args: u32,
649 span: Span,
650 ) -> ArgumentContext<'b, 'source> {
651 ArgumentContext {
652 args: args.iter(),
653 min_args,
654 args_used: 0,
655 total_args: args.len() as u32,
656 span,
657 }
658 }
659
660 fn binary_op_splat(
668 &mut self,
669 op: crate::BinaryOperator,
670 left: &mut Handle<crate::Expression>,
671 right: &mut Handle<crate::Expression>,
672 ) -> Result<(), Error<'source>> {
673 if matches!(
674 op,
675 crate::BinaryOperator::Add
676 | crate::BinaryOperator::Subtract
677 | crate::BinaryOperator::Divide
678 | crate::BinaryOperator::Modulo
679 ) {
680 match resolve_inner_binary!(self, *left, *right) {
681 (&crate::TypeInner::Vector { size, .. }, &crate::TypeInner::Scalar { .. }) => {
682 *right = self.append_expression(
683 crate::Expression::Splat {
684 size,
685 value: *right,
686 },
687 self.get_expression_span(*right),
688 )?;
689 }
690 (&crate::TypeInner::Scalar { .. }, &crate::TypeInner::Vector { size, .. }) => {
691 *left = self.append_expression(
692 crate::Expression::Splat { size, value: *left },
693 self.get_expression_span(*left),
694 )?;
695 }
696 _ => {}
697 }
698 }
699
700 Ok(())
701 }
702
703 fn interrupt_emitter(
708 &mut self,
709 expression: crate::Expression,
710 span: Span,
711 ) -> Result<Handle<crate::Expression>, Error<'source>> {
712 match self.expr_type {
713 ExpressionContextType::Runtime(ref mut rctx)
714 | ExpressionContextType::Constant(Some(ref mut rctx)) => {
715 rctx.block
716 .extend(rctx.emitter.finish(&rctx.function.expressions));
717 }
718 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {}
719 }
720 let result = self.append_expression(expression, span);
721 match self.expr_type {
722 ExpressionContextType::Runtime(ref mut rctx)
723 | ExpressionContextType::Constant(Some(ref mut rctx)) => {
724 rctx.emitter.start(&rctx.function.expressions);
725 }
726 ExpressionContextType::Constant(None) | ExpressionContextType::Override => {}
727 }
728 result
729 }
730
731 fn apply_load_rule(
736 &mut self,
737 expr: Typed<Handle<crate::Expression>>,
738 ) -> Result<Handle<crate::Expression>, Error<'source>> {
739 match expr {
740 Typed::Reference(pointer) => {
741 let load = crate::Expression::Load { pointer };
742 let span = self.get_expression_span(pointer);
743 self.append_expression(load, span)
744 }
745 Typed::Plain(handle) => Ok(handle),
746 }
747 }
748
749 fn ensure_type_exists(&mut self, inner: crate::TypeInner) -> Handle<crate::Type> {
750 self.as_global().ensure_type_exists(None, inner)
751 }
752}
753
754struct ArgumentContext<'ctx, 'source> {
755 args: std::slice::Iter<'ctx, Handle<ast::Expression<'source>>>,
756 min_args: u32,
757 args_used: u32,
758 total_args: u32,
759 span: Span,
760}
761
762impl<'source> ArgumentContext<'_, 'source> {
763 pub fn finish(self) -> Result<(), Error<'source>> {
764 if self.args.len() == 0 {
765 Ok(())
766 } else {
767 Err(Error::WrongArgumentCount {
768 found: self.total_args,
769 expected: self.min_args..self.args_used + 1,
770 span: self.span,
771 })
772 }
773 }
774
775 pub fn next(&mut self) -> Result<Handle<ast::Expression<'source>>, Error<'source>> {
776 match self.args.next().copied() {
777 Some(arg) => {
778 self.args_used += 1;
779 Ok(arg)
780 }
781 None => Err(Error::WrongArgumentCount {
782 found: self.total_args,
783 expected: self.min_args..self.args_used + 1,
784 span: self.span,
785 }),
786 }
787 }
788}
789
790#[derive(Debug, Copy, Clone)]
791enum Declared<T> {
792 Const(T),
794
795 Runtime(T),
797}
798
799impl<T> Declared<T> {
800 fn runtime(self) -> T {
801 match self {
802 Declared::Const(t) | Declared::Runtime(t) => t,
803 }
804 }
805
806 fn const_time(self) -> Option<T> {
807 match self {
808 Declared::Const(t) => Some(t),
809 Declared::Runtime(_) => None,
810 }
811 }
812}
813
814#[derive(Debug, Copy, Clone)]
836enum Typed<T> {
837 Reference(T),
839
840 Plain(T),
842}
843
844impl<T> Typed<T> {
845 fn map<U>(self, mut f: impl FnMut(T) -> U) -> Typed<U> {
846 match self {
847 Self::Reference(v) => Typed::Reference(f(v)),
848 Self::Plain(v) => Typed::Plain(f(v)),
849 }
850 }
851
852 fn try_map<U, E>(self, mut f: impl FnMut(T) -> Result<U, E>) -> Result<Typed<U>, E> {
853 Ok(match self {
854 Self::Reference(expr) => Typed::Reference(f(expr)?),
855 Self::Plain(expr) => Typed::Plain(f(expr)?),
856 })
857 }
858}
859
860enum Components {
866 Single(u32),
867 Swizzle {
868 size: crate::VectorSize,
869 pattern: [crate::SwizzleComponent; 4],
870 },
871}
872
873impl Components {
874 const fn letter_component(letter: char) -> Option<crate::SwizzleComponent> {
875 use crate::SwizzleComponent as Sc;
876 match letter {
877 'x' | 'r' => Some(Sc::X),
878 'y' | 'g' => Some(Sc::Y),
879 'z' | 'b' => Some(Sc::Z),
880 'w' | 'a' => Some(Sc::W),
881 _ => None,
882 }
883 }
884
885 fn single_component(name: &str, name_span: Span) -> Result<u32, Error> {
886 let ch = name.chars().next().ok_or(Error::BadAccessor(name_span))?;
887 match Self::letter_component(ch) {
888 Some(sc) => Ok(sc as u32),
889 None => Err(Error::BadAccessor(name_span)),
890 }
891 }
892
893 fn new(name: &str, name_span: Span) -> Result<Self, Error> {
897 let size = match name.len() {
898 1 => return Ok(Components::Single(Self::single_component(name, name_span)?)),
899 2 => crate::VectorSize::Bi,
900 3 => crate::VectorSize::Tri,
901 4 => crate::VectorSize::Quad,
902 _ => return Err(Error::BadAccessor(name_span)),
903 };
904
905 let mut pattern = [crate::SwizzleComponent::X; 4];
906 for (comp, ch) in pattern.iter_mut().zip(name.chars()) {
907 *comp = Self::letter_component(ch).ok_or(Error::BadAccessor(name_span))?;
908 }
909
910 if name.chars().all(|c| matches!(c, 'x' | 'y' | 'z' | 'w'))
911 || name.chars().all(|c| matches!(c, 'r' | 'g' | 'b' | 'a'))
912 {
913 Ok(Components::Swizzle { size, pattern })
914 } else {
915 Err(Error::BadAccessor(name_span))
916 }
917 }
918}
919
920enum LoweredGlobalDecl {
922 Function(Handle<crate::Function>),
923 Var(Handle<crate::GlobalVariable>),
924 Const(Handle<crate::Constant>),
925 Override(Handle<crate::Override>),
926 Type(Handle<crate::Type>),
927 EntryPoint,
928}
929
930enum Texture {
931 Gather,
932 GatherCompare,
933
934 Sample,
935 SampleBias,
936 SampleCompare,
937 SampleCompareLevel,
938 SampleGrad,
939 SampleLevel,
940 }
942
943impl Texture {
944 pub fn map(word: &str) -> Option<Self> {
945 Some(match word {
946 "textureGather" => Self::Gather,
947 "textureGatherCompare" => Self::GatherCompare,
948
949 "textureSample" => Self::Sample,
950 "textureSampleBias" => Self::SampleBias,
951 "textureSampleCompare" => Self::SampleCompare,
952 "textureSampleCompareLevel" => Self::SampleCompareLevel,
953 "textureSampleGrad" => Self::SampleGrad,
954 "textureSampleLevel" => Self::SampleLevel,
955 _ => return None,
957 })
958 }
959
960 pub const fn min_argument_count(&self) -> u32 {
961 match *self {
962 Self::Gather => 3,
963 Self::GatherCompare => 4,
964
965 Self::Sample => 3,
966 Self::SampleBias => 5,
967 Self::SampleCompare => 5,
968 Self::SampleCompareLevel => 5,
969 Self::SampleGrad => 6,
970 Self::SampleLevel => 5,
971 }
973 }
974}
975
976enum SubgroupGather {
977 BroadcastFirst,
978 Broadcast,
979 Shuffle,
980 ShuffleDown,
981 ShuffleUp,
982 ShuffleXor,
983}
984
985impl SubgroupGather {
986 pub fn map(word: &str) -> Option<Self> {
987 Some(match word {
988 "subgroupBroadcastFirst" => Self::BroadcastFirst,
989 "subgroupBroadcast" => Self::Broadcast,
990 "subgroupShuffle" => Self::Shuffle,
991 "subgroupShuffleDown" => Self::ShuffleDown,
992 "subgroupShuffleUp" => Self::ShuffleUp,
993 "subgroupShuffleXor" => Self::ShuffleXor,
994 _ => return None,
995 })
996 }
997}
998
999pub struct Lowerer<'source, 'temp> {
1000 index: &'temp Index<'source>,
1001 layouter: Layouter,
1002}
1003
1004impl<'source, 'temp> Lowerer<'source, 'temp> {
1005 pub fn new(index: &'temp Index<'source>) -> Self {
1006 Self {
1007 index,
1008 layouter: Layouter::default(),
1009 }
1010 }
1011
1012 pub fn lower(
1013 &mut self,
1014 tu: &'temp ast::TranslationUnit<'source>,
1015 ) -> Result<crate::Module, Error<'source>> {
1016 let mut module = crate::Module {
1017 diagnostic_filters: tu.diagnostic_filters.clone(),
1018 diagnostic_filter_leaf: tu.diagnostic_filter_leaf,
1019 ..Default::default()
1020 };
1021
1022 let mut ctx = GlobalContext {
1023 ast_expressions: &tu.expressions,
1024 globals: &mut FastHashMap::default(),
1025 types: &tu.types,
1026 module: &mut module,
1027 const_typifier: &mut Typifier::new(),
1028 global_expression_kind_tracker: &mut crate::proc::ExpressionKindTracker::new(),
1029 };
1030
1031 for decl_handle in self.index.visit_ordered() {
1032 let span = tu.decls.get_span(decl_handle);
1033 let decl = &tu.decls[decl_handle];
1034
1035 match decl.kind {
1036 ast::GlobalDeclKind::Fn(ref f) => {
1037 let lowered_decl = self.function(f, span, &mut ctx)?;
1038 ctx.globals.insert(f.name.name, lowered_decl);
1039 }
1040 ast::GlobalDeclKind::Var(ref v) => {
1041 let explicit_ty =
1042 v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx))
1043 .transpose()?;
1044
1045 let (ty, initializer) =
1046 self.type_and_init(v.name, v.init, explicit_ty, &mut ctx.as_override())?;
1047
1048 let binding = if let Some(ref binding) = v.binding {
1049 Some(crate::ResourceBinding {
1050 group: self.const_u32(binding.group, &mut ctx.as_const())?.0,
1051 binding: self.const_u32(binding.binding, &mut ctx.as_const())?.0,
1052 })
1053 } else {
1054 None
1055 };
1056
1057 let handle = ctx.module.global_variables.append(
1058 crate::GlobalVariable {
1059 name: Some(v.name.name.to_string()),
1060 space: v.space,
1061 binding,
1062 ty,
1063 init: initializer,
1064 },
1065 span,
1066 );
1067
1068 ctx.globals
1069 .insert(v.name.name, LoweredGlobalDecl::Var(handle));
1070 }
1071 ast::GlobalDeclKind::Const(ref c) => {
1072 let mut ectx = ctx.as_const();
1073 let mut init = self.expression_for_abstract(c.init, &mut ectx)?;
1074
1075 let ty;
1076 if let Some(explicit_ty) = c.ty {
1077 let explicit_ty =
1078 self.resolve_ast_type(explicit_ty, &mut ectx.as_global())?;
1079 let explicit_ty_res = crate::proc::TypeResolution::Handle(explicit_ty);
1080 init = ectx
1081 .try_automatic_conversions(init, &explicit_ty_res, c.name.span)
1082 .map_err(|error| match error {
1083 Error::AutoConversion(e) => Error::InitializationTypeMismatch {
1084 name: c.name.span,
1085 expected: e.dest_type,
1086 got: e.source_type,
1087 },
1088 other => other,
1089 })?;
1090 ty = explicit_ty;
1091 } else {
1092 init = ectx.concretize(init)?;
1093 ty = ectx.register_type(init)?;
1094 }
1095
1096 let handle = ctx.module.constants.append(
1097 crate::Constant {
1098 name: Some(c.name.name.to_string()),
1099 ty,
1100 init,
1101 },
1102 span,
1103 );
1104
1105 ctx.globals
1106 .insert(c.name.name, LoweredGlobalDecl::Const(handle));
1107 }
1108 ast::GlobalDeclKind::Override(ref o) => {
1109 let explicit_ty =
1110 o.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx))
1111 .transpose()?;
1112
1113 let mut ectx = ctx.as_override();
1114
1115 let (ty, init) = self.type_and_init(o.name, o.init, explicit_ty, &mut ectx)?;
1116
1117 let id =
1118 o.id.map(|id| self.const_u32(id, &mut ctx.as_const()))
1119 .transpose()?;
1120
1121 let id = if let Some((id, id_span)) = id {
1122 Some(
1123 u16::try_from(id)
1124 .map_err(|_| Error::PipelineConstantIDValue(id_span))?,
1125 )
1126 } else {
1127 None
1128 };
1129
1130 let handle = ctx.module.overrides.append(
1131 crate::Override {
1132 name: Some(o.name.name.to_string()),
1133 id,
1134 ty,
1135 init,
1136 },
1137 span,
1138 );
1139
1140 ctx.globals
1141 .insert(o.name.name, LoweredGlobalDecl::Override(handle));
1142 }
1143 ast::GlobalDeclKind::Struct(ref s) => {
1144 let handle = self.r#struct(s, span, &mut ctx)?;
1145 ctx.globals
1146 .insert(s.name.name, LoweredGlobalDecl::Type(handle));
1147 }
1148 ast::GlobalDeclKind::Type(ref alias) => {
1149 let ty = self.resolve_named_ast_type(
1150 alias.ty,
1151 Some(alias.name.name.to_string()),
1152 &mut ctx,
1153 )?;
1154 ctx.globals
1155 .insert(alias.name.name, LoweredGlobalDecl::Type(ty));
1156 }
1157 ast::GlobalDeclKind::ConstAssert(condition) => {
1158 let condition = self.expression(condition, &mut ctx.as_const())?;
1159
1160 let span = ctx.module.global_expressions.get_span(condition);
1161 match ctx
1162 .module
1163 .to_ctx()
1164 .eval_expr_to_bool_from(condition, &ctx.module.global_expressions)
1165 {
1166 Some(true) => Ok(()),
1167 Some(false) => Err(Error::ConstAssertFailed(span)),
1168 _ => Err(Error::NotBool(span)),
1169 }?;
1170 }
1171 }
1172 }
1173
1174 crate::compact::compact(&mut module);
1178
1179 Ok(module)
1180 }
1181
1182 fn type_and_init(
1184 &mut self,
1185 name: ast::Ident<'source>,
1186 init: Option<Handle<ast::Expression<'source>>>,
1187 explicit_ty: Option<Handle<crate::Type>>,
1188 ectx: &mut ExpressionContext<'source, '_, '_>,
1189 ) -> Result<(Handle<crate::Type>, Option<Handle<crate::Expression>>), Error<'source>> {
1190 let ty;
1191 let initializer;
1192 match (init, explicit_ty) {
1193 (Some(init), Some(explicit_ty)) => {
1194 let init = self.expression_for_abstract(init, ectx)?;
1195 let ty_res = crate::proc::TypeResolution::Handle(explicit_ty);
1196 let init = ectx
1197 .try_automatic_conversions(init, &ty_res, name.span)
1198 .map_err(|error| match error {
1199 Error::AutoConversion(e) => Error::InitializationTypeMismatch {
1200 name: name.span,
1201 expected: e.dest_type,
1202 got: e.source_type,
1203 },
1204 other => other,
1205 })?;
1206 ty = explicit_ty;
1207 initializer = Some(init);
1208 }
1209 (Some(init), None) => {
1210 let concretized = self.expression(init, ectx)?;
1211 ty = ectx.register_type(concretized)?;
1212 initializer = Some(concretized);
1213 }
1214 (None, Some(explicit_ty)) => {
1215 ty = explicit_ty;
1216 initializer = None;
1217 }
1218 (None, None) => return Err(Error::DeclMissingTypeAndInit(name.span)),
1219 }
1220 Ok((ty, initializer))
1221 }
1222
1223 fn function(
1224 &mut self,
1225 f: &ast::Function<'source>,
1226 span: Span,
1227 ctx: &mut GlobalContext<'source, '_, '_>,
1228 ) -> Result<LoweredGlobalDecl, Error<'source>> {
1229 let mut local_table = FastHashMap::default();
1230 let mut expressions = Arena::new();
1231 let mut named_expressions = FastIndexMap::default();
1232 let mut local_expression_kind_tracker = crate::proc::ExpressionKindTracker::new();
1233
1234 let arguments = f
1235 .arguments
1236 .iter()
1237 .enumerate()
1238 .map(|(i, arg)| -> Result<_, Error<'_>> {
1239 let ty = self.resolve_ast_type(arg.ty, ctx)?;
1240 let expr = expressions
1241 .append(crate::Expression::FunctionArgument(i as u32), arg.name.span);
1242 local_table.insert(arg.handle, Declared::Runtime(Typed::Plain(expr)));
1243 named_expressions.insert(expr, (arg.name.name.to_string(), arg.name.span));
1244 local_expression_kind_tracker.insert(expr, crate::proc::ExpressionKind::Runtime);
1245
1246 Ok(crate::FunctionArgument {
1247 name: Some(arg.name.name.to_string()),
1248 ty,
1249 binding: self.binding(&arg.binding, ty, ctx)?,
1250 })
1251 })
1252 .collect::<Result<Vec<_>, _>>()?;
1253
1254 let result = f
1255 .result
1256 .as_ref()
1257 .map(|res| -> Result<_, Error<'_>> {
1258 let ty = self.resolve_ast_type(res.ty, ctx)?;
1259 Ok(crate::FunctionResult {
1260 ty,
1261 binding: self.binding(&res.binding, ty, ctx)?,
1262 })
1263 })
1264 .transpose()?;
1265
1266 let mut function = crate::Function {
1267 name: Some(f.name.name.to_string()),
1268 arguments,
1269 result,
1270 local_variables: Arena::new(),
1271 expressions,
1272 named_expressions: crate::NamedExpressions::default(),
1273 body: crate::Block::default(),
1274 diagnostic_filter_leaf: f.diagnostic_filter_leaf,
1275 };
1276
1277 let mut typifier = Typifier::default();
1278 let mut stmt_ctx = StatementContext {
1279 local_table: &mut local_table,
1280 globals: ctx.globals,
1281 ast_expressions: ctx.ast_expressions,
1282 const_typifier: ctx.const_typifier,
1283 typifier: &mut typifier,
1284 function: &mut function,
1285 named_expressions: &mut named_expressions,
1286 types: ctx.types,
1287 module: ctx.module,
1288 local_expression_kind_tracker: &mut local_expression_kind_tracker,
1289 global_expression_kind_tracker: ctx.global_expression_kind_tracker,
1290 };
1291 let mut body = self.block(&f.body, false, &mut stmt_ctx)?;
1292 ensure_block_returns(&mut body);
1293
1294 function.body = body;
1295 function.named_expressions = named_expressions
1296 .into_iter()
1297 .map(|(key, (name, _))| (key, name))
1298 .collect();
1299
1300 if let Some(ref entry) = f.entry_point {
1301 let workgroup_size_info = if let Some(workgroup_size) = entry.workgroup_size {
1302 let mut workgroup_size_out = [1; 3];
1304 let mut workgroup_size_overrides_out = [None; 3];
1305 for (i, size) in workgroup_size.into_iter().enumerate() {
1306 if let Some(size_expr) = size {
1307 match self.const_u32(size_expr, &mut ctx.as_const()) {
1308 Ok(value) => {
1309 workgroup_size_out[i] = value.0;
1310 }
1311 err => {
1312 if let Err(Error::ConstantEvaluatorError(ref ty, _)) = err {
1313 match **ty {
1314 crate::proc::ConstantEvaluatorError::OverrideExpr => {
1315 workgroup_size_overrides_out[i] =
1316 Some(self.workgroup_size_override(
1317 size_expr,
1318 &mut ctx.as_override(),
1319 )?);
1320 }
1321 _ => {
1322 err?;
1323 }
1324 }
1325 } else {
1326 err?;
1327 }
1328 }
1329 }
1330 }
1331 }
1332 if workgroup_size_overrides_out.iter().all(|x| x.is_none()) {
1333 (workgroup_size_out, None)
1334 } else {
1335 (workgroup_size_out, Some(workgroup_size_overrides_out))
1336 }
1337 } else {
1338 ([0; 3], None)
1339 };
1340
1341 let (workgroup_size, workgroup_size_overrides) = workgroup_size_info;
1342 ctx.module.entry_points.push(crate::EntryPoint {
1343 name: f.name.name.to_string(),
1344 stage: entry.stage,
1345 early_depth_test: entry.early_depth_test,
1346 workgroup_size,
1347 workgroup_size_overrides,
1348 function,
1349 });
1350 Ok(LoweredGlobalDecl::EntryPoint)
1351 } else {
1352 let handle = ctx.module.functions.append(function, span);
1353 Ok(LoweredGlobalDecl::Function(handle))
1354 }
1355 }
1356
1357 fn workgroup_size_override(
1358 &mut self,
1359 size_expr: Handle<ast::Expression<'source>>,
1360 ctx: &mut ExpressionContext<'source, '_, '_>,
1361 ) -> Result<Handle<crate::Expression>, Error<'source>> {
1362 let span = ctx.ast_expressions.get_span(size_expr);
1363 let expr = self.expression(size_expr, ctx)?;
1364 match resolve_inner!(ctx, expr).scalar_kind().ok_or(0) {
1365 Ok(crate::ScalarKind::Sint) | Ok(crate::ScalarKind::Uint) => Ok(expr),
1366 _ => Err(Error::ExpectedConstExprConcreteIntegerScalar(span)),
1367 }
1368 }
1369
1370 fn block(
1371 &mut self,
1372 b: &ast::Block<'source>,
1373 is_inside_loop: bool,
1374 ctx: &mut StatementContext<'source, '_, '_>,
1375 ) -> Result<crate::Block, Error<'source>> {
1376 let mut block = crate::Block::default();
1377
1378 for stmt in b.stmts.iter() {
1379 self.statement(stmt, &mut block, is_inside_loop, ctx)?;
1380 }
1381
1382 Ok(block)
1383 }
1384
1385 fn statement(
1386 &mut self,
1387 stmt: &ast::Statement<'source>,
1388 block: &mut crate::Block,
1389 is_inside_loop: bool,
1390 ctx: &mut StatementContext<'source, '_, '_>,
1391 ) -> Result<(), Error<'source>> {
1392 let out = match stmt.kind {
1393 ast::StatementKind::Block(ref block) => {
1394 let block = self.block(block, is_inside_loop, ctx)?;
1395 crate::Statement::Block(block)
1396 }
1397 ast::StatementKind::LocalDecl(ref decl) => match *decl {
1398 ast::LocalDecl::Let(ref l) => {
1399 let mut emitter = Emitter::default();
1400 emitter.start(&ctx.function.expressions);
1401
1402 let value =
1403 self.expression(l.init, &mut ctx.as_expression(block, &mut emitter))?;
1404
1405 ctx.local_expression_kind_tracker.force_non_const(value);
1411
1412 let explicit_ty =
1413 l.ty.map(|ty| self.resolve_ast_type(ty, &mut ctx.as_global()))
1414 .transpose()?;
1415
1416 if let Some(ty) = explicit_ty {
1417 let mut ctx = ctx.as_expression(block, &mut emitter);
1418 let init_ty = ctx.register_type(value)?;
1419 if !ctx.module.types[ty]
1420 .inner
1421 .equivalent(&ctx.module.types[init_ty].inner, &ctx.module.types)
1422 {
1423 let gctx = &ctx.module.to_ctx();
1424 return Err(Error::InitializationTypeMismatch {
1425 name: l.name.span,
1426 expected: ty.to_wgsl(gctx).into(),
1427 got: init_ty.to_wgsl(gctx).into(),
1428 });
1429 }
1430 }
1431
1432 block.extend(emitter.finish(&ctx.function.expressions));
1433 ctx.local_table
1434 .insert(l.handle, Declared::Runtime(Typed::Plain(value)));
1435 ctx.named_expressions
1436 .insert(value, (l.name.name.to_string(), l.name.span));
1437
1438 return Ok(());
1439 }
1440 ast::LocalDecl::Var(ref v) => {
1441 let explicit_ty =
1442 v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_global()))
1443 .transpose()?;
1444
1445 let mut emitter = Emitter::default();
1446 emitter.start(&ctx.function.expressions);
1447 let mut ectx = ctx.as_expression(block, &mut emitter);
1448
1449 let ty;
1450 let initializer;
1451 match (v.init, explicit_ty) {
1452 (Some(init), Some(explicit_ty)) => {
1453 let init = self.expression_for_abstract(init, &mut ectx)?;
1454 let ty_res = crate::proc::TypeResolution::Handle(explicit_ty);
1455 let init = ectx
1456 .try_automatic_conversions(init, &ty_res, v.name.span)
1457 .map_err(|error| match error {
1458 Error::AutoConversion(e) => Error::InitializationTypeMismatch {
1459 name: v.name.span,
1460 expected: e.dest_type,
1461 got: e.source_type,
1462 },
1463 other => other,
1464 })?;
1465 ty = explicit_ty;
1466 initializer = Some(init);
1467 }
1468 (Some(init), None) => {
1469 let concretized = self.expression(init, &mut ectx)?;
1470 ty = ectx.register_type(concretized)?;
1471 initializer = Some(concretized);
1472 }
1473 (None, Some(explicit_ty)) => {
1474 ty = explicit_ty;
1475 initializer = None;
1476 }
1477 (None, None) => return Err(Error::DeclMissingTypeAndInit(v.name.span)),
1478 }
1479
1480 let (const_initializer, initializer) = {
1481 match initializer {
1482 Some(init) => {
1483 if is_inside_loop
1491 || !ctx.local_expression_kind_tracker.is_const_or_override(init)
1492 {
1493 (None, Some(init))
1494 } else {
1495 (Some(init), None)
1496 }
1497 }
1498 None => (None, None),
1499 }
1500 };
1501
1502 let var = ctx.function.local_variables.append(
1503 crate::LocalVariable {
1504 name: Some(v.name.name.to_string()),
1505 ty,
1506 init: const_initializer,
1507 },
1508 stmt.span,
1509 );
1510
1511 let handle = ctx.as_expression(block, &mut emitter).interrupt_emitter(
1512 crate::Expression::LocalVariable(var),
1513 Span::UNDEFINED,
1514 )?;
1515 block.extend(emitter.finish(&ctx.function.expressions));
1516 ctx.local_table
1517 .insert(v.handle, Declared::Runtime(Typed::Reference(handle)));
1518
1519 match initializer {
1520 Some(initializer) => crate::Statement::Store {
1521 pointer: handle,
1522 value: initializer,
1523 },
1524 None => return Ok(()),
1525 }
1526 }
1527 ast::LocalDecl::Const(ref c) => {
1528 let mut emitter = Emitter::default();
1529 emitter.start(&ctx.function.expressions);
1530
1531 let ectx = &mut ctx.as_const(block, &mut emitter);
1532
1533 let mut init = self.expression_for_abstract(c.init, ectx)?;
1534
1535 if let Some(explicit_ty) = c.ty {
1536 let explicit_ty =
1537 self.resolve_ast_type(explicit_ty, &mut ectx.as_global())?;
1538 let explicit_ty_res = crate::proc::TypeResolution::Handle(explicit_ty);
1539 init = ectx
1540 .try_automatic_conversions(init, &explicit_ty_res, c.name.span)
1541 .map_err(|error| match error {
1542 Error::AutoConversion(error) => Error::InitializationTypeMismatch {
1543 name: c.name.span,
1544 expected: error.dest_type,
1545 got: error.source_type,
1546 },
1547 other => other,
1548 })?;
1549 } else {
1550 init = ectx.concretize(init)?;
1551 ectx.register_type(init)?;
1552 }
1553
1554 block.extend(emitter.finish(&ctx.function.expressions));
1555 ctx.local_table
1556 .insert(c.handle, Declared::Const(Typed::Plain(init)));
1557 ctx.named_expressions
1558 .insert(init, (c.name.name.to_string(), c.name.span));
1559
1560 return Ok(());
1561 }
1562 },
1563 ast::StatementKind::If {
1564 condition,
1565 ref accept,
1566 ref reject,
1567 } => {
1568 let mut emitter = Emitter::default();
1569 emitter.start(&ctx.function.expressions);
1570
1571 let condition =
1572 self.expression(condition, &mut ctx.as_expression(block, &mut emitter))?;
1573 block.extend(emitter.finish(&ctx.function.expressions));
1574
1575 let accept = self.block(accept, is_inside_loop, ctx)?;
1576 let reject = self.block(reject, is_inside_loop, ctx)?;
1577
1578 crate::Statement::If {
1579 condition,
1580 accept,
1581 reject,
1582 }
1583 }
1584 ast::StatementKind::Switch {
1585 selector,
1586 ref cases,
1587 } => {
1588 let mut emitter = Emitter::default();
1589 emitter.start(&ctx.function.expressions);
1590
1591 let mut ectx = ctx.as_expression(block, &mut emitter);
1592 let selector = self.expression(selector, &mut ectx)?;
1593
1594 let uint =
1595 resolve_inner!(ectx, selector).scalar_kind() == Some(crate::ScalarKind::Uint);
1596 block.extend(emitter.finish(&ctx.function.expressions));
1597
1598 let cases = cases
1599 .iter()
1600 .map(|case| {
1601 Ok(crate::SwitchCase {
1602 value: match case.value {
1603 ast::SwitchValue::Expr(expr) => {
1604 let span = ctx.ast_expressions.get_span(expr);
1605 let expr =
1606 self.expression(expr, &mut ctx.as_global().as_const())?;
1607 match ctx.module.to_ctx().eval_expr_to_literal(expr) {
1608 Some(crate::Literal::I32(value)) if !uint => {
1609 crate::SwitchValue::I32(value)
1610 }
1611 Some(crate::Literal::U32(value)) if uint => {
1612 crate::SwitchValue::U32(value)
1613 }
1614 _ => {
1615 return Err(Error::InvalidSwitchValue { uint, span });
1616 }
1617 }
1618 }
1619 ast::SwitchValue::Default => crate::SwitchValue::Default,
1620 },
1621 body: self.block(&case.body, is_inside_loop, ctx)?,
1622 fall_through: case.fall_through,
1623 })
1624 })
1625 .collect::<Result<_, _>>()?;
1626
1627 crate::Statement::Switch { selector, cases }
1628 }
1629 ast::StatementKind::Loop {
1630 ref body,
1631 ref continuing,
1632 break_if,
1633 } => {
1634 let body = self.block(body, true, ctx)?;
1635 let mut continuing = self.block(continuing, true, ctx)?;
1636
1637 let mut emitter = Emitter::default();
1638 emitter.start(&ctx.function.expressions);
1639 let break_if = break_if
1640 .map(|expr| {
1641 self.expression(expr, &mut ctx.as_expression(&mut continuing, &mut emitter))
1642 })
1643 .transpose()?;
1644 continuing.extend(emitter.finish(&ctx.function.expressions));
1645
1646 crate::Statement::Loop {
1647 body,
1648 continuing,
1649 break_if,
1650 }
1651 }
1652 ast::StatementKind::Break => crate::Statement::Break,
1653 ast::StatementKind::Continue => crate::Statement::Continue,
1654 ast::StatementKind::Return { value } => {
1655 let mut emitter = Emitter::default();
1656 emitter.start(&ctx.function.expressions);
1657
1658 let value = value
1659 .map(|expr| self.expression(expr, &mut ctx.as_expression(block, &mut emitter)))
1660 .transpose()?;
1661 block.extend(emitter.finish(&ctx.function.expressions));
1662
1663 crate::Statement::Return { value }
1664 }
1665 ast::StatementKind::Kill => crate::Statement::Kill,
1666 ast::StatementKind::Call {
1667 ref function,
1668 ref arguments,
1669 } => {
1670 let mut emitter = Emitter::default();
1671 emitter.start(&ctx.function.expressions);
1672
1673 let _ = self.call(
1674 stmt.span,
1675 function,
1676 arguments,
1677 &mut ctx.as_expression(block, &mut emitter),
1678 true,
1679 )?;
1680 block.extend(emitter.finish(&ctx.function.expressions));
1681 return Ok(());
1682 }
1683 ast::StatementKind::Assign {
1684 target: ast_target,
1685 op,
1686 value,
1687 } => {
1688 let mut emitter = Emitter::default();
1689 emitter.start(&ctx.function.expressions);
1690 let target_span = ctx.ast_expressions.get_span(ast_target);
1691
1692 let mut ectx = ctx.as_expression(block, &mut emitter);
1693 let target = self.expression_for_reference(ast_target, &mut ectx)?;
1694 let target_handle = match target {
1695 Typed::Reference(handle) => handle,
1696 Typed::Plain(handle) => {
1697 let ty = ctx.invalid_assignment_type(handle);
1698 return Err(Error::InvalidAssignment {
1699 span: target_span,
1700 ty,
1701 });
1702 }
1703 };
1704
1705 let target_scalar = match op {
1710 Some(crate::BinaryOperator::ShiftLeft | crate::BinaryOperator::ShiftRight) => {
1711 Some(crate::Scalar::U32)
1712 }
1713 _ => resolve_inner!(ectx, target_handle)
1714 .pointer_automatically_convertible_scalar(&ectx.module.types),
1715 };
1716
1717 let value = self.expression_for_abstract(value, &mut ectx)?;
1718 let mut value = match target_scalar {
1719 Some(target_scalar) => ectx.try_automatic_conversion_for_leaf_scalar(
1720 value,
1721 target_scalar,
1722 target_span,
1723 )?,
1724 None => value,
1725 };
1726
1727 let value = match op {
1728 Some(op) => {
1729 let mut left = ectx.apply_load_rule(target)?;
1730 ectx.binary_op_splat(op, &mut left, &mut value)?;
1731 ectx.append_expression(
1732 crate::Expression::Binary {
1733 op,
1734 left,
1735 right: value,
1736 },
1737 stmt.span,
1738 )?
1739 }
1740 None => value,
1741 };
1742 block.extend(emitter.finish(&ctx.function.expressions));
1743
1744 crate::Statement::Store {
1745 pointer: target_handle,
1746 value,
1747 }
1748 }
1749 ast::StatementKind::Increment(value) | ast::StatementKind::Decrement(value) => {
1750 let mut emitter = Emitter::default();
1751 emitter.start(&ctx.function.expressions);
1752
1753 let op = match stmt.kind {
1754 ast::StatementKind::Increment(_) => crate::BinaryOperator::Add,
1755 ast::StatementKind::Decrement(_) => crate::BinaryOperator::Subtract,
1756 _ => unreachable!(),
1757 };
1758
1759 let value_span = ctx.ast_expressions.get_span(value);
1760 let target = self
1761 .expression_for_reference(value, &mut ctx.as_expression(block, &mut emitter))?;
1762 let target_handle = match target {
1763 Typed::Reference(handle) => handle,
1764 Typed::Plain(_) => return Err(Error::BadIncrDecrReferenceType(value_span)),
1765 };
1766
1767 let mut ectx = ctx.as_expression(block, &mut emitter);
1768 let scalar = match *resolve_inner!(ectx, target_handle) {
1769 crate::TypeInner::ValuePointer {
1770 size: None, scalar, ..
1771 } => scalar,
1772 crate::TypeInner::Pointer { base, .. } => match ectx.module.types[base].inner {
1773 crate::TypeInner::Scalar(scalar) => scalar,
1774 _ => return Err(Error::BadIncrDecrReferenceType(value_span)),
1775 },
1776 _ => return Err(Error::BadIncrDecrReferenceType(value_span)),
1777 };
1778 let literal = match scalar.kind {
1779 crate::ScalarKind::Sint | crate::ScalarKind::Uint => {
1780 crate::Literal::one(scalar)
1781 .ok_or(Error::BadIncrDecrReferenceType(value_span))?
1782 }
1783 _ => return Err(Error::BadIncrDecrReferenceType(value_span)),
1784 };
1785
1786 let right =
1787 ectx.interrupt_emitter(crate::Expression::Literal(literal), Span::UNDEFINED)?;
1788 let rctx = ectx.runtime_expression_ctx(stmt.span)?;
1789 let left = rctx.function.expressions.append(
1790 crate::Expression::Load {
1791 pointer: target_handle,
1792 },
1793 value_span,
1794 );
1795 let value = rctx
1796 .function
1797 .expressions
1798 .append(crate::Expression::Binary { op, left, right }, stmt.span);
1799 rctx.local_expression_kind_tracker
1800 .insert(left, crate::proc::ExpressionKind::Runtime);
1801 rctx.local_expression_kind_tracker
1802 .insert(value, crate::proc::ExpressionKind::Runtime);
1803
1804 block.extend(emitter.finish(&ctx.function.expressions));
1805 crate::Statement::Store {
1806 pointer: target_handle,
1807 value,
1808 }
1809 }
1810 ast::StatementKind::ConstAssert(condition) => {
1811 let mut emitter = Emitter::default();
1812 emitter.start(&ctx.function.expressions);
1813
1814 let condition =
1815 self.expression(condition, &mut ctx.as_const(block, &mut emitter))?;
1816
1817 let span = ctx.function.expressions.get_span(condition);
1818 match ctx
1819 .module
1820 .to_ctx()
1821 .eval_expr_to_bool_from(condition, &ctx.function.expressions)
1822 {
1823 Some(true) => Ok(()),
1824 Some(false) => Err(Error::ConstAssertFailed(span)),
1825 _ => Err(Error::NotBool(span)),
1826 }?;
1827
1828 block.extend(emitter.finish(&ctx.function.expressions));
1829
1830 return Ok(());
1831 }
1832 ast::StatementKind::Phony(expr) => {
1833 let mut emitter = Emitter::default();
1834 emitter.start(&ctx.function.expressions);
1835
1836 let value = self.expression(expr, &mut ctx.as_expression(block, &mut emitter))?;
1837 block.extend(emitter.finish(&ctx.function.expressions));
1838 ctx.named_expressions
1839 .insert(value, ("phony".to_string(), stmt.span));
1840 return Ok(());
1841 }
1842 };
1843
1844 block.push(out, stmt.span);
1845
1846 Ok(())
1847 }
1848
1849 fn expression(
1857 &mut self,
1858 expr: Handle<ast::Expression<'source>>,
1859 ctx: &mut ExpressionContext<'source, '_, '_>,
1860 ) -> Result<Handle<crate::Expression>, Error<'source>> {
1861 let expr = self.expression_for_abstract(expr, ctx)?;
1862 ctx.concretize(expr)
1863 }
1864
1865 fn expression_for_abstract(
1866 &mut self,
1867 expr: Handle<ast::Expression<'source>>,
1868 ctx: &mut ExpressionContext<'source, '_, '_>,
1869 ) -> Result<Handle<crate::Expression>, Error<'source>> {
1870 let expr = self.expression_for_reference(expr, ctx)?;
1871 ctx.apply_load_rule(expr)
1872 }
1873
1874 fn expression_for_reference(
1875 &mut self,
1876 expr: Handle<ast::Expression<'source>>,
1877 ctx: &mut ExpressionContext<'source, '_, '_>,
1878 ) -> Result<Typed<Handle<crate::Expression>>, Error<'source>> {
1879 let span = ctx.ast_expressions.get_span(expr);
1880 let expr = &ctx.ast_expressions[expr];
1881
1882 let expr: Typed<crate::Expression> = match *expr {
1883 ast::Expression::Literal(literal) => {
1884 let literal = match literal {
1885 ast::Literal::Number(Number::F32(f)) => crate::Literal::F32(f),
1886 ast::Literal::Number(Number::I32(i)) => crate::Literal::I32(i),
1887 ast::Literal::Number(Number::U32(u)) => crate::Literal::U32(u),
1888 ast::Literal::Number(Number::I64(i)) => crate::Literal::I64(i),
1889 ast::Literal::Number(Number::U64(u)) => crate::Literal::U64(u),
1890 ast::Literal::Number(Number::F64(f)) => crate::Literal::F64(f),
1891 ast::Literal::Number(Number::AbstractInt(i)) => crate::Literal::AbstractInt(i),
1892 ast::Literal::Number(Number::AbstractFloat(f)) => {
1893 crate::Literal::AbstractFloat(f)
1894 }
1895 ast::Literal::Bool(b) => crate::Literal::Bool(b),
1896 };
1897 let handle = ctx.interrupt_emitter(crate::Expression::Literal(literal), span)?;
1898 return Ok(Typed::Plain(handle));
1899 }
1900 ast::Expression::Ident(ast::IdentExpr::Local(local)) => {
1901 return ctx.local(&local, span);
1902 }
1903 ast::Expression::Ident(ast::IdentExpr::Unresolved(name)) => {
1904 let global = ctx
1905 .globals
1906 .get(name)
1907 .ok_or(Error::UnknownIdent(span, name))?;
1908 let expr = match *global {
1909 LoweredGlobalDecl::Var(handle) => {
1910 let expr = crate::Expression::GlobalVariable(handle);
1911 match ctx.module.global_variables[handle].space {
1912 crate::AddressSpace::Handle => Typed::Plain(expr),
1913 _ => Typed::Reference(expr),
1914 }
1915 }
1916 LoweredGlobalDecl::Const(handle) => {
1917 Typed::Plain(crate::Expression::Constant(handle))
1918 }
1919 LoweredGlobalDecl::Override(handle) => {
1920 Typed::Plain(crate::Expression::Override(handle))
1921 }
1922 LoweredGlobalDecl::Function(_)
1923 | LoweredGlobalDecl::Type(_)
1924 | LoweredGlobalDecl::EntryPoint => {
1925 return Err(Error::Unexpected(span, ExpectedToken::Variable));
1926 }
1927 };
1928
1929 return expr.try_map(|handle| ctx.interrupt_emitter(handle, span));
1930 }
1931 ast::Expression::Construct {
1932 ref ty,
1933 ty_span,
1934 ref components,
1935 } => {
1936 let handle = self.construct(span, ty, ty_span, components, ctx)?;
1937 return Ok(Typed::Plain(handle));
1938 }
1939 ast::Expression::Unary { op, expr } => {
1940 let expr = self.expression_for_abstract(expr, ctx)?;
1941 Typed::Plain(crate::Expression::Unary { op, expr })
1942 }
1943 ast::Expression::AddrOf(expr) => {
1944 match self.expression_for_reference(expr, ctx)? {
1947 Typed::Reference(handle) => {
1948 return Ok(Typed::Plain(handle));
1950 }
1951 Typed::Plain(_) => {
1952 return Err(Error::NotReference("the operand of the `&` operator", span));
1953 }
1954 }
1955 }
1956 ast::Expression::Deref(expr) => {
1957 let pointer = self.expression(expr, ctx)?;
1959
1960 if resolve_inner!(ctx, pointer).pointer_space().is_none() {
1961 return Err(Error::NotPointer(span));
1962 }
1963
1964 return Ok(Typed::Reference(pointer));
1966 }
1967 ast::Expression::Binary { op, left, right } => {
1968 self.binary(op, left, right, span, ctx)?
1969 }
1970 ast::Expression::Call {
1971 ref function,
1972 ref arguments,
1973 } => {
1974 let handle = self
1975 .call(span, function, arguments, ctx, false)?
1976 .ok_or(Error::FunctionReturnsVoid(function.span))?;
1977 return Ok(Typed::Plain(handle));
1978 }
1979 ast::Expression::Index { base, index } => {
1980 let lowered_base = self.expression_for_reference(base, ctx)?;
1981 let index = self.expression(index, ctx)?;
1982
1983 if let Typed::Plain(handle) = lowered_base {
1984 if resolve_inner!(ctx, handle).pointer_space().is_some() {
1985 return Err(Error::Pointer(
1986 "the value indexed by a `[]` subscripting expression",
1987 ctx.ast_expressions.get_span(base),
1988 ));
1989 }
1990 }
1991
1992 lowered_base.map(|base| match ctx.const_access(index) {
1993 Some(index) => crate::Expression::AccessIndex { base, index },
1994 None => crate::Expression::Access { base, index },
1995 })
1996 }
1997 ast::Expression::Member { base, ref field } => {
1998 let lowered_base = self.expression_for_reference(base, ctx)?;
1999
2000 let temp_inner;
2001 let composite_type: &crate::TypeInner = match lowered_base {
2002 Typed::Reference(handle) => {
2003 let inner = resolve_inner!(ctx, handle);
2004 match *inner {
2005 crate::TypeInner::Pointer { base, .. } => &ctx.module.types[base].inner,
2006 crate::TypeInner::ValuePointer {
2007 size: None, scalar, ..
2008 } => {
2009 temp_inner = crate::TypeInner::Scalar(scalar);
2010 &temp_inner
2011 }
2012 crate::TypeInner::ValuePointer {
2013 size: Some(size),
2014 scalar,
2015 ..
2016 } => {
2017 temp_inner = crate::TypeInner::Vector { size, scalar };
2018 &temp_inner
2019 }
2020 _ => unreachable!(
2021 "In Typed::Reference(handle), handle must be a Naga pointer"
2022 ),
2023 }
2024 }
2025
2026 Typed::Plain(handle) => {
2027 let inner = resolve_inner!(ctx, handle);
2028 if let crate::TypeInner::Pointer { .. }
2029 | crate::TypeInner::ValuePointer { .. } = *inner
2030 {
2031 return Err(Error::Pointer(
2032 "the value accessed by a `.member` expression",
2033 ctx.ast_expressions.get_span(base),
2034 ));
2035 }
2036 inner
2037 }
2038 };
2039
2040 let access = match *composite_type {
2041 crate::TypeInner::Struct { ref members, .. } => {
2042 let index = members
2043 .iter()
2044 .position(|m| m.name.as_deref() == Some(field.name))
2045 .ok_or(Error::BadAccessor(field.span))?
2046 as u32;
2047
2048 lowered_base.map(|base| crate::Expression::AccessIndex { base, index })
2049 }
2050 crate::TypeInner::Vector { .. } => {
2051 match Components::new(field.name, field.span)? {
2052 Components::Swizzle { size, pattern } => {
2053 Typed::Plain(crate::Expression::Swizzle {
2054 size,
2055 vector: ctx.apply_load_rule(lowered_base)?,
2056 pattern,
2057 })
2058 }
2059 Components::Single(index) => lowered_base
2060 .map(|base| crate::Expression::AccessIndex { base, index }),
2061 }
2062 }
2063 _ => return Err(Error::BadAccessor(field.span)),
2064 };
2065
2066 access
2067 }
2068 ast::Expression::Bitcast { expr, to, ty_span } => {
2069 let expr = self.expression(expr, ctx)?;
2070 let to_resolved = self.resolve_ast_type(to, &mut ctx.as_global())?;
2071
2072 let element_scalar = match ctx.module.types[to_resolved].inner {
2073 crate::TypeInner::Scalar(scalar) => scalar,
2074 crate::TypeInner::Vector { scalar, .. } => scalar,
2075 _ => {
2076 let ty = resolve!(ctx, expr);
2077 let gctx = &ctx.module.to_ctx();
2078 return Err(Error::BadTypeCast {
2079 from_type: ty.to_wgsl(gctx).into(),
2080 span: ty_span,
2081 to_type: to_resolved.to_wgsl(gctx).into(),
2082 });
2083 }
2084 };
2085
2086 Typed::Plain(crate::Expression::As {
2087 expr,
2088 kind: element_scalar.kind,
2089 convert: None,
2090 })
2091 }
2092 };
2093
2094 expr.try_map(|handle| ctx.append_expression(handle, span))
2095 }
2096
2097 fn binary(
2098 &mut self,
2099 op: crate::BinaryOperator,
2100 left: Handle<ast::Expression<'source>>,
2101 right: Handle<ast::Expression<'source>>,
2102 span: Span,
2103 ctx: &mut ExpressionContext<'source, '_, '_>,
2104 ) -> Result<Typed<crate::Expression>, Error<'source>> {
2105 let mut left = self.expression_for_abstract(left, ctx)?;
2107 let mut right = self.expression_for_abstract(right, ctx)?;
2108
2109 ctx.binary_op_splat(op, &mut left, &mut right)?;
2112
2113 match op {
2115 crate::BinaryOperator::ShiftLeft | crate::BinaryOperator::ShiftRight => {
2120 right =
2121 ctx.try_automatic_conversion_for_leaf_scalar(right, crate::Scalar::U32, span)?;
2122 }
2123
2124 _ => {
2129 ctx.grow_types(left)?;
2130 ctx.grow_types(right)?;
2131 if let Ok(consensus_scalar) =
2132 ctx.automatic_conversion_consensus([left, right].iter())
2133 {
2134 ctx.convert_to_leaf_scalar(&mut left, consensus_scalar)?;
2135 ctx.convert_to_leaf_scalar(&mut right, consensus_scalar)?;
2136 }
2137 }
2138 }
2139
2140 Ok(Typed::Plain(crate::Expression::Binary { op, left, right }))
2141 }
2142
2143 fn call(
2162 &mut self,
2163 span: Span,
2164 function: &ast::Ident<'source>,
2165 arguments: &[Handle<ast::Expression<'source>>],
2166 ctx: &mut ExpressionContext<'source, '_, '_>,
2167 is_statement: bool,
2168 ) -> Result<Option<Handle<crate::Expression>>, Error<'source>> {
2169 match ctx.globals.get(function.name) {
2170 Some(&LoweredGlobalDecl::Type(ty)) => {
2171 let handle = self.construct(
2172 span,
2173 &ast::ConstructorType::Type(ty),
2174 function.span,
2175 arguments,
2176 ctx,
2177 )?;
2178 Ok(Some(handle))
2179 }
2180 Some(
2181 &LoweredGlobalDecl::Const(_)
2182 | &LoweredGlobalDecl::Override(_)
2183 | &LoweredGlobalDecl::Var(_),
2184 ) => Err(Error::Unexpected(function.span, ExpectedToken::Function)),
2185 Some(&LoweredGlobalDecl::EntryPoint) => Err(Error::CalledEntryPoint(function.span)),
2186 Some(&LoweredGlobalDecl::Function(function)) => {
2187 let arguments = arguments
2188 .iter()
2189 .enumerate()
2190 .map(|(i, &arg)| {
2191 let Some(&crate::FunctionArgument {
2193 ty: parameter_ty, ..
2194 }) = ctx.module.functions[function].arguments.get(i)
2195 else {
2196 return self.expression(arg, ctx);
2199 };
2200
2201 let expr = self.expression_for_abstract(arg, ctx)?;
2202 ctx.try_automatic_conversions(
2203 expr,
2204 &crate::proc::TypeResolution::Handle(parameter_ty),
2205 ctx.ast_expressions.get_span(arg),
2206 )
2207 })
2208 .collect::<Result<Vec<_>, _>>()?;
2209
2210 let has_result = ctx.module.functions[function].result.is_some();
2211 let rctx = ctx.runtime_expression_ctx(span)?;
2212 rctx.block
2214 .extend(rctx.emitter.finish(&rctx.function.expressions));
2215 let result = has_result.then(|| {
2216 let result = rctx
2217 .function
2218 .expressions
2219 .append(crate::Expression::CallResult(function), span);
2220 rctx.local_expression_kind_tracker
2221 .insert(result, crate::proc::ExpressionKind::Runtime);
2222 result
2223 });
2224 rctx.emitter.start(&rctx.function.expressions);
2225 rctx.block.push(
2226 crate::Statement::Call {
2227 function,
2228 arguments,
2229 result,
2230 },
2231 span,
2232 );
2233
2234 Ok(result)
2235 }
2236 None => {
2237 let span = function.span;
2238 let expr = if let Some(fun) = conv::map_relational_fun(function.name) {
2239 let mut args = ctx.prepare_args(arguments, 1, span);
2240 let argument = self.expression(args.next()?, ctx)?;
2241 args.finish()?;
2242
2243 let argument_unmodified = matches!(
2245 fun,
2246 crate::RelationalFunction::All | crate::RelationalFunction::Any
2247 ) && {
2248 matches!(
2249 resolve_inner!(ctx, argument),
2250 &crate::TypeInner::Scalar(crate::Scalar {
2251 kind: crate::ScalarKind::Bool,
2252 ..
2253 })
2254 )
2255 };
2256
2257 if argument_unmodified {
2258 return Ok(Some(argument));
2259 } else {
2260 crate::Expression::Relational { fun, argument }
2261 }
2262 } else if let Some((axis, ctrl)) = conv::map_derivative(function.name) {
2263 let mut args = ctx.prepare_args(arguments, 1, span);
2264 let expr = self.expression(args.next()?, ctx)?;
2265 args.finish()?;
2266
2267 crate::Expression::Derivative { axis, ctrl, expr }
2268 } else if let Some(fun) = conv::map_standard_fun(function.name) {
2269 let expected = fun.argument_count() as _;
2270 let mut args = ctx.prepare_args(arguments, expected, span);
2271
2272 let arg = self.expression(args.next()?, ctx)?;
2273 let arg1 = args
2274 .next()
2275 .map(|x| self.expression(x, ctx))
2276 .ok()
2277 .transpose()?;
2278 let arg2 = args
2279 .next()
2280 .map(|x| self.expression(x, ctx))
2281 .ok()
2282 .transpose()?;
2283 let arg3 = args
2284 .next()
2285 .map(|x| self.expression(x, ctx))
2286 .ok()
2287 .transpose()?;
2288
2289 args.finish()?;
2290
2291 if fun == crate::MathFunction::Modf || fun == crate::MathFunction::Frexp {
2292 if let Some((size, scalar)) = match *resolve_inner!(ctx, arg) {
2293 crate::TypeInner::Scalar(scalar) => Some((None, scalar)),
2294 crate::TypeInner::Vector { size, scalar, .. } => {
2295 Some((Some(size), scalar))
2296 }
2297 _ => None,
2298 } {
2299 ctx.module.generate_predeclared_type(
2300 if fun == crate::MathFunction::Modf {
2301 crate::PredeclaredType::ModfResult { size, scalar }
2302 } else {
2303 crate::PredeclaredType::FrexpResult { size, scalar }
2304 },
2305 );
2306 }
2307 }
2308
2309 crate::Expression::Math {
2310 fun,
2311 arg,
2312 arg1,
2313 arg2,
2314 arg3,
2315 }
2316 } else if let Some(fun) = Texture::map(function.name) {
2317 self.texture_sample_helper(fun, arguments, span, ctx)?
2318 } else if let Some((op, cop)) = conv::map_subgroup_operation(function.name) {
2319 return Ok(Some(
2320 self.subgroup_operation_helper(span, op, cop, arguments, ctx)?,
2321 ));
2322 } else if let Some(mode) = SubgroupGather::map(function.name) {
2323 return Ok(Some(
2324 self.subgroup_gather_helper(span, mode, arguments, ctx)?,
2325 ));
2326 } else if let Some(fun) = crate::AtomicFunction::map(function.name) {
2327 return self.atomic_helper(span, fun, arguments, is_statement, ctx);
2328 } else {
2329 match function.name {
2330 "select" => {
2331 let mut args = ctx.prepare_args(arguments, 3, span);
2332
2333 let reject = self.expression(args.next()?, ctx)?;
2334 let accept = self.expression(args.next()?, ctx)?;
2335 let condition = self.expression(args.next()?, ctx)?;
2336
2337 args.finish()?;
2338
2339 crate::Expression::Select {
2340 reject,
2341 accept,
2342 condition,
2343 }
2344 }
2345 "arrayLength" => {
2346 let mut args = ctx.prepare_args(arguments, 1, span);
2347 let expr = self.expression(args.next()?, ctx)?;
2348 args.finish()?;
2349
2350 crate::Expression::ArrayLength(expr)
2351 }
2352 "atomicLoad" => {
2353 let mut args = ctx.prepare_args(arguments, 1, span);
2354 let pointer = self.atomic_pointer(args.next()?, ctx)?;
2355 args.finish()?;
2356
2357 crate::Expression::Load { pointer }
2358 }
2359 "atomicStore" => {
2360 let mut args = ctx.prepare_args(arguments, 2, span);
2361 let pointer = self.atomic_pointer(args.next()?, ctx)?;
2362 let value = self.expression(args.next()?, ctx)?;
2363 args.finish()?;
2364
2365 let rctx = ctx.runtime_expression_ctx(span)?;
2366 rctx.block
2367 .extend(rctx.emitter.finish(&rctx.function.expressions));
2368 rctx.emitter.start(&rctx.function.expressions);
2369 rctx.block
2370 .push(crate::Statement::Store { pointer, value }, span);
2371 return Ok(None);
2372 }
2373 "atomicCompareExchangeWeak" => {
2374 let mut args = ctx.prepare_args(arguments, 3, span);
2375
2376 let pointer = self.atomic_pointer(args.next()?, ctx)?;
2377
2378 let compare = self.expression(args.next()?, ctx)?;
2379
2380 let value = args.next()?;
2381 let value_span = ctx.ast_expressions.get_span(value);
2382 let value = self.expression(value, ctx)?;
2383
2384 args.finish()?;
2385
2386 let expression = match *resolve_inner!(ctx, value) {
2387 crate::TypeInner::Scalar(scalar) => {
2388 crate::Expression::AtomicResult {
2389 ty: ctx.module.generate_predeclared_type(
2390 crate::PredeclaredType::AtomicCompareExchangeWeakResult(
2391 scalar,
2392 ),
2393 ),
2394 comparison: true,
2395 }
2396 }
2397 _ => return Err(Error::InvalidAtomicOperandType(value_span)),
2398 };
2399
2400 let result = ctx.interrupt_emitter(expression, span)?;
2401 let rctx = ctx.runtime_expression_ctx(span)?;
2402 rctx.block.push(
2403 crate::Statement::Atomic {
2404 pointer,
2405 fun: crate::AtomicFunction::Exchange {
2406 compare: Some(compare),
2407 },
2408 value,
2409 result: Some(result),
2410 },
2411 span,
2412 );
2413 return Ok(Some(result));
2414 }
2415 "textureAtomicMin" | "textureAtomicMax" | "textureAtomicAdd"
2416 | "textureAtomicAnd" | "textureAtomicOr" | "textureAtomicXor" => {
2417 let mut args = ctx.prepare_args(arguments, 3, span);
2418
2419 let image = args.next()?;
2420 let image_span = ctx.ast_expressions.get_span(image);
2421 let image = self.expression(image, ctx)?;
2422
2423 let coordinate = self.expression(args.next()?, ctx)?;
2424
2425 let (_, arrayed) = ctx.image_data(image, image_span)?;
2426 let array_index = arrayed
2427 .then(|| {
2428 args.min_args += 1;
2429 self.expression(args.next()?, ctx)
2430 })
2431 .transpose()?;
2432
2433 let value = self.expression(args.next()?, ctx)?;
2434
2435 args.finish()?;
2436
2437 let rctx = ctx.runtime_expression_ctx(span)?;
2438 rctx.block
2439 .extend(rctx.emitter.finish(&rctx.function.expressions));
2440 rctx.emitter.start(&rctx.function.expressions);
2441 let stmt = crate::Statement::ImageAtomic {
2442 image,
2443 coordinate,
2444 array_index,
2445 fun: match function.name {
2446 "textureAtomicMin" => crate::AtomicFunction::Min,
2447 "textureAtomicMax" => crate::AtomicFunction::Max,
2448 "textureAtomicAdd" => crate::AtomicFunction::Add,
2449 "textureAtomicAnd" => crate::AtomicFunction::And,
2450 "textureAtomicOr" => crate::AtomicFunction::InclusiveOr,
2451 "textureAtomicXor" => crate::AtomicFunction::ExclusiveOr,
2452 _ => unreachable!(),
2453 },
2454 value,
2455 };
2456 rctx.block.push(stmt, span);
2457 return Ok(None);
2458 }
2459 "storageBarrier" => {
2460 ctx.prepare_args(arguments, 0, span).finish()?;
2461
2462 let rctx = ctx.runtime_expression_ctx(span)?;
2463 rctx.block
2464 .push(crate::Statement::Barrier(crate::Barrier::STORAGE), span);
2465 return Ok(None);
2466 }
2467 "workgroupBarrier" => {
2468 ctx.prepare_args(arguments, 0, span).finish()?;
2469
2470 let rctx = ctx.runtime_expression_ctx(span)?;
2471 rctx.block
2472 .push(crate::Statement::Barrier(crate::Barrier::WORK_GROUP), span);
2473 return Ok(None);
2474 }
2475 "subgroupBarrier" => {
2476 ctx.prepare_args(arguments, 0, span).finish()?;
2477
2478 let rctx = ctx.runtime_expression_ctx(span)?;
2479 rctx.block
2480 .push(crate::Statement::Barrier(crate::Barrier::SUB_GROUP), span);
2481 return Ok(None);
2482 }
2483 "workgroupUniformLoad" => {
2484 let mut args = ctx.prepare_args(arguments, 1, span);
2485 let expr = args.next()?;
2486 args.finish()?;
2487
2488 let pointer = self.expression(expr, ctx)?;
2489 let result_ty = match *resolve_inner!(ctx, pointer) {
2490 crate::TypeInner::Pointer {
2491 base,
2492 space: crate::AddressSpace::WorkGroup,
2493 } => base,
2494 ref other => {
2495 log::error!("Type {other:?} passed to workgroupUniformLoad");
2496 let span = ctx.ast_expressions.get_span(expr);
2497 return Err(Error::InvalidWorkGroupUniformLoad(span));
2498 }
2499 };
2500 let result = ctx.interrupt_emitter(
2501 crate::Expression::WorkGroupUniformLoadResult { ty: result_ty },
2502 span,
2503 )?;
2504 let rctx = ctx.runtime_expression_ctx(span)?;
2505 rctx.block.push(
2506 crate::Statement::WorkGroupUniformLoad { pointer, result },
2507 span,
2508 );
2509
2510 return Ok(Some(result));
2511 }
2512 "textureStore" => {
2513 let mut args = ctx.prepare_args(arguments, 3, span);
2514
2515 let image = args.next()?;
2516 let image_span = ctx.ast_expressions.get_span(image);
2517 let image = self.expression(image, ctx)?;
2518
2519 let coordinate = self.expression(args.next()?, ctx)?;
2520
2521 let (_, arrayed) = ctx.image_data(image, image_span)?;
2522 let array_index = arrayed
2523 .then(|| {
2524 args.min_args += 1;
2525 self.expression(args.next()?, ctx)
2526 })
2527 .transpose()?;
2528
2529 let value = self.expression(args.next()?, ctx)?;
2530
2531 args.finish()?;
2532
2533 let rctx = ctx.runtime_expression_ctx(span)?;
2534 rctx.block
2535 .extend(rctx.emitter.finish(&rctx.function.expressions));
2536 rctx.emitter.start(&rctx.function.expressions);
2537 let stmt = crate::Statement::ImageStore {
2538 image,
2539 coordinate,
2540 array_index,
2541 value,
2542 };
2543 rctx.block.push(stmt, span);
2544 return Ok(None);
2545 }
2546 "textureLoad" => {
2547 let mut args = ctx.prepare_args(arguments, 2, span);
2548
2549 let image = args.next()?;
2550 let image_span = ctx.ast_expressions.get_span(image);
2551 let image = self.expression(image, ctx)?;
2552
2553 let coordinate = self.expression(args.next()?, ctx)?;
2554
2555 let (class, arrayed) = ctx.image_data(image, image_span)?;
2556 let array_index = arrayed
2557 .then(|| {
2558 args.min_args += 1;
2559 self.expression(args.next()?, ctx)
2560 })
2561 .transpose()?;
2562
2563 let level = class
2564 .is_mipmapped()
2565 .then(|| {
2566 args.min_args += 1;
2567 self.expression(args.next()?, ctx)
2568 })
2569 .transpose()?;
2570
2571 let sample = class
2572 .is_multisampled()
2573 .then(|| self.expression(args.next()?, ctx))
2574 .transpose()?;
2575
2576 args.finish()?;
2577
2578 crate::Expression::ImageLoad {
2579 image,
2580 coordinate,
2581 array_index,
2582 level,
2583 sample,
2584 }
2585 }
2586 "textureDimensions" => {
2587 let mut args = ctx.prepare_args(arguments, 1, span);
2588 let image = self.expression(args.next()?, ctx)?;
2589 let level = args
2590 .next()
2591 .map(|arg| self.expression(arg, ctx))
2592 .ok()
2593 .transpose()?;
2594 args.finish()?;
2595
2596 crate::Expression::ImageQuery {
2597 image,
2598 query: crate::ImageQuery::Size { level },
2599 }
2600 }
2601 "textureNumLevels" => {
2602 let mut args = ctx.prepare_args(arguments, 1, span);
2603 let image = self.expression(args.next()?, ctx)?;
2604 args.finish()?;
2605
2606 crate::Expression::ImageQuery {
2607 image,
2608 query: crate::ImageQuery::NumLevels,
2609 }
2610 }
2611 "textureNumLayers" => {
2612 let mut args = ctx.prepare_args(arguments, 1, span);
2613 let image = self.expression(args.next()?, ctx)?;
2614 args.finish()?;
2615
2616 crate::Expression::ImageQuery {
2617 image,
2618 query: crate::ImageQuery::NumLayers,
2619 }
2620 }
2621 "textureNumSamples" => {
2622 let mut args = ctx.prepare_args(arguments, 1, span);
2623 let image = self.expression(args.next()?, ctx)?;
2624 args.finish()?;
2625
2626 crate::Expression::ImageQuery {
2627 image,
2628 query: crate::ImageQuery::NumSamples,
2629 }
2630 }
2631 "rayQueryInitialize" => {
2632 let mut args = ctx.prepare_args(arguments, 3, span);
2633 let query = self.ray_query_pointer(args.next()?, ctx)?;
2634 let acceleration_structure = self.expression(args.next()?, ctx)?;
2635 let descriptor = self.expression(args.next()?, ctx)?;
2636 args.finish()?;
2637
2638 let _ = ctx.module.generate_ray_desc_type();
2639 let fun = crate::RayQueryFunction::Initialize {
2640 acceleration_structure,
2641 descriptor,
2642 };
2643
2644 let rctx = ctx.runtime_expression_ctx(span)?;
2645 rctx.block
2646 .extend(rctx.emitter.finish(&rctx.function.expressions));
2647 rctx.emitter.start(&rctx.function.expressions);
2648 rctx.block
2649 .push(crate::Statement::RayQuery { query, fun }, span);
2650 return Ok(None);
2651 }
2652 "rayQueryProceed" => {
2653 let mut args = ctx.prepare_args(arguments, 1, span);
2654 let query = self.ray_query_pointer(args.next()?, ctx)?;
2655 args.finish()?;
2656
2657 let result = ctx.interrupt_emitter(
2658 crate::Expression::RayQueryProceedResult,
2659 span,
2660 )?;
2661 let fun = crate::RayQueryFunction::Proceed { result };
2662 let rctx = ctx.runtime_expression_ctx(span)?;
2663 rctx.block
2664 .push(crate::Statement::RayQuery { query, fun }, span);
2665 return Ok(Some(result));
2666 }
2667 "rayQueryGetCommittedIntersection" => {
2668 let mut args = ctx.prepare_args(arguments, 1, span);
2669 let query = self.ray_query_pointer(args.next()?, ctx)?;
2670 args.finish()?;
2671
2672 let _ = ctx.module.generate_ray_intersection_type();
2673 crate::Expression::RayQueryGetIntersection {
2674 query,
2675 committed: true,
2676 }
2677 }
2678 "rayQueryGetCandidateIntersection" => {
2679 let mut args = ctx.prepare_args(arguments, 1, span);
2680 let query = self.ray_query_pointer(args.next()?, ctx)?;
2681 args.finish()?;
2682
2683 let _ = ctx.module.generate_ray_intersection_type();
2684 crate::Expression::RayQueryGetIntersection {
2685 query,
2686 committed: false,
2687 }
2688 }
2689 "RayDesc" => {
2690 let ty = ctx.module.generate_ray_desc_type();
2691 let handle = self.construct(
2692 span,
2693 &ast::ConstructorType::Type(ty),
2694 function.span,
2695 arguments,
2696 ctx,
2697 )?;
2698 return Ok(Some(handle));
2699 }
2700 "subgroupBallot" => {
2701 let mut args = ctx.prepare_args(arguments, 0, span);
2702 let predicate = if arguments.len() == 1 {
2703 Some(self.expression(args.next()?, ctx)?)
2704 } else {
2705 None
2706 };
2707 args.finish()?;
2708
2709 let result = ctx
2710 .interrupt_emitter(crate::Expression::SubgroupBallotResult, span)?;
2711 let rctx = ctx.runtime_expression_ctx(span)?;
2712 rctx.block
2713 .push(crate::Statement::SubgroupBallot { result, predicate }, span);
2714 return Ok(Some(result));
2715 }
2716 _ => return Err(Error::UnknownIdent(function.span, function.name)),
2717 }
2718 };
2719
2720 let expr = ctx.append_expression(expr, span)?;
2721 Ok(Some(expr))
2722 }
2723 }
2724 }
2725
2726 fn atomic_pointer(
2727 &mut self,
2728 expr: Handle<ast::Expression<'source>>,
2729 ctx: &mut ExpressionContext<'source, '_, '_>,
2730 ) -> Result<Handle<crate::Expression>, Error<'source>> {
2731 let span = ctx.ast_expressions.get_span(expr);
2732 let pointer = self.expression(expr, ctx)?;
2733
2734 match *resolve_inner!(ctx, pointer) {
2735 crate::TypeInner::Pointer { base, .. } => match ctx.module.types[base].inner {
2736 crate::TypeInner::Atomic { .. } => Ok(pointer),
2737 ref other => {
2738 log::error!("Pointer type to {:?} passed to atomic op", other);
2739 Err(Error::InvalidAtomicPointer(span))
2740 }
2741 },
2742 ref other => {
2743 log::error!("Type {:?} passed to atomic op", other);
2744 Err(Error::InvalidAtomicPointer(span))
2745 }
2746 }
2747 }
2748
2749 fn atomic_helper(
2750 &mut self,
2751 span: Span,
2752 fun: crate::AtomicFunction,
2753 args: &[Handle<ast::Expression<'source>>],
2754 is_statement: bool,
2755 ctx: &mut ExpressionContext<'source, '_, '_>,
2756 ) -> Result<Option<Handle<crate::Expression>>, Error<'source>> {
2757 let mut args = ctx.prepare_args(args, 2, span);
2758
2759 let pointer = self.atomic_pointer(args.next()?, ctx)?;
2760 let value = self.expression(args.next()?, ctx)?;
2761 let value_inner = resolve_inner!(ctx, value);
2762 args.finish()?;
2763
2764 let is_64_bit_min_max =
2769 matches!(fun, crate::AtomicFunction::Min | crate::AtomicFunction::Max)
2770 && matches!(
2771 *value_inner,
2772 crate::TypeInner::Scalar(crate::Scalar { width: 8, .. })
2773 );
2774 let result = if is_64_bit_min_max && is_statement {
2775 let rctx = ctx.runtime_expression_ctx(span)?;
2776 rctx.block
2777 .extend(rctx.emitter.finish(&rctx.function.expressions));
2778 rctx.emitter.start(&rctx.function.expressions);
2779 None
2780 } else {
2781 let ty = ctx.register_type(value)?;
2782 Some(ctx.interrupt_emitter(
2783 crate::Expression::AtomicResult {
2784 ty,
2785 comparison: false,
2786 },
2787 span,
2788 )?)
2789 };
2790 let rctx = ctx.runtime_expression_ctx(span)?;
2791 rctx.block.push(
2792 crate::Statement::Atomic {
2793 pointer,
2794 fun,
2795 value,
2796 result,
2797 },
2798 span,
2799 );
2800 Ok(result)
2801 }
2802
2803 fn texture_sample_helper(
2804 &mut self,
2805 fun: Texture,
2806 args: &[Handle<ast::Expression<'source>>],
2807 span: Span,
2808 ctx: &mut ExpressionContext<'source, '_, '_>,
2809 ) -> Result<crate::Expression, Error<'source>> {
2810 let mut args = ctx.prepare_args(args, fun.min_argument_count(), span);
2811
2812 fn get_image_and_span<'source>(
2813 lowerer: &mut Lowerer<'source, '_>,
2814 args: &mut ArgumentContext<'_, 'source>,
2815 ctx: &mut ExpressionContext<'source, '_, '_>,
2816 ) -> Result<(Handle<crate::Expression>, Span), Error<'source>> {
2817 let image = args.next()?;
2818 let image_span = ctx.ast_expressions.get_span(image);
2819 let image = lowerer.expression(image, ctx)?;
2820 Ok((image, image_span))
2821 }
2822
2823 let (image, image_span, gather) = match fun {
2824 Texture::Gather => {
2825 let image_or_component = args.next()?;
2826 let image_or_component_span = ctx.ast_expressions.get_span(image_or_component);
2827 let lowered_image_or_component = self.expression(image_or_component, ctx)?;
2829
2830 match *resolve_inner!(ctx, lowered_image_or_component) {
2831 crate::TypeInner::Image {
2832 class: crate::ImageClass::Depth { .. },
2833 ..
2834 } => (
2835 lowered_image_or_component,
2836 image_or_component_span,
2837 Some(crate::SwizzleComponent::X),
2838 ),
2839 _ => {
2840 let (image, image_span) = get_image_and_span(self, &mut args, ctx)?;
2841 (
2842 image,
2843 image_span,
2844 Some(ctx.gather_component(
2845 lowered_image_or_component,
2846 image_or_component_span,
2847 span,
2848 )?),
2849 )
2850 }
2851 }
2852 }
2853 Texture::GatherCompare => {
2854 let (image, image_span) = get_image_and_span(self, &mut args, ctx)?;
2855 (image, image_span, Some(crate::SwizzleComponent::X))
2856 }
2857
2858 _ => {
2859 let (image, image_span) = get_image_and_span(self, &mut args, ctx)?;
2860 (image, image_span, None)
2861 }
2862 };
2863
2864 let sampler = self.expression(args.next()?, ctx)?;
2865
2866 let coordinate = self.expression(args.next()?, ctx)?;
2867
2868 let (_, arrayed) = ctx.image_data(image, image_span)?;
2869 let array_index = arrayed
2870 .then(|| self.expression(args.next()?, ctx))
2871 .transpose()?;
2872
2873 let (level, depth_ref) = match fun {
2874 Texture::Gather => (crate::SampleLevel::Zero, None),
2875 Texture::GatherCompare => {
2876 let reference = self.expression(args.next()?, ctx)?;
2877 (crate::SampleLevel::Zero, Some(reference))
2878 }
2879
2880 Texture::Sample => (crate::SampleLevel::Auto, None),
2881 Texture::SampleBias => {
2882 let bias = self.expression(args.next()?, ctx)?;
2883 (crate::SampleLevel::Bias(bias), None)
2884 }
2885 Texture::SampleCompare => {
2886 let reference = self.expression(args.next()?, ctx)?;
2887 (crate::SampleLevel::Auto, Some(reference))
2888 }
2889 Texture::SampleCompareLevel => {
2890 let reference = self.expression(args.next()?, ctx)?;
2891 (crate::SampleLevel::Zero, Some(reference))
2892 }
2893 Texture::SampleGrad => {
2894 let x = self.expression(args.next()?, ctx)?;
2895 let y = self.expression(args.next()?, ctx)?;
2896 (crate::SampleLevel::Gradient { x, y }, None)
2897 }
2898 Texture::SampleLevel => {
2899 let level = self.expression(args.next()?, ctx)?;
2900 (crate::SampleLevel::Exact(level), None)
2901 }
2902 };
2903
2904 let offset = args
2905 .next()
2906 .map(|arg| self.expression(arg, &mut ctx.as_const()))
2907 .ok()
2908 .transpose()?;
2909
2910 args.finish()?;
2911
2912 Ok(crate::Expression::ImageSample {
2913 image,
2914 sampler,
2915 gather,
2916 coordinate,
2917 array_index,
2918 offset,
2919 level,
2920 depth_ref,
2921 })
2922 }
2923
2924 fn subgroup_operation_helper(
2925 &mut self,
2926 span: Span,
2927 op: crate::SubgroupOperation,
2928 collective_op: crate::CollectiveOperation,
2929 arguments: &[Handle<ast::Expression<'source>>],
2930 ctx: &mut ExpressionContext<'source, '_, '_>,
2931 ) -> Result<Handle<crate::Expression>, Error<'source>> {
2932 let mut args = ctx.prepare_args(arguments, 1, span);
2933
2934 let argument = self.expression(args.next()?, ctx)?;
2935 args.finish()?;
2936
2937 let ty = ctx.register_type(argument)?;
2938
2939 let result =
2940 ctx.interrupt_emitter(crate::Expression::SubgroupOperationResult { ty }, span)?;
2941 let rctx = ctx.runtime_expression_ctx(span)?;
2942 rctx.block.push(
2943 crate::Statement::SubgroupCollectiveOperation {
2944 op,
2945 collective_op,
2946 argument,
2947 result,
2948 },
2949 span,
2950 );
2951 Ok(result)
2952 }
2953
2954 fn subgroup_gather_helper(
2955 &mut self,
2956 span: Span,
2957 mode: SubgroupGather,
2958 arguments: &[Handle<ast::Expression<'source>>],
2959 ctx: &mut ExpressionContext<'source, '_, '_>,
2960 ) -> Result<Handle<crate::Expression>, Error<'source>> {
2961 let mut args = ctx.prepare_args(arguments, 2, span);
2962
2963 let argument = self.expression(args.next()?, ctx)?;
2964
2965 use SubgroupGather as Sg;
2966 let mode = if let Sg::BroadcastFirst = mode {
2967 crate::GatherMode::BroadcastFirst
2968 } else {
2969 let index = self.expression(args.next()?, ctx)?;
2970 match mode {
2971 Sg::Broadcast => crate::GatherMode::Broadcast(index),
2972 Sg::Shuffle => crate::GatherMode::Shuffle(index),
2973 Sg::ShuffleDown => crate::GatherMode::ShuffleDown(index),
2974 Sg::ShuffleUp => crate::GatherMode::ShuffleUp(index),
2975 Sg::ShuffleXor => crate::GatherMode::ShuffleXor(index),
2976 Sg::BroadcastFirst => unreachable!(),
2977 }
2978 };
2979
2980 args.finish()?;
2981
2982 let ty = ctx.register_type(argument)?;
2983
2984 let result =
2985 ctx.interrupt_emitter(crate::Expression::SubgroupOperationResult { ty }, span)?;
2986 let rctx = ctx.runtime_expression_ctx(span)?;
2987 rctx.block.push(
2988 crate::Statement::SubgroupGather {
2989 mode,
2990 argument,
2991 result,
2992 },
2993 span,
2994 );
2995 Ok(result)
2996 }
2997
2998 fn r#struct(
2999 &mut self,
3000 s: &ast::Struct<'source>,
3001 span: Span,
3002 ctx: &mut GlobalContext<'source, '_, '_>,
3003 ) -> Result<Handle<crate::Type>, Error<'source>> {
3004 let mut offset = 0;
3005 let mut struct_alignment = Alignment::ONE;
3006 let mut members = Vec::with_capacity(s.members.len());
3007
3008 for member in s.members.iter() {
3009 let ty = self.resolve_ast_type(member.ty, ctx)?;
3010
3011 self.layouter.update(ctx.module.to_ctx()).unwrap();
3012
3013 let member_min_size = self.layouter[ty].size;
3014 let member_min_alignment = self.layouter[ty].alignment;
3015
3016 let member_size = if let Some(size_expr) = member.size {
3017 let (size, span) = self.const_u32(size_expr, &mut ctx.as_const())?;
3018 if size < member_min_size {
3019 return Err(Error::SizeAttributeTooLow(span, member_min_size));
3020 } else {
3021 size
3022 }
3023 } else {
3024 member_min_size
3025 };
3026
3027 let member_alignment = if let Some(align_expr) = member.align {
3028 let (align, span) = self.const_u32(align_expr, &mut ctx.as_const())?;
3029 if let Some(alignment) = Alignment::new(align) {
3030 if alignment < member_min_alignment {
3031 return Err(Error::AlignAttributeTooLow(span, member_min_alignment));
3032 } else {
3033 alignment
3034 }
3035 } else {
3036 return Err(Error::NonPowerOfTwoAlignAttribute(span));
3037 }
3038 } else {
3039 member_min_alignment
3040 };
3041
3042 let binding = self.binding(&member.binding, ty, ctx)?;
3043
3044 offset = member_alignment.round_up(offset);
3045 struct_alignment = struct_alignment.max(member_alignment);
3046
3047 members.push(crate::StructMember {
3048 name: Some(member.name.name.to_owned()),
3049 ty,
3050 binding,
3051 offset,
3052 });
3053
3054 offset += member_size;
3055 }
3056
3057 let size = struct_alignment.round_up(offset);
3058 let inner = crate::TypeInner::Struct {
3059 members,
3060 span: size,
3061 };
3062
3063 let handle = ctx.module.types.insert(
3064 crate::Type {
3065 name: Some(s.name.name.to_string()),
3066 inner,
3067 },
3068 span,
3069 );
3070 Ok(handle)
3071 }
3072
3073 fn const_u32(
3074 &mut self,
3075 expr: Handle<ast::Expression<'source>>,
3076 ctx: &mut ExpressionContext<'source, '_, '_>,
3077 ) -> Result<(u32, Span), Error<'source>> {
3078 let span = ctx.ast_expressions.get_span(expr);
3079 let expr = self.expression(expr, ctx)?;
3080 let value = ctx
3081 .module
3082 .to_ctx()
3083 .eval_expr_to_u32(expr)
3084 .map_err(|err| match err {
3085 crate::proc::U32EvalError::NonConst => {
3086 Error::ExpectedConstExprConcreteIntegerScalar(span)
3087 }
3088 crate::proc::U32EvalError::Negative => Error::ExpectedNonNegative(span),
3089 })?;
3090 Ok((value, span))
3091 }
3092
3093 fn array_size(
3094 &mut self,
3095 size: ast::ArraySize<'source>,
3096 ctx: &mut GlobalContext<'source, '_, '_>,
3097 ) -> Result<crate::ArraySize, Error<'source>> {
3098 Ok(match size {
3099 ast::ArraySize::Constant(expr) => {
3100 let span = ctx.ast_expressions.get_span(expr);
3101 let const_expr = self.expression(expr, &mut ctx.as_const());
3102 match const_expr {
3103 Ok(value) => {
3104 let len =
3105 ctx.module.to_ctx().eval_expr_to_u32(value).map_err(
3106 |err| match err {
3107 crate::proc::U32EvalError::NonConst => {
3108 Error::ExpectedConstExprConcreteIntegerScalar(span)
3109 }
3110 crate::proc::U32EvalError::Negative => {
3111 Error::ExpectedPositiveArrayLength(span)
3112 }
3113 },
3114 )?;
3115 let size =
3116 NonZeroU32::new(len).ok_or(Error::ExpectedPositiveArrayLength(span))?;
3117 crate::ArraySize::Constant(size)
3118 }
3119 err => {
3120 if let Err(Error::ConstantEvaluatorError(ref ty, _)) = err {
3121 match **ty {
3122 crate::proc::ConstantEvaluatorError::OverrideExpr => {
3123 crate::ArraySize::Pending(self.array_size_override(
3124 expr,
3125 &mut ctx.as_override(),
3126 span,
3127 )?)
3128 }
3129 _ => {
3130 err?;
3131 unreachable!()
3132 }
3133 }
3134 } else {
3135 err?;
3136 unreachable!()
3137 }
3138 }
3139 }
3140 }
3141 ast::ArraySize::Dynamic => crate::ArraySize::Dynamic,
3142 })
3143 }
3144
3145 fn array_size_override(
3146 &mut self,
3147 size_expr: Handle<ast::Expression<'source>>,
3148 ctx: &mut ExpressionContext<'source, '_, '_>,
3149 span: Span,
3150 ) -> Result<crate::PendingArraySize, Error<'source>> {
3151 let expr = self.expression(size_expr, ctx)?;
3152 match resolve_inner!(ctx, expr).scalar_kind().ok_or(0) {
3153 Ok(crate::ScalarKind::Sint) | Ok(crate::ScalarKind::Uint) => Ok({
3154 if let crate::Expression::Override(handle) = ctx.module.global_expressions[expr] {
3155 crate::PendingArraySize::Override(handle)
3156 } else {
3157 crate::PendingArraySize::Expression(expr)
3158 }
3159 }),
3160 _ => Err(Error::ExpectedConstExprConcreteIntegerScalar(span)),
3161 }
3162 }
3163
3164 fn resolve_named_ast_type(
3174 &mut self,
3175 handle: Handle<ast::Type<'source>>,
3176 name: Option<String>,
3177 ctx: &mut GlobalContext<'source, '_, '_>,
3178 ) -> Result<Handle<crate::Type>, Error<'source>> {
3179 let inner = match ctx.types[handle] {
3180 ast::Type::Scalar(scalar) => scalar.to_inner_scalar(),
3181 ast::Type::Vector { size, ty, ty_span } => {
3182 let ty = self.resolve_ast_type(ty, ctx)?;
3183 let scalar = match ctx.module.types[ty].inner {
3184 crate::TypeInner::Scalar(sc) => sc,
3185 _ => return Err(Error::UnknownScalarType(ty_span)),
3186 };
3187 crate::TypeInner::Vector { size, scalar }
3188 }
3189 ast::Type::Matrix {
3190 rows,
3191 columns,
3192 ty,
3193 ty_span,
3194 } => {
3195 let ty = self.resolve_ast_type(ty, ctx)?;
3196 let scalar = match ctx.module.types[ty].inner {
3197 crate::TypeInner::Scalar(sc) => sc,
3198 _ => return Err(Error::UnknownScalarType(ty_span)),
3199 };
3200 match scalar.kind {
3201 crate::ScalarKind::Float => crate::TypeInner::Matrix {
3202 columns,
3203 rows,
3204 scalar,
3205 },
3206 _ => return Err(Error::BadMatrixScalarKind(ty_span, scalar)),
3207 }
3208 }
3209 ast::Type::Atomic(scalar) => scalar.to_inner_atomic(),
3210 ast::Type::Pointer { base, space } => {
3211 let base = self.resolve_ast_type(base, ctx)?;
3212 crate::TypeInner::Pointer { base, space }
3213 }
3214 ast::Type::Array { base, size } => {
3215 let base = self.resolve_ast_type(base, ctx)?;
3216 let size = self.array_size(size, ctx)?;
3217
3218 self.layouter.update(ctx.module.to_ctx()).unwrap();
3219 let stride = self.layouter[base].to_stride();
3220
3221 crate::TypeInner::Array { base, size, stride }
3222 }
3223 ast::Type::Image {
3224 dim,
3225 arrayed,
3226 class,
3227 } => crate::TypeInner::Image {
3228 dim,
3229 arrayed,
3230 class,
3231 },
3232 ast::Type::Sampler { comparison } => crate::TypeInner::Sampler { comparison },
3233 ast::Type::AccelerationStructure => crate::TypeInner::AccelerationStructure,
3234 ast::Type::RayQuery => crate::TypeInner::RayQuery,
3235 ast::Type::BindingArray { base, size } => {
3236 let base = self.resolve_ast_type(base, ctx)?;
3237 let size = self.array_size(size, ctx)?;
3238 crate::TypeInner::BindingArray { base, size }
3239 }
3240 ast::Type::RayDesc => {
3241 return Ok(ctx.module.generate_ray_desc_type());
3242 }
3243 ast::Type::RayIntersection => {
3244 return Ok(ctx.module.generate_ray_intersection_type());
3245 }
3246 ast::Type::User(ref ident) => {
3247 return match ctx.globals.get(ident.name) {
3248 Some(&LoweredGlobalDecl::Type(handle)) => Ok(handle),
3249 Some(_) => Err(Error::Unexpected(ident.span, ExpectedToken::Type)),
3250 None => Err(Error::UnknownType(ident.span)),
3251 }
3252 }
3253 };
3254
3255 Ok(ctx.ensure_type_exists(name, inner))
3256 }
3257
3258 fn resolve_ast_type(
3260 &mut self,
3261 handle: Handle<ast::Type<'source>>,
3262 ctx: &mut GlobalContext<'source, '_, '_>,
3263 ) -> Result<Handle<crate::Type>, Error<'source>> {
3264 self.resolve_named_ast_type(handle, None, ctx)
3265 }
3266
3267 fn binding(
3268 &mut self,
3269 binding: &Option<ast::Binding<'source>>,
3270 ty: Handle<crate::Type>,
3271 ctx: &mut GlobalContext<'source, '_, '_>,
3272 ) -> Result<Option<crate::Binding>, Error<'source>> {
3273 Ok(match *binding {
3274 Some(ast::Binding::BuiltIn(b)) => Some(crate::Binding::BuiltIn(b)),
3275 Some(ast::Binding::Location {
3276 location,
3277 second_blend_source,
3278 interpolation,
3279 sampling,
3280 }) => {
3281 let mut binding = crate::Binding::Location {
3282 location: self.const_u32(location, &mut ctx.as_const())?.0,
3283 second_blend_source,
3284 interpolation,
3285 sampling,
3286 };
3287 binding.apply_default_interpolation(&ctx.module.types[ty].inner);
3288 Some(binding)
3289 }
3290 None => None,
3291 })
3292 }
3293
3294 fn ray_query_pointer(
3295 &mut self,
3296 expr: Handle<ast::Expression<'source>>,
3297 ctx: &mut ExpressionContext<'source, '_, '_>,
3298 ) -> Result<Handle<crate::Expression>, Error<'source>> {
3299 let span = ctx.ast_expressions.get_span(expr);
3300 let pointer = self.expression(expr, ctx)?;
3301
3302 match *resolve_inner!(ctx, pointer) {
3303 crate::TypeInner::Pointer { base, .. } => match ctx.module.types[base].inner {
3304 crate::TypeInner::RayQuery => Ok(pointer),
3305 ref other => {
3306 log::error!("Pointer type to {:?} passed to ray query op", other);
3307 Err(Error::InvalidRayQueryPointer(span))
3308 }
3309 },
3310 ref other => {
3311 log::error!("Type {:?} passed to ray query op", other);
3312 Err(Error::InvalidRayQueryPointer(span))
3313 }
3314 }
3315 }
3316}
3317
3318impl crate::AtomicFunction {
3319 pub fn map(word: &str) -> Option<Self> {
3320 Some(match word {
3321 "atomicAdd" => crate::AtomicFunction::Add,
3322 "atomicSub" => crate::AtomicFunction::Subtract,
3323 "atomicAnd" => crate::AtomicFunction::And,
3324 "atomicOr" => crate::AtomicFunction::InclusiveOr,
3325 "atomicXor" => crate::AtomicFunction::ExclusiveOr,
3326 "atomicMin" => crate::AtomicFunction::Min,
3327 "atomicMax" => crate::AtomicFunction::Max,
3328 "atomicExchange" => crate::AtomicFunction::Exchange { compare: None },
3329 _ => return None,
3330 })
3331 }
3332}