Skip to content

Fix constant evaluation of indexing by an i32 override #8814

New issue

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

Sign up for GitHub

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

Already on GitHub? Sign in to your account

Merged
teoxoy merged 3 commits into from
Jan 8, 2026
Merged

Fix constant evaluation of indexing by an i32 override #8814

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
c91edea
refactor(naga): Rename U32EvalError to ConstValueError
andyleiserson Jan 5, 2026
3b9db6f
refactor(naga): Consolidate some redundant constant "evaluation" helpers
andyleiserson Jan 5, 2026
5694172
fix(naga): Fix constant evaluation of indexing by an i32 override
andyleiserson Jan 5, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions cts_runner/test.lst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -209,6 +209,7 @@ webgpu:shader,execution,flow_control,return:*
// Fails on Metal in CI only, not when running locally.
fails-if(metal) webgpu:shader,execution,robust_access_vertex:vertex_buffer_access:indexed=true;indirect=false;drawCallTestParameter="baseVertex";type="float32x4";additionalBuffers=4;partialLastNumber=false;offsetVertexBuffer=true
webgpu:shader,validation,const_assert,const_assert:*
webgpu:shader,validation,expression,access,array:early_eval_errors:case="override_in_bounds"
webgpu:shader,validation,expression,binary,short_circuiting_and_or:array_override:op="%26%26";a_val=1;b_val=1
webgpu:shader,validation,expression,binary,short_circuiting_and_or:invalid_types:*
webgpu:shader,validation,expression,binary,short_circuiting_and_or:scalar_vector:op="%26%26";lhs="bool";rhs="bool"
Expand Down
6 changes: 3 additions & 3 deletions naga/src/back/pipeline_constants.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -280,7 +280,7 @@ fn process_workgroup_size_override(
Some(h) => {
ep.workgroup_size[i] = module
.to_ctx()
.eval_expr_to_u32(adjusted_global_expressions[h])
.get_const_val(adjusted_global_expressions[h])
.map(|n| {
if n == 0 {
Err(PipelineConstantError::NegativeWorkgroupSize)
Expand Down Expand Up @@ -309,13 +309,13 @@ fn process_mesh_shader_overrides(
if let Some(r#override) = mesh_info.max_vertices_override {
mesh_info.max_vertices = module
.to_ctx()
.eval_expr_to_u32(adjusted_global_expressions[r#override])
.get_const_val(adjusted_global_expressions[r#override])
.map_err(|_| PipelineConstantError::NegativeMeshOutputMax)?;
}
if let Some(r#override) = mesh_info.max_primitives_override {
mesh_info.max_primitives = module
.to_ctx()
.eval_expr_to_u32(adjusted_global_expressions[r#override])
.get_const_val(adjusted_global_expressions[r#override])
.map_err(|_| PipelineConstantError::NegativeMeshOutputMax)?;
}
}
Expand Down
2 changes: 1 addition & 1 deletion naga/src/front/glsl/context.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -561,7 +561,7 @@ impl<'a> Context<'a> {
_ => self
.module
.to_ctx()
.eval_expr_to_u32_from(index, &self.expressions)
.get_const_val_from(index, &self.expressions)
.ok(),
};

Expand Down
8 changes: 4 additions & 4 deletions naga/src/front/glsl/parser.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@ use super::{
variables::{GlobalOrConstant, VarDeclaration},
Frontend, Result,
};
use crate::{arena::Handle, proc::U32EvalError, Expression, Module, Span, Type};
use crate::{arena::Handle, proc::ConstValueError, Expression, Module, Span, Type};

mod declarations;
mod expressions;
Expand Down Expand Up @@ -211,15 +211,15 @@ impl<'source> ParsingContext<'source> {
ctx.global_expression_kind_tracker,
)?;

let res = ctx.module.to_ctx().eval_expr_to_u32(const_expr);
let res = ctx.module.to_ctx().get_const_val(const_expr);

let int = match res {
Ok(value) => Ok(value),
Err(U32EvalError::Negative) => Err(Error {
Err(ConstValueError::Negative) => Err(Error {
kind: ErrorKind::SemanticError("int constant overflows".into()),
meta,
}),
Err(U32EvalError::NonConst) => Err(Error {
Err(ConstValueError::NonConst | ConstValueError::InvalidType) => Err(Error {
kind: ErrorKind::SemanticError("Expected a uint constant".into()),
meta,
}),
Expand Down
2 changes: 1 addition & 1 deletion naga/src/front/spv/next_block.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -277,7 +277,7 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
let index_maybe = match *index_expr_data {
crate::Expression::Constant(const_handle) => Some(
ctx.gctx()
.eval_expr_to_u32(ctx.module.constants[const_handle].init)
.get_const_val(ctx.module.constants[const_handle].init)
.map_err(|_| {
Error::InvalidAccess(crate::Expression::Constant(
const_handle,
Expand Down
95 changes: 40 additions & 55 deletions naga/src/front/wgsl/lower/mod.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -535,50 +535,28 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
.map_err(|e| Box::new(Error::ConstantEvaluatorError(e.into(), span)))
}

fn const_eval_expr_to_u32(
fn get_const_val<T: TryFrom<crate::Literal, Error = proc::ConstValueError>>(
&self,
handle: Handle<ir::Expression>,
) -> core::result::Result<u32, proc::U32EvalError> {
) -> core::result::Result<T, proc::ConstValueError> {
match self.expr_type {
ExpressionContextType::Runtime(ref ctx) => {
if !ctx.local_expression_kind_tracker.is_const(handle) {
return Err(proc::U32EvalError::NonConst);
return Err(proc::ConstValueError::NonConst);
}

self.module
.to_ctx()
.eval_expr_to_u32_from(handle, &ctx.function.expressions)
.get_const_val_from(handle, &ctx.function.expressions)
}
ExpressionContextType::Constant(Some(ref ctx)) => {
assert!(ctx.local_expression_kind_tracker.is_const(handle));
self.module
.to_ctx()
.eval_expr_to_u32_from(handle, &ctx.function.expressions)
.get_const_val_from(handle, &ctx.function.expressions)
}
ExpressionContextType::Constant(None) => self.module.to_ctx().eval_expr_to_u32(handle),
ExpressionContextType::Override => Err(proc::U32EvalError::NonConst),
}
}

fn const_eval_expr_to_bool(&self, handle: Handle<ir::Expression>) -> Option<bool> {
match self.expr_type {
ExpressionContextType::Runtime(ref ctx) => {
if !ctx.local_expression_kind_tracker.is_const(handle) {
return None;
}

self.module
.to_ctx()
.eval_expr_to_bool_from(handle, &ctx.function.expressions)
}
ExpressionContextType::Constant(Some(ref ctx)) => {
assert!(ctx.local_expression_kind_tracker.is_const(handle));
self.module
.to_ctx()
.eval_expr_to_bool_from(handle, &ctx.function.expressions)
}
ExpressionContextType::Constant(None) => self.module.to_ctx().eval_expr_to_bool(handle),
ExpressionContextType::Override => None,
ExpressionContextType::Constant(None) => self.module.to_ctx().get_const_val(handle),
ExpressionContextType::Override => Err(proc::ConstValueError::NonConst),
}
}

Expand Down Expand Up @@ -716,12 +694,14 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
let index = self
.module
.to_ctx()
.eval_expr_to_u32_from(expr, &rctx.function.expressions)
.get_const_val_from::<u32, _>(expr, &rctx.function.expressions)
.map_err(|err| match err {
proc::U32EvalError::NonConst => {
proc::ConstValueError::NonConst | proc::ConstValueError::InvalidType => {
Error::ExpectedConstExprConcreteIntegerScalar(component_span)
}
proc::U32EvalError::Negative => Error::ExpectedNonNegative(component_span),
proc::ConstValueError::Negative => {
Error::ExpectedNonNegative(component_span)
}
})?;
ir::SwizzleComponent::XYZW
.get(index as usize)
Expand Down Expand Up @@ -1417,11 +1397,14 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
match ctx
.module
.to_ctx()
.eval_expr_to_bool_from(condition, &ctx.module.global_expressions)
.get_const_val_from(condition, &ctx.module.global_expressions)
{
Some(true) => Ok(()),
Some(false) => Err(Error::ConstAssertFailed(span)),
_ => Err(Error::NotBool(span)),
Ok(true) => Ok(()),
Ok(false) => Err(Error::ConstAssertFailed(span)),
Err(proc::ConstValueError::NonConst | proc::ConstValueError::Negative) => {
unreachable!()
}
Err(proc::ConstValueError::InvalidType) => Err(Error::NotBool(span)),
}?;
}
}
Expand Down Expand Up @@ -1940,14 +1923,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
match ctx
.module
.to_ctx()
.eval_expr_to_literal_from(expr, &ctx.function.expressions)
.get_const_val_from(expr, &ctx.function.expressions)
{
Some(ir::Literal::I32(value)) => {
ir::SwitchValue::I32(value)
}
Some(ir::Literal::U32(value)) => {
ir::SwitchValue::U32(value)
}
Ok(ir::Literal::I32(value)) => ir::SwitchValue::I32(value),
Ok(ir::Literal::U32(value)) => ir::SwitchValue::U32(value),
_ => {
return Err(Box::new(Error::InvalidSwitchCase {
span,
Expand Down Expand Up @@ -2168,11 +2147,14 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
match ctx
.module
.to_ctx()
.eval_expr_to_bool_from(condition, &ctx.function.expressions)
.get_const_val_from(condition, &ctx.function.expressions)
{
Some(true) => Ok(()),
Some(false) => Err(Error::ConstAssertFailed(span)),
_ => Err(Error::NotBool(span)),
Ok(true) => Ok(()),
Ok(false) => Err(Error::ConstAssertFailed(span)),
Err(proc::ConstValueError::NonConst | proc::ConstValueError::Negative) => {
unreachable!()
}
Err(proc::ConstValueError::InvalidType) => Err(Error::NotBool(span)),
}?;

block.extend(emitter.finish(&ctx.function.expressions));
Expand Down Expand Up @@ -2370,7 +2352,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
}
}

lowered_base.try_map(|base| match ctx.const_eval_expr_to_u32(index).ok() {
lowered_base.try_map(|base| match ctx.get_const_val(index).ok() {
Some(index) => Ok::<_, Box<Error>>(ir::Expression::AccessIndex { base, index }),
None => {
// When an abstract array value e is indexed by an expression
Expand Down Expand Up @@ -2565,7 +2547,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
result_var,
)))
} else {
let left_val = ctx.const_eval_expr_to_bool(left);
let left_val: Option<bool> = ctx.get_const_val(left).ok();

if left_val.is_some_and(|left_val| {
op == crate::BinaryOperator::LogicalAnd && !left_val
Expand Down Expand Up @@ -4201,10 +4183,12 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
let value = ctx
.module
.to_ctx()
.eval_expr_to_u32(expr)
.get_const_val(expr)
.map_err(|err| match err {
proc::U32EvalError::NonConst => Error::ExpectedConstExprConcreteIntegerScalar(span),
proc::U32EvalError::Negative => Error::ExpectedNonNegative(span),
proc::ConstValueError::NonConst | proc::ConstValueError::InvalidType => {
Error::ExpectedConstExprConcreteIntegerScalar(span)
}
proc::ConstValueError::Negative => Error::ExpectedNonNegative(span),
})?;
Ok((value, span))
}
Expand All @@ -4220,12 +4204,13 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
let const_expr = self.expression(expr, &mut ctx.as_const());
match const_expr {
Ok(value) => {
let len = ctx.const_eval_expr_to_u32(value).map_err(|err| {
let len = ctx.get_const_val(value).map_err(|err| {
Box::new(match err {
proc::U32EvalError::NonConst => {
proc::ConstValueError::NonConst
| proc::ConstValueError::InvalidType => {
Error::ExpectedConstExprConcreteIntegerScalar(span)
}
proc::U32EvalError::Negative => {
proc::ConstValueError::Negative => {
Error::ExpectedPositiveArrayLength(span)
}
})
Expand Down
24 changes: 5 additions & 19 deletions naga/src/proc/constant_evaluator.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1268,7 +1268,11 @@ impl<'a> ConstantEvaluator<'a> {
let base = self.check_and_get(base)?;
let index = self.check_and_get(index)?;

self.access(base, self.constant_index(index)?, span)
let index_val: u32 = self
.to_ctx()
.get_const_val_from(index, self.expressions)
.map_err(|_| ConstantEvaluatorError::InvalidAccessIndexTy)?;
self.access(base, index_val as usize, span)
}
Expression::Swizzle {
size,
Expand Down Expand Up @@ -2116,24 +2120,6 @@ impl<'a> ConstantEvaluator<'a> {
}
}

fn constant_index(&self, expr: Handle<Expression>) -> Result<usize, ConstantEvaluatorError> {
match self.expressions[expr] {
Expression::ZeroValue(ty)
if matches!(
self.types[ty].inner,
TypeInner::Scalar(crate::Scalar {
kind: ScalarKind::Uint,
..
})
) =>
{
Ok(0)
}
Expression::Literal(Literal::U32(index)) => Ok(index as usize),
_ => Err(ConstantEvaluatorError::InvalidAccessIndexTy),
}
}

/// Lower [`ZeroValue`] and [`Splat`] expressions to [`Literal`] and [`Compose`] expressions.
///
/// [`ZeroValue`]: Expression::ZeroValue
Expand Down
2 changes: 1 addition & 1 deletion naga/src/proc/index.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -483,7 +483,7 @@ impl GuardedIndex {
expressions: &crate::Arena<crate::Expression>,
module: &crate::Module,
) -> Self {
match module.to_ctx().eval_expr_to_u32_from(expr, expressions) {
match module.to_ctx().get_const_val_from(expr, expressions) {
Ok(value) => Self::Known(value),
Err(_) => Self::Expression(expr),
}
Expand Down
Loading
Loading