diff --git a/naga/src/back/hlsl/mod.rs b/naga/src/back/hlsl/mod.rs index dd9a5bf6e5..7f55398040 100644 --- a/naga/src/back/hlsl/mod.rs +++ b/naga/src/back/hlsl/mod.rs @@ -13,11 +13,17 @@ type should be stored in `uniform` and `storage` buffers. The HLSL we generate must access values in that form, even when it is not what HLSL would use normally. -The rules described here only apply to WGSL `uniform` variables. WGSL -`storage` buffers are translated as HLSL `ByteAddressBuffers`, for -which we generate `Load` and `Store` method calls with explicit byte -offsets. WGSL pipeline inputs must be scalars or vectors; they cannot -be matrices, which is where the interesting problems arise. +Matching the WGSL memory layout is a concern only for `uniform` +variables. WGSL `storage` buffers are translated as HLSL +`ByteAddressBuffers`, for which we generate `Load` and `Store` method +calls with explicit byte offsets. WGSL pipeline inputs must be scalars +or vectors; they cannot be matrices, which is where the interesting +problems arise. However, when an affected type appears in a struct +definition, the transformations described here are applied without +consideration of where the struct is used. + +Access to storage buffers is implemented in `storage.rs`. Access to +uniform buffers is implemented where applicable in `writer.rs`. ## Row- and column-major ordering for matrices @@ -57,10 +63,9 @@ that the columns of a `matKx2` need only be [aligned as required for `vec2`][ilov], which is [eight-byte alignment][8bb]. To compensate for this, any time a `matKx2` appears in a WGSL -`uniform` variable, whether directly as the variable's type or as part -of a struct/array, we actually emit `K` separate `float2` members, and -assemble/disassemble the matrix from its columns (in WGSL; rows in -HLSL) upon load and store. +`uniform` value or as part of a struct/array, we actually emit `K` +separate `float2` members, and assemble/disassemble the matrix from its +columns (in WGSL; rows in HLSL) upon load and store. For example, the following WGSL struct type: diff --git a/naga/src/back/hlsl/storage.rs b/naga/src/back/hlsl/storage.rs index 9f92d86639..5d66af0008 100644 --- a/naga/src/back/hlsl/storage.rs +++ b/naga/src/back/hlsl/storage.rs @@ -108,6 +108,13 @@ pub(super) enum StoreValue { base: Handle, member_index: u32, }, + // Access to a single column of a Cx2 matrix within a struct + TempColumnAccess { + depth: usize, + base: Handle, + member_index: u32, + column: u32, + }, } impl super::Writer<'_, W> { @@ -290,6 +297,15 @@ impl super::Writer<'_, W> { let name = &self.names[&NameKey::StructMember(base, member_index)]; write!(self.out, "{STORE_TEMP_NAME}{depth}.{name}")? } + StoreValue::TempColumnAccess { + depth, + base, + member_index, + column, + } => { + let name = &self.names[&NameKey::StructMember(base, member_index)]; + write!(self.out, "{STORE_TEMP_NAME}{depth}.{name}_{column}")? + } } Ok(()) } @@ -302,6 +318,7 @@ impl super::Writer<'_, W> { value: StoreValue, func_ctx: &FunctionCtx, level: crate::back::Level, + within_struct: Option>, ) -> BackendResult { let temp_resolution; let ty_resolution = match value { @@ -325,6 +342,9 @@ impl super::Writer<'_, W> { temp_resolution = TypeResolution::Handle(ty_handle); &temp_resolution } + StoreValue::TempColumnAccess { .. } => { + unreachable!("attempting write_storage_store for TempColumnAccess"); + } }; match *ty_resolution.inner_with(&module.types) { crate::TypeInner::Scalar(scalar) => { @@ -372,37 +392,92 @@ impl super::Writer<'_, W> { rows, scalar, } => { - // first, assign the value to a temporary - writeln!(self.out, "{level}{{")?; - let depth = level.0 + 1; - write!( - self.out, - "{}{}{}x{} {}{} = ", - level.next(), - scalar.to_hlsl_str()?, - columns as u8, - rows as u8, - STORE_TEMP_NAME, - depth, - )?; - self.write_store_value(module, &value, func_ctx)?; - writeln!(self.out, ";")?; - // Note: Matrices containing vec3s, due to padding, act like they contain vec4s. let row_stride = Alignment::from(rows) * scalar.width as u32; - // then iterate the stores - for i in 0..columns as u32 { - self.temp_access_chain - .push(SubAccess::Offset(i * row_stride)); - let ty_inner = crate::TypeInner::Vector { size: rows, scalar }; - let sv = StoreValue::TempIndex { - depth, - index: i, - ty: TypeResolution::Value(ty_inner), - }; - self.write_storage_store(module, var_handle, sv, func_ctx, level.next())?; - self.temp_access_chain.pop(); + writeln!(self.out, "{level}{{")?; + + match within_struct { + Some(containing_struct) if rows == crate::VectorSize::Bi => { + // If we are within a struct, then the struct was already assigned to + // a temporary, we don't need to make another. + let mut chain = mem::take(&mut self.temp_access_chain); + for i in 0..columns as u32 { + chain.push(SubAccess::Offset(i * row_stride)); + // working around the borrow checker in `self.write_expr` + let var_name = &self.names[&NameKey::GlobalVariable(var_handle)]; + let StoreValue::TempAccess { member_index, .. } = value else { + unreachable!( + "write_storage_store within_struct but not TempAccess" + ); + }; + let column_value = StoreValue::TempColumnAccess { + depth: level.0, // note not incrementing, b/c no temp + base: containing_struct, + member_index, + column: i, + }; + // See note about DXC and Load/Store in the module's documentation. + if scalar.width == 4 { + write!( + self.out, + "{}{}.Store{}(", + level.next(), + var_name, + rows as u8 + )?; + self.write_storage_address(module, &chain, func_ctx)?; + write!(self.out, ", asuint(")?; + self.write_store_value(module, &column_value, func_ctx)?; + writeln!(self.out, "));")?; + } else { + write!(self.out, "{}{var_name}.Store(", level.next())?; + self.write_storage_address(module, &chain, func_ctx)?; + write!(self.out, ", ")?; + self.write_store_value(module, &column_value, func_ctx)?; + writeln!(self.out, ");")?; + } + chain.pop(); + } + self.temp_access_chain = chain; + } + _ => { + // first, assign the value to a temporary + let depth = level.0 + 1; + write!( + self.out, + "{}{}{}x{} {}{} = ", + level.next(), + scalar.to_hlsl_str()?, + columns as u8, + rows as u8, + STORE_TEMP_NAME, + depth, + )?; + self.write_store_value(module, &value, func_ctx)?; + writeln!(self.out, ";")?; + + // then iterate the stores + for i in 0..columns as u32 { + self.temp_access_chain + .push(SubAccess::Offset(i * row_stride)); + let ty_inner = crate::TypeInner::Vector { size: rows, scalar }; + let sv = StoreValue::TempIndex { + depth, + index: i, + ty: TypeResolution::Value(ty_inner), + }; + self.write_storage_store( + module, + var_handle, + sv, + func_ctx, + level.next(), + None, + )?; + self.temp_access_chain.pop(); + } + } } // done writeln!(self.out, "{level}}}")?; @@ -415,7 +490,7 @@ impl super::Writer<'_, W> { // first, assign the value to a temporary writeln!(self.out, "{level}{{")?; write!(self.out, "{}", level.next())?; - self.write_value_type(module, &module.types[base].inner)?; + self.write_type(module, base)?; let depth = level.next().0; write!(self.out, " {STORE_TEMP_NAME}{depth}")?; self.write_array_size(module, base, crate::ArraySize::Constant(size))?; @@ -430,7 +505,7 @@ impl super::Writer<'_, W> { index: i, ty: TypeResolution::Handle(base), }; - self.write_storage_store(module, var_handle, sv, func_ctx, level.next())?; + self.write_storage_store(module, var_handle, sv, func_ctx, level.next(), None)?; self.temp_access_chain.pop(); } // done @@ -461,7 +536,14 @@ impl super::Writer<'_, W> { base: struct_ty, member_index: i as u32, }; - self.write_storage_store(module, var_handle, sv, func_ctx, level.next())?; + self.write_storage_store( + module, + var_handle, + sv, + func_ctx, + level.next(), + Some(struct_ty), + )?; self.temp_access_chain.pop(); } // done diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index 2096aff0f9..75e890d272 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -45,6 +45,11 @@ pub(crate) const F2U64_FUNCTION: &str = "naga_f2u64"; pub(crate) const IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION: &str = "nagaTextureSampleBaseClampToEdge"; +enum Index { + Expression(Handle), + Static(u32), +} + struct EpStructMember { name: String, ty: Handle, @@ -1797,6 +1802,23 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Ok(()) } + fn write_index( + &mut self, + module: &Module, + index: Index, + func_ctx: &back::FunctionCtx<'_>, + ) -> BackendResult { + match index { + Index::Static(index) => { + write!(self.out, "{index}")?; + } + Index::Expression(index) => { + self.write_expr(module, index, func_ctx)?; + } + } + Ok(()) + } + /// Helper method used to write statements /// /// # Notes @@ -1945,6 +1967,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { StoreValue::Expression(value), func_ctx, level, + None, )?; } else { // We treat matrices of the form `matCx2` as a sequence of C `vec2`s. @@ -1952,13 +1975,15 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // // We handle matrix Stores here directly (including sub accesses for Vectors and Scalars). // Loads are handled by `Expression::AccessIndex` (since sub accesses work fine for Loads). - struct MatrixAccess { - base: Handle, - index: u32, - } - enum Index { - Expression(Handle), - Static(u32), + enum MatrixAccess { + Direct { + base: Handle, + index: u32, + }, + Struct { + columns: crate::VectorSize, + base: Handle, + }, } let get_members = |expr: Handle| { @@ -1972,187 +1997,28 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } }; - let mut matrix = None; - let mut vector = None; - let mut scalar = None; - - let mut current_expr = pointer; - for _ in 0..3 { - let resolved = func_ctx.resolve_type(current_expr, &module.types); - - match (resolved, &func_ctx.expressions[current_expr]) { - ( - &TypeInner::Pointer { base: ty, .. }, - &crate::Expression::AccessIndex { base, index }, - ) if matches!( - module.types[ty].inner, - TypeInner::Matrix { - rows: crate::VectorSize::Bi, - .. - } - ) && get_members(base) - .map(|members| members[index as usize].binding.is_none()) - == Some(true) => - { - matrix = Some(MatrixAccess { base, index }); - break; - } - ( - &TypeInner::ValuePointer { - size: Some(crate::VectorSize::Bi), - .. - }, - &crate::Expression::Access { base, index }, - ) => { - vector = Some(Index::Expression(index)); - current_expr = base; - } - ( - &TypeInner::ValuePointer { - size: Some(crate::VectorSize::Bi), - .. - }, - &crate::Expression::AccessIndex { base, index }, - ) => { - vector = Some(Index::Static(index)); - current_expr = base; - } - ( - &TypeInner::ValuePointer { size: None, .. }, - &crate::Expression::Access { base, index }, - ) => { - scalar = Some(Index::Expression(index)); - current_expr = base; - } - ( - &TypeInner::ValuePointer { size: None, .. }, - &crate::Expression::AccessIndex { base, index }, - ) => { - scalar = Some(Index::Static(index)); - current_expr = base; - } - _ => break, - } - } - write!(self.out, "{level}")?; - if let Some(MatrixAccess { index, base }) = matrix { - let base_ty_res = &func_ctx.info[base].ty; - let resolved = base_ty_res.inner_with(&module.types); - let ty = match *resolved { - TypeInner::Pointer { base, .. } => base, - _ => base_ty_res.handle().unwrap(), - }; - - if let Some(Index::Static(vec_index)) = vector { - self.write_expr(module, base, func_ctx)?; - write!( - self.out, - ".{}_{}", - &self.names[&NameKey::StructMember(ty, index)], - vec_index - )?; - - if let Some(scalar_index) = scalar { - write!(self.out, "[")?; - match scalar_index { - Index::Static(index) => { - write!(self.out, "{index}")?; - } - Index::Expression(index) => { - self.write_expr(module, index, func_ctx)?; - } - } - write!(self.out, "]")?; - } - - write!(self.out, " = ")?; - self.write_expr(module, value, func_ctx)?; - writeln!(self.out, ";")?; - } else { - let access = WrappedStructMatrixAccess { ty, index }; - match (&vector, &scalar) { - (&Some(_), &Some(_)) => { - self.write_wrapped_struct_matrix_set_scalar_function_name( - access, - )?; - } - (&Some(_), &None) => { - self.write_wrapped_struct_matrix_set_vec_function_name(access)?; - } - (&None, _) => { - self.write_wrapped_struct_matrix_set_function_name(access)?; - } - } - - write!(self.out, "(")?; - self.write_expr(module, base, func_ctx)?; - write!(self.out, ", ")?; - self.write_expr(module, value, func_ctx)?; - - if let Some(Index::Expression(vec_index)) = vector { - write!(self.out, ", ")?; - self.write_expr(module, vec_index, func_ctx)?; - - if let Some(scalar_index) = scalar { - write!(self.out, ", ")?; - match scalar_index { - Index::Static(index) => { - write!(self.out, "{index}")?; - } - Index::Expression(index) => { - self.write_expr(module, index, func_ctx)?; - } - } - } - } - writeln!(self.out, ");")?; - } - } else { - // We handle `Store`s to __matCx2 column vectors and scalar elements via - // the previously injected functions __set_col_of_matCx2 / __set_el_of_matCx2. - struct MatrixData { - columns: crate::VectorSize, - base: Handle, - } - - enum Index { - Expression(Handle), - Static(u32), - } - - let mut matrix = None; - let mut vector = None; - let mut scalar = None; - - let mut current_expr = pointer; - for _ in 0..3 { - let resolved = func_ctx.resolve_type(current_expr, &module.types); - match (resolved, &func_ctx.expressions[current_expr]) { - ( - &TypeInner::ValuePointer { - size: Some(crate::VectorSize::Bi), - .. - }, - &crate::Expression::Access { base, index }, - ) => { - vector = Some(index); - current_expr = base; - } - ( - &TypeInner::ValuePointer { size: None, .. }, - &crate::Expression::Access { base, index }, - ) => { - scalar = Some(Index::Expression(index)); - current_expr = base; - } + let matrix_access_on_lhs = + find_matrix_in_access_chain(module, pointer, func_ctx).and_then( + |(matrix_expr, vector, scalar)| match ( + func_ctx.resolve_type(matrix_expr, &module.types), + &func_ctx.expressions[matrix_expr], + ) { ( - &TypeInner::ValuePointer { size: None, .. }, + &TypeInner::Pointer { base: ty, .. }, &crate::Expression::AccessIndex { base, index }, - ) => { - scalar = Some(Index::Static(index)); - current_expr = base; + ) if matches!( + module.types[ty].inner, + TypeInner::Matrix { + rows: crate::VectorSize::Bi, + .. + } + ) && get_members(base) + .map(|members| members[index as usize].binding.is_none()) + == Some(true) => + { + Some((MatrixAccess::Direct { base, index }, vector, scalar)) } _ => { if let Some(MatrixType { @@ -2161,24 +2027,95 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { width: 4, }) = get_inner_matrix_of_struct_array_member( module, - current_expr, + matrix_expr, func_ctx, true, ) { - matrix = Some(MatrixData { - columns, - base: current_expr, - }); + Some(( + MatrixAccess::Struct { + columns, + base: matrix_expr, + }, + vector, + scalar, + )) + } else { + None + } + } + }, + ); + + match matrix_access_on_lhs { + Some((MatrixAccess::Direct { index, base }, vector, scalar)) => { + let base_ty_res = &func_ctx.info[base].ty; + let resolved = base_ty_res.inner_with(&module.types); + let ty = match *resolved { + TypeInner::Pointer { base, .. } => base, + _ => base_ty_res.handle().unwrap(), + }; + + if let Some(Index::Static(vec_index)) = vector { + self.write_expr(module, base, func_ctx)?; + write!( + self.out, + ".{}_{}", + &self.names[&NameKey::StructMember(ty, index)], + vec_index + )?; + + if let Some(scalar_index) = scalar { + write!(self.out, "[")?; + self.write_index(module, scalar_index, func_ctx)?; + write!(self.out, "]")?; + } + + write!(self.out, " = ")?; + self.write_expr(module, value, func_ctx)?; + writeln!(self.out, ";")?; + } else { + let access = WrappedStructMatrixAccess { ty, index }; + match (&vector, &scalar) { + (&Some(_), &Some(_)) => { + self.write_wrapped_struct_matrix_set_scalar_function_name( + access, + )?; + } + (&Some(_), &None) => { + self.write_wrapped_struct_matrix_set_vec_function_name( + access, + )?; } + (&None, _) => { + self.write_wrapped_struct_matrix_set_function_name(access)?; + } + } + + write!(self.out, "(")?; + self.write_expr(module, base, func_ctx)?; + write!(self.out, ", ")?; + self.write_expr(module, value, func_ctx)?; - break; + if let Some(Index::Expression(vec_index)) = vector { + write!(self.out, ", ")?; + self.write_expr(module, vec_index, func_ctx)?; + + if let Some(scalar_index) = scalar { + write!(self.out, ", ")?; + self.write_index(module, scalar_index, func_ctx)?; + } } + writeln!(self.out, ");")?; } } + Some(( + MatrixAccess::Struct { columns, base }, + Some(Index::Expression(vec_index)), + scalar, + )) => { + // We handle `Store`s to __matCx2 column vectors and scalar elements via + // the previously injected functions __set_col_of_matCx2 / __set_el_of_matCx2. - if let (Some(MatrixData { columns, base }), Some(vec_index)) = - (matrix, vector) - { if scalar.is_some() { write!(self.out, "__set_el_of_mat{}x2", columns as u8)?; } else { @@ -2191,21 +2128,17 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { if let Some(scalar_index) = scalar { write!(self.out, ", ")?; - match scalar_index { - Index::Static(index) => { - write!(self.out, "{index}")?; - } - Index::Expression(index) => { - self.write_expr(module, index, func_ctx)?; - } - } + self.write_index(module, scalar_index, func_ctx)?; } write!(self.out, ", ")?; self.write_expr(module, value, func_ctx)?; writeln!(self.out, ");")?; - } else { + } + Some((MatrixAccess::Struct { .. }, Some(Index::Static(_)), _)) + | Some((MatrixAccess::Struct { .. }, None, _)) + | None => { self.write_expr(module, pointer, func_ctx)?; write!(self.out, " = ")?; @@ -2963,6 +2896,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { rows: crate::VectorSize::Bi, width: 4, }) = get_inner_matrix_of_struct_array_member(module, base, func_ctx, true) + .or_else(|| get_global_uniform_matrix(module, base, func_ctx)) { write!(self.out, "__get_col_of_mat{}x2(", columns as u8)?; self.write_expr(module, base, func_ctx)?; @@ -3075,13 +3009,15 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { { // do nothing, the chain is written on `Load`/`Store` } else { - // We write the matrix column access in a special way since - // the type of `base` is our special __matCx2 struct. + // See if we need to write the matrix column access in a + // special way since the type of `base` is our special + // __matCx2 struct. if let Some(MatrixType { rows: crate::VectorSize::Bi, width: 4, .. }) = get_inner_matrix_of_struct_array_member(module, base, func_ctx, true) + .or_else(|| get_global_uniform_matrix(module, base, func_ctx)) { self.write_expr(module, base, func_ctx)?; write!(self.out, "._{index}")?; @@ -3381,8 +3317,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { .or_else(|| get_inner_matrix_of_global_uniform(module, pointer, func_ctx)) { let mut resolved = func_ctx.resolve_type(pointer, &module.types); - if let TypeInner::Pointer { base, .. } = *resolved { - resolved = &module.types[base].inner; + let ptr_tr = resolved.pointer_base_type(); + if let Some(ptr_ty) = + ptr_tr.as_ref().map(|tr| tr.inner_with(&module.types)) + { + resolved = ptr_ty; } write!(self.out, "((")?; @@ -4416,6 +4355,44 @@ pub(super) fn get_inner_matrix_data( } } +/// If `base` is an access chain of the form `mat`, `mat[col]`, or `mat[col][row]`, +/// returns a tuple of the matrix, the column (vector) index (if present), and +/// the row (scalar) index (if present). +fn find_matrix_in_access_chain( + module: &Module, + base: Handle, + func_ctx: &back::FunctionCtx<'_>, +) -> Option<(Handle, Option, Option)> { + let mut current_base = base; + let mut vector = None; + let mut scalar = None; + loop { + let resolved_tr = func_ctx + .resolve_type(current_base, &module.types) + .pointer_base_type(); + let resolved = resolved_tr.as_ref()?.inner_with(&module.types); + + match *resolved { + TypeInner::Matrix { .. } => return Some((current_base, vector, scalar)), + TypeInner::Scalar(_) | TypeInner::Vector { .. } => {} + _ => return None, + } + + let index; + (current_base, index) = match func_ctx.expressions[current_base] { + crate::Expression::Access { base, index } => (base, Index::Expression(index)), + crate::Expression::AccessIndex { base, index } => (base, Index::Static(index)), + _ => return None, + }; + + match *resolved { + TypeInner::Scalar(_) => scalar = Some(index), + TypeInner::Vector { .. } => vector = Some(index), + _ => unreachable!(), + } + } +} + /// Returns the matrix data if the access chain starting at `base`: /// - starts with an expression with resolved type of [`TypeInner::Matrix`] if `direct = true` /// - contains one or more expressions with resolved type of [`TypeInner::Array`] of [`TypeInner::Matrix`] @@ -4474,6 +4451,36 @@ pub(super) fn get_inner_matrix_of_struct_array_member( None } +/// Simpler version of get_inner_matrix_of_global_uniform that only looks at the +/// immediate expression, rather than traversing an access chain. +fn get_global_uniform_matrix( + module: &Module, + base: Handle, + func_ctx: &back::FunctionCtx<'_>, +) -> Option { + let base_tr = func_ctx + .resolve_type(base, &module.types) + .pointer_base_type(); + let base_ty = base_tr.as_ref().map(|tr| tr.inner_with(&module.types)); + match (&func_ctx.expressions[base], base_ty) { + ( + &crate::Expression::GlobalVariable(handle), + Some(&TypeInner::Matrix { + columns, + rows, + scalar, + }), + ) if module.global_variables[handle].space == crate::AddressSpace::Uniform => { + Some(MatrixType { + columns, + rows, + width: scalar.width, + }) + } + _ => None, + } +} + /// Returns the matrix data if the access chain starting at `base`: /// - starts with an expression with resolved type of [`TypeInner::Matrix`] /// - contains zero or more expressions with resolved type of [`TypeInner::Array`] of [`TypeInner::Matrix`] diff --git a/naga/tests/in/wgsl/access.wgsl b/naga/tests/in/wgsl/access.wgsl index d93f2b30a1..0e89bcfb0e 100644 --- a/naga/tests/in/wgsl/access.wgsl +++ b/naga/tests/in/wgsl/access.wgsl @@ -35,6 +35,9 @@ var baz: Baz; var qux: vec2; fn test_matrix_within_struct_accesses() { + // Test HLSL accesses to Cx2 matrices. There are additional tests + // in `hlsl_mat_cx2.wgsl`. + var idx = 1; idx--; diff --git a/naga/tests/in/wgsl/hlsl_mat_cx2.toml b/naga/tests/in/wgsl/hlsl_mat_cx2.toml new file mode 100644 index 0000000000..3ca0b52f4e --- /dev/null +++ b/naga/tests/in/wgsl/hlsl_mat_cx2.toml @@ -0,0 +1 @@ +targets = "HLSL" diff --git a/naga/tests/in/wgsl/hlsl_mat_cx2.wgsl b/naga/tests/in/wgsl/hlsl_mat_cx2.wgsl new file mode 100644 index 0000000000..50bc188794 --- /dev/null +++ b/naga/tests/in/wgsl/hlsl_mat_cx2.wgsl @@ -0,0 +1,177 @@ +// Test HLSL handling of N-by-2 matrices. +// See the doc comment on `naga::back::hlsl` for details. +// +// There are additional tests in `access.wgsl`. +// +// Tests that we don't apply this handling to other sizes are in hlsl_mat_cx3.wgsl. + +// Access type (3rd item in variable names) +// S = Struct +// M = Matrix +// C = Column +// E = Element + +// Index type (4th item in variable names) +// C = Constant +// V = Variable + +alias Mat = mat2x2; + +@group(0) @binding(0) +var s_m: Mat; + +@group(0) @binding(1) +var u_m: Mat; + +fn access_m() { + var idx = 1; + idx--; + + // loads from storage + let l_s_m = s_m; + let l_s_c_c = s_m[0]; + let l_s_c_v = s_m[idx]; + let l_s_e_cc = s_m[0][0]; + let l_s_e_cv = s_m[0][idx]; + let l_s_e_vc = s_m[idx][0]; + let l_s_e_vv = s_m[idx][idx]; + + // loads from uniform + let l_u_m = u_m; + let l_u_c_c = u_m[0]; + let l_u_c_v = u_m[idx]; + let l_u_e_cc = u_m[0][0]; + let l_u_e_cv = u_m[0][idx]; + let l_u_e_vc = u_m[idx][0]; + let l_u_e_vv = u_m[idx][idx]; + + // stores to storage + s_m = l_u_m; + s_m[0] = l_u_c_c; + s_m[idx] = l_u_c_v; + s_m[0][0] = l_u_e_cc; + s_m[0][idx] = l_u_e_cv; + s_m[idx][0] = l_u_e_vc; + s_m[idx][idx] = l_u_e_vv; +} + +struct StructWithMat { + m: Mat, +} + +@group(1) @binding(0) +var s_sm: StructWithMat; + +@group(1) @binding(1) +var u_sm: StructWithMat; + +fn access_sm() { + var idx = 1; + idx--; + + // loads from storage + let l_s_s = s_sm; + let l_s_m = s_sm.m; + let l_s_c_c = s_sm.m[0]; + let l_s_c_v = s_sm.m[idx]; + let l_s_e_cc = s_sm.m[0][0]; + let l_s_e_cv = s_sm.m[0][idx]; + let l_s_e_vc = s_sm.m[idx][0]; + let l_s_e_vv = s_sm.m[idx][idx]; + + // loads from uniform + let l_u_s = u_sm; + let l_u_m = u_sm.m; + let l_u_c_c = u_sm.m[0]; + let l_u_c_v = u_sm.m[idx]; + let l_u_e_cc = u_sm.m[0][0]; + let l_u_e_cv = u_sm.m[0][idx]; + let l_u_e_vc = u_sm.m[idx][0]; + let l_u_e_vv = u_sm.m[idx][idx]; + + // stores to storage + s_sm = l_u_s; + s_sm.m = l_u_m; + s_sm.m[0] = l_u_c_c; + s_sm.m[idx] = l_u_c_v; + s_sm.m[0][0] = l_u_e_cc; + s_sm.m[0][idx] = l_u_e_cv; + s_sm.m[idx][0] = l_u_e_vc; + s_sm.m[idx][idx] = l_u_e_vv; +} + +struct StructWithArrayOfStructOfMat { + a: array, +} + +@group(2) @binding(0) +var s_sasm: StructWithArrayOfStructOfMat; + +@group(2) @binding(1) +var u_sasm: StructWithArrayOfStructOfMat; + +fn access_sasm() { + var idx = 1; + idx--; + + // loads from storage + let l_s_s = s_sasm; + let l_s_a = s_sasm.a; + let l_s_m_c = s_sasm.a[0].m; + let l_s_m_v = s_sasm.a[idx].m; + let l_s_c_cc = s_sasm.a[0].m[0]; + let l_s_c_cv = s_sasm.a[0].m[idx]; + let l_s_c_vc = s_sasm.a[idx].m[0]; + let l_s_c_vv = s_sasm.a[idx].m[idx]; + let l_s_e_ccc = s_sasm.a[0].m[0][0]; + let l_s_e_ccv = s_sasm.a[0].m[0][idx]; + let l_s_e_cvc = s_sasm.a[0].m[idx][0]; + let l_s_e_cvv = s_sasm.a[0].m[idx][idx]; + let l_s_e_vcc = s_sasm.a[idx].m[0][0]; + let l_s_e_vcv = s_sasm.a[idx].m[0][idx]; + let l_s_e_vvc = s_sasm.a[idx].m[idx][0]; + let l_s_e_vvv = s_sasm.a[idx].m[idx][idx]; + + // loads from uniform + let l_u_s = u_sasm; + let l_u_a = u_sasm.a; + let l_u_m_c = u_sasm.a[0].m; + let l_u_m_v = u_sasm.a[idx].m; + let l_u_c_cc = u_sasm.a[0].m[0]; + let l_u_c_cv = u_sasm.a[0].m[idx]; + let l_u_c_vc = u_sasm.a[idx].m[0]; + let l_u_c_vv = u_sasm.a[idx].m[idx]; + let l_u_e_ccc = u_sasm.a[0].m[0][0]; + let l_u_e_ccv = u_sasm.a[0].m[0][idx]; + let l_u_e_cvc = u_sasm.a[0].m[idx][0]; + let l_u_e_cvv = u_sasm.a[0].m[idx][idx]; + let l_u_e_vcc = u_sasm.a[idx].m[0][0]; + let l_u_e_vcv = u_sasm.a[idx].m[0][idx]; + let l_u_e_vvc = u_sasm.a[idx].m[idx][0]; + let l_u_e_vvv = u_sasm.a[idx].m[idx][idx]; + + // stores to storage + s_sasm = l_u_s; + s_sasm.a = l_u_a; + s_sasm.a[0].m = l_u_m_c; + s_sasm.a[idx].m = l_u_m_v; + s_sasm.a[0].m[0] = l_u_c_cc; + s_sasm.a[0].m[idx] = l_u_c_cv; + s_sasm.a[idx].m[0] = l_u_c_vc; + s_sasm.a[idx].m[idx] = l_u_c_vv; + s_sasm.a[0].m[0][0] = l_u_e_ccc; + s_sasm.a[0].m[0][idx] = l_u_e_ccv; + s_sasm.a[0].m[idx][0] = l_u_e_cvc; + s_sasm.a[0].m[idx][idx] = l_u_e_cvv; + s_sasm.a[idx].m[0][0] = l_u_e_vcc; + s_sasm.a[idx].m[0][idx] = l_u_e_vcv; + s_sasm.a[idx].m[idx][0] = l_u_e_vvc; + s_sasm.a[idx].m[idx][idx] = l_u_e_vvv; +} + +@compute @workgroup_size(1) +fn main() { + access_m(); + access_sm(); + access_sasm(); +} diff --git a/naga/tests/in/wgsl/hlsl_mat_cx3.toml b/naga/tests/in/wgsl/hlsl_mat_cx3.toml new file mode 100644 index 0000000000..3ca0b52f4e --- /dev/null +++ b/naga/tests/in/wgsl/hlsl_mat_cx3.toml @@ -0,0 +1 @@ +targets = "HLSL" diff --git a/naga/tests/in/wgsl/hlsl_mat_cx3.wgsl b/naga/tests/in/wgsl/hlsl_mat_cx3.wgsl new file mode 100644 index 0000000000..e33f10fc9c --- /dev/null +++ b/naga/tests/in/wgsl/hlsl_mat_cx3.wgsl @@ -0,0 +1,173 @@ +// Test HLSL handling of N-by-3 matrices. These should not receive the special +// treatment that N-by-2 matrices receive (which is tested in hlsl_mat_cx2). + +// Access type (3rd item in variable names) +// S = Struct +// M = Matrix +// C = Column +// E = Element + +// Index type (4th item in variable names) +// C = Constant +// V = Variable + +alias Mat = mat3x3; + +@group(0) @binding(0) +var s_m: Mat; + +@group(0) @binding(1) +var u_m: Mat; + +fn access_m() { + var idx = 1; + idx--; + + // loads from storage + let l_s_m = s_m; + let l_s_c_c = s_m[0]; + let l_s_c_v = s_m[idx]; + let l_s_e_cc = s_m[0][0]; + let l_s_e_cv = s_m[0][idx]; + let l_s_e_vc = s_m[idx][0]; + let l_s_e_vv = s_m[idx][idx]; + + // loads from uniform + let l_u_m = u_m; + let l_u_c_c = u_m[0]; + let l_u_c_v = u_m[idx]; + let l_u_e_cc = u_m[0][0]; + let l_u_e_cv = u_m[0][idx]; + let l_u_e_vc = u_m[idx][0]; + let l_u_e_vv = u_m[idx][idx]; + + // stores to storage + s_m = l_u_m; + s_m[0] = l_u_c_c; + s_m[idx] = l_u_c_v; + s_m[0][0] = l_u_e_cc; + s_m[0][idx] = l_u_e_cv; + s_m[idx][0] = l_u_e_vc; + s_m[idx][idx] = l_u_e_vv; +} + +struct StructWithMat { + m: Mat, +} + +@group(1) @binding(0) +var s_sm: StructWithMat; + +@group(1) @binding(1) +var u_sm: StructWithMat; + +fn access_sm() { + var idx = 1; + idx--; + + // loads from storage + let l_s_s = s_sm; + let l_s_m = s_sm.m; + let l_s_c_c = s_sm.m[0]; + let l_s_c_v = s_sm.m[idx]; + let l_s_e_cc = s_sm.m[0][0]; + let l_s_e_cv = s_sm.m[0][idx]; + let l_s_e_vc = s_sm.m[idx][0]; + let l_s_e_vv = s_sm.m[idx][idx]; + + // loads from uniform + let l_u_s = u_sm; + let l_u_m = u_sm.m; + let l_u_c_c = u_sm.m[0]; + let l_u_c_v = u_sm.m[idx]; + let l_u_e_cc = u_sm.m[0][0]; + let l_u_e_cv = u_sm.m[0][idx]; + let l_u_e_vc = u_sm.m[idx][0]; + let l_u_e_vv = u_sm.m[idx][idx]; + + // stores to storage + s_sm = l_u_s; + s_sm.m = l_u_m; + s_sm.m[0] = l_u_c_c; + s_sm.m[idx] = l_u_c_v; + s_sm.m[0][0] = l_u_e_cc; + s_sm.m[0][idx] = l_u_e_cv; + s_sm.m[idx][0] = l_u_e_vc; + s_sm.m[idx][idx] = l_u_e_vv; +} + +struct StructWithArrayOfStructOfMat { + a: array, +} + +@group(2) @binding(0) +var s_sasm: StructWithArrayOfStructOfMat; + +@group(2) @binding(1) +var u_sasm: StructWithArrayOfStructOfMat; + +fn access_sasm() { + var idx = 1; + idx--; + + // loads from storage + let l_s_s = s_sasm; + let l_s_a = s_sasm.a; + let l_s_m_c = s_sasm.a[0].m; + let l_s_m_v = s_sasm.a[idx].m; + let l_s_c_cc = s_sasm.a[0].m[0]; + let l_s_c_cv = s_sasm.a[0].m[idx]; + let l_s_c_vc = s_sasm.a[idx].m[0]; + let l_s_c_vv = s_sasm.a[idx].m[idx]; + let l_s_e_ccc = s_sasm.a[0].m[0][0]; + let l_s_e_ccv = s_sasm.a[0].m[0][idx]; + let l_s_e_cvc = s_sasm.a[0].m[idx][0]; + let l_s_e_cvv = s_sasm.a[0].m[idx][idx]; + let l_s_e_vcc = s_sasm.a[idx].m[0][0]; + let l_s_e_vcv = s_sasm.a[idx].m[0][idx]; + let l_s_e_vvc = s_sasm.a[idx].m[idx][0]; + let l_s_e_vvv = s_sasm.a[idx].m[idx][idx]; + + // loads from uniform + let l_u_s = u_sasm; + let l_u_a = u_sasm.a; + let l_u_m_c = u_sasm.a[0].m; + let l_u_m_v = u_sasm.a[idx].m; + let l_u_c_cc = u_sasm.a[0].m[0]; + let l_u_c_cv = u_sasm.a[0].m[idx]; + let l_u_c_vc = u_sasm.a[idx].m[0]; + let l_u_c_vv = u_sasm.a[idx].m[idx]; + let l_u_e_ccc = u_sasm.a[0].m[0][0]; + let l_u_e_ccv = u_sasm.a[0].m[0][idx]; + let l_u_e_cvc = u_sasm.a[0].m[idx][0]; + let l_u_e_cvv = u_sasm.a[0].m[idx][idx]; + let l_u_e_vcc = u_sasm.a[idx].m[0][0]; + let l_u_e_vcv = u_sasm.a[idx].m[0][idx]; + let l_u_e_vvc = u_sasm.a[idx].m[idx][0]; + let l_u_e_vvv = u_sasm.a[idx].m[idx][idx]; + + // stores to storage + s_sasm = l_u_s; + s_sasm.a = l_u_a; + s_sasm.a[0].m = l_u_m_c; + s_sasm.a[idx].m = l_u_m_v; + s_sasm.a[0].m[0] = l_u_c_cc; + s_sasm.a[0].m[idx] = l_u_c_cv; + s_sasm.a[idx].m[0] = l_u_c_vc; + s_sasm.a[idx].m[idx] = l_u_c_vv; + s_sasm.a[0].m[0][0] = l_u_e_ccc; + s_sasm.a[0].m[0][idx] = l_u_e_ccv; + s_sasm.a[0].m[idx][0] = l_u_e_cvc; + s_sasm.a[0].m[idx][idx] = l_u_e_cvv; + s_sasm.a[idx].m[0][0] = l_u_e_vcc; + s_sasm.a[idx].m[0][idx] = l_u_e_vcv; + s_sasm.a[idx].m[idx][0] = l_u_e_vvc; + s_sasm.a[idx].m[idx][idx] = l_u_e_vvv; +} + +@compute @workgroup_size(1) +fn main() { + access_m(); + access_sm(); + access_sasm(); +} diff --git a/naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.hlsl b/naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.hlsl new file mode 100644 index 0000000000..1de9e069cd --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.hlsl @@ -0,0 +1,372 @@ +typedef struct { float2 _0; float2 _1; } __mat2x2; +float2 __get_col_of_mat2x2(__mat2x2 mat, uint idx) { + switch(idx) { + case 0: { return mat._0; } + case 1: { return mat._1; } + default: { return (float2)0; } + } +} +void __set_col_of_mat2x2(__mat2x2 mat, uint idx, float2 value) { + switch(idx) { + case 0: { mat._0 = value; break; } + case 1: { mat._1 = value; break; } + } +} +void __set_el_of_mat2x2(__mat2x2 mat, uint idx, uint vec_idx, float value) { + switch(idx) { + case 0: { mat._0[vec_idx] = value; break; } + case 1: { mat._1[vec_idx] = value; break; } + } +} + +struct StructWithMat { + float2 m_0; float2 m_1; +}; + +struct StructWithArrayOfStructOfMat { + StructWithMat a[4]; +}; + +RWByteAddressBuffer s_m : register(u0); +cbuffer u_m : register(b1) { __mat2x2 u_m; } +RWByteAddressBuffer s_sm : register(u0, space1); +cbuffer u_sm : register(b1, space1) { StructWithMat u_sm; } +RWByteAddressBuffer s_sasm : register(u0, space2); +cbuffer u_sasm : register(b1, space2) { StructWithArrayOfStructOfMat u_sasm; } + +void access_m() +{ + int idx = int(1); + + int _e3 = idx; + idx = asint(asuint(_e3) - asuint(int(1))); + float2x2 l_s_m = float2x2(asfloat(s_m.Load2(0)), asfloat(s_m.Load2(8))); + float2 l_s_c_c = asfloat(s_m.Load2(0)); + int _e11 = idx; + float2 l_s_c_v = asfloat(s_m.Load2(_e11*8)); + float l_s_e_cc = asfloat(s_m.Load(0+0)); + int _e20 = idx; + float l_s_e_cv = asfloat(s_m.Load(_e20*4+0)); + int _e24 = idx; + float l_s_e_vc = asfloat(s_m.Load(0+_e24*8)); + int _e29 = idx; + int _e31 = idx; + float l_s_e_vv = asfloat(s_m.Load(_e31*4+_e29*8)); + float2x2 l_u_m = ((float2x2)u_m); + float2 l_u_c_c = u_m._0; + int _e40 = idx; + float2 l_u_c_v = __get_col_of_mat2x2(u_m, _e40); + float l_u_e_cc = u_m._0.x; + int _e49 = idx; + float l_u_e_cv = u_m._0[_e49]; + int _e53 = idx; + float l_u_e_vc = __get_col_of_mat2x2(u_m, _e53).x; + int _e58 = idx; + int _e60 = idx; + float l_u_e_vv = __get_col_of_mat2x2(u_m, _e58)[_e60]; + { + float2x2 _value2 = l_u_m; + s_m.Store2(0, asuint(_value2[0])); + s_m.Store2(8, asuint(_value2[1])); + } + s_m.Store2(0, asuint(l_u_c_c)); + int _e67 = idx; + s_m.Store2(_e67*8, asuint(l_u_c_v)); + s_m.Store(0+0, asuint(l_u_e_cc)); + int _e74 = idx; + s_m.Store(_e74*4+0, asuint(l_u_e_cv)); + int _e77 = idx; + s_m.Store(0+_e77*8, asuint(l_u_e_vc)); + int _e81 = idx; + int _e83 = idx; + s_m.Store(_e83*4+_e81*8, asuint(l_u_e_vv)); + return; +} + +StructWithMat ConstructStructWithMat(float2x2 arg0) { + StructWithMat ret = (StructWithMat)0; + ret.m_0 = arg0[0]; + ret.m_1 = arg0[1]; + return ret; +} + +float2x2 GetMatmOnStructWithMat(StructWithMat obj) { + return float2x2(obj.m_0, obj.m_1); +} + +void SetMatmOnStructWithMat(StructWithMat obj, float2x2 mat) { + obj.m_0 = mat[0]; + obj.m_1 = mat[1]; +} + +void SetMatVecmOnStructWithMat(StructWithMat obj, float2 vec, uint mat_idx) { + switch(mat_idx) { + case 0: { obj.m_0 = vec; break; } + case 1: { obj.m_1 = vec; break; } + } +} + +void SetMatScalarmOnStructWithMat(StructWithMat obj, float scalar, uint mat_idx, uint vec_idx) { + switch(mat_idx) { + case 0: { obj.m_0[vec_idx] = scalar; break; } + case 1: { obj.m_1[vec_idx] = scalar; break; } + } +} + +void access_sm() +{ + int idx_1 = int(1); + + int _e3 = idx_1; + idx_1 = asint(asuint(_e3) - asuint(int(1))); + StructWithMat l_s_s = ConstructStructWithMat(float2x2(asfloat(s_sm.Load2(0+0)), asfloat(s_sm.Load2(0+8)))); + float2x2 l_s_m_1 = float2x2(asfloat(s_sm.Load2(0+0)), asfloat(s_sm.Load2(0+8))); + float2 l_s_c_c_1 = asfloat(s_sm.Load2(0+0)); + int _e16 = idx_1; + float2 l_s_c_v_1 = asfloat(s_sm.Load2(_e16*8+0)); + float l_s_e_cc_1 = asfloat(s_sm.Load(0+0+0)); + int _e27 = idx_1; + float l_s_e_cv_1 = asfloat(s_sm.Load(_e27*4+0+0)); + int _e32 = idx_1; + float l_s_e_vc_1 = asfloat(s_sm.Load(0+_e32*8+0)); + int _e38 = idx_1; + int _e40 = idx_1; + float l_s_e_vv_1 = asfloat(s_sm.Load(_e40*4+_e38*8+0)); + StructWithMat l_u_s = u_sm; + float2x2 l_u_m_1 = GetMatmOnStructWithMat(u_sm); + float2 l_u_c_c_1 = GetMatmOnStructWithMat(u_sm)[0]; + int _e54 = idx_1; + float2 l_u_c_v_1 = GetMatmOnStructWithMat(u_sm)[_e54]; + float l_u_e_cc_1 = GetMatmOnStructWithMat(u_sm)[0].x; + int _e65 = idx_1; + float l_u_e_cv_1 = GetMatmOnStructWithMat(u_sm)[0][_e65]; + int _e70 = idx_1; + float l_u_e_vc_1 = GetMatmOnStructWithMat(u_sm)[_e70].x; + int _e76 = idx_1; + int _e78 = idx_1; + float l_u_e_vv_1 = GetMatmOnStructWithMat(u_sm)[_e76][_e78]; + { + StructWithMat _value2 = l_u_s; + { + s_sm.Store2(0+0, asuint(_value2.m_0)); + s_sm.Store2(0+8, asuint(_value2.m_1)); + } + } + { + float2x2 _value2 = l_u_m_1; + s_sm.Store2(0+0, asuint(_value2[0])); + s_sm.Store2(0+8, asuint(_value2[1])); + } + s_sm.Store2(0+0, asuint(l_u_c_c_1)); + int _e89 = idx_1; + s_sm.Store2(_e89*8+0, asuint(l_u_c_v_1)); + s_sm.Store(0+0+0, asuint(l_u_e_cc_1)); + int _e98 = idx_1; + s_sm.Store(_e98*4+0+0, asuint(l_u_e_cv_1)); + int _e102 = idx_1; + s_sm.Store(0+_e102*8+0, asuint(l_u_e_vc_1)); + int _e107 = idx_1; + int _e109 = idx_1; + s_sm.Store(_e109*4+_e107*8+0, asuint(l_u_e_vv_1)); + return; +} + +typedef StructWithMat ret_Constructarray4_StructWithMat_[4]; +ret_Constructarray4_StructWithMat_ Constructarray4_StructWithMat_(StructWithMat arg0, StructWithMat arg1, StructWithMat arg2, StructWithMat arg3) { + StructWithMat ret[4] = { arg0, arg1, arg2, arg3 }; + return ret; +} + +StructWithArrayOfStructOfMat ConstructStructWithArrayOfStructOfMat(StructWithMat arg0[4]) { + StructWithArrayOfStructOfMat ret = (StructWithArrayOfStructOfMat)0; + ret.a = arg0; + return ret; +} + +void access_sasm() +{ + int idx_2 = int(1); + + int _e3 = idx_2; + idx_2 = asint(asuint(_e3) - asuint(int(1))); + StructWithArrayOfStructOfMat l_s_s_1 = ConstructStructWithArrayOfStructOfMat(Constructarray4_StructWithMat_(ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+0+0+0)), asfloat(s_sasm.Load2(0+0+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+16+0+0)), asfloat(s_sasm.Load2(0+16+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+32+0+0)), asfloat(s_sasm.Load2(0+32+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+48+0+0)), asfloat(s_sasm.Load2(0+48+0+8)))))); + StructWithMat l_s_a[4] = Constructarray4_StructWithMat_(ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+0+0+0)), asfloat(s_sasm.Load2(0+0+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+16+0+0)), asfloat(s_sasm.Load2(0+16+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+32+0+0)), asfloat(s_sasm.Load2(0+32+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+48+0+0)), asfloat(s_sasm.Load2(0+48+0+8))))); + float2x2 l_s_m_c = float2x2(asfloat(s_sasm.Load2(0+0+0+0)), asfloat(s_sasm.Load2(0+0+0+8))); + int _e17 = idx_2; + float2x2 l_s_m_v = float2x2(asfloat(s_sasm.Load2(0+_e17*16+0+0)), asfloat(s_sasm.Load2(0+_e17*16+0+8))); + float2 l_s_c_cc = asfloat(s_sasm.Load2(0+0+0+0)); + int _e31 = idx_2; + float2 l_s_c_cv = asfloat(s_sasm.Load2(_e31*8+0+0+0)); + int _e36 = idx_2; + float2 l_s_c_vc = asfloat(s_sasm.Load2(0+0+_e36*16+0)); + int _e43 = idx_2; + int _e46 = idx_2; + float2 l_s_c_vv = asfloat(s_sasm.Load2(_e46*8+0+_e43*16+0)); + float l_s_e_ccc = asfloat(s_sasm.Load(0+0+0+0+0)); + int _e61 = idx_2; + float l_s_e_ccv = asfloat(s_sasm.Load(_e61*4+0+0+0+0)); + int _e68 = idx_2; + float l_s_e_cvc = asfloat(s_sasm.Load(0+_e68*8+0+0+0)); + int _e76 = idx_2; + int _e78 = idx_2; + float l_s_e_cvv = asfloat(s_sasm.Load(_e78*4+_e76*8+0+0+0)); + int _e83 = idx_2; + float l_s_e_vcc = asfloat(s_sasm.Load(0+0+0+_e83*16+0)); + int _e91 = idx_2; + int _e95 = idx_2; + float l_s_e_vcv = asfloat(s_sasm.Load(_e95*4+0+0+_e91*16+0)); + int _e100 = idx_2; + int _e103 = idx_2; + float l_s_e_vvc = asfloat(s_sasm.Load(0+_e103*8+0+_e100*16+0)); + int _e109 = idx_2; + int _e112 = idx_2; + int _e114 = idx_2; + float l_s_e_vvv = asfloat(s_sasm.Load(_e114*4+_e112*8+0+_e109*16+0)); + StructWithArrayOfStructOfMat l_u_s_1 = u_sasm; + StructWithMat l_u_a[4] = u_sasm.a; + float2x2 l_u_m_c = GetMatmOnStructWithMat(u_sasm.a[0]); + int _e129 = idx_2; + float2x2 l_u_m_v = GetMatmOnStructWithMat(u_sasm.a[_e129]); + float2 l_u_c_cc = GetMatmOnStructWithMat(u_sasm.a[0])[0]; + int _e143 = idx_2; + float2 l_u_c_cv = GetMatmOnStructWithMat(u_sasm.a[0])[_e143]; + int _e148 = idx_2; + float2 l_u_c_vc = GetMatmOnStructWithMat(u_sasm.a[_e148])[0]; + int _e155 = idx_2; + int _e158 = idx_2; + float2 l_u_c_vv = GetMatmOnStructWithMat(u_sasm.a[_e155])[_e158]; + float l_u_e_ccc = GetMatmOnStructWithMat(u_sasm.a[0])[0].x; + int _e173 = idx_2; + float l_u_e_ccv = GetMatmOnStructWithMat(u_sasm.a[0])[0][_e173]; + int _e180 = idx_2; + float l_u_e_cvc = GetMatmOnStructWithMat(u_sasm.a[0])[_e180].x; + int _e188 = idx_2; + int _e190 = idx_2; + float l_u_e_cvv = GetMatmOnStructWithMat(u_sasm.a[0])[_e188][_e190]; + int _e195 = idx_2; + float l_u_e_vcc = GetMatmOnStructWithMat(u_sasm.a[_e195])[0].x; + int _e203 = idx_2; + int _e207 = idx_2; + float l_u_e_vcv = GetMatmOnStructWithMat(u_sasm.a[_e203])[0][_e207]; + int _e212 = idx_2; + int _e215 = idx_2; + float l_u_e_vvc = GetMatmOnStructWithMat(u_sasm.a[_e212])[_e215].x; + int _e221 = idx_2; + int _e224 = idx_2; + int _e226 = idx_2; + float l_u_e_vvv = GetMatmOnStructWithMat(u_sasm.a[_e221])[_e224][_e226]; + { + StructWithArrayOfStructOfMat _value2 = l_u_s_1; + { + StructWithMat _value3[4] = _value2.a; + { + StructWithMat _value4 = _value3[0]; + { + s_sasm.Store2(0+0+0+0, asuint(_value4.m_0)); + s_sasm.Store2(0+0+0+8, asuint(_value4.m_1)); + } + } + { + StructWithMat _value4 = _value3[1]; + { + s_sasm.Store2(0+16+0+0, asuint(_value4.m_0)); + s_sasm.Store2(0+16+0+8, asuint(_value4.m_1)); + } + } + { + StructWithMat _value4 = _value3[2]; + { + s_sasm.Store2(0+32+0+0, asuint(_value4.m_0)); + s_sasm.Store2(0+32+0+8, asuint(_value4.m_1)); + } + } + { + StructWithMat _value4 = _value3[3]; + { + s_sasm.Store2(0+48+0+0, asuint(_value4.m_0)); + s_sasm.Store2(0+48+0+8, asuint(_value4.m_1)); + } + } + } + } + { + StructWithMat _value2[4] = l_u_a; + { + StructWithMat _value3 = _value2[0]; + { + s_sasm.Store2(0+0+0+0, asuint(_value3.m_0)); + s_sasm.Store2(0+0+0+8, asuint(_value3.m_1)); + } + } + { + StructWithMat _value3 = _value2[1]; + { + s_sasm.Store2(0+16+0+0, asuint(_value3.m_0)); + s_sasm.Store2(0+16+0+8, asuint(_value3.m_1)); + } + } + { + StructWithMat _value3 = _value2[2]; + { + s_sasm.Store2(0+32+0+0, asuint(_value3.m_0)); + s_sasm.Store2(0+32+0+8, asuint(_value3.m_1)); + } + } + { + StructWithMat _value3 = _value2[3]; + { + s_sasm.Store2(0+48+0+0, asuint(_value3.m_0)); + s_sasm.Store2(0+48+0+8, asuint(_value3.m_1)); + } + } + } + { + float2x2 _value2 = l_u_m_c; + s_sasm.Store2(0+0+0+0, asuint(_value2[0])); + s_sasm.Store2(0+0+0+8, asuint(_value2[1])); + } + int _e238 = idx_2; + { + float2x2 _value2 = l_u_m_v; + s_sasm.Store2(0+_e238*16+0+0, asuint(_value2[0])); + s_sasm.Store2(0+_e238*16+0+8, asuint(_value2[1])); + } + s_sasm.Store2(0+0+0+0, asuint(l_u_c_cc)); + int _e250 = idx_2; + s_sasm.Store2(_e250*8+0+0+0, asuint(l_u_c_cv)); + int _e254 = idx_2; + s_sasm.Store2(0+0+_e254*16+0, asuint(l_u_c_vc)); + int _e260 = idx_2; + int _e263 = idx_2; + s_sasm.Store2(_e263*8+0+_e260*16+0, asuint(l_u_c_vv)); + s_sasm.Store(0+0+0+0+0, asuint(l_u_e_ccc)); + int _e276 = idx_2; + s_sasm.Store(_e276*4+0+0+0+0, asuint(l_u_e_ccv)); + int _e282 = idx_2; + s_sasm.Store(0+_e282*8+0+0+0, asuint(l_u_e_cvc)); + int _e289 = idx_2; + int _e291 = idx_2; + s_sasm.Store(_e291*4+_e289*8+0+0+0, asuint(l_u_e_cvv)); + int _e295 = idx_2; + s_sasm.Store(0+0+0+_e295*16+0, asuint(l_u_e_vcc)); + int _e302 = idx_2; + int _e306 = idx_2; + s_sasm.Store(_e306*4+0+0+_e302*16+0, asuint(l_u_e_vcv)); + int _e310 = idx_2; + int _e313 = idx_2; + s_sasm.Store(0+_e313*8+0+_e310*16+0, asuint(l_u_e_vvc)); + int _e318 = idx_2; + int _e321 = idx_2; + int _e323 = idx_2; + s_sasm.Store(_e323*4+_e321*8+0+_e318*16+0, asuint(l_u_e_vvv)); + return; +} + +[numthreads(1, 1, 1)] +void main() +{ + access_m(); + access_sm(); + access_sasm(); + return; +} diff --git a/naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.ron b/naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.ron new file mode 100644 index 0000000000..a07b03300b --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_5_1", + ), + ], +) diff --git a/naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.hlsl b/naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.hlsl new file mode 100644 index 0000000000..f90cdff6e1 --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.hlsl @@ -0,0 +1,350 @@ +struct StructWithMat { + row_major float3x3 m; + int _end_pad_0; +}; + +struct StructWithArrayOfStructOfMat { + StructWithMat a[4]; +}; + +RWByteAddressBuffer s_m : register(u0); +cbuffer u_m : register(b1) { row_major float3x3 u_m; } +RWByteAddressBuffer s_sm : register(u0, space1); +cbuffer u_sm : register(b1, space1) { StructWithMat u_sm; } +RWByteAddressBuffer s_sasm : register(u0, space2); +cbuffer u_sasm : register(b1, space2) { StructWithArrayOfStructOfMat u_sasm; } + +void access_m() +{ + int idx = int(1); + + int _e3 = idx; + idx = asint(asuint(_e3) - asuint(int(1))); + float3x3 l_s_m = float3x3(asfloat(s_m.Load3(0)), asfloat(s_m.Load3(16)), asfloat(s_m.Load3(32))); + float3 l_s_c_c = asfloat(s_m.Load3(0)); + int _e11 = idx; + float3 l_s_c_v = asfloat(s_m.Load3(_e11*16)); + float l_s_e_cc = asfloat(s_m.Load(0+0)); + int _e20 = idx; + float l_s_e_cv = asfloat(s_m.Load(_e20*4+0)); + int _e24 = idx; + float l_s_e_vc = asfloat(s_m.Load(0+_e24*16)); + int _e29 = idx; + int _e31 = idx; + float l_s_e_vv = asfloat(s_m.Load(_e31*4+_e29*16)); + float3x3 l_u_m = u_m; + float3 l_u_c_c = u_m[0]; + int _e40 = idx; + float3 l_u_c_v = u_m[_e40]; + float l_u_e_cc = u_m[0].x; + int _e49 = idx; + float l_u_e_cv = u_m[0][_e49]; + int _e53 = idx; + float l_u_e_vc = u_m[_e53].x; + int _e58 = idx; + int _e60 = idx; + float l_u_e_vv = u_m[_e58][_e60]; + { + float3x3 _value2 = l_u_m; + s_m.Store3(0, asuint(_value2[0])); + s_m.Store3(16, asuint(_value2[1])); + s_m.Store3(32, asuint(_value2[2])); + } + s_m.Store3(0, asuint(l_u_c_c)); + int _e67 = idx; + s_m.Store3(_e67*16, asuint(l_u_c_v)); + s_m.Store(0+0, asuint(l_u_e_cc)); + int _e74 = idx; + s_m.Store(_e74*4+0, asuint(l_u_e_cv)); + int _e77 = idx; + s_m.Store(0+_e77*16, asuint(l_u_e_vc)); + int _e81 = idx; + int _e83 = idx; + s_m.Store(_e83*4+_e81*16, asuint(l_u_e_vv)); + return; +} + +StructWithMat ConstructStructWithMat(float3x3 arg0) { + StructWithMat ret = (StructWithMat)0; + ret.m = arg0; + return ret; +} + +void access_sm() +{ + int idx_1 = int(1); + + int _e3 = idx_1; + idx_1 = asint(asuint(_e3) - asuint(int(1))); + StructWithMat l_s_s = ConstructStructWithMat(float3x3(asfloat(s_sm.Load3(0+0)), asfloat(s_sm.Load3(0+16)), asfloat(s_sm.Load3(0+32)))); + float3x3 l_s_m_1 = float3x3(asfloat(s_sm.Load3(0+0)), asfloat(s_sm.Load3(0+16)), asfloat(s_sm.Load3(0+32))); + float3 l_s_c_c_1 = asfloat(s_sm.Load3(0+0)); + int _e16 = idx_1; + float3 l_s_c_v_1 = asfloat(s_sm.Load3(_e16*16+0)); + float l_s_e_cc_1 = asfloat(s_sm.Load(0+0+0)); + int _e27 = idx_1; + float l_s_e_cv_1 = asfloat(s_sm.Load(_e27*4+0+0)); + int _e32 = idx_1; + float l_s_e_vc_1 = asfloat(s_sm.Load(0+_e32*16+0)); + int _e38 = idx_1; + int _e40 = idx_1; + float l_s_e_vv_1 = asfloat(s_sm.Load(_e40*4+_e38*16+0)); + StructWithMat l_u_s = u_sm; + float3x3 l_u_m_1 = u_sm.m; + float3 l_u_c_c_1 = u_sm.m[0]; + int _e54 = idx_1; + float3 l_u_c_v_1 = u_sm.m[_e54]; + float l_u_e_cc_1 = u_sm.m[0].x; + int _e65 = idx_1; + float l_u_e_cv_1 = u_sm.m[0][_e65]; + int _e70 = idx_1; + float l_u_e_vc_1 = u_sm.m[_e70].x; + int _e76 = idx_1; + int _e78 = idx_1; + float l_u_e_vv_1 = u_sm.m[_e76][_e78]; + { + StructWithMat _value2 = l_u_s; + { + float3x3 _value3 = _value2.m; + s_sm.Store3(0+0, asuint(_value3[0])); + s_sm.Store3(0+16, asuint(_value3[1])); + s_sm.Store3(0+32, asuint(_value3[2])); + } + } + { + float3x3 _value2 = l_u_m_1; + s_sm.Store3(0+0, asuint(_value2[0])); + s_sm.Store3(0+16, asuint(_value2[1])); + s_sm.Store3(0+32, asuint(_value2[2])); + } + s_sm.Store3(0+0, asuint(l_u_c_c_1)); + int _e89 = idx_1; + s_sm.Store3(_e89*16+0, asuint(l_u_c_v_1)); + s_sm.Store(0+0+0, asuint(l_u_e_cc_1)); + int _e98 = idx_1; + s_sm.Store(_e98*4+0+0, asuint(l_u_e_cv_1)); + int _e102 = idx_1; + s_sm.Store(0+_e102*16+0, asuint(l_u_e_vc_1)); + int _e107 = idx_1; + int _e109 = idx_1; + s_sm.Store(_e109*4+_e107*16+0, asuint(l_u_e_vv_1)); + return; +} + +typedef StructWithMat ret_Constructarray4_StructWithMat_[4]; +ret_Constructarray4_StructWithMat_ Constructarray4_StructWithMat_(StructWithMat arg0, StructWithMat arg1, StructWithMat arg2, StructWithMat arg3) { + StructWithMat ret[4] = { arg0, arg1, arg2, arg3 }; + return ret; +} + +StructWithArrayOfStructOfMat ConstructStructWithArrayOfStructOfMat(StructWithMat arg0[4]) { + StructWithArrayOfStructOfMat ret = (StructWithArrayOfStructOfMat)0; + ret.a = arg0; + return ret; +} + +void access_sasm() +{ + int idx_2 = int(1); + + int _e3 = idx_2; + idx_2 = asint(asuint(_e3) - asuint(int(1))); + StructWithArrayOfStructOfMat l_s_s_1 = ConstructStructWithArrayOfStructOfMat(Constructarray4_StructWithMat_(ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+0+0+0)), asfloat(s_sasm.Load3(0+0+0+16)), asfloat(s_sasm.Load3(0+0+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+48+0+0)), asfloat(s_sasm.Load3(0+48+0+16)), asfloat(s_sasm.Load3(0+48+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+96+0+0)), asfloat(s_sasm.Load3(0+96+0+16)), asfloat(s_sasm.Load3(0+96+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+144+0+0)), asfloat(s_sasm.Load3(0+144+0+16)), asfloat(s_sasm.Load3(0+144+0+32)))))); + StructWithMat l_s_a[4] = Constructarray4_StructWithMat_(ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+0+0+0)), asfloat(s_sasm.Load3(0+0+0+16)), asfloat(s_sasm.Load3(0+0+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+48+0+0)), asfloat(s_sasm.Load3(0+48+0+16)), asfloat(s_sasm.Load3(0+48+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+96+0+0)), asfloat(s_sasm.Load3(0+96+0+16)), asfloat(s_sasm.Load3(0+96+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+144+0+0)), asfloat(s_sasm.Load3(0+144+0+16)), asfloat(s_sasm.Load3(0+144+0+32))))); + float3x3 l_s_m_c = float3x3(asfloat(s_sasm.Load3(0+0+0+0)), asfloat(s_sasm.Load3(0+0+0+16)), asfloat(s_sasm.Load3(0+0+0+32))); + int _e17 = idx_2; + float3x3 l_s_m_v = float3x3(asfloat(s_sasm.Load3(0+_e17*48+0+0)), asfloat(s_sasm.Load3(0+_e17*48+0+16)), asfloat(s_sasm.Load3(0+_e17*48+0+32))); + float3 l_s_c_cc = asfloat(s_sasm.Load3(0+0+0+0)); + int _e31 = idx_2; + float3 l_s_c_cv = asfloat(s_sasm.Load3(_e31*16+0+0+0)); + int _e36 = idx_2; + float3 l_s_c_vc = asfloat(s_sasm.Load3(0+0+_e36*48+0)); + int _e43 = idx_2; + int _e46 = idx_2; + float3 l_s_c_vv = asfloat(s_sasm.Load3(_e46*16+0+_e43*48+0)); + float l_s_e_ccc = asfloat(s_sasm.Load(0+0+0+0+0)); + int _e61 = idx_2; + float l_s_e_ccv = asfloat(s_sasm.Load(_e61*4+0+0+0+0)); + int _e68 = idx_2; + float l_s_e_cvc = asfloat(s_sasm.Load(0+_e68*16+0+0+0)); + int _e76 = idx_2; + int _e78 = idx_2; + float l_s_e_cvv = asfloat(s_sasm.Load(_e78*4+_e76*16+0+0+0)); + int _e83 = idx_2; + float l_s_e_vcc = asfloat(s_sasm.Load(0+0+0+_e83*48+0)); + int _e91 = idx_2; + int _e95 = idx_2; + float l_s_e_vcv = asfloat(s_sasm.Load(_e95*4+0+0+_e91*48+0)); + int _e100 = idx_2; + int _e103 = idx_2; + float l_s_e_vvc = asfloat(s_sasm.Load(0+_e103*16+0+_e100*48+0)); + int _e109 = idx_2; + int _e112 = idx_2; + int _e114 = idx_2; + float l_s_e_vvv = asfloat(s_sasm.Load(_e114*4+_e112*16+0+_e109*48+0)); + StructWithArrayOfStructOfMat l_u_s_1 = u_sasm; + StructWithMat l_u_a[4] = u_sasm.a; + float3x3 l_u_m_c = u_sasm.a[0].m; + int _e129 = idx_2; + float3x3 l_u_m_v = u_sasm.a[_e129].m; + float3 l_u_c_cc = u_sasm.a[0].m[0]; + int _e143 = idx_2; + float3 l_u_c_cv = u_sasm.a[0].m[_e143]; + int _e148 = idx_2; + float3 l_u_c_vc = u_sasm.a[_e148].m[0]; + int _e155 = idx_2; + int _e158 = idx_2; + float3 l_u_c_vv = u_sasm.a[_e155].m[_e158]; + float l_u_e_ccc = u_sasm.a[0].m[0].x; + int _e173 = idx_2; + float l_u_e_ccv = u_sasm.a[0].m[0][_e173]; + int _e180 = idx_2; + float l_u_e_cvc = u_sasm.a[0].m[_e180].x; + int _e188 = idx_2; + int _e190 = idx_2; + float l_u_e_cvv = u_sasm.a[0].m[_e188][_e190]; + int _e195 = idx_2; + float l_u_e_vcc = u_sasm.a[_e195].m[0].x; + int _e203 = idx_2; + int _e207 = idx_2; + float l_u_e_vcv = u_sasm.a[_e203].m[0][_e207]; + int _e212 = idx_2; + int _e215 = idx_2; + float l_u_e_vvc = u_sasm.a[_e212].m[_e215].x; + int _e221 = idx_2; + int _e224 = idx_2; + int _e226 = idx_2; + float l_u_e_vvv = u_sasm.a[_e221].m[_e224][_e226]; + { + StructWithArrayOfStructOfMat _value2 = l_u_s_1; + { + StructWithMat _value3[4] = _value2.a; + { + StructWithMat _value4 = _value3[0]; + { + float3x3 _value5 = _value4.m; + s_sasm.Store3(0+0+0+0, asuint(_value5[0])); + s_sasm.Store3(0+0+0+16, asuint(_value5[1])); + s_sasm.Store3(0+0+0+32, asuint(_value5[2])); + } + } + { + StructWithMat _value4 = _value3[1]; + { + float3x3 _value5 = _value4.m; + s_sasm.Store3(0+48+0+0, asuint(_value5[0])); + s_sasm.Store3(0+48+0+16, asuint(_value5[1])); + s_sasm.Store3(0+48+0+32, asuint(_value5[2])); + } + } + { + StructWithMat _value4 = _value3[2]; + { + float3x3 _value5 = _value4.m; + s_sasm.Store3(0+96+0+0, asuint(_value5[0])); + s_sasm.Store3(0+96+0+16, asuint(_value5[1])); + s_sasm.Store3(0+96+0+32, asuint(_value5[2])); + } + } + { + StructWithMat _value4 = _value3[3]; + { + float3x3 _value5 = _value4.m; + s_sasm.Store3(0+144+0+0, asuint(_value5[0])); + s_sasm.Store3(0+144+0+16, asuint(_value5[1])); + s_sasm.Store3(0+144+0+32, asuint(_value5[2])); + } + } + } + } + { + StructWithMat _value2[4] = l_u_a; + { + StructWithMat _value3 = _value2[0]; + { + float3x3 _value4 = _value3.m; + s_sasm.Store3(0+0+0+0, asuint(_value4[0])); + s_sasm.Store3(0+0+0+16, asuint(_value4[1])); + s_sasm.Store3(0+0+0+32, asuint(_value4[2])); + } + } + { + StructWithMat _value3 = _value2[1]; + { + float3x3 _value4 = _value3.m; + s_sasm.Store3(0+48+0+0, asuint(_value4[0])); + s_sasm.Store3(0+48+0+16, asuint(_value4[1])); + s_sasm.Store3(0+48+0+32, asuint(_value4[2])); + } + } + { + StructWithMat _value3 = _value2[2]; + { + float3x3 _value4 = _value3.m; + s_sasm.Store3(0+96+0+0, asuint(_value4[0])); + s_sasm.Store3(0+96+0+16, asuint(_value4[1])); + s_sasm.Store3(0+96+0+32, asuint(_value4[2])); + } + } + { + StructWithMat _value3 = _value2[3]; + { + float3x3 _value4 = _value3.m; + s_sasm.Store3(0+144+0+0, asuint(_value4[0])); + s_sasm.Store3(0+144+0+16, asuint(_value4[1])); + s_sasm.Store3(0+144+0+32, asuint(_value4[2])); + } + } + } + { + float3x3 _value2 = l_u_m_c; + s_sasm.Store3(0+0+0+0, asuint(_value2[0])); + s_sasm.Store3(0+0+0+16, asuint(_value2[1])); + s_sasm.Store3(0+0+0+32, asuint(_value2[2])); + } + int _e238 = idx_2; + { + float3x3 _value2 = l_u_m_v; + s_sasm.Store3(0+_e238*48+0+0, asuint(_value2[0])); + s_sasm.Store3(0+_e238*48+0+16, asuint(_value2[1])); + s_sasm.Store3(0+_e238*48+0+32, asuint(_value2[2])); + } + s_sasm.Store3(0+0+0+0, asuint(l_u_c_cc)); + int _e250 = idx_2; + s_sasm.Store3(_e250*16+0+0+0, asuint(l_u_c_cv)); + int _e254 = idx_2; + s_sasm.Store3(0+0+_e254*48+0, asuint(l_u_c_vc)); + int _e260 = idx_2; + int _e263 = idx_2; + s_sasm.Store3(_e263*16+0+_e260*48+0, asuint(l_u_c_vv)); + s_sasm.Store(0+0+0+0+0, asuint(l_u_e_ccc)); + int _e276 = idx_2; + s_sasm.Store(_e276*4+0+0+0+0, asuint(l_u_e_ccv)); + int _e282 = idx_2; + s_sasm.Store(0+_e282*16+0+0+0, asuint(l_u_e_cvc)); + int _e289 = idx_2; + int _e291 = idx_2; + s_sasm.Store(_e291*4+_e289*16+0+0+0, asuint(l_u_e_cvv)); + int _e295 = idx_2; + s_sasm.Store(0+0+0+_e295*48+0, asuint(l_u_e_vcc)); + int _e302 = idx_2; + int _e306 = idx_2; + s_sasm.Store(_e306*4+0+0+_e302*48+0, asuint(l_u_e_vcv)); + int _e310 = idx_2; + int _e313 = idx_2; + s_sasm.Store(0+_e313*16+0+_e310*48+0, asuint(l_u_e_vvc)); + int _e318 = idx_2; + int _e321 = idx_2; + int _e323 = idx_2; + s_sasm.Store(_e323*4+_e321*16+0+_e318*48+0, asuint(l_u_e_vvv)); + return; +} + +[numthreads(1, 1, 1)] +void main() +{ + access_m(); + access_sm(); + access_sasm(); + return; +} diff --git a/naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.ron b/naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.ron new file mode 100644 index 0000000000..a07b03300b --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_5_1", + ), + ], +) diff --git a/naga/tests/out/spv/wgsl-access.spvasm b/naga/tests/out/spv/wgsl-access.spvasm index 8bbfd03edb..ae6a1d1c5f 100644 --- a/naga/tests/out/spv/wgsl-access.spvasm +++ b/naga/tests/out/spv/wgsl-access.spvasm @@ -49,6 +49,9 @@ var baz: Baz; var qux: vec2; fn test_matrix_within_struct_accesses() { + // Test HLSL accesses to Cx2 matrices. There are additional tests + // in `hlsl_mat_cx2.wgsl`. + var idx = 1; idx--; @@ -576,91 +579,91 @@ OpDecorate %388 Location 0 %69 = OpAccessChain %68 %56 %48 OpBranch %98 %98 = OpLabel -OpLine %3 40 5 +OpLine %3 43 5 %99 = OpLoad %6 %94 %100 = OpISub %6 %99 %70 -OpLine %3 40 5 +OpLine %3 43 5 OpStore %94 %100 -OpLine %3 43 14 +OpLine %3 46 14 %102 = OpAccessChain %101 %69 %48 %103 = OpLoad %22 %102 -OpLine %3 44 14 -OpLine %3 44 14 +OpLine %3 47 14 +OpLine %3 47 14 %105 = OpAccessChain %104 %69 %48 %48 %106 = OpLoad %13 %105 -OpLine %3 45 14 +OpLine %3 48 14 %107 = OpLoad %6 %94 %108 = OpAccessChain %104 %69 %48 %107 %109 = OpLoad %13 %108 -OpLine %3 46 14 -OpLine %3 46 14 -OpLine %3 46 14 +OpLine %3 49 14 +OpLine %3 49 14 +OpLine %3 49 14 %111 = OpAccessChain %110 %69 %48 %48 %44 %112 = OpLoad %9 %111 -OpLine %3 47 14 -OpLine %3 47 14 +OpLine %3 50 14 +OpLine %3 50 14 %113 = OpLoad %6 %94 %114 = OpAccessChain %110 %69 %48 %48 %113 %115 = OpLoad %9 %114 -OpLine %3 48 14 +OpLine %3 51 14 %116 = OpLoad %6 %94 -OpLine %3 48 14 +OpLine %3 51 14 %117 = OpAccessChain %110 %69 %48 %116 %44 %118 = OpLoad %9 %117 -OpLine %3 49 14 +OpLine %3 52 14 %119 = OpLoad %6 %94 %120 = OpLoad %6 %94 %121 = OpAccessChain %110 %69 %48 %119 %120 %122 = OpLoad %9 %121 -OpLine %3 51 29 -OpLine %3 51 45 -OpLine %3 51 13 -OpLine %3 53 5 +OpLine %3 54 29 +OpLine %3 54 45 +OpLine %3 54 13 +OpLine %3 56 5 %123 = OpLoad %6 %94 %124 = OpIAdd %6 %123 %70 -OpLine %3 53 5 -OpStore %94 %124 -OpLine %3 56 5 -OpLine %3 56 23 -OpLine %3 56 39 -OpLine %3 56 11 OpLine %3 56 5 +OpStore %94 %124 +OpLine %3 59 5 +OpLine %3 59 23 +OpLine %3 59 39 +OpLine %3 59 11 +OpLine %3 59 5 %126 = OpAccessChain %125 %96 %48 OpStore %126 %85 -OpLine %3 57 5 -OpLine %3 57 5 -OpLine %3 57 14 -OpLine %3 57 5 +OpLine %3 60 5 +OpLine %3 60 5 +OpLine %3 60 14 +OpLine %3 60 5 %128 = OpAccessChain %127 %96 %48 %48 OpStore %128 %87 -OpLine %3 58 5 +OpLine %3 61 5 %129 = OpLoad %6 %94 -OpLine %3 58 16 -OpLine %3 58 5 +OpLine %3 61 16 +OpLine %3 61 5 %130 = OpAccessChain %127 %96 %48 %129 OpStore %130 %89 -OpLine %3 59 5 -OpLine %3 59 5 -OpLine %3 59 5 -OpLine %3 59 5 +OpLine %3 62 5 +OpLine %3 62 5 +OpLine %3 62 5 +OpLine %3 62 5 %131 = OpAccessChain %28 %96 %48 %48 %44 OpStore %131 %90 -OpLine %3 60 5 -OpLine %3 60 5 +OpLine %3 63 5 +OpLine %3 63 5 %132 = OpLoad %6 %94 -OpLine %3 60 5 +OpLine %3 63 5 %133 = OpAccessChain %28 %96 %48 %48 %132 OpStore %133 %91 -OpLine %3 61 5 +OpLine %3 64 5 %134 = OpLoad %6 %94 -OpLine %3 61 5 -OpLine %3 61 5 +OpLine %3 64 5 +OpLine %3 64 5 %135 = OpAccessChain %28 %96 %48 %134 %44 OpStore %135 %92 -OpLine %3 62 5 +OpLine %3 65 5 %136 = OpLoad %6 %94 %137 = OpLoad %6 %94 -OpLine %3 62 5 +OpLine %3 65 5 %138 = OpAccessChain %28 %96 %48 %136 %137 OpStore %138 %93 OpReturn @@ -672,111 +675,111 @@ OpFunctionEnd %142 = OpAccessChain %141 %62 %48 OpBranch %153 %153 = OpLabel -OpLine %3 75 5 +OpLine %3 78 5 %154 = OpLoad %6 %150 %155 = OpISub %6 %154 %70 -OpLine %3 75 5 +OpLine %3 78 5 OpStore %150 %155 -OpLine %3 78 14 +OpLine %3 81 14 %157 = OpAccessChain %156 %142 %48 %158 = OpLoad %26 %157 -OpLine %3 79 14 -OpLine %3 79 14 +OpLine %3 82 14 +OpLine %3 82 14 %160 = OpAccessChain %159 %142 %48 %48 %161 = OpLoad %25 %160 -OpLine %3 80 14 -OpLine %3 80 14 -OpLine %3 80 14 +OpLine %3 83 14 +OpLine %3 83 14 +OpLine %3 83 14 %162 = OpAccessChain %104 %142 %48 %48 %48 %163 = OpLoad %13 %162 -OpLine %3 81 14 -OpLine %3 81 14 +OpLine %3 84 14 +OpLine %3 84 14 %164 = OpLoad %6 %150 %165 = OpAccessChain %104 %142 %48 %48 %164 %166 = OpLoad %13 %165 -OpLine %3 82 14 -OpLine %3 82 14 -OpLine %3 82 14 -OpLine %3 82 14 +OpLine %3 85 14 +OpLine %3 85 14 +OpLine %3 85 14 +OpLine %3 85 14 %167 = OpAccessChain %110 %142 %48 %48 %48 %44 %168 = OpLoad %9 %167 -OpLine %3 83 14 -OpLine %3 83 14 -OpLine %3 83 14 +OpLine %3 86 14 +OpLine %3 86 14 +OpLine %3 86 14 %169 = OpLoad %6 %150 %170 = OpAccessChain %110 %142 %48 %48 %48 %169 %171 = OpLoad %9 %170 -OpLine %3 84 14 -OpLine %3 84 14 +OpLine %3 87 14 +OpLine %3 87 14 %172 = OpLoad %6 %150 -OpLine %3 84 14 +OpLine %3 87 14 %173 = OpAccessChain %110 %142 %48 %48 %172 %44 %174 = OpLoad %9 %173 -OpLine %3 85 14 -OpLine %3 85 14 +OpLine %3 88 14 +OpLine %3 88 14 %175 = OpLoad %6 %150 %176 = OpLoad %6 %150 %177 = OpAccessChain %110 %142 %48 %48 %175 %176 %178 = OpLoad %9 %177 -OpLine %3 87 13 -OpLine %3 89 5 +OpLine %3 90 13 +OpLine %3 92 5 %179 = OpLoad %6 %150 %180 = OpIAdd %6 %179 %70 -OpLine %3 89 5 -OpStore %150 %180 -OpLine %3 92 5 OpLine %3 92 5 +OpStore %150 %180 +OpLine %3 95 5 +OpLine %3 95 5 %182 = OpAccessChain %181 %151 %48 OpStore %182 %143 -OpLine %3 93 5 -OpLine %3 93 5 -OpLine %3 93 27 -OpLine %3 93 43 -OpLine %3 93 59 -OpLine %3 93 15 -OpLine %3 93 5 +OpLine %3 96 5 +OpLine %3 96 5 +OpLine %3 96 27 +OpLine %3 96 43 +OpLine %3 96 59 +OpLine %3 96 15 +OpLine %3 96 5 %184 = OpAccessChain %183 %151 %48 %48 OpStore %184 %149 -OpLine %3 94 5 -OpLine %3 94 5 -OpLine %3 94 5 -OpLine %3 94 18 -OpLine %3 94 5 +OpLine %3 97 5 +OpLine %3 97 5 +OpLine %3 97 5 +OpLine %3 97 18 +OpLine %3 97 5 %185 = OpAccessChain %127 %151 %48 %48 %48 OpStore %185 %87 -OpLine %3 95 5 -OpLine %3 95 5 +OpLine %3 98 5 +OpLine %3 98 5 %186 = OpLoad %6 %150 -OpLine %3 95 20 -OpLine %3 95 5 +OpLine %3 98 20 +OpLine %3 98 5 %187 = OpAccessChain %127 %151 %48 %48 %186 OpStore %187 %89 -OpLine %3 96 5 -OpLine %3 96 5 -OpLine %3 96 5 -OpLine %3 96 5 -OpLine %3 96 5 +OpLine %3 99 5 +OpLine %3 99 5 +OpLine %3 99 5 +OpLine %3 99 5 +OpLine %3 99 5 %188 = OpAccessChain %28 %151 %48 %48 %48 %44 OpStore %188 %90 -OpLine %3 97 5 -OpLine %3 97 5 -OpLine %3 97 5 +OpLine %3 100 5 +OpLine %3 100 5 +OpLine %3 100 5 %189 = OpLoad %6 %150 -OpLine %3 97 5 +OpLine %3 100 5 %190 = OpAccessChain %28 %151 %48 %48 %48 %189 OpStore %190 %91 -OpLine %3 98 5 -OpLine %3 98 5 +OpLine %3 101 5 +OpLine %3 101 5 %191 = OpLoad %6 %150 -OpLine %3 98 5 -OpLine %3 98 5 +OpLine %3 101 5 +OpLine %3 101 5 %192 = OpAccessChain %28 %151 %48 %48 %191 %44 OpStore %192 %92 -OpLine %3 99 5 -OpLine %3 99 5 +OpLine %3 102 5 +OpLine %3 102 5 %193 = OpLoad %6 %150 %194 = OpLoad %6 %150 -OpLine %3 99 5 +OpLine %3 102 5 %195 = OpAccessChain %28 %151 %48 %48 %193 %194 OpStore %195 %93 OpReturn @@ -786,7 +789,7 @@ OpFunctionEnd %196 = OpLabel OpBranch %200 %200 = OpLabel -OpLine %3 102 22 +OpLine %3 105 22 %201 = OpLoad %9 %197 OpReturnValue %201 OpFunctionEnd @@ -795,9 +798,9 @@ OpFunctionEnd %202 = OpLabel OpBranch %206 %206 = OpLabel -OpLine %3 107 12 +OpLine %3 110 12 %207 = OpCompositeExtract %29 %203 4 -OpLine %3 107 12 +OpLine %3 110 12 %208 = OpCompositeExtract %9 %207 9 OpReturnValue %208 OpFunctionEnd @@ -806,7 +809,7 @@ OpFunctionEnd %209 = OpLabel OpBranch %214 %214 = OpLabel -OpLine %3 156 5 +OpLine %3 159 5 OpStore %210 %213 OpReturn OpFunctionEnd @@ -815,11 +818,11 @@ OpFunctionEnd %215 = OpLabel OpBranch %222 %222 = OpLabel -OpLine %3 160 32 -OpLine %3 160 43 -OpLine %3 160 32 -OpLine %3 160 12 -OpLine %3 160 5 +OpLine %3 163 32 +OpLine %3 163 43 +OpLine %3 163 32 +OpLine %3 163 12 +OpLine %3 163 5 OpStore %216 %221 OpReturn OpFunctionEnd @@ -829,13 +832,13 @@ OpFunctionEnd %230 = OpVariable %36 Function %228 OpBranch %231 %231 = OpLabel -OpLine %3 165 5 -%232 = OpFunctionCall %2 %211 %229 -OpLine %3 167 35 -OpLine %3 167 46 -OpLine %3 167 35 -OpLine %3 167 15 OpLine %3 168 5 +%232 = OpFunctionCall %2 %211 %229 +OpLine %3 170 35 +OpLine %3 170 46 +OpLine %3 170 35 +OpLine %3 170 15 +OpLine %3 171 5 %233 = OpFunctionCall %2 %217 %230 OpReturn OpFunctionEnd @@ -844,7 +847,7 @@ OpFunctionEnd %234 = OpLabel OpBranch %238 %238 = OpLabel -OpLine %3 176 10 +OpLine %3 179 10 %239 = OpAccessChain %34 %235 %48 %240 = OpLoad %4 %239 OpReturnValue %240 @@ -854,8 +857,8 @@ OpFunctionEnd %241 = OpLabel OpBranch %245 %245 = OpLabel -OpLine %3 180 3 -OpLine %3 180 3 +OpLine %3 183 3 +OpLine %3 183 3 %246 = OpAccessChain %34 %242 %48 OpStore %246 %17 OpReturn @@ -865,7 +868,7 @@ OpFunctionEnd %247 = OpLabel OpBranch %251 %251 = OpLabel -OpLine %3 184 10 +OpLine %3 187 10 %252 = OpAccessChain %34 %248 %44 %253 = OpLoad %4 %252 OpReturnValue %253 @@ -875,8 +878,8 @@ OpFunctionEnd %254 = OpLabel OpBranch %258 %258 = OpLabel -OpLine %3 188 3 -OpLine %3 188 3 +OpLine %3 191 3 +OpLine %3 191 3 %259 = OpAccessChain %34 %255 %44 OpStore %259 %17 OpReturn @@ -887,13 +890,13 @@ OpFunctionEnd %264 = OpVariable %41 Function %265 OpBranch %266 %266 = OpLabel -OpLine %3 193 4 +OpLine %3 196 4 %267 = OpFunctionCall %2 %243 %262 -OpLine %3 194 4 -%268 = OpFunctionCall %4 %236 %262 OpLine %3 197 4 +%268 = OpFunctionCall %4 %236 %262 +OpLine %3 200 4 %269 = OpFunctionCall %2 %256 %264 -OpLine %3 198 4 +OpLine %3 201 4 %270 = OpFunctionCall %4 %249 %264 OpReturn OpFunctionEnd @@ -903,11 +906,11 @@ OpFunctionEnd %275 = OpVariable %276 Function %277 OpBranch %278 %278 = OpLabel -OpLine %3 202 13 +OpLine %3 205 13 %279 = OpCompositeConstruct %43 %272 -OpLine %3 202 5 +OpLine %3 205 5 OpStore %275 %279 -OpLine %3 204 12 +OpLine %3 207 12 %281 = OpAccessChain %280 %275 %48 %282 = OpLoad %42 %281 OpReturnValue %282 @@ -917,8 +920,8 @@ OpFunctionEnd %288 = OpVariable %289 Function %287 OpBranch %290 %290 = OpLabel -OpLine %3 210 16 -OpLine %3 212 12 +OpLine %3 213 16 +OpLine %3 215 12 %291 = OpAccessChain %95 %288 %48 %292 = OpLoad %6 %291 OpReturnValue %292 @@ -927,19 +930,19 @@ OpFunctionEnd %293 = OpLabel OpBranch %296 %296 = OpLabel -OpLine %3 222 17 +OpLine %3 225 17 %297 = OpCompositeExtract %46 %295 0 -OpLine %3 223 20 +OpLine %3 226 20 %298 = OpCompositeExtract %6 %297 0 -OpLine %3 225 9 +OpLine %3 228 9 %299 = OpCompositeExtract %4 %295 1 %300 = OpBitcast %4 %298 %301 = OpINotEqual %42 %299 %300 -OpLine %3 225 5 +OpLine %3 228 5 OpSelectionMerge %302 None OpBranchConditional %301 %302 %302 %302 = OpLabel -OpLine %3 229 12 +OpLine %3 232 12 %303 = OpCompositeExtract %46 %295 0 %304 = OpCompositeExtract %6 %303 0 OpReturnValue %304 @@ -951,27 +954,27 @@ OpFunctionEnd %312 = OpVariable %95 Function %313 OpBranch %314 %314 = OpLabel -OpLine %3 235 17 +OpLine %3 238 17 %315 = OpAccessChain %310 %307 %48 %316 = OpLoad %46 %315 -OpLine %3 235 5 +OpLine %3 238 5 OpStore %309 %316 -OpLine %3 236 20 +OpLine %3 239 20 %317 = OpAccessChain %95 %309 %48 %318 = OpLoad %6 %317 -OpLine %3 236 5 +OpLine %3 239 5 OpStore %312 %318 -OpLine %3 238 9 +OpLine %3 241 9 %319 = OpAccessChain %34 %307 %44 %320 = OpLoad %4 %319 %321 = OpLoad %6 %312 %322 = OpBitcast %4 %321 %323 = OpINotEqual %42 %320 %322 -OpLine %3 238 5 +OpLine %3 241 5 OpSelectionMerge %324 None OpBranchConditional %323 %324 %324 %324 = OpLabel -OpLine %3 242 12 +OpLine %3 245 12 %325 = OpAccessChain %95 %307 %48 %48 %326 = OpLoad %6 %325 OpReturnValue %326 @@ -988,58 +991,58 @@ OpBranch %348 %348 = OpLabel OpLine %3 1 1 %349 = OpLoad %9 %344 -OpLine %3 115 5 +OpLine %3 118 5 OpStore %344 %71 -OpLine %3 117 9 +OpLine %3 120 9 %350 = OpLoad %7 %52 -OpLine %3 118 5 +OpLine %3 121 5 %351 = OpFunctionCall %2 %66 -OpLine %3 119 5 +OpLine %3 122 5 %352 = OpFunctionCall %2 %140 -OpLine %3 122 19 +OpLine %3 125 19 %354 = OpAccessChain %353 %54 %48 %355 = OpLoad %10 %354 -OpLine %3 123 15 +OpLine %3 126 15 %357 = OpAccessChain %356 %54 %40 %358 = OpLoad %19 %357 -OpLine %3 125 13 +OpLine %3 128 13 %361 = OpAccessChain %360 %54 %48 %339 %48 %362 = OpLoad %9 %361 -OpLine %3 126 13 -OpLine %3 126 22 +OpLine %3 129 13 +OpLine %3 129 22 %364 = OpArrayLength %4 %54 5 -OpLine %3 126 13 +OpLine %3 129 13 %365 = OpISub %4 %364 %15 %368 = OpAccessChain %367 %54 %31 %365 %48 %369 = OpLoad %6 %368 -OpLine %3 127 13 +OpLine %3 130 13 %370 = OpLoad %24 %336 -OpLine %3 130 56 -OpLine %3 130 56 -OpLine %3 131 21 +OpLine %3 133 56 +OpLine %3 133 56 +OpLine %3 134 21 %371 = OpFunctionCall %9 %198 %344 -OpLine %3 134 31 +OpLine %3 137 31 %374 = OpExtInst %9 %1 FClamp %362 %372 %373 %375 = OpConvertFToS %6 %374 -OpLine %3 134 14 +OpLine %3 137 14 %376 = OpCompositeConstruct %33 %369 %375 %340 %341 %342 -OpLine %3 134 5 +OpLine %3 137 5 OpStore %345 %376 -OpLine %3 135 5 +OpLine %3 138 5 %377 = OpIAdd %4 %330 %44 -OpLine %3 135 5 +OpLine %3 138 5 %378 = OpAccessChain %95 %345 %377 OpStore %378 %286 -OpLine %3 136 17 +OpLine %3 139 17 %379 = OpAccessChain %95 %345 %330 %380 = OpLoad %6 %379 -OpLine %3 138 5 +OpLine %3 141 5 %381 = OpFunctionCall %9 %204 %343 -OpLine %3 140 22 +OpLine %3 143 22 %383 = OpCompositeConstruct %382 %380 %380 %380 %380 %384 = OpConvertSToF %32 %383 %385 = OpMatrixTimesVector %11 %355 %384 -OpLine %3 140 12 +OpLine %3 143 12 %386 = OpCompositeConstruct %32 %385 %73 OpStore %331 %386 OpReturn @@ -1049,33 +1052,33 @@ OpFunctionEnd %390 = OpAccessChain %335 %59 %48 OpBranch %401 %401 = OpLabel -OpLine %3 146 5 -OpLine %3 146 5 -OpLine %3 146 5 +OpLine %3 149 5 +OpLine %3 149 5 +OpLine %3 149 5 %402 = OpAccessChain %360 %54 %48 %44 %15 OpStore %402 %71 -OpLine %3 147 5 -OpLine %3 147 31 -OpLine %3 147 47 -OpLine %3 147 63 -OpLine %3 147 19 -OpLine %3 147 5 +OpLine %3 150 5 +OpLine %3 150 31 +OpLine %3 150 47 +OpLine %3 150 63 +OpLine %3 150 19 +OpLine %3 150 5 %403 = OpAccessChain %353 %54 %48 OpStore %403 %395 -OpLine %3 148 5 -OpLine %3 148 35 -OpLine %3 148 15 -OpLine %3 148 5 +OpLine %3 151 5 +OpLine %3 151 35 +OpLine %3 151 15 +OpLine %3 151 5 %404 = OpAccessChain %356 %54 %40 OpStore %404 %398 -OpLine %3 149 5 -OpLine %3 149 5 -OpLine %3 149 5 +OpLine %3 152 5 +OpLine %3 152 5 +OpLine %3 152 5 %405 = OpAccessChain %367 %54 %31 %44 %48 OpStore %405 %70 -OpLine %3 150 5 +OpLine %3 153 5 OpStore %390 %399 -OpLine %3 152 12 +OpLine %3 155 12 OpStore %388 %400 OpReturn OpFunctionEnd @@ -1083,17 +1086,17 @@ OpFunctionEnd %406 = OpLabel OpBranch %409 %409 = OpLabel -OpLine %3 247 5 +OpLine %3 250 5 %410 = OpFunctionCall %2 %224 -OpLine %3 248 5 +OpLine %3 251 5 %411 = OpFunctionCall %2 %261 -OpLine %3 249 5 +OpLine %3 252 5 %412 = OpFunctionCall %42 %273 %408 -OpLine %3 250 5 +OpLine %3 253 5 %413 = OpFunctionCall %6 %284 -OpLine %3 251 5 +OpLine %3 254 5 %414 = OpFunctionCall %6 %294 -OpLine %3 252 5 +OpLine %3 255 5 %415 = OpFunctionCall %6 %306 OpReturn OpFunctionEnd \ No newline at end of file