nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-02-17 01:12:41 +00:00
Code Issues Packages Projects Releases Wiki Activity

hlsl: storage reads support, array length

Browse Source
This commit is contained in:
Dzmitry Malyshau 2021-07-24 02:21:54 -04:00 committed by Dzmitry Malyshau
parent 526a29e2b8
commit 4b6846d5da
20 changed files with 1204 additions and 699 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

123
src/back/hlsl/conv.rs Normal file
View File

@ -0,0 +1,123 @@
use super::Error;
impl crate::ScalarKind {
pub(super) fn to_hlsl_cast(self) -> &'static str {
match self {
Self::Float => "asfloat",
Self::Sint => "asint",
Self::Uint => "asuint",
Self::Bool => unreachable!(),
}
}
/// Helper function that returns scalar related strings
/// https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-scalar
pub(super) fn to_hlsl_str(self, width: crate::Bytes) -> Result<&'static str, Error> {
match self {
Self::Sint => Ok("int"),
Self::Uint => Ok("uint"),
Self::Float => match width {
2 => Ok("half"),
4 => Ok("float"),
8 => Ok("double"),
_ => Err(Error::UnsupportedScalar(self, width)),
},
Self::Bool => Ok("bool"),
}
}
}
impl crate::TypeInner {
pub(super) fn is_matrix(&self) -> bool {
match *self {
Self::Matrix { .. } => true,
_ => false,
}
}
}
impl crate::StorageFormat {
pub(super) fn to_hlsl_str(self) -> &'static str {
match self {
Self::R16Float => "float",
Self::R8Unorm => "unorm float",
Self::R8Snorm => "snorm float",
Self::R8Uint | Self::R16Uint => "uint",
Self::R8Sint | Self::R16Sint => "int",
Self::Rg16Float => "float2",
Self::Rg8Unorm => "unorm float2",
Self::Rg8Snorm => "snorm float2",
Self::Rg8Sint | Self::Rg16Sint => "int2",
Self::Rg8Uint | Self::Rg16Uint => "uint2",
Self::Rg11b10Float => "float3",
Self::Rgba16Float | Self::R32Float | Self::Rg32Float | Self::Rgba32Float => "float4",
Self::Rgba8Unorm | Self::Rgb10a2Unorm => "unorm float4",
Self::Rgba8Snorm => "snorm float4",
Self::Rgba8Uint
| Self::Rgba16Uint
| Self::R32Uint
| Self::Rg32Uint
| Self::Rgba32Uint => "uint4",
Self::Rgba8Sint
| Self::Rgba16Sint
| Self::R32Sint
| Self::Rg32Sint
| Self::Rgba32Sint => "int4",
}
}
}
impl crate::BuiltIn {
pub(super) fn to_hlsl_str(self) -> Result<&'static str, Error> {
Ok(match self {
Self::Position => "SV_Position",
// vertex
Self::ClipDistance => "SV_ClipDistance",
Self::CullDistance => "SV_CullDistance",
Self::InstanceIndex => "SV_InstanceID",
// based on this page https://docs.microsoft.com/en-us/windows/uwp/gaming/glsl-to-hlsl-reference#comparing-opengl-es-20-with-direct3d-11
// No meaning unless you target Direct3D 9
Self::PointSize => "PSIZE",
Self::VertexIndex => "SV_VertexID",
// fragment
Self::FragDepth => "SV_Depth",
Self::FrontFacing => "SV_IsFrontFace",
Self::PrimitiveIndex => "SV_PrimitiveID",
Self::SampleIndex => "SV_SampleIndex",
Self::SampleMask => "SV_Coverage",
// compute
Self::GlobalInvocationId => "SV_DispatchThreadID",
Self::LocalInvocationId => "SV_GroupThreadID",
Self::LocalInvocationIndex => "SV_GroupIndex",
Self::WorkGroupId => "SV_GroupID",
_ => return Err(Error::Unimplemented(format!("builtin {:?}", self))),
})
}
}
impl crate::Interpolation {
/// Helper function that returns the string corresponding to the HLSL interpolation qualifier
pub(super) fn to_hlsl_str(self) -> &'static str {
match self {
Self::Perspective => "linear",
Self::Linear => "noperspective",
Self::Flat => "nointerpolation",
}
}
}
impl crate::Sampling {
/// Return the HLSL auxiliary qualifier for the given sampling value.
pub(super) fn to_hlsl_str(self) -> Option<&'static str> {
match self {
Self::Center => None,
Self::Centroid => Some("centroid"),
Self::Sample => Some("sample"),
}
}
}

425
src/back/hlsl/help.rs Normal file
View File

