Merge branch 'trunk' into new-get-vertex

Cargo.toml

@@ -89,7 +89,7 @@ arrayvec = { version = "0.7.6", default-features = false }
 bincode = "1"
 bit-set = { version = "0.8", default-features = false }
 bit-vec = { version = "0.8", default-features = false }
-bitflags = "2.7"
+bitflags = "2.9"
 bytemuck = { version = "1.21", features = [
     "derive",
     "extern_crate_alloc",

naga/src/front/wgsl/lower/mod.rs

@@ -2035,10 +2035,18 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                     }
                 }
 
-                lowered_base.map(|base| match ctx.const_eval_expr_to_u32(index).ok() {
-                    Some(index) => crate::Expression::AccessIndex { base, index },
-                    None => crate::Expression::Access { base, index },
-                })
+                lowered_base.try_map(|base| match ctx.const_eval_expr_to_u32(index).ok() {
+                    Some(index) => Ok::<_, Error>(crate::Expression::AccessIndex { base, index }),
+                    None => {
+                        // When an abstract array value e is indexed by an expression
+                        // that is not a const-expression, then the array is concretized
+                        // before the index is applied.
+                        // https://www.w3.org/TR/WGSL/#array-access-expr
+                        // Also applies to vectors and matrices.
+                        let base = ctx.concretize(base)?;
+                        Ok(crate::Expression::Access { base, index })
+                    }
+                })?
             }
             ast::Expression::Member { base, ref field } => {
                 let mut lowered_base = self.expression_for_reference(base, ctx)?;

naga/tests/in/wgsl/const-exprs.wgsl

@@ -115,3 +115,18 @@ fn test_local_const() {
     const local_const = 2;
 	var arr: array<f32, local_const>;
 }
+
+const ABSTRACT_ARRAY = array(1, 2, 3, 4, 5, 6, 7, 8, 9);
+const ABSTRACT_VECTOR = vec4(1, 2, 3, 4);
+
+fn abstract_access(i: u32) {
+    // Constant indexing of abstract types is allowed, therefore we can assign
+    // to f32 or u32 vars just fine.
+    var a: f32 = ABSTRACT_ARRAY[0];
+    var b: u32 = ABSTRACT_VECTOR.x;
+
+    // For non constant indices the base type is concretized prior to indexing,
+    // therefore we can only assign to i32 in this case.
+    var c: i32 = ABSTRACT_ARRAY[i];
+    var d: i32 = ABSTRACT_VECTOR[i];
+}

naga/tests/out/glsl/const-exprs.main.Compute.glsl

@@ -112,6 +112,16 @@ void relational() {
     return;
 }
 
+void abstract_access(uint i) {
+    float a_1 = 1.0;
+    uint b_1 = 1u;
+    int c_1 = 0;
+    int d = 0;
+    c_1 = int[9](1, 2, 3, 4, 5, 6, 7, 8, 9)[i];
+    d = ivec4(1, 2, 3, 4)[i];
+    return;
+}
+
 void main() {
     swizzle_of_compose();
     index_of_compose();

naga/tests/out/hlsl/const-exprs.hlsl

@@ -125,6 +125,24 @@ void relational()
     return;
 }
 
+typedef int ret_Constructarray9_int_[9];
+ret_Constructarray9_int_ Constructarray9_int_(int arg0, int arg1, int arg2, int arg3, int arg4, int arg5, int arg6, int arg7, int arg8) {
+    int ret[9] = { arg0, arg1, arg2, arg3, arg4, arg5, arg6, arg7, arg8 };
+    return ret;
+}
+
+void abstract_access(uint i)
+{
+    float a_1 = 1.0;
+    uint b_1 = 1u;
+    int c_1 = (int)0;
+    int d = (int)0;
+
+    c_1 = Constructarray9_int_(int(1), int(2), int(3), int(4), int(5), int(6), int(7), int(8), int(9))[min(uint(i), 8u)];
+    d = int4(int(1), int(2), int(3), int(4))[min(uint(i), 3u)];
+    return;
+}
+
 [numthreads(2, 3, 1)]
 void main()
 {

naga/tests/out/msl/const-exprs.msl

@@ -7,6 +7,9 @@ using metal::uint;
 struct type_6 {
     float inner[2];
 };
+struct type_10 {
+    int inner[9];
+};
 constant uint TWO = 2u;
 constant int THREE = 3;
 constant bool TRUE = true;
@@ -125,6 +128,18 @@ void relational(
     return;
 }
 
+void abstract_access(
+    uint i
+) {
+    float a_1 = 1.0;
+    uint b_1 = 1u;
+    int c_1 = {};
+    int d = {};
+    c_1 = type_10 {1, 2, 3, 4, 5, 6, 7, 8, 9}.inner[i];
+    d = metal::int4(1, 2, 3, 4)[i];
+    return;
+}
+
 kernel void main_(
 ) {
     swizzle_of_compose();