mirror of
https://github.com/gfx-rs/wgpu.git
synced 2026-04-22 03:02:01 -04:00
[naga wgsl-in] Ensure correct array stride when converting types during const evaluation (#7112)
When converting the underlying scalar type of an array during const evaluation, we currently use the resulting base type's size as the array stride for the resulting type. For certain types, this may not match the required alignment and will therefore result in a validation error. For example, `array<vec3<f32>, N>` should have a stride of 16. But if declared with an abstract initializer, eg `array(vec3(0.0))` we will incorrectly determine the stride to be 12. To solve this, we use the proc::Layouter struct to determine the required array stride during const evaluation. To avoid repeating layouting work, we reuse the Lowerer's layouter, passing it through the various *Contexts through to the ConstantEvaluator.
This commit is contained in:
@@ -75,6 +75,7 @@ pub fn process_overrides<'a>(
|
||||
let mut adjusted_constant_initializers = HashSet::with_capacity(module.constants.len());
|
||||
|
||||
let mut global_expression_kind_tracker = crate::proc::ExpressionKindTracker::new();
|
||||
let mut layouter = crate::proc::Layouter::default();
|
||||
|
||||
// An iterator through the original overrides table, consumed in
|
||||
// approximate tandem with the global expressions.
|
||||
@@ -147,6 +148,7 @@ pub fn process_overrides<'a>(
|
||||
let mut evaluator = ConstantEvaluator::for_wgsl_module(
|
||||
&mut module,
|
||||
&mut global_expression_kind_tracker,
|
||||
&mut layouter,
|
||||
false,
|
||||
);
|
||||
adjust_expr(&adjusted_global_expressions, &mut expr);
|
||||
@@ -186,13 +188,13 @@ pub fn process_overrides<'a>(
|
||||
|
||||
let mut functions = mem::take(&mut module.functions);
|
||||
for (_, function) in functions.iter_mut() {
|
||||
process_function(&mut module, &override_map, function)?;
|
||||
process_function(&mut module, &override_map, &mut layouter, function)?;
|
||||
}
|
||||
module.functions = functions;
|
||||
|
||||
let mut entry_points = mem::take(&mut module.entry_points);
|
||||
for ep in entry_points.iter_mut() {
|
||||
process_function(&mut module, &override_map, &mut ep.function)?;
|
||||
process_function(&mut module, &override_map, &mut layouter, &mut ep.function)?;
|
||||
process_workgroup_size_override(&mut module, &adjusted_global_expressions, ep)?;
|
||||
}
|
||||
module.entry_points = entry_points;
|
||||
@@ -365,6 +367,7 @@ fn process_override(
|
||||
fn process_function(
|
||||
module: &mut Module,
|
||||
override_map: &HandleVec<Override, Handle<Constant>>,
|
||||
layouter: &mut crate::proc::Layouter,
|
||||
function: &mut Function,
|
||||
) -> Result<(), ConstantEvaluatorError> {
|
||||
// A map from original local expression handles to
|
||||
@@ -389,6 +392,7 @@ fn process_function(
|
||||
module,
|
||||
&mut function.expressions,
|
||||
&mut local_expression_kind_tracker,
|
||||
layouter,
|
||||
&mut emitter,
|
||||
&mut block,
|
||||
false,
|
||||
|
||||
@@ -8,9 +8,9 @@ use super::{
|
||||
Frontend, Result,
|
||||
};
|
||||
use crate::{
|
||||
front::Typifier, proc::Emitter, AddressSpace, Arena, BinaryOperator, Block, Expression,
|
||||
FastHashMap, FunctionArgument, Handle, Literal, LocalVariable, RelationalFunction, Scalar,
|
||||
Span, Statement, Type, TypeInner, VectorSize,
|
||||
front::Typifier, proc::Emitter, proc::Layouter, AddressSpace, Arena, BinaryOperator, Block,
|
||||
Expression, FastHashMap, FunctionArgument, Handle, Literal, LocalVariable, RelationalFunction,
|
||||
Scalar, Span, Statement, Type, TypeInner, VectorSize,
|
||||
};
|
||||
use std::ops::Index;
|
||||
|
||||
@@ -72,6 +72,7 @@ pub struct Context<'a> {
|
||||
|
||||
pub const_typifier: Typifier,
|
||||
pub typifier: Typifier,
|
||||
layouter: Layouter,
|
||||
emitter: Emitter,
|
||||
stmt_ctx: Option<StmtContext>,
|
||||
pub body: Block,
|
||||
@@ -103,6 +104,7 @@ impl<'a> Context<'a> {
|
||||
|
||||
const_typifier: Typifier::new(),
|
||||
typifier: Typifier::new(),
|
||||
layouter: Layouter::default(),
|
||||
emitter: Emitter::default(),
|
||||
stmt_ctx: Some(StmtContext::new()),
|
||||
body: Block::new(),
|
||||
@@ -260,12 +262,14 @@ impl<'a> Context<'a> {
|
||||
crate::proc::ConstantEvaluator::for_glsl_module(
|
||||
self.module,
|
||||
self.global_expression_kind_tracker,
|
||||
&mut self.layouter,
|
||||
)
|
||||
} else {
|
||||
crate::proc::ConstantEvaluator::for_glsl_function(
|
||||
self.module,
|
||||
&mut self.expressions,
|
||||
&mut self.local_expression_kind_tracker,
|
||||
&mut self.layouter,
|
||||
&mut self.emitter,
|
||||
&mut self.body,
|
||||
)
|
||||
|
||||
@@ -490,8 +490,8 @@ impl<'source> Lowerer<'source, '_> {
|
||||
NonZeroU32::new(u32::try_from(components.len()).unwrap()).unwrap(),
|
||||
),
|
||||
stride: {
|
||||
self.layouter.update(ctx.module.to_ctx()).unwrap();
|
||||
self.layouter[base].to_stride()
|
||||
ctx.layouter.update(ctx.module.to_ctx()).unwrap();
|
||||
ctx.layouter[base].to_stride()
|
||||
},
|
||||
};
|
||||
let ty = ctx.ensure_type_exists(inner);
|
||||
@@ -616,8 +616,8 @@ impl<'source> Lowerer<'source, '_> {
|
||||
let base = self.resolve_ast_type(base, &mut ctx.as_global())?;
|
||||
let size = self.array_size(size, &mut ctx.as_global())?;
|
||||
|
||||
self.layouter.update(ctx.module.to_ctx()).unwrap();
|
||||
let stride = self.layouter[base].to_stride();
|
||||
ctx.layouter.update(ctx.module.to_ctx()).unwrap();
|
||||
let stride = ctx.layouter[base].to_stride();
|
||||
|
||||
let ty = ctx.ensure_type_exists(crate::TypeInner::Array { base, size, stride });
|
||||
Constructor::Type(ty)
|
||||
|
||||
@@ -87,6 +87,8 @@ pub struct GlobalContext<'source, 'temp, 'out> {
|
||||
|
||||
const_typifier: &'temp mut Typifier,
|
||||
|
||||
layouter: &'temp mut Layouter,
|
||||
|
||||
global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
|
||||
}
|
||||
|
||||
@@ -98,6 +100,7 @@ impl<'source> GlobalContext<'source, '_, '_> {
|
||||
types: self.types,
|
||||
module: self.module,
|
||||
const_typifier: self.const_typifier,
|
||||
layouter: self.layouter,
|
||||
expr_type: ExpressionContextType::Constant(None),
|
||||
global_expression_kind_tracker: self.global_expression_kind_tracker,
|
||||
}
|
||||
@@ -110,6 +113,7 @@ impl<'source> GlobalContext<'source, '_, '_> {
|
||||
types: self.types,
|
||||
module: self.module,
|
||||
const_typifier: self.const_typifier,
|
||||
layouter: self.layouter,
|
||||
expr_type: ExpressionContextType::Override,
|
||||
global_expression_kind_tracker: self.global_expression_kind_tracker,
|
||||
}
|
||||
@@ -165,6 +169,7 @@ pub struct StatementContext<'source, 'temp, 'out> {
|
||||
|
||||
const_typifier: &'temp mut Typifier,
|
||||
typifier: &'temp mut Typifier,
|
||||
layouter: &'temp mut Layouter,
|
||||
function: &'out mut crate::Function,
|
||||
/// Stores the names of expressions that are assigned in `let` statement
|
||||
/// Also stores the spans of the names, for use in errors.
|
||||
@@ -198,6 +203,7 @@ impl<'a, 'temp> StatementContext<'a, 'temp, '_> {
|
||||
types: self.types,
|
||||
ast_expressions: self.ast_expressions,
|
||||
const_typifier: self.const_typifier,
|
||||
layouter: self.layouter,
|
||||
global_expression_kind_tracker: self.global_expression_kind_tracker,
|
||||
module: self.module,
|
||||
expr_type: ExpressionContextType::Constant(Some(LocalExpressionContext {
|
||||
@@ -224,6 +230,7 @@ impl<'a, 'temp> StatementContext<'a, 'temp, '_> {
|
||||
types: self.types,
|
||||
ast_expressions: self.ast_expressions,
|
||||
const_typifier: self.const_typifier,
|
||||
layouter: self.layouter,
|
||||
global_expression_kind_tracker: self.global_expression_kind_tracker,
|
||||
module: self.module,
|
||||
expr_type: ExpressionContextType::Runtime(LocalExpressionContext {
|
||||
@@ -244,6 +251,7 @@ impl<'a, 'temp> StatementContext<'a, 'temp, '_> {
|
||||
types: self.types,
|
||||
module: self.module,
|
||||
const_typifier: self.const_typifier,
|
||||
layouter: self.layouter,
|
||||
global_expression_kind_tracker: self.global_expression_kind_tracker,
|
||||
}
|
||||
}
|
||||
@@ -364,6 +372,7 @@ pub struct ExpressionContext<'source, 'temp, 'out> {
|
||||
///
|
||||
/// [`module::global_expressions`]: crate::Module::global_expressions
|
||||
const_typifier: &'temp mut Typifier,
|
||||
layouter: &'temp mut Layouter,
|
||||
global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
|
||||
|
||||
/// Whether we are lowering a constant expression or a general
|
||||
@@ -379,6 +388,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
|
||||
types: self.types,
|
||||
ast_expressions: self.ast_expressions,
|
||||
const_typifier: self.const_typifier,
|
||||
layouter: self.layouter,
|
||||
module: self.module,
|
||||
expr_type: ExpressionContextType::Constant(match self.expr_type {
|
||||
ExpressionContextType::Runtime(ref mut local_expression_context)
|
||||
@@ -406,6 +416,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
|
||||
types: self.types,
|
||||
module: self.module,
|
||||
const_typifier: self.const_typifier,
|
||||
layouter: self.layouter,
|
||||
global_expression_kind_tracker: self.global_expression_kind_tracker,
|
||||
}
|
||||
}
|
||||
@@ -416,6 +427,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
|
||||
self.module,
|
||||
&mut rctx.function.expressions,
|
||||
rctx.local_expression_kind_tracker,
|
||||
self.layouter,
|
||||
rctx.emitter,
|
||||
rctx.block,
|
||||
false,
|
||||
@@ -425,6 +437,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
|
||||
self.module,
|
||||
&mut rctx.function.expressions,
|
||||
rctx.local_expression_kind_tracker,
|
||||
self.layouter,
|
||||
rctx.emitter,
|
||||
rctx.block,
|
||||
true,
|
||||
@@ -433,11 +446,13 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
|
||||
ExpressionContextType::Constant(None) => ConstantEvaluator::for_wgsl_module(
|
||||
self.module,
|
||||
self.global_expression_kind_tracker,
|
||||
self.layouter,
|
||||
false,
|
||||
),
|
||||
ExpressionContextType::Override => ConstantEvaluator::for_wgsl_module(
|
||||
self.module,
|
||||
self.global_expression_kind_tracker,
|
||||
self.layouter,
|
||||
true,
|
||||
),
|
||||
}
|
||||
@@ -1016,15 +1031,11 @@ impl SubgroupGather {
|
||||
|
||||
pub struct Lowerer<'source, 'temp> {
|
||||
index: &'temp Index<'source>,
|
||||
layouter: Layouter,
|
||||
}
|
||||
|
||||
impl<'source, 'temp> Lowerer<'source, 'temp> {
|
||||
pub fn new(index: &'temp Index<'source>) -> Self {
|
||||
Self {
|
||||
index,
|
||||
layouter: Layouter::default(),
|
||||
}
|
||||
pub const fn new(index: &'temp Index<'source>) -> Self {
|
||||
Self { index }
|
||||
}
|
||||
|
||||
pub fn lower(
|
||||
@@ -1043,6 +1054,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
|
||||
types: &tu.types,
|
||||
module: &mut module,
|
||||
const_typifier: &mut Typifier::new(),
|
||||
layouter: &mut Layouter::default(),
|
||||
global_expression_kind_tracker: &mut crate::proc::ExpressionKindTracker::new(),
|
||||
};
|
||||
|
||||
@@ -1299,6 +1311,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
|
||||
ast_expressions: ctx.ast_expressions,
|
||||
const_typifier: ctx.const_typifier,
|
||||
typifier: &mut typifier,
|
||||
layouter: ctx.layouter,
|
||||
function: &mut function,
|
||||
named_expressions: &mut named_expressions,
|
||||
types: ctx.types,
|
||||
@@ -3051,10 +3064,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
|
||||
for member in s.members.iter() {
|
||||
let ty = self.resolve_ast_type(member.ty, ctx)?;
|
||||
|
||||
self.layouter.update(ctx.module.to_ctx()).unwrap();
|
||||
ctx.layouter.update(ctx.module.to_ctx()).unwrap();
|
||||
|
||||
let member_min_size = self.layouter[ty].size;
|
||||
let member_min_alignment = self.layouter[ty].alignment;
|
||||
let member_min_size = ctx.layouter[ty].size;
|
||||
let member_min_alignment = ctx.layouter[ty].alignment;
|
||||
|
||||
let member_size = if let Some(size_expr) = member.size {
|
||||
let (size, span) = self.const_u32(size_expr, &mut ctx.as_const())?;
|
||||
@@ -3258,8 +3271,8 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
|
||||
let base = self.resolve_ast_type(base, ctx)?;
|
||||
let size = self.array_size(size, ctx)?;
|
||||
|
||||
self.layouter.update(ctx.module.to_ctx()).unwrap();
|
||||
let stride = self.layouter[base].to_stride();
|
||||
ctx.layouter.update(ctx.module.to_ctx()).unwrap();
|
||||
let stride = ctx.layouter[base].to_stride();
|
||||
|
||||
crate::TypeInner::Array { base, size, stride }
|
||||
}
|
||||
|
||||
@@ -312,6 +312,8 @@ pub struct ConstantEvaluator<'a> {
|
||||
|
||||
/// Tracks the constness of expressions residing in [`Self::expressions`]
|
||||
expression_kind_tracker: &'a mut ExpressionKindTracker,
|
||||
|
||||
layouter: &'a mut crate::proc::Layouter,
|
||||
}
|
||||
|
||||
#[derive(Debug)]
|
||||
@@ -594,6 +596,7 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
pub fn for_wgsl_module(
|
||||
module: &'a mut crate::Module,
|
||||
global_expression_kind_tracker: &'a mut ExpressionKindTracker,
|
||||
layouter: &'a mut crate::proc::Layouter,
|
||||
in_override_ctx: bool,
|
||||
) -> Self {
|
||||
Self::for_module(
|
||||
@@ -604,6 +607,7 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
}),
|
||||
module,
|
||||
global_expression_kind_tracker,
|
||||
layouter,
|
||||
)
|
||||
}
|
||||
|
||||
@@ -614,11 +618,13 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
pub fn for_glsl_module(
|
||||
module: &'a mut crate::Module,
|
||||
global_expression_kind_tracker: &'a mut ExpressionKindTracker,
|
||||
layouter: &'a mut crate::proc::Layouter,
|
||||
) -> Self {
|
||||
Self::for_module(
|
||||
Behavior::Glsl(GlslRestrictions::Const),
|
||||
module,
|
||||
global_expression_kind_tracker,
|
||||
layouter,
|
||||
)
|
||||
}
|
||||
|
||||
@@ -626,6 +632,7 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
behavior: Behavior<'a>,
|
||||
module: &'a mut crate::Module,
|
||||
global_expression_kind_tracker: &'a mut ExpressionKindTracker,
|
||||
layouter: &'a mut crate::proc::Layouter,
|
||||
) -> Self {
|
||||
Self {
|
||||
behavior,
|
||||
@@ -634,6 +641,7 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
overrides: &module.overrides,
|
||||
expressions: &mut module.global_expressions,
|
||||
expression_kind_tracker: global_expression_kind_tracker,
|
||||
layouter,
|
||||
}
|
||||
}
|
||||
|
||||
@@ -645,6 +653,7 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
module: &'a mut crate::Module,
|
||||
expressions: &'a mut Arena<Expression>,
|
||||
local_expression_kind_tracker: &'a mut ExpressionKindTracker,
|
||||
layouter: &'a mut crate::proc::Layouter,
|
||||
emitter: &'a mut super::Emitter,
|
||||
block: &'a mut crate::Block,
|
||||
is_const: bool,
|
||||
@@ -665,6 +674,7 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
overrides: &module.overrides,
|
||||
expressions,
|
||||
expression_kind_tracker: local_expression_kind_tracker,
|
||||
layouter,
|
||||
}
|
||||
}
|
||||
|
||||
@@ -676,6 +686,7 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
module: &'a mut crate::Module,
|
||||
expressions: &'a mut Arena<Expression>,
|
||||
local_expression_kind_tracker: &'a mut ExpressionKindTracker,
|
||||
layouter: &'a mut crate::proc::Layouter,
|
||||
emitter: &'a mut super::Emitter,
|
||||
block: &'a mut crate::Block,
|
||||
) -> Self {
|
||||
@@ -690,6 +701,7 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
overrides: &module.overrides,
|
||||
expressions,
|
||||
expression_kind_tracker: local_expression_kind_tracker,
|
||||
layouter,
|
||||
}
|
||||
}
|
||||
|
||||
@@ -1718,7 +1730,11 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
self.types.insert(Type { name: None, inner }, span)
|
||||
}
|
||||
};
|
||||
let new_base_stride = self.types[new_base].inner.size(self.to_ctx());
|
||||
let mut layouter = std::mem::take(self.layouter);
|
||||
layouter.update(self.to_ctx()).unwrap();
|
||||
*self.layouter = layouter;
|
||||
|
||||
let new_base_stride = self.layouter[new_base].to_stride();
|
||||
let new_array_ty = self.types.insert(
|
||||
Type {
|
||||
name: None,
|
||||
@@ -2567,6 +2583,7 @@ mod tests {
|
||||
overrides: &overrides,
|
||||
expressions: &mut global_expressions,
|
||||
expression_kind_tracker,
|
||||
layouter: &mut crate::proc::Layouter::default(),
|
||||
};
|
||||
|
||||
let res1 = solver
|
||||
@@ -2653,6 +2670,7 @@ mod tests {
|
||||
overrides: &overrides,
|
||||
expressions: &mut global_expressions,
|
||||
expression_kind_tracker,
|
||||
layouter: &mut crate::proc::Layouter::default(),
|
||||
};
|
||||
|
||||
let res = solver
|
||||
@@ -2771,6 +2789,7 @@ mod tests {
|
||||
overrides: &overrides,
|
||||
expressions: &mut global_expressions,
|
||||
expression_kind_tracker,
|
||||
layouter: &mut crate::proc::Layouter::default(),
|
||||
};
|
||||
|
||||
let root1 = Expression::AccessIndex { base, index: 1 };
|
||||
@@ -2864,6 +2883,7 @@ mod tests {
|
||||
overrides: &overrides,
|
||||
expressions: &mut global_expressions,
|
||||
expression_kind_tracker,
|
||||
layouter: &mut crate::proc::Layouter::default(),
|
||||
};
|
||||
|
||||
let solved_compose = solver
|
||||
@@ -2946,6 +2966,7 @@ mod tests {
|
||||
overrides: &overrides,
|
||||
expressions: &mut global_expressions,
|
||||
expression_kind_tracker,
|
||||
layouter: &mut crate::proc::Layouter::default(),
|
||||
};
|
||||
|
||||
let solved_compose = solver
|
||||
@@ -3034,6 +3055,7 @@ mod tests {
|
||||
overrides: &overrides,
|
||||
expressions: &mut global_expressions,
|
||||
expression_kind_tracker,
|
||||
layouter: &mut crate::proc::Layouter::default(),
|
||||
};
|
||||
|
||||
let solved_add = solver
|
||||
|
||||
@@ -43,6 +43,10 @@ var<private> xafpaiaf: array<f32, 2> = array(1, 2.0);
|
||||
var<private> xafpafai: array<f32, 2> = array(1.0, 2);
|
||||
var<private> xafpafaf: array<f32, 2> = array(1.0, 2.0);
|
||||
|
||||
var<private> xavipai: array<vec3<i32>, 1> = array(vec3(1));
|
||||
var<private> xavfpai: array<vec3<f32>, 1> = array(vec3(1));
|
||||
var<private> xavfpaf: array<vec3<f32>, 1> = array(vec3(1.0));
|
||||
|
||||
var<private> ivispai = vec2(1);
|
||||
var<private> ivfspaf = vec2(1.0);
|
||||
var<private> ivis_ai = vec2<i32>(1);
|
||||
@@ -58,6 +62,10 @@ var<private> iafpafaf = array(1.0, 2.0);
|
||||
var<private> iafpaiaf = array(1, 2.0);
|
||||
var<private> iafpafai = array(1.0, 2);
|
||||
|
||||
var<private> iavipai = array(vec3(1));
|
||||
var<private> iavfpai = array(vec3(1));
|
||||
var<private> iavfpaf = array(vec3(1.0));
|
||||
|
||||
fn all_constant_arguments() {
|
||||
var xvipaiai: vec2<i32> = vec2(42, 43);
|
||||
var xvupaiai: vec2<u32> = vec2(44, 45);
|
||||
@@ -100,6 +108,10 @@ fn all_constant_arguments() {
|
||||
var xafpafai: array<f32, 2> = array(1.0, 2);
|
||||
var xafpafaf: array<f32, 2> = array(1.0, 2.0);
|
||||
|
||||
var xavipai: array<vec3<i32>, 1> = array(vec3(1));
|
||||
var xavfpai: array<vec3<f32>, 1> = array(vec3(1));
|
||||
var xavfpaf: array<vec3<f32>, 1> = array(vec3(1.0));
|
||||
|
||||
var iaipaiai = array(1, 2);
|
||||
var iafpaiaf = array(1, 2.0);
|
||||
var iafpafai = array(1.0, 2);
|
||||
@@ -147,6 +159,10 @@ fn all_constant_arguments() {
|
||||
xafpafai = array(1.0, 2);
|
||||
xafpafaf = array(1.0, 2.0);
|
||||
|
||||
xavipai = array(vec3(1));
|
||||
xavfpai = array(vec3(1));
|
||||
xavfpaf = array(vec3(1.0));
|
||||
|
||||
iaipaiai = array(1, 2);
|
||||
iafpaiaf = array(1, 2.0);
|
||||
iafpafai = array(1.0, 2);
|
||||
|
||||
@@ -10,6 +10,12 @@ struct type_7 {
|
||||
struct type_8 {
|
||||
int inner[2];
|
||||
};
|
||||
struct type_10 {
|
||||
metal::int3 inner[1];
|
||||
};
|
||||
struct type_12 {
|
||||
metal::float3 inner[1];
|
||||
};
|
||||
|
||||
void all_constant_arguments(
|
||||
) {
|
||||
@@ -46,6 +52,9 @@ void all_constant_arguments(
|
||||
type_7 xafpaiaf = type_7 {1.0, 2.0};
|
||||
type_7 xafpafai = type_7 {1.0, 2.0};
|
||||
type_7 xafpafaf = type_7 {1.0, 2.0};
|
||||
type_10 xavipai = type_10 {metal::int3(1)};
|
||||
type_12 xavfpai = type_12 {metal::float3(1.0)};
|
||||
type_12 xavfpaf = type_12 {metal::float3(1.0)};
|
||||
type_8 iaipaiai = type_8 {1, 2};
|
||||
type_7 iafpaiaf = type_7 {1.0, 2.0};
|
||||
type_7 iafpafai = type_7 {1.0, 2.0};
|
||||
@@ -83,6 +92,9 @@ void all_constant_arguments(
|
||||
xafpaiaf = type_7 {1.0, 2.0};
|
||||
xafpafai = type_7 {1.0, 2.0};
|
||||
xafpafaf = type_7 {1.0, 2.0};
|
||||
xavipai = type_10 {metal::int3(1)};
|
||||
xavfpai = type_12 {metal::float3(1.0)};
|
||||
xavfpaf = type_12 {metal::float3(1.0)};
|
||||
iaipaiai = type_8 {1, 2};
|
||||
iafpaiaf = type_7 {1.0, 2.0};
|
||||
iafpafai = type_7 {1.0, 2.0};
|
||||
|
||||
@@ -1,13 +1,15 @@
|
||||
; SPIR-V
|
||||
; Version: 1.1
|
||||
; Generator: rspirv
|
||||
; Bound: 269
|
||||
; Bound: 290
|
||||
OpCapability Shader
|
||||
OpCapability Linkage
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpDecorate %10 ArrayStride 4
|
||||
OpDecorate %12 ArrayStride 4
|
||||
OpDecorate %14 ArrayStride 16
|
||||
OpDecorate %17 ArrayStride 16
|
||||
%2 = OpTypeVoid
|
||||
%3 = OpTypeInt 32 1
|
||||
%4 = OpTypeVector %3 2
|
||||
@@ -19,342 +21,366 @@ OpDecorate %12 ArrayStride 4
|
||||
%11 = OpConstant %5 2
|
||||
%10 = OpTypeArray %7 %11
|
||||
%12 = OpTypeArray %3 %11
|
||||
%13 = OpConstant %3 42
|
||||
%14 = OpConstant %3 43
|
||||
%15 = OpConstantComposite %4 %13 %14
|
||||
%16 = OpConstant %5 44
|
||||
%17 = OpConstant %5 45
|
||||
%18 = OpConstantComposite %6 %16 %17
|
||||
%19 = OpConstant %7 46.0
|
||||
%20 = OpConstant %7 47.0
|
||||
%21 = OpConstantComposite %8 %19 %20
|
||||
%22 = OpConstant %5 42
|
||||
%23 = OpConstant %5 43
|
||||
%24 = OpConstantComposite %6 %22 %23
|
||||
%25 = OpConstant %7 1.0
|
||||
%26 = OpConstant %7 2.0
|
||||
%27 = OpConstantComposite %8 %25 %26
|
||||
%28 = OpConstant %7 3.0
|
||||
%29 = OpConstant %7 4.0
|
||||
%30 = OpConstantComposite %8 %28 %29
|
||||
%31 = OpConstantComposite %9 %27 %30
|
||||
%32 = OpConstant %3 1
|
||||
%33 = OpConstantComposite %4 %32 %32
|
||||
%34 = OpConstantComposite %8 %25 %25
|
||||
%35 = OpConstant %5 1
|
||||
%36 = OpConstantComposite %6 %35 %35
|
||||
%37 = OpConstantComposite %10 %25 %26
|
||||
%38 = OpConstant %3 2
|
||||
%39 = OpConstantComposite %12 %32 %38
|
||||
%41 = OpTypePointer Private %4
|
||||
%40 = OpVariable %41 Private %15
|
||||
%43 = OpTypePointer Private %6
|
||||
%42 = OpVariable %43 Private %18
|
||||
%45 = OpTypePointer Private %8
|
||||
%44 = OpVariable %45 Private %21
|
||||
%46 = OpVariable %43 Private %24
|
||||
%47 = OpVariable %43 Private %24
|
||||
%48 = OpVariable %43 Private %24
|
||||
%49 = OpVariable %43 Private %24
|
||||
%51 = OpTypePointer Private %9
|
||||
%50 = OpVariable %51 Private %31
|
||||
%52 = OpVariable %51 Private %31
|
||||
%53 = OpVariable %51 Private %31
|
||||
%54 = OpVariable %51 Private %31
|
||||
%55 = OpVariable %51 Private %31
|
||||
%56 = OpVariable %41 Private %33
|
||||
%57 = OpVariable %45 Private %34
|
||||
%58 = OpVariable %41 Private %33
|
||||
%59 = OpVariable %43 Private %36
|
||||
%60 = OpVariable %45 Private %34
|
||||
%61 = OpVariable %45 Private %34
|
||||
%63 = OpTypePointer Private %10
|
||||
%62 = OpVariable %63 Private %37
|
||||
%64 = OpVariable %63 Private %37
|
||||
%66 = OpTypePointer Private %12
|
||||
%65 = OpVariable %66 Private %39
|
||||
%67 = OpVariable %63 Private %37
|
||||
%68 = OpVariable %63 Private %37
|
||||
%69 = OpVariable %63 Private %37
|
||||
%70 = OpVariable %41 Private %33
|
||||
%71 = OpVariable %45 Private %34
|
||||
%72 = OpVariable %41 Private %33
|
||||
%73 = OpVariable %43 Private %36
|
||||
%74 = OpVariable %45 Private %34
|
||||
%75 = OpVariable %45 Private %34
|
||||
%76 = OpVariable %63 Private %37
|
||||
%77 = OpVariable %63 Private %37
|
||||
%78 = OpVariable %66 Private %39
|
||||
%79 = OpVariable %63 Private %37
|
||||
%80 = OpVariable %63 Private %37
|
||||
%81 = OpVariable %63 Private %37
|
||||
%84 = OpTypeFunction %2
|
||||
%86 = OpTypePointer Function %4
|
||||
%88 = OpTypePointer Function %6
|
||||
%90 = OpTypePointer Function %8
|
||||
%96 = OpTypePointer Function %9
|
||||
%112 = OpTypePointer Function %10
|
||||
%117 = OpTypePointer Function %12
|
||||
%132 = OpTypePointer Function %5
|
||||
%133 = OpConstantNull %5
|
||||
%135 = OpTypePointer Function %3
|
||||
%136 = OpConstantNull %3
|
||||
%138 = OpTypePointer Function %7
|
||||
%139 = OpConstantNull %7
|
||||
%141 = OpConstantNull %6
|
||||
%143 = OpConstantNull %6
|
||||
%145 = OpConstantNull %6
|
||||
%147 = OpConstantNull %6
|
||||
%149 = OpConstantNull %9
|
||||
%151 = OpConstantNull %9
|
||||
%153 = OpConstantNull %9
|
||||
%155 = OpConstantNull %9
|
||||
%157 = OpConstantNull %10
|
||||
%159 = OpConstantNull %10
|
||||
%161 = OpConstantNull %10
|
||||
%163 = OpConstantNull %10
|
||||
%165 = OpConstantNull %12
|
||||
%167 = OpConstantNull %12
|
||||
%169 = OpConstantNull %10
|
||||
%171 = OpConstantNull %10
|
||||
%173 = OpConstantNull %10
|
||||
%175 = OpConstantNull %10
|
||||
%177 = OpConstantNull %12
|
||||
%179 = OpConstantNull %12
|
||||
%83 = OpFunction %2 None %84
|
||||
%82 = OpLabel
|
||||
%127 = OpVariable %112 Function %37
|
||||
%124 = OpVariable %117 Function %39
|
||||
%121 = OpVariable %112 Function %37
|
||||
%118 = OpVariable %117 Function %39
|
||||
%114 = OpVariable %112 Function %37
|
||||
%110 = OpVariable %90 Function %34
|
||||
%107 = OpVariable %86 Function %33
|
||||
%104 = OpVariable %96 Function %31
|
||||
%101 = OpVariable %96 Function %31
|
||||
%98 = OpVariable %96 Function %31
|
||||
%94 = OpVariable %88 Function %24
|
||||
%91 = OpVariable %88 Function %24
|
||||
%85 = OpVariable %86 Function %15
|
||||
%125 = OpVariable %112 Function %37
|
||||
%122 = OpVariable %112 Function %37
|
||||
%119 = OpVariable %117 Function %39
|
||||
%115 = OpVariable %112 Function %37
|
||||
%111 = OpVariable %112 Function %37
|
||||
%108 = OpVariable %88 Function %36
|
||||
%105 = OpVariable %86 Function %33
|
||||
%102 = OpVariable %96 Function %31
|
||||
%99 = OpVariable %96 Function %31
|
||||
%95 = OpVariable %96 Function %31
|
||||
%92 = OpVariable %88 Function %24
|
||||
%87 = OpVariable %88 Function %18
|
||||
%126 = OpVariable %112 Function %37
|
||||
%123 = OpVariable %112 Function %37
|
||||
%120 = OpVariable %112 Function %37
|
||||
%116 = OpVariable %117 Function %39
|
||||
%113 = OpVariable %112 Function %37
|
||||
%109 = OpVariable %90 Function %34
|
||||
%106 = OpVariable %90 Function %34
|
||||
%103 = OpVariable %96 Function %31
|
||||
%100 = OpVariable %96 Function %31
|
||||
%97 = OpVariable %96 Function %31
|
||||
%93 = OpVariable %88 Function %24
|
||||
%89 = OpVariable %90 Function %21
|
||||
OpBranch %128
|
||||
%128 = OpLabel
|
||||
OpStore %85 %15
|
||||
OpStore %87 %18
|
||||
OpStore %89 %21
|
||||
OpStore %91 %24
|
||||
OpStore %92 %24
|
||||
OpStore %93 %24
|
||||
OpStore %94 %24
|
||||
OpStore %95 %31
|
||||
OpStore %97 %31
|
||||
OpStore %98 %31
|
||||
OpStore %99 %31
|
||||
OpStore %100 %31
|
||||
OpStore %101 %31
|
||||
OpStore %102 %31
|
||||
OpStore %103 %31
|
||||
OpStore %104 %31
|
||||
OpStore %105 %33
|
||||
OpStore %106 %34
|
||||
OpStore %107 %33
|
||||
OpStore %108 %36
|
||||
OpStore %109 %34
|
||||
OpStore %110 %34
|
||||
OpStore %111 %37
|
||||
OpStore %113 %37
|
||||
OpStore %114 %37
|
||||
OpStore %115 %37
|
||||
OpStore %116 %39
|
||||
OpStore %118 %39
|
||||
OpStore %119 %39
|
||||
OpStore %120 %37
|
||||
OpStore %121 %37
|
||||
OpStore %122 %37
|
||||
OpStore %123 %37
|
||||
OpStore %124 %39
|
||||
OpStore %125 %37
|
||||
OpStore %126 %37
|
||||
OpStore %127 %37
|
||||
%13 = OpTypeVector %3 3
|
||||
%15 = OpConstant %5 1
|
||||
%14 = OpTypeArray %13 %15
|
||||
%16 = OpTypeVector %7 3
|
||||
%17 = OpTypeArray %16 %15
|
||||
%18 = OpConstant %3 42
|
||||
%19 = OpConstant %3 43
|
||||
%20 = OpConstantComposite %4 %18 %19
|
||||
%21 = OpConstant %5 44
|
||||
%22 = OpConstant %5 45
|
||||
%23 = OpConstantComposite %6 %21 %22
|
||||
%24 = OpConstant %7 46.0
|
||||
%25 = OpConstant %7 47.0
|
||||
%26 = OpConstantComposite %8 %24 %25
|
||||
%27 = OpConstant %5 42
|
||||
%28 = OpConstant %5 43
|
||||
%29 = OpConstantComposite %6 %27 %28
|
||||
%30 = OpConstant %7 1.0
|
||||
%31 = OpConstant %7 2.0
|
||||
%32 = OpConstantComposite %8 %30 %31
|
||||
%33 = OpConstant %7 3.0
|
||||
%34 = OpConstant %7 4.0
|
||||
%35 = OpConstantComposite %8 %33 %34
|
||||
%36 = OpConstantComposite %9 %32 %35
|
||||
%37 = OpConstant %3 1
|
||||
%38 = OpConstantComposite %4 %37 %37
|
||||
%39 = OpConstantComposite %8 %30 %30
|
||||
%40 = OpConstantComposite %6 %15 %15
|
||||
%41 = OpConstantComposite %10 %30 %31
|
||||
%42 = OpConstant %3 2
|
||||
%43 = OpConstantComposite %12 %37 %42
|
||||
%44 = OpConstantComposite %13 %37 %37 %37
|
||||
%45 = OpConstantComposite %14 %44
|
||||
%46 = OpConstantComposite %16 %30 %30 %30
|
||||
%47 = OpConstantComposite %17 %46
|
||||
%49 = OpTypePointer Private %4
|
||||
%48 = OpVariable %49 Private %20
|
||||
%51 = OpTypePointer Private %6
|
||||
%50 = OpVariable %51 Private %23
|
||||
%53 = OpTypePointer Private %8
|
||||
%52 = OpVariable %53 Private %26
|
||||
%54 = OpVariable %51 Private %29
|
||||
%55 = OpVariable %51 Private %29
|
||||
%56 = OpVariable %51 Private %29
|
||||
%57 = OpVariable %51 Private %29
|
||||
%59 = OpTypePointer Private %9
|
||||
%58 = OpVariable %59 Private %36
|
||||
%60 = OpVariable %59 Private %36
|
||||
%61 = OpVariable %59 Private %36
|
||||
%62 = OpVariable %59 Private %36
|
||||
%63 = OpVariable %59 Private %36
|
||||
%64 = OpVariable %49 Private %38
|
||||
%65 = OpVariable %53 Private %39
|
||||
%66 = OpVariable %49 Private %38
|
||||
%67 = OpVariable %51 Private %40
|
||||
%68 = OpVariable %53 Private %39
|
||||
%69 = OpVariable %53 Private %39
|
||||
%71 = OpTypePointer Private %10
|
||||
%70 = OpVariable %71 Private %41
|
||||
%72 = OpVariable %71 Private %41
|
||||
%74 = OpTypePointer Private %12
|
||||
%73 = OpVariable %74 Private %43
|
||||
%75 = OpVariable %71 Private %41
|
||||
%76 = OpVariable %71 Private %41
|
||||
%77 = OpVariable %71 Private %41
|
||||
%79 = OpTypePointer Private %14
|
||||
%78 = OpVariable %79 Private %45
|
||||
%81 = OpTypePointer Private %17
|
||||
%80 = OpVariable %81 Private %47
|
||||
%82 = OpVariable %81 Private %47
|
||||
%83 = OpVariable %49 Private %38
|
||||
%84 = OpVariable %53 Private %39
|
||||
%85 = OpVariable %49 Private %38
|
||||
%86 = OpVariable %51 Private %40
|
||||
%87 = OpVariable %53 Private %39
|
||||
%88 = OpVariable %53 Private %39
|
||||
%89 = OpVariable %71 Private %41
|
||||
%90 = OpVariable %71 Private %41
|
||||
%91 = OpVariable %74 Private %43
|
||||
%92 = OpVariable %71 Private %41
|
||||
%93 = OpVariable %71 Private %41
|
||||
%94 = OpVariable %71 Private %41
|
||||
%95 = OpVariable %79 Private %45
|
||||
%96 = OpVariable %79 Private %45
|
||||
%97 = OpVariable %81 Private %47
|
||||
%100 = OpTypeFunction %2
|
||||
%102 = OpTypePointer Function %4
|
||||
%104 = OpTypePointer Function %6
|
||||
%106 = OpTypePointer Function %8
|
||||
%112 = OpTypePointer Function %9
|
||||
%128 = OpTypePointer Function %10
|
||||
%133 = OpTypePointer Function %12
|
||||
%141 = OpTypePointer Function %14
|
||||
%143 = OpTypePointer Function %17
|
||||
%153 = OpTypePointer Function %5
|
||||
%154 = OpConstantNull %5
|
||||
%156 = OpTypePointer Function %3
|
||||
%157 = OpConstantNull %3
|
||||
%159 = OpTypePointer Function %7
|
||||
%160 = OpConstantNull %7
|
||||
%162 = OpConstantNull %6
|
||||
%164 = OpConstantNull %6
|
||||
%166 = OpConstantNull %6
|
||||
%168 = OpConstantNull %6
|
||||
%170 = OpConstantNull %9
|
||||
%172 = OpConstantNull %9
|
||||
%174 = OpConstantNull %9
|
||||
%176 = OpConstantNull %9
|
||||
%178 = OpConstantNull %10
|
||||
%180 = OpConstantNull %10
|
||||
%182 = OpConstantNull %10
|
||||
%184 = OpConstantNull %10
|
||||
%186 = OpConstantNull %12
|
||||
%188 = OpConstantNull %12
|
||||
%190 = OpConstantNull %10
|
||||
%192 = OpConstantNull %10
|
||||
%194 = OpConstantNull %10
|
||||
%196 = OpConstantNull %10
|
||||
%198 = OpConstantNull %12
|
||||
%200 = OpConstantNull %12
|
||||
%99 = OpFunction %2 None %100
|
||||
%98 = OpLabel
|
||||
%148 = OpVariable %128 Function %41
|
||||
%145 = OpVariable %133 Function %43
|
||||
%140 = OpVariable %141 Function %45
|
||||
%137 = OpVariable %128 Function %41
|
||||
%134 = OpVariable %133 Function %43
|
||||
%130 = OpVariable %128 Function %41
|
||||
%126 = OpVariable %106 Function %39
|
||||
%123 = OpVariable %102 Function %38
|
||||
%120 = OpVariable %112 Function %36
|
||||
%117 = OpVariable %112 Function %36
|
||||
%114 = OpVariable %112 Function %36
|
||||
%110 = OpVariable %104 Function %29
|
||||
%107 = OpVariable %104 Function %29
|
||||
%101 = OpVariable %102 Function %20
|
||||
%146 = OpVariable %128 Function %41
|
||||
%142 = OpVariable %143 Function %47
|
||||
%138 = OpVariable %128 Function %41
|
||||
%135 = OpVariable %133 Function %43
|
||||
%131 = OpVariable %128 Function %41
|
||||
%127 = OpVariable %128 Function %41
|
||||
%124 = OpVariable %104 Function %40
|
||||
%121 = OpVariable %102 Function %38
|
||||
%118 = OpVariable %112 Function %36
|
||||
%115 = OpVariable %112 Function %36
|
||||
%111 = OpVariable %112 Function %36
|
||||
%108 = OpVariable %104 Function %29
|
||||
%103 = OpVariable %104 Function %23
|
||||
%147 = OpVariable %128 Function %41
|
||||
%144 = OpVariable %143 Function %47
|
||||
%139 = OpVariable %128 Function %41
|
||||
%136 = OpVariable %128 Function %41
|
||||
%132 = OpVariable %133 Function %43
|
||||
%129 = OpVariable %128 Function %41
|
||||
%125 = OpVariable %106 Function %39
|
||||
%122 = OpVariable %106 Function %39
|
||||
%119 = OpVariable %112 Function %36
|
||||
%116 = OpVariable %112 Function %36
|
||||
%113 = OpVariable %112 Function %36
|
||||
%109 = OpVariable %104 Function %29
|
||||
%105 = OpVariable %106 Function %26
|
||||
OpBranch %149
|
||||
%149 = OpLabel
|
||||
OpStore %101 %20
|
||||
OpStore %103 %23
|
||||
OpStore %105 %26
|
||||
OpStore %107 %29
|
||||
OpStore %108 %29
|
||||
OpStore %109 %29
|
||||
OpStore %110 %29
|
||||
OpStore %111 %36
|
||||
OpStore %113 %36
|
||||
OpStore %114 %36
|
||||
OpStore %115 %36
|
||||
OpStore %116 %36
|
||||
OpStore %117 %36
|
||||
OpStore %118 %36
|
||||
OpStore %119 %36
|
||||
OpStore %120 %36
|
||||
OpStore %121 %38
|
||||
OpStore %122 %39
|
||||
OpStore %123 %38
|
||||
OpStore %124 %40
|
||||
OpStore %125 %39
|
||||
OpStore %126 %39
|
||||
OpStore %127 %41
|
||||
OpStore %129 %41
|
||||
OpStore %130 %41
|
||||
OpStore %131 %41
|
||||
OpStore %132 %43
|
||||
OpStore %134 %43
|
||||
OpStore %135 %43
|
||||
OpStore %136 %41
|
||||
OpStore %137 %41
|
||||
OpStore %138 %41
|
||||
OpStore %139 %41
|
||||
OpStore %140 %45
|
||||
OpStore %142 %47
|
||||
OpStore %144 %47
|
||||
OpStore %145 %43
|
||||
OpStore %146 %41
|
||||
OpStore %147 %41
|
||||
OpStore %148 %41
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%130 = OpFunction %2 None %84
|
||||
%129 = OpLabel
|
||||
%178 = OpVariable %117 Function %179
|
||||
%172 = OpVariable %112 Function %173
|
||||
%166 = OpVariable %117 Function %167
|
||||
%160 = OpVariable %112 Function %161
|
||||
%154 = OpVariable %96 Function %155
|
||||
%148 = OpVariable %96 Function %149
|
||||
%142 = OpVariable %88 Function %143
|
||||
%134 = OpVariable %135 Function %136
|
||||
%176 = OpVariable %117 Function %177
|
||||
%170 = OpVariable %112 Function %171
|
||||
%164 = OpVariable %117 Function %165
|
||||
%158 = OpVariable %112 Function %159
|
||||
%152 = OpVariable %96 Function %153
|
||||
%146 = OpVariable %88 Function %147
|
||||
%140 = OpVariable %88 Function %141
|
||||
%131 = OpVariable %132 Function %133
|
||||
%174 = OpVariable %112 Function %175
|
||||
%168 = OpVariable %112 Function %169
|
||||
%162 = OpVariable %112 Function %163
|
||||
%156 = OpVariable %112 Function %157
|
||||
%150 = OpVariable %96 Function %151
|
||||
%144 = OpVariable %88 Function %145
|
||||
%137 = OpVariable %138 Function %139
|
||||
OpBranch %180
|
||||
%180 = OpLabel
|
||||
%181 = OpLoad %5 %131
|
||||
%182 = OpCompositeConstruct %6 %181 %23
|
||||
OpStore %140 %182
|
||||
%183 = OpLoad %5 %131
|
||||
%184 = OpCompositeConstruct %6 %22 %183
|
||||
OpStore %142 %184
|
||||
%185 = OpLoad %5 %131
|
||||
%186 = OpCompositeConstruct %6 %185 %23
|
||||
OpStore %144 %186
|
||||
%187 = OpLoad %5 %131
|
||||
%188 = OpCompositeConstruct %6 %22 %187
|
||||
OpStore %146 %188
|
||||
%189 = OpLoad %7 %137
|
||||
%190 = OpCompositeConstruct %8 %189 %26
|
||||
%191 = OpCompositeConstruct %9 %190 %30
|
||||
OpStore %148 %191
|
||||
%192 = OpLoad %7 %137
|
||||
%193 = OpCompositeConstruct %8 %25 %192
|
||||
%194 = OpCompositeConstruct %9 %193 %30
|
||||
OpStore %150 %194
|
||||
%195 = OpLoad %7 %137
|
||||
%196 = OpCompositeConstruct %8 %195 %29
|
||||
%197 = OpCompositeConstruct %9 %27 %196
|
||||
OpStore %152 %197
|
||||
%198 = OpLoad %7 %137
|
||||
%199 = OpCompositeConstruct %8 %28 %198
|
||||
%200 = OpCompositeConstruct %9 %27 %199
|
||||
OpStore %154 %200
|
||||
%201 = OpLoad %7 %137
|
||||
%202 = OpCompositeConstruct %10 %201 %26
|
||||
OpStore %156 %202
|
||||
%203 = OpLoad %7 %137
|
||||
%204 = OpCompositeConstruct %10 %25 %203
|
||||
OpStore %158 %204
|
||||
%205 = OpLoad %7 %137
|
||||
%206 = OpCompositeConstruct %10 %205 %26
|
||||
OpStore %160 %206
|
||||
%207 = OpLoad %7 %137
|
||||
%208 = OpCompositeConstruct %10 %25 %207
|
||||
OpStore %162 %208
|
||||
%209 = OpLoad %3 %134
|
||||
%210 = OpCompositeConstruct %12 %209 %38
|
||||
OpStore %164 %210
|
||||
%211 = OpLoad %3 %134
|
||||
%212 = OpCompositeConstruct %12 %32 %211
|
||||
OpStore %166 %212
|
||||
%213 = OpLoad %7 %137
|
||||
%214 = OpCompositeConstruct %10 %213 %26
|
||||
OpStore %168 %214
|
||||
%215 = OpLoad %7 %137
|
||||
%216 = OpCompositeConstruct %10 %25 %215
|
||||
OpStore %170 %216
|
||||
%217 = OpLoad %7 %137
|
||||
%218 = OpCompositeConstruct %10 %217 %26
|
||||
OpStore %172 %218
|
||||
%219 = OpLoad %7 %137
|
||||
%220 = OpCompositeConstruct %10 %25 %219
|
||||
OpStore %174 %220
|
||||
%221 = OpLoad %3 %134
|
||||
%222 = OpCompositeConstruct %12 %221 %38
|
||||
OpStore %176 %222
|
||||
%223 = OpLoad %3 %134
|
||||
%224 = OpCompositeConstruct %12 %32 %223
|
||||
OpStore %178 %224
|
||||
%225 = OpLoad %5 %131
|
||||
%226 = OpCompositeConstruct %6 %225 %23
|
||||
OpStore %140 %226
|
||||
%227 = OpLoad %5 %131
|
||||
%228 = OpCompositeConstruct %6 %22 %227
|
||||
OpStore %142 %228
|
||||
%229 = OpLoad %5 %131
|
||||
%230 = OpCompositeConstruct %6 %229 %23
|
||||
OpStore %144 %230
|
||||
%231 = OpLoad %5 %131
|
||||
%232 = OpCompositeConstruct %6 %22 %231
|
||||
OpStore %146 %232
|
||||
%233 = OpLoad %7 %137
|
||||
%234 = OpCompositeConstruct %8 %233 %26
|
||||
%235 = OpCompositeConstruct %9 %234 %30
|
||||
OpStore %148 %235
|
||||
%236 = OpLoad %7 %137
|
||||
%237 = OpCompositeConstruct %8 %25 %236
|
||||
%238 = OpCompositeConstruct %9 %237 %30
|
||||
OpStore %150 %238
|
||||
%239 = OpLoad %7 %137
|
||||
%240 = OpCompositeConstruct %8 %239 %29
|
||||
%241 = OpCompositeConstruct %9 %27 %240
|
||||
OpStore %152 %241
|
||||
%242 = OpLoad %7 %137
|
||||
%243 = OpCompositeConstruct %8 %28 %242
|
||||
%244 = OpCompositeConstruct %9 %27 %243
|
||||
OpStore %154 %244
|
||||
%245 = OpLoad %7 %137
|
||||
%246 = OpCompositeConstruct %10 %245 %26
|
||||
OpStore %156 %246
|
||||
%247 = OpLoad %7 %137
|
||||
%248 = OpCompositeConstruct %10 %25 %247
|
||||
OpStore %158 %248
|
||||
%249 = OpLoad %7 %137
|
||||
%250 = OpCompositeConstruct %10 %249 %26
|
||||
OpStore %160 %250
|
||||
%251 = OpLoad %7 %137
|
||||
%252 = OpCompositeConstruct %10 %25 %251
|
||||
OpStore %162 %252
|
||||
%253 = OpLoad %3 %134
|
||||
%254 = OpCompositeConstruct %12 %253 %38
|
||||
OpStore %164 %254
|
||||
%255 = OpLoad %3 %134
|
||||
%256 = OpCompositeConstruct %12 %32 %255
|
||||
OpStore %166 %256
|
||||
%257 = OpLoad %7 %137
|
||||
%258 = OpCompositeConstruct %10 %257 %26
|
||||
OpStore %168 %258
|
||||
%259 = OpLoad %7 %137
|
||||
%260 = OpCompositeConstruct %10 %25 %259
|
||||
OpStore %170 %260
|
||||
%261 = OpLoad %7 %137
|
||||
%262 = OpCompositeConstruct %10 %261 %26
|
||||
OpStore %172 %262
|
||||
%263 = OpLoad %7 %137
|
||||
%264 = OpCompositeConstruct %10 %25 %263
|
||||
OpStore %174 %264
|
||||
%265 = OpLoad %3 %134
|
||||
%266 = OpCompositeConstruct %12 %265 %38
|
||||
OpStore %176 %266
|
||||
%267 = OpLoad %3 %134
|
||||
%268 = OpCompositeConstruct %12 %32 %267
|
||||
OpStore %178 %268
|
||||
%151 = OpFunction %2 None %100
|
||||
%150 = OpLabel
|
||||
%199 = OpVariable %133 Function %200
|
||||
%193 = OpVariable %128 Function %194
|
||||
%187 = OpVariable %133 Function %188
|
||||
%181 = OpVariable %128 Function %182
|
||||
%175 = OpVariable %112 Function %176
|
||||
%169 = OpVariable %112 Function %170
|
||||
%163 = OpVariable %104 Function %164
|
||||
%155 = OpVariable %156 Function %157
|
||||
%197 = OpVariable %133 Function %198
|
||||
%191 = OpVariable %128 Function %192
|
||||
%185 = OpVariable %133 Function %186
|
||||
%179 = OpVariable %128 Function %180
|
||||
%173 = OpVariable %112 Function %174
|
||||
%167 = OpVariable %104 Function %168
|
||||
%161 = OpVariable %104 Function %162
|
||||
%152 = OpVariable %153 Function %154
|
||||
%195 = OpVariable %128 Function %196
|
||||
%189 = OpVariable %128 Function %190
|
||||
%183 = OpVariable %128 Function %184
|
||||
%177 = OpVariable %128 Function %178
|
||||
%171 = OpVariable %112 Function %172
|
||||
%165 = OpVariable %104 Function %166
|
||||
%158 = OpVariable %159 Function %160
|
||||
OpBranch %201
|
||||
%201 = OpLabel
|
||||
%202 = OpLoad %5 %152
|
||||
%203 = OpCompositeConstruct %6 %202 %28
|
||||
OpStore %161 %203
|
||||
%204 = OpLoad %5 %152
|
||||
%205 = OpCompositeConstruct %6 %27 %204
|
||||
OpStore %163 %205
|
||||
%206 = OpLoad %5 %152
|
||||
%207 = OpCompositeConstruct %6 %206 %28
|
||||
OpStore %165 %207
|
||||
%208 = OpLoad %5 %152
|
||||
%209 = OpCompositeConstruct %6 %27 %208
|
||||
OpStore %167 %209
|
||||
%210 = OpLoad %7 %158
|
||||
%211 = OpCompositeConstruct %8 %210 %31
|
||||
%212 = OpCompositeConstruct %9 %211 %35
|
||||
OpStore %169 %212
|
||||
%213 = OpLoad %7 %158
|
||||
%214 = OpCompositeConstruct %8 %30 %213
|
||||
%215 = OpCompositeConstruct %9 %214 %35
|
||||
OpStore %171 %215
|
||||
%216 = OpLoad %7 %158
|
||||
%217 = OpCompositeConstruct %8 %216 %34
|
||||
%218 = OpCompositeConstruct %9 %32 %217
|
||||
OpStore %173 %218
|
||||
%219 = OpLoad %7 %158
|
||||
%220 = OpCompositeConstruct %8 %33 %219
|
||||
%221 = OpCompositeConstruct %9 %32 %220
|
||||
OpStore %175 %221
|
||||
%222 = OpLoad %7 %158
|
||||
%223 = OpCompositeConstruct %10 %222 %31
|
||||
OpStore %177 %223
|
||||
%224 = OpLoad %7 %158
|
||||
%225 = OpCompositeConstruct %10 %30 %224
|
||||
OpStore %179 %225
|
||||
%226 = OpLoad %7 %158
|
||||
%227 = OpCompositeConstruct %10 %226 %31
|
||||
OpStore %181 %227
|
||||
%228 = OpLoad %7 %158
|
||||
%229 = OpCompositeConstruct %10 %30 %228
|
||||
OpStore %183 %229
|
||||
%230 = OpLoad %3 %155
|
||||
%231 = OpCompositeConstruct %12 %230 %42
|
||||
OpStore %185 %231
|
||||
%232 = OpLoad %3 %155
|
||||
%233 = OpCompositeConstruct %12 %37 %232
|
||||
OpStore %187 %233
|
||||
%234 = OpLoad %7 %158
|
||||
%235 = OpCompositeConstruct %10 %234 %31
|
||||
OpStore %189 %235
|
||||
%236 = OpLoad %7 %158
|
||||
%237 = OpCompositeConstruct %10 %30 %236
|
||||
OpStore %191 %237
|
||||
%238 = OpLoad %7 %158
|
||||
%239 = OpCompositeConstruct %10 %238 %31
|
||||
OpStore %193 %239
|
||||
%240 = OpLoad %7 %158
|
||||
%241 = OpCompositeConstruct %10 %30 %240
|
||||
OpStore %195 %241
|
||||
%242 = OpLoad %3 %155
|
||||
%243 = OpCompositeConstruct %12 %242 %42
|
||||
OpStore %197 %243
|
||||
%244 = OpLoad %3 %155
|
||||
%245 = OpCompositeConstruct %12 %37 %244
|
||||
OpStore %199 %245
|
||||
%246 = OpLoad %5 %152
|
||||
%247 = OpCompositeConstruct %6 %246 %28
|
||||
OpStore %161 %247
|
||||
%248 = OpLoad %5 %152
|
||||
%249 = OpCompositeConstruct %6 %27 %248
|
||||
OpStore %163 %249
|
||||
%250 = OpLoad %5 %152
|
||||
%251 = OpCompositeConstruct %6 %250 %28
|
||||
OpStore %165 %251
|
||||
%252 = OpLoad %5 %152
|
||||
%253 = OpCompositeConstruct %6 %27 %252
|
||||
OpStore %167 %253
|
||||
%254 = OpLoad %7 %158
|
||||
%255 = OpCompositeConstruct %8 %254 %31
|
||||
%256 = OpCompositeConstruct %9 %255 %35
|
||||
OpStore %169 %256
|
||||
%257 = OpLoad %7 %158
|
||||
%258 = OpCompositeConstruct %8 %30 %257
|
||||
%259 = OpCompositeConstruct %9 %258 %35
|
||||
OpStore %171 %259
|
||||
%260 = OpLoad %7 %158
|
||||
%261 = OpCompositeConstruct %8 %260 %34
|
||||
%262 = OpCompositeConstruct %9 %32 %261
|
||||
OpStore %173 %262
|
||||
%263 = OpLoad %7 %158
|
||||
%264 = OpCompositeConstruct %8 %33 %263
|
||||
%265 = OpCompositeConstruct %9 %32 %264
|
||||
OpStore %175 %265
|
||||
%266 = OpLoad %7 %158
|
||||
%267 = OpCompositeConstruct %10 %266 %31
|
||||
OpStore %177 %267
|
||||
%268 = OpLoad %7 %158
|
||||
%269 = OpCompositeConstruct %10 %30 %268
|
||||
OpStore %179 %269
|
||||
%270 = OpLoad %7 %158
|
||||
%271 = OpCompositeConstruct %10 %270 %31
|
||||
OpStore %181 %271
|
||||
%272 = OpLoad %7 %158
|
||||
%273 = OpCompositeConstruct %10 %30 %272
|
||||
OpStore %183 %273
|
||||
%274 = OpLoad %3 %155
|
||||
%275 = OpCompositeConstruct %12 %274 %42
|
||||
OpStore %185 %275
|
||||
%276 = OpLoad %3 %155
|
||||
%277 = OpCompositeConstruct %12 %37 %276
|
||||
OpStore %187 %277
|
||||
%278 = OpLoad %7 %158
|
||||
%279 = OpCompositeConstruct %10 %278 %31
|
||||
OpStore %189 %279
|
||||
%280 = OpLoad %7 %158
|
||||
%281 = OpCompositeConstruct %10 %30 %280
|
||||
OpStore %191 %281
|
||||
%282 = OpLoad %7 %158
|
||||
%283 = OpCompositeConstruct %10 %282 %31
|
||||
OpStore %193 %283
|
||||
%284 = OpLoad %7 %158
|
||||
%285 = OpCompositeConstruct %10 %30 %284
|
||||
OpStore %195 %285
|
||||
%286 = OpLoad %3 %155
|
||||
%287 = OpCompositeConstruct %12 %286 %42
|
||||
OpStore %197 %287
|
||||
%288 = OpLoad %3 %155
|
||||
%289 = OpCompositeConstruct %12 %37 %288
|
||||
OpStore %199 %289
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -22,6 +22,9 @@ var<private> xafpaiai_1: array<i32, 2> = array<i32, 2>(1i, 2i);
|
||||
var<private> xafpaiaf_1: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
var<private> xafpafai_1: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
var<private> xafpafaf_1: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
var<private> xavipai_1: array<vec3<i32>, 1> = array<vec3<i32>, 1>(vec3(1i));
|
||||
var<private> xavfpai_1: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
|
||||
var<private> xavfpaf_1: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
|
||||
var<private> ivispai: vec2<i32> = vec2(1i);
|
||||
var<private> ivfspaf: vec2<f32> = vec2(1f);
|
||||
var<private> ivis_ai: vec2<i32> = vec2(1i);
|
||||
@@ -34,6 +37,9 @@ var<private> iaipaiai_1: array<i32, 2> = array<i32, 2>(1i, 2i);
|
||||
var<private> iafpafaf_1: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
var<private> iafpaiaf_1: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
var<private> iafpafai_1: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
var<private> iavipai: array<vec3<i32>, 1> = array<vec3<i32>, 1>(vec3(1i));
|
||||
var<private> iavfpai: array<vec3<i32>, 1> = array<vec3<i32>, 1>(vec3(1i));
|
||||
var<private> iavfpaf: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
|
||||
|
||||
fn all_constant_arguments() {
|
||||
var xvipaiai: vec2<i32> = vec2<i32>(42i, 43i);
|
||||
@@ -69,6 +75,9 @@ fn all_constant_arguments() {
|
||||
var xafpaiaf: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
var xafpafai: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
var xafpafaf: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
var xavipai: array<vec3<i32>, 1> = array<vec3<i32>, 1>(vec3(1i));
|
||||
var xavfpai: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
|
||||
var xavfpaf: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
|
||||
var iaipaiai: array<i32, 2> = array<i32, 2>(1i, 2i);
|
||||
var iafpaiaf: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
var iafpafai: array<f32, 2> = array<f32, 2>(1f, 2f);
|
||||
@@ -107,6 +116,9 @@ fn all_constant_arguments() {
|
||||
xafpaiaf = array<f32, 2>(1f, 2f);
|
||||
xafpafai = array<f32, 2>(1f, 2f);
|
||||
xafpafaf = array<f32, 2>(1f, 2f);
|
||||
xavipai = array<vec3<i32>, 1>(vec3(1i));
|
||||
xavfpai = array<vec3<f32>, 1>(vec3(1f));
|
||||
xavfpaf = array<vec3<f32>, 1>(vec3(1f));
|
||||
iaipaiai = array<i32, 2>(1i, 2i);
|
||||
iafpaiaf = array<f32, 2>(1f, 2f);
|
||||
iafpafai = array<f32, 2>(1f, 2f);
|
||||
|
||||
Reference in New Issue
Block a user