@ -0,0 +1,425 @@
// Important note about `Expression::ImageQuery`/`Expression::ArrayLength` and hlsl backend:
// Due to implementation of `GetDimensions` function in hlsl (https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions)
// backend can't work with it as an expression.
// Instead, it generates a unique wrapped function per `Expression::ImageQuery`, based on texure info and query function.
// See `WrappedImageQuery` struct that represents a unique function and will be generated before writing all statements and expressions.
// This allowed to works with `Expression::ImageQuery` as expression and write wrapped function.
//
// For example:
// ```wgsl
// let dim_1d = textureDimensions(image_1d);
// ```
//
// ```hlsl
// int NagaDimensions1D(Texture1D<float4>)
// {
// uint4 ret;
// image_1d.GetDimensions(ret.x);
// return ret.x;
// }
//
// int dim_1d = NagaDimensions1D(image_1d);
// ```
use super::{super::FunctionCtx, writer::BackendResult};
use crate::arena::Handle;
use std::fmt::Write;
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
pub(super) struct WrappedArrayLength {
pub(super) writable: bool,
}
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
pub(super) struct WrappedImageQuery {
pub(super) dim: crate::ImageDimension,
pub(super) arrayed: bool,
pub(super) class: crate::ImageClass,
pub(super) query: ImageQuery,
}
// HLSL backend requires its own `ImageQuery` enum.
// It is used inside `WrappedImageQuery` and should be unique per ImageQuery function.
// IR version can't be unique per function, because it's store mipmap level as an expression.
//
// For example:
// ```wgsl
// let dim_cube_array_lod = textureDimensions(image_cube_array, 1);
// let dim_cube_array_lod2 = textureDimensions(image_cube_array, 1);
// ```
//
// ```ir
// ImageQuery {
// image: [1],
// query: Size {
// level: Some(
// [1],
// ),
// },
// },
// ImageQuery {
// image: [1],
// query: Size {
// level: Some(
// [2],
// ),
// },
// },
// ```
//
// HLSL should generate only 1 function for this case.
//
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
pub(super) enum ImageQuery {
Size,
SizeLevel,
NumLevels,
NumLayers,
NumSamples,
}
impl From<crate::ImageQuery> for ImageQuery {
fn from(q: crate::ImageQuery) -> Self {
use crate::ImageQuery as Iq;
match q {
Iq::Size { level: Some(_) } => ImageQuery::SizeLevel,
Iq::Size { level: None } => ImageQuery::Size,
Iq::NumLevels => ImageQuery::NumLevels,
Iq::NumLayers => ImageQuery::NumLayers,
Iq::NumSamples => ImageQuery::NumSamples,
}
}
}
#[derive(Clone, Copy, PartialEq)]
pub(super) enum MipLevelCoordinate {
NotApplicable,
Zero,
Expression(Handle<crate::Expression>),
}
impl<'a, W: Write> super::Writer<'a, W> {
pub(super) fn write_image_type(
&mut self,
dim: crate::ImageDimension,
arrayed: bool,
class: crate::ImageClass,
) -> BackendResult {
let dim_str = dim.to_hlsl_str();
let arrayed_str = if arrayed { "Array" } else { "" };
write!(self.out, "Texture{}{}", dim_str, arrayed_str)?;
match class {
crate::ImageClass::Depth { multi } => {
let multi_str = if multi { "MS" } else { "" };
write!(self.out, "{}<float>", multi_str)?
}
crate::ImageClass::Sampled { kind, multi } => {
let multi_str = if multi { "MS" } else { "" };
let scalar_kind_str = kind.to_hlsl_str(4)?;
write!(self.out, "{}<{}4>", multi_str, scalar_kind_str)?
}
crate::ImageClass::Storage(format) => {
let storage_format_str = format.to_hlsl_str();
write!(self.out, "<{}>", storage_format_str)?
}
}
Ok(())
}
pub(super) fn write_wrapped_array_length_function_name(
&mut self,
query: WrappedArrayLength,
) -> BackendResult {
let access_str = if query.writable { "RW" } else { "" };
write!(self.out, "NagaBufferLength{}", access_str,)?;
Ok(())
}
/// Helper function that write wrapped function for `Expression::ArrayLength`
/// https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/sm5-object-rwbyteaddressbuffer-getdimensions
pub(super) fn write_wrapped_array_length_function(
&mut self,
module: &crate::Module,
wal: WrappedArrayLength,
expr_handle: Handle<crate::Expression>,
func_ctx: &FunctionCtx,
) -> BackendResult {
use crate::back::INDENT;
const ARGUMENT_VARIABLE_NAME: &str = "buffer";
const RETURN_VARIABLE_NAME: &str = "ret";
// Write function return type and name
let ret_ty = func_ctx.info[expr_handle].ty.inner_with(&module.types);
self.write_value_type(module, ret_ty)?;
write!(self.out, " ")?;
self.write_wrapped_array_length_function_name(wal)?;
// Write function parameters
write!(self.out, "(")?;
let access_str = if wal.writable { "RW" } else { "" };
writeln!(
self.out,
"{}ByteAddressBuffer {})",
access_str, ARGUMENT_VARIABLE_NAME
)?;
// Write function body
writeln!(self.out, "{{")?;
// Write `GetDimensions` function.
writeln!(self.out, "{}uint {};", INDENT, RETURN_VARIABLE_NAME)?;
writeln!(
self.out,
"{}{}.GetDimensions({});",
INDENT, ARGUMENT_VARIABLE_NAME, RETURN_VARIABLE_NAME
)?;
// Write return value
writeln!(self.out, "{}return {};", INDENT, RETURN_VARIABLE_NAME)?;
// End of function body
writeln!(self.out, "}}")?;
// Write extra new line
writeln!(self.out)?;
Ok(())
}
pub(super) fn write_wrapped_image_query_function_name(
&mut self,
query: WrappedImageQuery,
) -> BackendResult {
let dim_str = query.dim.to_hlsl_str();
let class_str = match query.class {
crate::ImageClass::Sampled { multi: true, .. } => "MS",
crate::ImageClass::Depth { multi: true } => "DepthMS",
crate::ImageClass::Depth { multi: false } => "Depth",
crate::ImageClass::Sampled { multi: false, .. } | crate::ImageClass::Storage { .. } => {
""
}
};
let arrayed_str = if query.arrayed { "Array" } else { "" };
let query_str = match query.query {
ImageQuery::Size => "Dimensions",
ImageQuery::SizeLevel => "MipDimensions",
ImageQuery::NumLevels => "NumLevels",
ImageQuery::NumLayers => "NumLayers",
ImageQuery::NumSamples => "NumSamples",
};
write!(
self.out,
"Naga{}{}{}{}",
class_str, query_str, dim_str, arrayed_str
)?;
Ok(())
}
/// Helper function that write wrapped function for `Expression::ImageQuery`
/// https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions
pub(super) fn write_wrapped_image_query_function(
&mut self,
module: &crate::Module,
wiq: WrappedImageQuery,
expr_handle: Handle<crate::Expression>,
func_ctx: &FunctionCtx,
) -> BackendResult {
use crate::{back::INDENT, ImageDimension as IDim};
const ARGUMENT_VARIABLE_NAME: &str = "texture";
const RETURN_VARIABLE_NAME: &str = "ret";
const MIP_LEVEL_PARAM: &str = "mip_level";
// Write function return type and name
let ret_ty = func_ctx.info[expr_handle].ty.inner_with(&module.types);
self.write_value_type(module, ret_ty)?;
write!(self.out, " ")?;
self.write_wrapped_image_query_function_name(wiq)?;
// Write function parameters
write!(self.out, "(")?;
// Texture always first parameter
self.write_image_type(wiq.dim, wiq.arrayed, wiq.class)?;
write!(self.out, " {}", ARGUMENT_VARIABLE_NAME)?;
// Mipmap is a second parameter if exists
if let ImageQuery::SizeLevel = wiq.query {
write!(self.out, ", uint {}", MIP_LEVEL_PARAM)?;
}
writeln!(self.out, ")")?;
// Write function body
writeln!(self.out, "{{")?;
let array_coords = if wiq.arrayed { 1 } else { 0 };
// GetDimensions Overloaded Methods
// https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions#overloaded-methods
let (ret_swizzle, number_of_params) = match wiq.query {
ImageQuery::Size | ImageQuery::SizeLevel => match wiq.dim {
IDim::D1 => ("x", 1 + array_coords),
IDim::D2 => ("xy", 3 + array_coords),
IDim::D3 => ("xyz", 4),
IDim::Cube => ("xy", 3 + array_coords),
},
ImageQuery::NumLevels | ImageQuery::NumSamples | ImageQuery::NumLayers => {
if wiq.arrayed || wiq.dim == IDim::D3 {
("w", 4)
} else {
("z", 3)
}
}
};
// Write `GetDimensions` function.
writeln!(self.out, "{}uint4 {};", INDENT, RETURN_VARIABLE_NAME)?;
write!(
self.out,
"{}{}.GetDimensions(",
INDENT, ARGUMENT_VARIABLE_NAME
)?;
match wiq.query {
ImageQuery::SizeLevel => {
write!(self.out, "{}, ", MIP_LEVEL_PARAM)?;
}
_ => match wiq.class {
crate::ImageClass::Sampled { multi: true, .. }
| crate::ImageClass::Depth { multi: true } => {}
_ => match wiq.dim {
// Write zero mipmap level for supported types
IDim::D2 | IDim::D3 | IDim::Cube => {
write!(self.out, "0, ")?;
}
IDim::D1 => {}
},
},
}
for component in crate::back::COMPONENTS[..number_of_params - 1].iter() {
write!(self.out, "{}.{}, ", RETURN_VARIABLE_NAME, component)?;
}
// write last parameter without comma and space for last parameter
write!(
self.out,
"{}.{}",
RETURN_VARIABLE_NAME,
crate::back::COMPONENTS[number_of_params - 1]
)?;
writeln!(self.out, ");")?;
// Write return value
writeln!(
self.out,
"{}return {}.{};",
INDENT, RETURN_VARIABLE_NAME, ret_swizzle
)?;
// End of function body
writeln!(self.out, "}}")?;
// Write extra new line
writeln!(self.out)?;
Ok(())
}
/// Helper function that write wrapped function for `Expression::ImageQuery` and `Expression::ArrayLength`
/// https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions
pub(super) fn write_wrapped_functions(
&mut self,
module: &crate::Module,
func_ctx: &FunctionCtx,
) -> BackendResult {
for (handle, _) in func_ctx.expressions.iter() {
match func_ctx.expressions[handle] {
crate::Expression::ArrayLength(expr) => {
let global_expr = match func_ctx.expressions[expr] {
crate::Expression::AccessIndex { base, index: _ } => base,
ref other => unreachable!("Array length of {:?}", other),
};
let global_var = match func_ctx.expressions[global_expr] {
crate::Expression::GlobalVariable(var_handle) => {
&module.global_variables[var_handle]
}
ref other => unreachable!("Array length of base {:?}", other),
};
let wal = WrappedArrayLength {
writable: global_var
.storage_access
.contains(crate::StorageAccess::STORE),
};
if !self.wrapped_array_lengths.contains(&wal) {
self.write_wrapped_array_length_function(module, wal, handle, func_ctx)?;
self.wrapped_array_lengths.insert(wal);
}
}
crate::Expression::ImageQuery { image, query } => {
let wiq = match *func_ctx.info[image].ty.inner_with(&module.types) {
crate::TypeInner::Image {
dim,
arrayed,
class,
} => WrappedImageQuery {
dim,
arrayed,
class,
query: query.into(),
},
_ => unreachable!("we only query images"),
};
if !self.wrapped_image_queries.contains(&wiq) {
self.write_wrapped_image_query_function(module, wiq, handle, func_ctx)?;
self.wrapped_image_queries.insert(wiq);
}
}
_ => {}
};
}
Ok(())
}
pub(super) fn write_texture_coordinates(
&mut self,
kind: &str,
coordinate: Handle<crate::Expression>,
array_index: Option<Handle<crate::Expression>>,
mip_level: MipLevelCoordinate,
module: &crate::Module,
func_ctx: &FunctionCtx,
) -> BackendResult {
// HLSL expects the array index to be merged with the coordinate
let extra = array_index.is_some() as usize
+ (mip_level != MipLevelCoordinate::NotApplicable) as usize;
if extra == 0 {
self.write_expr(module, coordinate, func_ctx)?;
} else {
let num_coords = match *func_ctx.info[coordinate].ty.inner_with(&module.types) {
crate::TypeInner::Scalar { .. } => 1,
crate::TypeInner::Vector { size, .. } => size as usize,
_ => unreachable!(),
};
write!(self.out, "{}{}(", kind, num_coords + extra)?;
self.write_expr(module, coordinate, func_ctx)?;
if let Some(expr) = array_index {
write!(self.out, ", ")?;
self.write_expr(module, expr, func_ctx)?;
}
match mip_level {
MipLevelCoordinate::NotApplicable => {}
MipLevelCoordinate::Zero => {
write!(self.out, ", 0")?;
}
MipLevelCoordinate::Expression(expr) => {
write!(self.out, ", ")?;
self.write_expr(module, expr, func_ctx)?;
}
}
write!(self.out, ")")?;
}
Ok(())
}
}

299
src/back/hlsl/image.rs
View File

@ -1,299 +0,0 @@
// Important note about `Expression::ImageQuery` and hlsl backend:
// Due to implementation of `GetDimensions` function in hlsl (https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions)
// backend can't work with it as an expression.
// Instead, it generates a unique wrapped function per `Expression::ImageQuery`, based on texure info and query function.
// See `WrappedImageQuery` struct that represents a unique function and will be generated before writing all statements and expressions.
// This allowed to works with `Expression::ImageQuery` as expression and write wrapped function.
//
// For example:
// ```wgsl
// let dim_1d = textureDimensions(image_1d);
// ```
//
// ```hlsl
// int NagaDimensions1D(Texture1D<float4>)
// {
// uint4 ret;
// image_1d.GetDimensions(ret.x);
// return ret.x;
// }
//
// int dim_1d = NagaDimensions1D(image_1d);
// ```
use super::{super::FunctionCtx, writer::BackendResult};
use crate::arena::Handle;
use std::fmt::Write;
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
pub(super) struct WrappedImageQuery {
pub(super) dim: crate::ImageDimension,
pub(super) arrayed: bool,
pub(super) class: crate::ImageClass,
pub(super) query: ImageQuery,
}
// HLSL backend requires its own `ImageQuery` enum.
// It is used inside `WrappedImageQuery` and should be unique per ImageQuery function.
// IR version can't be unique per function, because it's store mipmap level as an expression.
//
// For example:
// ```wgsl
// let dim_cube_array_lod = textureDimensions(image_cube_array, 1);
// let dim_cube_array_lod2 = textureDimensions(image_cube_array, 1);
// ```
//
// ```ir
// ImageQuery {
// image: [1],
// query: Size {
// level: Some(
// [1],
// ),
// },
// },
// ImageQuery {
// image: [1],
// query: Size {
// level: Some(
// [2],
// ),
// },
// },
// ```
//
// HLSL should generate only 1 function for this case.
//
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
pub(super) enum ImageQuery {
Size,
SizeLevel,
NumLevels,
NumLayers,
NumSamples,
}
impl From<crate::ImageQuery> for ImageQuery {
fn from(q: crate::ImageQuery) -> Self {
use crate::ImageQuery as Iq;
match q {
Iq::Size { level: Some(_) } => ImageQuery::SizeLevel,
Iq::Size { level: None } => ImageQuery::Size,
Iq::NumLevels => ImageQuery::NumLevels,
Iq::NumLayers => ImageQuery::NumLayers,
Iq::NumSamples => ImageQuery::NumSamples,
}
}
}
#[derive(Clone, Copy, PartialEq)]
pub(super) enum MipLevelCoordinate {
NotApplicable,
Zero,
Expression(Handle<crate::Expression>),
}
impl<'a, W: Write> super::Writer<'a, W> {
pub(super) fn write_wrapped_image_query_function_name(
&mut self,
query: WrappedImageQuery,
) -> BackendResult {
let dim_str = query.dim.to_hlsl_str();
let class_str = match query.class {
crate::ImageClass::Sampled { multi: true, .. } => "MS",
crate::ImageClass::Depth { multi: true } => "DepthMS",
crate::ImageClass::Depth { multi: false } => "Depth",
crate::ImageClass::Sampled { multi: false, .. } | crate::ImageClass::Storage { .. } => {
""
}
};
let arrayed_str = if query.arrayed { "Array" } else { "" };
let query_str = match query.query {
ImageQuery::Size => "Dimensions",
ImageQuery::SizeLevel => "MipDimensions",
ImageQuery::NumLevels => "NumLevels",
ImageQuery::NumLayers => "NumLayers",
ImageQuery::NumSamples => "NumSamples",
};
write!(
self.out,
"Naga{}{}{}{}",
class_str, query_str, dim_str, arrayed_str
)?;
Ok(())
}
/// Helper function that write wrapped function for `Expression::ImageQuery`
/// https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions
pub(super) fn write_wrapped_image_query_functions(
&mut self,
module: &crate::Module,
func_ctx: &FunctionCtx,
) -> BackendResult {
use crate::{back::INDENT, ImageDimension as IDim, ImageQuery as Iq};
const RETURN_VARIABLE_NAME: &str = "ret";
const MIP_LEVEL_PARAM: &str = "MipLevel";
for (handle, _) in func_ctx.expressions.iter() {
let (image, query) = match func_ctx.expressions[handle] {
crate::Expression::ImageQuery { image, query } => (image, query),
_ => continue,
};
let image_ty = func_ctx.info[image].ty.inner_with(&module.types);
let ret_ty = func_ctx.info[handle].ty.inner_with(&module.types);
let (dim, arrayed, class) = match *image_ty {
crate::TypeInner::Image {
dim,
arrayed,
class,
} => (dim, arrayed, class),
_ => unreachable!("we only query images"),
};
let wrapped_image_query = WrappedImageQuery {
dim,
arrayed,
class,
query: query.into(),
};
if self.wrapped_image_queries.contains(&wrapped_image_query) {
continue;
}
// Write function return type and name
self.write_value_type(module, ret_ty)?;
write!(self.out, " ")?;
self.write_wrapped_image_query_function_name(wrapped_image_query)?;
// Write function parameters
write!(self.out, "(")?;
// Texture always first parameter
self.write_value_type(module, image_ty)?;
// Mipmap is a second parameter if exists
if let crate::ImageQuery::Size { level: Some(_) } = query {
write!(self.out, ", uint {}", MIP_LEVEL_PARAM)?;
}
writeln!(self.out, ")")?;
// Write function body
writeln!(self.out, "{{")?;
let array_coords = if arrayed { 1 } else { 0 };
// GetDimensions Overloaded Methods
// https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions#overloaded-methods
let (ret_swizzle, number_of_params) = match query {
Iq::Size { .. } => match dim {
IDim::D1 => ("x", 1 + array_coords),
IDim::D2 => ("xy", 3 + array_coords),
IDim::D3 => ("xyz", 4),
IDim::Cube => ("xy", 3 + array_coords),
},
Iq::NumLevels | Iq::NumSamples | Iq::NumLayers => {
if arrayed || dim == IDim::D3 {
("w", 4)
} else {
("z", 3)
}
}
};
// Write `GetDimensions` function.
writeln!(self.out, "{}uint4 {};", INDENT, RETURN_VARIABLE_NAME)?;
write!(self.out, "{}", INDENT)?;
self.write_expr(module, image, func_ctx)?;
write!(self.out, ".GetDimensions(")?;
match query {
Iq::Size { level: Some(_) } => {
write!(self.out, "{}, ", MIP_LEVEL_PARAM)?;
}
_ => match class {
crate::ImageClass::Sampled { multi: true, .. }
| crate::ImageClass::Depth { multi: true } => {}
_ => match dim {
// Write zero mipmap level for supported types
IDim::D2 | IDim::D3 | IDim::Cube => {
write!(self.out, "0, ")?;
}
IDim::D1 => {}
},
},
}
for component in crate::back::COMPONENTS[..number_of_params - 1].iter() {
write!(self.out, "{}.{}, ", RETURN_VARIABLE_NAME, component)?;
}
// write last parameter without comma and space for last parameter
write!(
self.out,
"{}.{}",
RETURN_VARIABLE_NAME,
crate::back::COMPONENTS[number_of_params - 1]
)?;
writeln!(self.out, ");")?;
// Write return value
writeln!(
self.out,
"{}return {}.{};",
INDENT, RETURN_VARIABLE_NAME, ret_swizzle
)?;
// End of function body
writeln!(self.out, "}}")?;
// Write extra new line
writeln!(self.out)?;
self.wrapped_image_queries.insert(wrapped_image_query);
}
Ok(())
}
pub(super) fn write_texture_coordinates(
&mut self,
kind: &str,
coordinate: Handle<crate::Expression>,
array_index: Option<Handle<crate::Expression>>,
mip_level: MipLevelCoordinate,
module: &crate::Module,
func_ctx: &FunctionCtx,
) -> BackendResult {
// HLSL expects the array index to be merged with the coordinate
let extra = array_index.is_some() as usize
+ (mip_level != MipLevelCoordinate::NotApplicable) as usize;
if extra == 0 {
self.write_expr(module, coordinate, func_ctx)?;
} else {
let num_coords = match *func_ctx.info[coordinate].ty.inner_with(&module.types) {
crate::TypeInner::Scalar { .. } => 1,
crate::TypeInner::Vector { size, .. } => size as usize,
_ => unreachable!(),
};
write!(self.out, "{}{}(", kind, num_coords + extra)?;
self.write_expr(module, coordinate, func_ctx)?;
if let Some(expr) = array_index {
write!(self.out, ", ")?;
self.write_expr(module, expr, func_ctx)?;
}
match mip_level {
MipLevelCoordinate::NotApplicable => {}
MipLevelCoordinate::Zero => {
write!(self.out, ", 0")?;
}
MipLevelCoordinate::Expression(expr) => {
write!(self.out, ", ")?;
self.write_expr(module, expr, func_ctx)?;
}
}
write!(self.out, ")")?;
}
Ok(())
}
}

3
src/back/hlsl/mod.rs
View File

@ -6,7 +6,8 @@
//! - 6.0
//!
mod image;
mod conv;
mod help;
mod keywords;
mod writer;

640
src/back/hlsl/writer.rs
View File

@ -1,5 +1,5 @@
use super::{
image::{MipLevelCoordinate, WrappedImageQuery},
help::{MipLevelCoordinate, WrappedArrayLength, WrappedImageQuery},
Error, Options,
};
use crate::{
@ -7,20 +7,25 @@ use crate::{
proc::{self, NameKey},
valid, Handle, Module, ShaderStage, TypeInner,
};
use std::fmt::Write;
use std::{fmt, mem};
const LOCATION_SEMANTIC: &str = "LOC";
/// Shorthand result used internally by the backend
pub(super) type BackendResult = Result<(), Error>;
impl TypeInner {
fn is_matrix(&self) -> bool {
match *self {
Self::Matrix { .. } => true,
_ => false,
}
}
#[derive(Copy, Clone, PartialEq)]
enum Io {
Input,
Output,
}
enum SubAccess {
Offset(u32),
Index {
value: Handle<crate::Expression>,
stride: u32,
},
}
/// Structure contains information required for generating
@ -48,10 +53,12 @@ pub struct Writer<'a, W> {
ep_inputs: Vec<Option<EntryPointBinding>>,
/// Set of expressions that have associated temporary variables
named_expressions: crate::NamedExpressions,
pub(super) wrapped_array_lengths: crate::FastHashSet<WrappedArrayLength>,
pub(super) wrapped_image_queries: crate::FastHashSet<WrappedImageQuery>,
temp_access_chain: Vec<SubAccess>,
}
impl<'a, W: Write> Writer<'a, W> {
impl<'a, W: fmt::Write> Writer<'a, W> {
pub fn new(out: W, options: &'a Options) -> Self {
Self {
out,
@ -60,7 +67,9 @@ impl<'a, W: Write> Writer<'a, W> {
options,
ep_inputs: Vec::new(),
named_expressions: crate::NamedExpressions::default(),
wrapped_array_lengths: crate::FastHashSet::default(),
wrapped_image_queries: crate::FastHashSet::default(),
temp_access_chain: Vec::new(),
}
}
@ -68,8 +77,9 @@ impl<'a, W: Write> Writer<'a, W> {
self.names.clear();
self.namer
.reset(module, super::keywords::RESERVED, &[], &mut self.names);
self.named_expressions.clear();
self.ep_inputs.clear();
self.named_expressions.clear();
self.wrapped_array_lengths.clear();
self.wrapped_image_queries.clear();
}
@ -118,6 +128,17 @@ impl<'a, W: Write> Writer<'a, W> {
..
} = ty.inner
{
if let Some(member) = members.last() {
if let TypeInner::Array {
size: crate::ArraySize::Dynamic,
..
} = module.types[member.ty].inner
{
// unsized arrays can only be in storage buffers, for which we use `ByteAddressBuffer` anyway.
continue;
}
}
let ep_result = ep_results.iter().find(|e| {
if let Some(ref result) = e.1 {
result.ty == handle
@ -126,18 +147,13 @@ impl<'a, W: Write> Writer<'a, W> {
}
});
if let Some(result) = ep_result {
self.write_struct(
module,
handle,
top_level,
members,
Some(result.0),
Some(true),
)?;
} else {
self.write_struct(module, handle, top_level, members, None, None)?;
}
self.write_struct(
module,
handle,
top_level,
members,
ep_result.map(|r| (r.0, Io::Output)),
)?;
writeln!(self.out)?;
}
}
@ -192,8 +208,9 @@ impl<'a, W: Write> Writer<'a, W> {
};
let name = self.names[&NameKey::Function(handle)].clone();
// Write wrapped function for `Expression::ImageQuery` before writing all statements and expressions
self.write_wrapped_image_query_functions(module, &ctx)?;
// Write wrapped function for `Expression::ImageQuery` and `Expressions::ArrayLength`
// before writing all statements and expressions.
self.write_wrapped_functions(module, &ctx)?;
self.write_function(module, name.as_str(), function, &ctx)?;
@ -232,8 +249,9 @@ impl<'a, W: Write> Writer<'a, W> {
named_expressions: &ep.function.named_expressions,
};
// Write wrapped function for `Expression::ImageQuery` before writing all statements and expressions
self.write_wrapped_image_query_functions(module, &ctx)?;
// Write wrapped function for `Expression::ImageQuery` and `Expressions::ArrayLength`
// before writing all statements and expressions.
self.write_wrapped_functions(module, &ctx)?;
if ep.stage == ShaderStage::Compute {
// HLSL is calling workgroup size "num threads"
@ -261,15 +279,15 @@ impl<'a, W: Write> Writer<'a, W> {
fn write_semantic(
&mut self,
binding: &crate::Binding,
stage: Option<ShaderStage>,
output: Option<bool>,
stage: Option<(ShaderStage, Io)>,
) -> BackendResult {
match *binding {
crate::Binding::BuiltIn(builtin) => {
write!(self.out, " : {}", builtin_str(builtin))?;
let builtin_str = builtin.to_hlsl_str()?;
write!(self.out, " : {}", builtin_str)?;
}
crate::Binding::Location { location, .. } => {
if stage == Some(crate::ShaderStage::Fragment) && output == Some(true) {
if stage == Some((crate::ShaderStage::Fragment, Io::Output)) {
write!(self.out, " : SV_Target{}", location)?;
} else {
write!(self.out, " : {}{}", LOCATION_SEMANTIC, location)?;
@ -316,7 +334,7 @@ impl<'a, W: Write> Writer<'a, W> {
self.write_type(module, member.ty)?;
write!(self.out, " {}", &member.name)?;
if let Some(ref binding) = member.binding {
self.write_semantic(binding, Some(stage), Some(false))?;
self.write_semantic(binding, Some((stage, Io::Input)))?;
}
write!(self.out, ";")?;
writeln!(self.out)?;
@ -389,9 +407,7 @@ impl<'a, W: Write> Writer<'a, W> {
} else {
("", "t")
};
write!(self.out, "{}StructuredBuffer<", prefix)?;
self.write_type(module, global.ty)?;
write!(self.out, ">")?;
write!(self.out, "{}ByteAddressBuffer", prefix)?;
register
}
crate::StorageClass::Handle => {
@ -502,22 +518,13 @@ impl<'a, W: Write> Writer<'a, W> {
// Panics if `ArraySize::Constant` has a constant that isn't an sint or uint
match size {
crate::ArraySize::Constant(const_handle) => {
match module.constants[const_handle].inner {
crate::ConstantInner::Scalar {
width: _,
value: crate::ScalarValue::Uint(size),
} => write!(self.out, "{}", size)?,
crate::ConstantInner::Scalar {
width: _,
value: crate::ScalarValue::Sint(size),
} => write!(self.out, "{}", size)?,
_ => unreachable!(),
}
let size = module.constants[const_handle].to_array_length().unwrap();
write!(self.out, "{}", size)?;
}
crate::ArraySize::Dynamic => {
//TODO: https://github.com/gfx-rs/naga/issues/1127
log::warn!("Dynamically sized arrays are not properly supported yet");
write!(self.out, "1")?
write!(self.out, "1")?;
}
}
@ -535,8 +542,7 @@ impl<'a, W: Write> Writer<'a, W> {
handle: Handle<crate::Type>,
_block: bool,
members: &[crate::StructMember],
shader_stage: Option<ShaderStage>,
out: Option<bool>,
shader_stage: Option<(ShaderStage, Io)>,
) -> BackendResult {
// Write struct name
write!(self.out, "struct {}", self.names[&NameKey::Type(handle)])?;
@ -553,15 +559,22 @@ impl<'a, W: Write> Writer<'a, W> {
stride: _,
} => {
// HLSL arrays are written as `type name[size]`
let ty_name = match module.types[base].inner {
let (ty_name, vec_size) = match module.types[base].inner {
// Write scalar type by backend so as not to depend on the front-end implementation
// Name returned from frontend can be generated (type1, float1, etc.)
TypeInner::Scalar { kind, width } => scalar_kind_str(kind, width)?,
_ => &self.names[&NameKey::Type(base)],
TypeInner::Scalar { kind, width } => (kind.to_hlsl_str(width)?, None),
// Similarly, write vector types directly.
TypeInner::Vector { size, kind, width } => {
(kind.to_hlsl_str(width)?, Some(size))
}
_ => (self.names[&NameKey::Type(base)].as_str(), None),
};
// Write `type` and `name`
write!(self.out, "{}", ty_name)?;
if let Some(s) = vec_size {
write!(self.out, "{}", s as usize)?;
}
write!(
self.out,
" {}",
@ -579,12 +592,12 @@ impl<'a, W: Write> Writer<'a, W> {
}) = member.binding
{
if let Some(interpolation) = interpolation {
write!(self.out, "{} ", interpolation_str(interpolation))?
write!(self.out, "{} ", interpolation.to_hlsl_str())?
}
if let Some(sampling) = sampling {
if let Some(str) = sampling_str(sampling) {
write!(self.out, "{} ", str)?
if let Some(string) = sampling.to_hlsl_str() {
write!(self.out, "{} ", string)?
}
}
}
@ -600,7 +613,7 @@ impl<'a, W: Write> Writer<'a, W> {
}
if let Some(ref binding) = member.binding {
self.write_semantic(binding, shader_stage, out)?;
self.write_semantic(binding, shader_stage)?;
};
write!(self.out, ";")?;
writeln!(self.out)?;
@ -633,13 +646,13 @@ impl<'a, W: Write> Writer<'a, W> {
pub(super) fn write_value_type(&mut self, module: &Module, inner: &TypeInner) -> BackendResult {
match *inner {
TypeInner::Scalar { kind, width } => {
write!(self.out, "{}", scalar_kind_str(kind, width)?)?;
write!(self.out, "{}", kind.to_hlsl_str(width)?)?;
}
TypeInner::Vector { size, kind, width } => {
write!(
self.out,
"{}{}",
scalar_kind_str(kind, width)?,
kind.to_hlsl_str(width)?,
back::vector_size_str(size)
)?;
}
@ -653,7 +666,7 @@ impl<'a, W: Write> Writer<'a, W> {
write!(
self.out,
"{}{}x{}",
scalar_kind_str(crate::ScalarKind::Float, width)?,
crate::ScalarKind::Float.to_hlsl_str(width)?,
back::vector_size_str(columns),
back::vector_size_str(rows),
)?;
@ -663,26 +676,7 @@ impl<'a, W: Write> Writer<'a, W> {
arrayed,
class,
} => {
use crate::ImageClass as Ic;
let dim_str = dim.to_hlsl_str();
let arrayed_str = if arrayed { "Array" } else { "" };
write!(self.out, "Texture{}{}", dim_str, arrayed_str)?;
match class {
Ic::Depth { multi } => {
let multi_str = if multi { "MS" } else { "" };
write!(self.out, "{}<float>", multi_str)?
}
Ic::Sampled { kind, multi } => {
let multi_str = if multi { "MS" } else { "" };
let scalar_kind_str = scalar_kind_str(kind, 4)?;
write!(self.out, "{}<{}4>", multi_str, scalar_kind_str)?
}
Ic::Storage(format) => {
let storage_format_str = storage_format_to_texture_type(format);
write!(self.out, "<{}>", storage_format_str)?
}
}
self.write_image_type(dim, arrayed, class)?;
}
TypeInner::Sampler { comparison } => {
let sampler = if comparison {
@ -782,8 +776,7 @@ impl<'a, W: Write> Writer<'a, W> {
};
if let Some(ref result) = func.result {
if let Some(ref binding) = result.binding {
let output = stage.is_some();
self.write_semantic(binding, stage, Some(output))?;
self.write_semantic(binding, stage.map(|s| (s, Io::Output)))?;
}
}
@ -955,11 +948,48 @@ impl<'a, W: Write> Writer<'a, W> {
}
}
Statement::Store { pointer, value } => {
write!(self.out, "{}", INDENT.repeat(indent))?;
self.write_expr(module, pointer, func_ctx)?;
write!(self.out, " = ")?;
self.write_expr(module, value, func_ctx)?;
writeln!(self.out, ";")?
let array_info = match *func_ctx.info[pointer].ty.inner_with(&module.types) {
TypeInner::Pointer { base, .. } => match module.types[base].inner {
crate::TypeInner::Array {
size: crate::ArraySize::Constant(ch),
..
} => Some((ch, base)),
_ => None,
},
_ => None,
};
if func_ctx.info[pointer]
.ty
.inner_with(&module.types)
.pointer_class()
== Some(crate::StorageClass::Storage)
{
return Err(Error::Unimplemented("Storage stores".to_string()));
} else if let Some((const_handle, base_ty)) = array_info {
let size = module.constants[const_handle].to_array_length().unwrap();
writeln!(self.out, "{}{{", INDENT.repeat(indent))?;
write!(self.out, "{}", INDENT.repeat(indent + 1))?;
self.write_type(module, base_ty)?;
write!(self.out, " _result[{}]=", size)?;
self.write_expr(module, value, func_ctx)?;
writeln!(self.out, ";")?;
write!(
self.out,
"{}for(int _i=0; _i<{}; ++_i) ",
INDENT.repeat(indent + 1),
size
)?;
self.write_expr(module, pointer, func_ctx)?;
writeln!(self.out, "[_i] = _result[_i];")?;
writeln!(self.out, "{}}}", INDENT.repeat(indent))?;
} else {
write!(self.out, "{}", INDENT.repeat(indent))?;
self.write_expr(module, pointer, func_ctx)?;
write!(self.out, " = ")?;
self.write_expr(module, value, func_ctx)?;
writeln!(self.out, ";")?
}
}
Statement::Call {
function,
@ -1134,40 +1164,65 @@ impl<'a, W: Write> Writer<'a, W> {
self.write_expr(module, right, func_ctx)?;
write!(self.out, ")")?;
}
// TODO: copy-paste from glsl-out
Expression::Access { base, index } => {
if func_ctx.info[expr]
.ty
.inner_with(&module.types)
.pointer_class()
== Some(crate::StorageClass::Storage)
{
// do nothing, the chain is written on `Load`/`Store`
} else {
self.write_expr(module, base, func_ctx)?;
write!(self.out, "[")?;
self.write_expr(module, index, func_ctx)?;
write!(self.out, "]")?;
}
}
Expression::AccessIndex { base, index } => {
self.write_expr(module, base, func_ctx)?;
if func_ctx.info[expr]
.ty
.inner_with(&module.types)
.pointer_class()
== Some(crate::StorageClass::Storage)
{
// do nothing, the chain is written on `Load`/`Store`
} else {
self.write_expr(module, base, func_ctx)?;
let base_ty_res = &func_ctx.info[base].ty;
let mut resolved = base_ty_res.inner_with(&module.types);
let base_ty_handle = match *resolved {
TypeInner::Pointer { base, class: _ } => {
resolved = &module.types[base].inner;
Some(base)
}
_ => base_ty_res.handle(),
};
let base_ty_res = &func_ctx.info[base].ty;
let mut resolved = base_ty_res.inner_with(&module.types);
let base_ty_handle = match *resolved {
TypeInner::Pointer { base, class: _ } => {
resolved = &module.types[base].inner;
Some(base)
}
_ => base_ty_res.handle(),
};
match *resolved {
TypeInner::Vector { .. } => {
// Write vector access as a swizzle
write!(self.out, ".{}", back::COMPONENTS[index as usize])?
}
TypeInner::Matrix { .. }
| TypeInner::Array { .. }
| TypeInner::ValuePointer { .. } => write!(self.out, "[{}]", index)?,
TypeInner::Struct { .. } => {
// This will never panic in case the type is a `Struct`, this is not true
// for other types so we can only check while inside this match arm
let ty = base_ty_handle.unwrap();
match *resolved {
TypeInner::Vector { .. } => {
// Write vector access as a swizzle
write!(self.out, ".{}", back::COMPONENTS[index as usize])?
}
TypeInner::Matrix { .. }
| TypeInner::Array { .. }
| TypeInner::ValuePointer { .. } => write!(self.out, "[{}]", index)?,
TypeInner::Struct { .. } => {
// This will never panic in case the type is a `Struct`, this is not true
// for other types so we can only check while inside this match arm
let ty = base_ty_handle.unwrap();
write!(
self.out,
".{}",
&self.names[&NameKey::StructMember(ty, index)]
)?
write!(
self.out,
".{}",
&self.names[&NameKey::StructMember(ty, index)]
)?
}
ref other => {
return Err(Error::Custom(format!("Cannot index {:?}", other)))
}
}
ref other => return Err(Error::Custom(format!("Cannot index {:?}", other))),
}
}
Expression::FunctionArgument(pos) => {
@ -1339,24 +1394,27 @@ impl<'a, W: Write> Writer<'a, W> {
}
}
Expression::GlobalVariable(handle) => {
let name = &self.names[&NameKey::GlobalVariable(handle)];
let postfix = match module.global_variables[handle].class {
crate::StorageClass::Storage => "[0]",
_ => "",
};
write!(self.out, "{}{}", name, postfix)?;
if module.global_variables[handle].class != crate::StorageClass::Storage {
let name = &self.names[&NameKey::GlobalVariable(handle)];
write!(self.out, "{}", name)?;
}
}
Expression::LocalVariable(handle) => {
write!(self.out, "{}", self.names[&func_ctx.name_key(handle)])?
}
Expression::Load { pointer } => {
self.write_expr(module, pointer, func_ctx)?;
}
Expression::Access { base, index } => {
self.write_expr(module, base, func_ctx)?;
write!(self.out, "[")?;
self.write_expr(module, index, func_ctx)?;
write!(self.out, "]")?;
if func_ctx.info[pointer]
.ty
.inner_with(&module.types)
.pointer_class()
== Some(crate::StorageClass::Storage)
{
let var_handle = self.fill_access_chain(module, pointer, func_ctx)?;
let result_ty = func_ctx.info[expr].ty.clone();
self.write_storage_load(module, var_handle, result_ty, func_ctx)?;
} else {
self.write_expr(module, pointer, func_ctx)?;
}
}
Expression::Unary { op, expr } => {
// https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-operators#unary-operators
@ -1374,12 +1432,12 @@ impl<'a, W: Write> Writer<'a, W> {
write!(
self.out,
"{}{}",
scalar_kind_str(kind, width)?,
kind.to_hlsl_str(width)?,
back::vector_size_str(size),
)?;
}
TypeInner::Scalar { width, .. } => {
write!(self.out, "{}", scalar_kind_str(kind, width)?)?
write!(self.out, "{}", kind.to_hlsl_str(width)?)?
}
_ => {
return Err(Error::Unimplemented(format!(
@ -1482,10 +1540,44 @@ impl<'a, W: Write> Writer<'a, W> {
self.out.write_char(back::COMPONENTS[sc as usize])?;
}
}
// `ArrayLength` is written as `expr.length()`
Expression::ArrayLength(expr) => {
self.write_expr(module, expr, func_ctx)?;
write!(self.out, ".length()")?
let var_handle = match func_ctx.expressions[expr] {
Expression::AccessIndex { base, index: _ } => {
match func_ctx.expressions[base] {
Expression::GlobalVariable(handle) => handle,
_ => unreachable!(),
}
}
Expression::GlobalVariable(handle) => handle,
_ => unreachable!(),
};
let var = &module.global_variables[var_handle];
let (offset, stride) = match module.types[var.ty].inner {
TypeInner::Array { stride, .. } => (0, stride),
TypeInner::Struct {
top_level: true,
ref members,
..
} => {
let last = members.last().unwrap();
let stride = match module.types[last.ty].inner {
TypeInner::Array { stride, .. } => stride,
_ => unreachable!(),
};
(last.offset, stride)
}
_ => unreachable!(),
};
let wrapped_array_length = WrappedArrayLength {
writable: var.storage_access.contains(crate::StorageAccess::STORE),
};
write!(self.out, "((")?;
self.write_wrapped_array_length_function_name(wrapped_array_length)?;
let var_name = &self.names[&NameKey::GlobalVariable(var_handle)];
write!(self.out, "({}) - {}) / {})", var_name, offset, stride)?
}
Expression::Derivative { axis, expr } => {
use crate::DerivativeAxis as Da;
@ -1654,109 +1746,193 @@ impl<'a, W: Write> Writer<'a, W> {
/// Helper function that write default zero initialization
fn write_default_init(&mut self, module: &Module, ty: Handle<crate::Type>) -> BackendResult {
write!(self.out, "(")?;
self.write_type(module, ty)?;
if let TypeInner::Array { size, .. } = module.types[ty].inner {
self.write_array_size(module, size)?;
match module.types[ty].inner {
TypeInner::Array {
size: crate::ArraySize::Constant(const_handle),
base,
..
} => {
write!(self.out, "{{")?;
let count = module.constants[const_handle].to_array_length().unwrap();
for i in 0..count {
if i != 0 {
write!(self.out, ",")?;
}
self.write_default_init(module, base)?;
}
write!(self.out, "}}")?;
}
_ => {
write!(self.out, "(")?;
self.write_type(module, ty)?;
write!(self.out, ")0")?;
}
}
write!(self.out, ")0")?;
Ok(())
}
}
fn builtin_str(built_in: crate::BuiltIn) -> &'static str {
use crate::BuiltIn as Bi;
fn write_storage_address(
&mut self,
module: &Module,
chain: &[SubAccess],
func_ctx: &back::FunctionCtx<'_>,
) -> BackendResult {
for (i, access) in chain.iter().enumerate() {
if i != 0 {
write!(self.out, "+")?;
}
match *access {
SubAccess::Offset(offset) => {
write!(self.out, "{}", offset)?;
}
SubAccess::Index { value, stride } => {
self.write_expr(module, value, func_ctx)?;
write!(self.out, "*{}", stride)?;
}
}
}
Ok(())
}
match built_in {
Bi::Position => "SV_Position",
// vertex
Bi::ClipDistance => "SV_ClipDistance",
Bi::CullDistance => "SV_CullDistance",
Bi::InstanceIndex => "SV_InstanceID",
// based on this page https://docs.microsoft.com/en-us/windows/uwp/gaming/glsl-to-hlsl-reference#comparing-opengl-es-20-with-direct3d-11
// No meaning unless you target Direct3D 9
Bi::PointSize => "PSIZE",
Bi::VertexIndex => "SV_VertexID",
// fragment
Bi::FragDepth => "SV_Depth",
Bi::FrontFacing => "SV_IsFrontFace",
Bi::PrimitiveIndex => "SV_PrimitiveID",
Bi::SampleIndex => "SV_SampleIndex",
Bi::SampleMask => "SV_Coverage",
// compute
Bi::GlobalInvocationId => "SV_DispatchThreadID",
Bi::LocalInvocationId => "SV_GroupThreadID",
Bi::LocalInvocationIndex => "SV_GroupIndex",
Bi::WorkGroupId => "SV_GroupID",
_ => todo!("builtin_str {:?}", built_in),
}
}
/// Helper function that returns scalar related strings
/// https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-scalar
fn scalar_kind_str(kind: crate::ScalarKind, width: crate::Bytes) -> Result<&'static str, Error> {
use crate::ScalarKind as Sk;
match kind {
Sk::Sint => Ok("int"),
Sk::Uint => Ok("uint"),
Sk::Float => match width {
2 => Ok("half"),
4 => Ok("float"),
8 => Ok("double"),
_ => Err(Error::UnsupportedScalar(kind, width)),
},
Sk::Bool => Ok("bool"),
}
}
/// Helper function that returns the string corresponding to the HLSL interpolation qualifier
fn interpolation_str(interpolation: crate::Interpolation) -> &'static str {
use crate::Interpolation as I;
match interpolation {
I::Perspective => "linear",
I::Linear => "noperspective",
I::Flat => "nointerpolation",
}
}
/// Return the HLSL auxiliary qualifier for the given sampling value.
fn sampling_str(sampling: crate::Sampling) -> Option<&'static str> {
use crate::Sampling as S;
match sampling {
S::Center => None,
S::Centroid => Some("centroid"),
S::Sample => Some("sample"),
}
}
fn storage_format_to_texture_type(format: crate::StorageFormat) -> &'static str {
use crate::StorageFormat as Sf;
match format {
Sf::R16Float => "float",
Sf::R8Unorm => "unorm float",
Sf::R8Snorm => "snorm float",
Sf::R8Uint | Sf::R16Uint => "uint",
Sf::R8Sint | Sf::R16Sint => "int",
Sf::Rg16Float => "float2",
Sf::Rg8Unorm => "unorm float2",
Sf::Rg8Snorm => "snorm float2",
Sf::Rg8Sint | Sf::Rg16Sint => "int2",
Sf::Rg8Uint | Sf::Rg16Uint => "uint2",
Sf::Rg11b10Float => "float3",
Sf::Rgba16Float | Sf::R32Float | Sf::Rg32Float | Sf::Rgba32Float => "float4",
Sf::Rgba8Unorm | Sf::Rgb10a2Unorm => "unorm float4",
Sf::Rgba8Snorm => "snorm float4",
Sf::Rgba8Uint | Sf::Rgba16Uint | Sf::R32Uint | Sf::Rg32Uint | Sf::Rgba32Uint => "uint4",
Sf::Rgba8Sint | Sf::Rgba16Sint | Sf::R32Sint | Sf::Rg32Sint | Sf::Rgba32Sint => "int4",
fn write_storage_load_sequence<I: Iterator<Item = (proc::TypeResolution, u32)>>(
&mut self,
module: &Module,
var_handle: Handle<crate::GlobalVariable>,
sequence: I,
func_ctx: &back::FunctionCtx<'_>,
) -> BackendResult {
for (i, (ty_resolution, offset)) in sequence.enumerate() {
// add the index temporarily
self.temp_access_chain.push(SubAccess::Offset(offset));
if i != 0 {
write!(self.out, ", ")?;
};
self.write_storage_load(module, var_handle, ty_resolution, func_ctx)?;
self.temp_access_chain.pop();
}
Ok(())
}
/// Helper function to write down the load operation on a `ByteAddressBuffer`.
fn write_storage_load(
&mut self,
module: &Module,
var_handle: Handle<crate::GlobalVariable>,
result_ty: proc::TypeResolution,
func_ctx: &back::FunctionCtx<'_>,
) -> BackendResult {
match *result_ty.inner_with(&module.types) {
TypeInner::Scalar { kind, width: _ } => {
// working around the borrow checker in `self.write_expr`
let chain = mem::take(&mut self.temp_access_chain);
let var_name = &self.names[&NameKey::GlobalVariable(var_handle)];
let cast = kind.to_hlsl_cast();
write!(self.out, "{}({}.Load(", cast, var_name)?;
self.write_storage_address(module, &chain, func_ctx)?;
write!(self.out, "))")?;
self.temp_access_chain = chain;
}
TypeInner::Vector {
size,
kind,
width: _,
} => {
// working around the borrow checker in `self.write_expr`
let chain = mem::take(&mut self.temp_access_chain);
let var_name = &self.names[&NameKey::GlobalVariable(var_handle)];
let cast = kind.to_hlsl_cast();
write!(self.out, "{}({}.Load{}(", cast, var_name, size as u8)?;
self.write_storage_address(module, &chain, func_ctx)?;
write!(self.out, "))")?;
self.temp_access_chain = chain;
}
TypeInner::Matrix {
columns,
rows,
width,
} => {
write!(
self.out,
"{}{}x{}(",
crate::ScalarKind::Float.to_hlsl_str(width)?,
columns as u8,
rows as u8,
)?;
let row_stride = width as u32 * rows as u32;
let iter = (0..columns as u32).map(|i| {
let ty_inner = TypeInner::Vector {
size: rows,
kind: crate::ScalarKind::Float,
width,
};
(proc::TypeResolution::Value(ty_inner), i * row_stride)
});
self.write_storage_load_sequence(module, var_handle, iter, func_ctx)?;
write!(self.out, ")")?;
}
TypeInner::Array {
base,
size: crate::ArraySize::Constant(const_handle),
..
} => {
write!(self.out, "{{")?;
let count = module.constants[const_handle].to_array_length().unwrap();
let stride = module.types[base].inner.span(&module.constants);
let iter = (0..count).map(|i| (proc::TypeResolution::Handle(base), stride * i));
self.write_storage_load_sequence(module, var_handle, iter, func_ctx)?;
write!(self.out, "}}")?;
}
TypeInner::Struct { ref members, .. } => {
write!(self.out, "{{")?;
let iter = members
.iter()
.map(|m| (proc::TypeResolution::Handle(m.ty), m.offset));
self.write_storage_load_sequence(module, var_handle, iter, func_ctx)?;
write!(self.out, "}}")?;
}
_ => unreachable!(),
}
Ok(())
}
fn fill_access_chain(
&mut self,
module: &Module,
mut cur_expr: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx<'_>,
) -> Result<Handle<crate::GlobalVariable>, Error> {
self.temp_access_chain.clear();
loop {
let stride = func_ctx.info[cur_expr]
.ty
.inner_with(&module.types)
.span(&module.constants);
let (next_expr, sub) = match func_ctx.expressions[cur_expr] {
crate::Expression::GlobalVariable(handle) => return Ok(handle),
crate::Expression::Access { base, index } => (
base,
SubAccess::Index {
value: index,
stride,
},
),
crate::Expression::AccessIndex { base, index } => {
match *func_ctx.info[base].ty.inner_with(&module.types) {
TypeInner::Struct { ref members, .. } => {
(base, SubAccess::Offset(members[index as usize].offset))
}
_ => (base, SubAccess::Offset(index * stride)),
}
}
ref other => {
return Err(Error::Unimplemented(format!(
"Pointer access of {:?}",
other
)))
}
};
self.temp_access_chain.push(sub);
cur_expr = next_expr;
}
}
}

7
src/back/spv/block.rs
View File

@ -199,12 +199,7 @@ impl<'w> BlockContext<'w> {
crate::Expression::GlobalVariable(_) | crate::Expression::LocalVariable(_) => true,
crate::Expression::FunctionArgument(index) => {
let arg = &self.ir_function.arguments[index as usize];
match self.ir_module.types[arg.ty].inner {
crate::TypeInner::Pointer { .. } | crate::TypeInner::ValuePointer { .. } => {
true
}
_ => false,
}
self.ir_module.types[arg.ty].inner.pointer_class().is_some()
}
// The chain rule: if this `Access...`'s `base` operand was

14
src/front/glsl/ast.rs
View File

@ -492,14 +492,12 @@ impl<'function> Context<'function> {
});
if !lhs {
match *program.resolve_type(self, pointer, meta)? {
TypeInner::Pointer { .. } | TypeInner::ValuePointer { .. } => {
return Ok((
Some(self.add_expression(Expression::Load { pointer }, body)),
meta,
));
}
_ => {}
let resolved = program.resolve_type(self, pointer, meta)?;
if resolved.pointer_class().is_some() {
return Ok((
Some(self.add_expression(Expression::Load { pointer }, body)),
meta,
));
}
}

21
src/front/spv/mod.rs
View File

@ -2906,19 +2906,18 @@ impl<I: Iterator<Item = u32>> Parser<I> {
let decor = self.future_decor.remove(&id);
let base_lookup_ty = self.lookup_type.lookup(type_id)?;
let class = match module.types[base_lookup_ty.handle].inner {
crate::TypeInner::Pointer { class, .. }
| crate::TypeInner::ValuePointer { class, .. } => class,
_ if self
.lookup_storage_buffer_types
.contains_key(&base_lookup_ty.handle) =>
{
crate::StorageClass::Storage
}
_ => match map_storage_class(storage_class)? {
let class = if let Some(class) = module.types[base_lookup_ty.handle].inner.pointer_class() {
class
} else if self
.lookup_storage_buffer_types
.contains_key(&base_lookup_ty.handle)
{
crate::StorageClass::Storage
} else {
match map_storage_class(storage_class)? {
ExtendedClass::Global(class) => class,
ExtendedClass::Input | ExtendedClass::Output => crate::StorageClass::Private,
},
}
};
// Don't bother with pointer stuff for `Handle` types.

8
src/proc/mod.rs
View File

@ -85,6 +85,14 @@ impl super::TypeInner {
}
}
pub fn pointer_class(&self) -> Option<crate::StorageClass> {
match *self {
Self::Pointer { class, .. } => Some(class),
Self::ValuePointer { class, .. } => Some(class),
_ => None,
}
}
pub fn span(&self, constants: &super::Arena<super::Constant>) -> u32 {
match *self {
Self::Scalar { kind: _, width } => width as u32,

6
tests/in/access.wgsl
View File

@ -3,6 +3,7 @@
[[block]]
struct Bar {
matrix: mat4x4<f32>;
arr: [[stride(8)]] array<vec2<u32>, 2>;
data: [[stride(4)]] array<i32>;
};
@ -16,6 +17,9 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
let baz: f32 = foo;
foo = 1.0;
let matrix = bar.matrix;
let arr = bar.arr;
let index = 3u;
let b = bar.matrix[index].x;
@ -25,5 +29,5 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
c[vi + 1u] = 42;
let value = c[vi];
return vec4<f32>(vec4<i32>(value));
return matrix * vec4<f32>(vec4<i32>(value));
}

34
tests/out/hlsl/access.hlsl Normal file
View File

@ -0,0 +1,34 @@
RWByteAddressBuffer bar : register(u0);
struct VertexInput_foo {
uint vi1 : SV_VertexID;
};
uint NagaBufferLengthRW(RWByteAddressBuffer buffer)
{
uint ret;
buffer.GetDimensions(ret);
return ret;
}
float4 foo(VertexInput_foo vertexinput_foo) : SV_Position
{
float foo1 = 0.0;
int c[5] = {(int)0,(int)0,(int)0,(int)0,(int)0};
float baz = foo1;
foo1 = 1.0;
float4x4 matrix1 = float4x4(asfloat(bar.Load4(0+0)), asfloat(bar.Load4(0+16)), asfloat(bar.Load4(0+32)), asfloat(bar.Load4(0+48)));
uint2 arr[2] = {asuint(bar.Load2(4+0)), asuint(bar.Load2(4+8))};
float4 _expr13 = asfloat(bar.Load4(12+0));
float b = _expr13.x;
int a = asint(bar.Load((((NagaBufferLengthRW(bar) - 80) / 4) - 2u)*4+8));
{
int _result[5]={ a, int(b), 3, 4, 5 };
for(int _i=0; _i<5; ++_i) c[_i] = _result[_i];
}
c[(vertexinput_foo.vi1 + 1u)] = 42;
int value = c[vertexinput_foo.vi1];
return mul(matrix1, float4(int4(value.xxxx)));
}

3
tests/out/hlsl/access.hlsl.config Normal file
View File

@ -0,0 +1,3 @@
vertex=(foo:vs_5_0 )
fragment=()
compute=()

80
tests/out/hlsl/image.hlsl
View File

@ -19,10 +19,10 @@ struct ComputeInput_main {
uint3 local_id1 : SV_GroupThreadID;
};
int2 NagaDimensions2D(Texture2D<uint4>)
int2 NagaDimensions2D(Texture2D<uint4> texture)
{
uint4 ret;
image_storage_src.GetDimensions(0, ret.x, ret.y, ret.z);
texture.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.xy;
}
@ -39,136 +39,136 @@ void main(ComputeInput_main computeinput_main)
return;
}
int NagaDimensions1D(Texture1D<float4>)
int NagaDimensions1D(Texture1D<float4> texture)
{
uint4 ret;
image_1d.GetDimensions(ret.x);
texture.GetDimensions(ret.x);
return ret.x;
}
int2 NagaDimensions2D(Texture2D<float4>)
int2 NagaDimensions2D(Texture2D<float4> texture)
{
uint4 ret;
image_2d.GetDimensions(0, ret.x, ret.y, ret.z);
texture.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.xy;
}
int NagaNumLevels2D(Texture2D<float4>)
int NagaNumLevels2D(Texture2D<float4> texture)
{
uint4 ret;
image_2d.GetDimensions(0, ret.x, ret.y, ret.z);
texture.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.z;
}
int2 NagaMipDimensions2D(Texture2D<float4>, uint MipLevel)
int2 NagaMipDimensions2D(Texture2D<float4> texture, uint mip_level)
{
uint4 ret;
image_2d.GetDimensions(MipLevel, ret.x, ret.y, ret.z);
texture.GetDimensions(mip_level, ret.x, ret.y, ret.z);
return ret.xy;
}
int2 NagaDimensions2DArray(Texture2DArray<float4>)
int2 NagaDimensions2DArray(Texture2DArray<float4> texture)
{
uint4 ret;
image_2d_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.xy;
}
int NagaNumLevels2DArray(Texture2DArray<float4>)
int NagaNumLevels2DArray(Texture2DArray<float4> texture)
{
uint4 ret;
image_2d_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int2 NagaMipDimensions2DArray(Texture2DArray<float4>, uint MipLevel)
int2 NagaMipDimensions2DArray(Texture2DArray<float4> texture, uint mip_level)
{
uint4 ret;
image_2d_array.GetDimensions(MipLevel, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(mip_level, ret.x, ret.y, ret.z, ret.w);
return ret.xy;
}
int NagaNumLayers2DArray(Texture2DArray<float4>)
int NagaNumLayers2DArray(Texture2DArray<float4> texture)
{
uint4 ret;
image_2d_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int2 NagaDimensionsCube(TextureCube<float4>)
int2 NagaDimensionsCube(TextureCube<float4> texture)
{
uint4 ret;
image_cube.GetDimensions(0, ret.x, ret.y, ret.z);
texture.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.xy;
}
int NagaNumLevelsCube(TextureCube<float4>)
int NagaNumLevelsCube(TextureCube<float4> texture)
{
uint4 ret;
image_cube.GetDimensions(0, ret.x, ret.y, ret.z);
texture.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.z;
}
int2 NagaMipDimensionsCube(TextureCube<float4>, uint MipLevel)
int2 NagaMipDimensionsCube(TextureCube<float4> texture, uint mip_level)
{
uint4 ret;
image_cube.GetDimensions(MipLevel, ret.x, ret.y, ret.z);
texture.GetDimensions(mip_level, ret.x, ret.y, ret.z);
return ret.xy;
}
int2 NagaDimensionsCubeArray(TextureCubeArray<float4>)
int2 NagaDimensionsCubeArray(TextureCubeArray<float4> texture)
{
uint4 ret;
image_cube_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.xy;
}
int NagaNumLevelsCubeArray(TextureCubeArray<float4>)
int NagaNumLevelsCubeArray(TextureCubeArray<float4> texture)
{
uint4 ret;
image_cube_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int2 NagaMipDimensionsCubeArray(TextureCubeArray<float4>, uint MipLevel)
int2 NagaMipDimensionsCubeArray(TextureCubeArray<float4> texture, uint mip_level)
{
uint4 ret;
image_cube_array.GetDimensions(MipLevel, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(mip_level, ret.x, ret.y, ret.z, ret.w);
return ret.xy;
}
int NagaNumLayersCubeArray(TextureCubeArray<float4>)
int NagaNumLayersCubeArray(TextureCubeArray<float4> texture)
{
uint4 ret;
image_cube_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int3 NagaDimensions3D(Texture3D<float4>)
int3 NagaDimensions3D(Texture3D<float4> texture)
{
uint4 ret;
image_3d.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.xyz;
}
int NagaNumLevels3D(Texture3D<float4>)
int NagaNumLevels3D(Texture3D<float4> texture)
{
uint4 ret;
image_3d.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(0, ret.x, ret.y, ret.z, ret.w);
return ret.w;
}
int3 NagaMipDimensions3D(Texture3D<float4>, uint MipLevel)
int3 NagaMipDimensions3D(Texture3D<float4> texture, uint mip_level)
{
uint4 ret;
image_3d.GetDimensions(MipLevel, ret.x, ret.y, ret.z, ret.w);
texture.GetDimensions(mip_level, ret.x, ret.y, ret.z, ret.w);
return ret.xyz;
}
int NagaMSNumSamples2D(Texture2DMS<float4>)
int NagaMSNumSamples2D(Texture2DMS<float4> texture)
{
uint4 ret;
image_aa.GetDimensions(ret.x, ret.y, ret.z);
texture.GetDimensions(ret.x, ret.y, ret.z);
return ret.z;
}

8
tests/out/hlsl/shadow.hlsl
View File

@ -11,12 +11,8 @@ struct Light {
float4 color;
};
struct Lights {
Light data[1];
};
cbuffer u_globals : register(b0) { Globals u_globals; }
StructuredBuffer<Lights> s_lights : register(t1);
ByteAddressBuffer s_lights : register(t1);
Texture2DArray<float> t_shadow : register(t2);
SamplerComparisonState sampler_shadow : register(s3);
@ -49,7 +45,7 @@ float4 fs_main(FragmentInput_fs_main fragmentinput_fs_main) : SV_Target0
break;
}
uint _expr19 = i;
Light light = s_lights[0].data[_expr19];
Light light = {float4x4(asfloat(s_lights.Load4(_expr19*4+0+0+0)), asfloat(s_lights.Load4(_expr19*4+0+0+16)), asfloat(s_lights.Load4(_expr19*4+0+0+32)), asfloat(s_lights.Load4(_expr19*4+0+0+48))), asfloat(s_lights.Load4(_expr19*4+0+64)), asfloat(s_lights.Load4(_expr19*4+0+80))};
uint _expr22 = i;
const float _e25 = fetch_shadow(_expr22, mul(light.proj, fragmentinput_fs_main.position1));
float3 light_dir = normalize((light.pos.xyz - fragmentinput_fs_main.position1.xyz));

15
tests/out/hlsl/texture-arg.hlsl Normal file
View File

@ -0,0 +1,15 @@
Texture2D<float4> Texture : register(t0);
SamplerState Sampler : register(s1);
float4 test(Texture2D<float4> Passed_Texture, SamplerState Passed_Sampler)
{
float4 _expr7 = Passed_Texture.Sample(Passed_Sampler, float2(0.0, 0.0));
return _expr7;
}
float4 main() : SV_Target0
{
const float4 _e2 = test(Texture, Sampler);
return _e2;
}

3
tests/out/hlsl/texture-arg.hlsl.config Normal file
View File

@ -0,0 +1,3 @@
vertex=()
fragment=(main:ps_5_0 )
compute=()

24
tests/out/msl/access.msl
View File

@ -6,12 +6,16 @@ struct _mslBufferSizes {
metal::uint size0;
};
typedef int type2[1];
struct type2 {
metal::uint2 inner[2];
};
typedef int type4[1];
struct Bar {
metal::float4x4 matrix;
type2 data;
type2 arr;
type4 data;
};
struct type6 {
struct type8 {
int inner[5];
};
@ -26,14 +30,16 @@ vertex fooOutput foo(
, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
) {
float foo1 = 0.0;
type6 c;
type8 c;
float baz = foo1;
foo1 = 1.0;
metal::float4 _e9 = bar.matrix[3];
float b = _e9.x;
int a = bar.data[(1 + (_buffer_sizes.size0 - 64 - 4) / 4) - 2u];
for(int _i=0; _i<5; ++_i) c.inner[_i] = type6 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
metal::float4x4 matrix = bar.matrix;
type2 arr = bar.arr;
metal::float4 _e13 = bar.matrix[3];
float b = _e13.x;
int a = bar.data[(1 + (_buffer_sizes.size0 - 80 - 4) / 4) - 2u];
for(int _i=0; _i<5; ++_i) c.inner[_i] = type8 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
c.inner[vi + 1u] = 42;
int value = c.inner[vi];
return fooOutput { static_cast<float4>(metal::int4(value)) };
return fooOutput { matrix * static_cast<float4>(metal::int4(value)) };
}

174
tests/out/spv/access.spvasm
View File

@ -1,92 +1,104 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 58
; Bound: 67
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %32 "foo" %27 %30
OpEntryPoint Vertex %35 "foo" %30 %33
OpSource GLSL 450
OpName %18 "Bar"
OpMemberName %18 0 "matrix"
OpMemberName %18 1 "data"
OpName %20 "bar"
OpName %22 "foo"
OpName %24 "c"
OpName %27 "vi"
OpName %32 "foo"
OpDecorate %17 ArrayStride 4
OpDecorate %18 Block
OpMemberDecorate %18 0 Offset 0
OpMemberDecorate %18 0 ColMajor
OpMemberDecorate %18 0 MatrixStride 16
OpMemberDecorate %18 1 Offset 64
OpDecorate %19 ArrayStride 4
OpDecorate %20 DescriptorSet 0
OpDecorate %20 Binding 0
OpDecorate %27 BuiltIn VertexIndex
OpDecorate %30 BuiltIn Position
OpName %21 "Bar"
OpMemberName %21 0 "matrix"
OpMemberName %21 1 "arr"
OpMemberName %21 2 "data"
OpName %23 "bar"
OpName %25 "foo"
OpName %27 "c"
OpName %30 "vi"
OpName %35 "foo"
OpDecorate %19 ArrayStride 8
OpDecorate %20 ArrayStride 4
OpDecorate %21 Block
OpMemberDecorate %21 0 Offset 0
OpMemberDecorate %21 0 ColMajor
OpMemberDecorate %21 0 MatrixStride 16
OpMemberDecorate %21 1 Offset 64
OpMemberDecorate %21 2 Offset 80
OpDecorate %22 ArrayStride 4
OpDecorate %23 DescriptorSet 0
OpDecorate %23 Binding 0
OpDecorate %30 BuiltIn VertexIndex
OpDecorate %33 BuiltIn Position
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpConstant %4 0.0
%5 = OpConstant %4 1.0
%7 = OpTypeInt 32 0
%6 = OpConstant %7 3
%8 = OpConstant %7 2
%10 = OpTypeInt 32 1
%9 = OpConstant %10 5
%11 = OpConstant %10 3
%12 = OpConstant %10 4
%13 = OpConstant %7 1
%14 = OpConstant %10 42
%16 = OpTypeVector %4 4
%15 = OpTypeMatrix %16 4
%17 = OpTypeRuntimeArray %10
%18 = OpTypeStruct %15 %17
%19 = OpTypeArray %10 %9
%21 = OpTypePointer StorageBuffer %18
%20 = OpVariable %21 StorageBuffer
%23 = OpTypePointer Function %4
%25 = OpTypePointer Function %19
%28 = OpTypePointer Input %7
%27 = OpVariable %28 Input
%31 = OpTypePointer Output %16
%30 = OpVariable %31 Output
%33 = OpTypeFunction %2
%36 = OpTypePointer StorageBuffer %15
%37 = OpTypePointer StorageBuffer %16
%38 = OpConstant %7 0
%42 = OpTypePointer StorageBuffer %17
%45 = OpTypePointer StorageBuffer %10
%51 = OpTypePointer Function %10
%55 = OpTypeVector %10 4
%32 = OpFunction %2 None %33
%26 = OpLabel
%22 = OpVariable %23 Function %3
%24 = OpVariable %25 Function
%29 = OpLoad %7 %27
OpBranch %34
%34 = OpLabel
%35 = OpLoad %4 %22
OpStore %22 %5
%39 = OpAccessChain %37 %20 %38 %6
%40 = OpLoad %16 %39
%41 = OpCompositeExtract %4 %40 0
%43 = OpArrayLength %7 %20 1
%44 = OpISub %7 %43 %8
%46 = OpAccessChain %45 %20 %13 %44
%47 = OpLoad %10 %46
%48 = OpConvertFToS %10 %41
%49 = OpCompositeConstruct %19 %47 %48 %11 %12 %9
OpStore %24 %49
%50 = OpIAdd %7 %29 %13
%52 = OpAccessChain %51 %24 %50
OpStore %52 %14
%53 = OpAccessChain %51 %24 %29
%54 = OpLoad %10 %53
%56 = OpCompositeConstruct %55 %54 %54 %54 %54
%57 = OpConvertSToF %16 %56
OpStore %30 %57
%4 = OpTypeInt 32 1
%3 = OpConstant %4 2
%6 = OpTypeFloat 32
%5 = OpConstant %6 0.0
%7 = OpConstant %6 1.0
%9 = OpTypeInt 32 0
%8 = OpConstant %9 3
%10 = OpConstant %9 2
%11 = OpConstant %4 5
%12 = OpConstant %4 3
%13 = OpConstant %4 4
%14 = OpConstant %9 1
%15 = OpConstant %4 42
%17 = OpTypeVector %6 4
%16 = OpTypeMatrix %17 4
%18 = OpTypeVector %9 2
%19 = OpTypeArray %18 %3
%20 = OpTypeRuntimeArray %4
%21 = OpTypeStruct %16 %19 %20
%22 = OpTypeArray %4 %11
%24 = OpTypePointer StorageBuffer %21
%23 = OpVariable %24 StorageBuffer
%26 = OpTypePointer Function %6
%28 = OpTypePointer Function %22
%31 = OpTypePointer Input %9
%30 = OpVariable %31 Input
%34 = OpTypePointer Output %17
%33 = OpVariable %34 Output
%36 = OpTypeFunction %2
%39 = OpTypePointer StorageBuffer %16
%40 = OpConstant %9 0
%43 = OpTypePointer StorageBuffer %19
%46 = OpTypePointer StorageBuffer %17
%50 = OpTypePointer StorageBuffer %20
%53 = OpTypePointer StorageBuffer %4
%59 = OpTypePointer Function %4
%63 = OpTypeVector %4 4
%35 = OpFunction %2 None %36
%29 = OpLabel
%25 = OpVariable %26 Function %5
%27 = OpVariable %28 Function
%32 = OpLoad %9 %30
OpBranch %37
%37 = OpLabel
%38 = OpLoad %6 %25
OpStore %25 %7
%41 = OpAccessChain %39 %23 %40
%42 = OpLoad %16 %41
%44 = OpAccessChain %43 %23 %14
%45 = OpLoad %19 %44
%47 = OpAccessChain %46 %23 %40 %8
%48 = OpLoad %17 %47
%49 = OpCompositeExtract %6 %48 0
%51 = OpArrayLength %9 %23 2
%52 = OpISub %9 %51 %10
%54 = OpAccessChain %53 %23 %10 %52
%55 = OpLoad %4 %54
%56 = OpConvertFToS %4 %49
%57 = OpCompositeConstruct %22 %55 %56 %12 %13 %11
OpStore %27 %57
%58 = OpIAdd %9 %32 %14
%60 = OpAccessChain %59 %27 %58
OpStore %60 %15
%61 = OpAccessChain %59 %27 %32
%62 = OpLoad %4 %61
%64 = OpCompositeConstruct %63 %62 %62 %62 %62
%65 = OpConvertSToF %17 %64
%66 = OpMatrixTimesVector %17 %42 %65
OpStore %33 %66
OpReturn
OpFunctionEnd

9
tests/out/wgsl/access.wgsl
View File

@ -1,6 +1,7 @@
[[block]]
struct Bar {
matrix: mat4x4<f32>;
arr: [[stride(8)]] array<vec2<u32>,2>;
data: [[stride(4)]] array<i32>;
};
@ -14,11 +15,13 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
let baz: f32 = foo1;
foo1 = 1.0;
let _e9: vec4<f32> = bar.matrix[3];
let b: f32 = _e9.x;
let matrix: mat4x4<f32> = bar.matrix;
let arr: array<vec2<u32>,2> = bar.arr;
let _e13: vec4<f32> = bar.matrix[3];
let b: f32 = _e13.x;
let a: i32 = bar.data[(arrayLength(&bar.data) - 2u)];
c = array<i32,5>(a, i32(b), 3, 4, 5);
c[(vi + 1u)] = 42;
let value: i32 = c[vi];
return vec4<f32>(vec4<i32>(value));
return (matrix * vec4<f32>(vec4<i32>(value)));
}

7
tests/snapshots.rs
View File

@ -423,7 +423,10 @@ fn convert_wgsl() {
"interpolate",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
),
("access", Targets::SPIRV | Targets::METAL | Targets::WGSL),
(
"access",
Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL,
),
(
"control-flow",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
@ -444,7 +447,7 @@ fn convert_wgsl() {
("bounds-check-zero", Targets::SPIRV),
(
"texture-arg",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
),
];