From d558a57668b5ac35307ea69ecba004e42c2af76c Mon Sep 17 00:00:00 2001 From: Gordon-F Date: Thu, 8 Jul 2021 23:55:14 +0300 Subject: [PATCH 1/4] [hlsl-out] Implement barrier statement --- src/back/hlsl/writer.rs | 17 +++++++++++++++++ tests/out/hlsl/control-flow.hlsl | 11 +++++++++++ tests/out/hlsl/control-flow.hlsl.config | 2 ++ tests/snapshots.rs | 2 +- 4 files changed, 31 insertions(+), 1 deletion(-) create mode 100644 tests/out/hlsl/control-flow.hlsl create mode 100644 tests/out/hlsl/control-flow.hlsl.config diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 2e8ca079c7..f61fc3750d 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -876,6 +876,23 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, "{}", INDENT.repeat(indent))?; writeln!(self.out, "continue;")? } + Statement::Barrier(barrier) => { + if barrier.contains(crate::Barrier::STORAGE) { + writeln!( + self.out, + "{}DeviceMemoryBarrierWithGroupSync();", + INDENT.repeat(indent) + )?; + } + + if barrier.contains(crate::Barrier::WORK_GROUP) { + writeln!( + self.out, + "{}GroupMemoryBarrierWithGroupSync();", + INDENT.repeat(indent) + )?; + } + } _ => return Err(Error::Unimplemented(format!("write_stmt {:?}", stmt))), } diff --git a/tests/out/hlsl/control-flow.hlsl b/tests/out/hlsl/control-flow.hlsl new file mode 100644 index 0000000000..1661ebab9e --- /dev/null +++ b/tests/out/hlsl/control-flow.hlsl @@ -0,0 +1,11 @@ +struct ComputeInput_main { + uint3 global_id1 : SV_DispatchThreadID; +}; + +[numthreads(1, 1, 1)] +void main(ComputeInput_main computeinput_main) +{ + DeviceMemoryBarrierWithGroupSync(); + GroupMemoryBarrierWithGroupSync(); + return; +} diff --git a/tests/out/hlsl/control-flow.hlsl.config b/tests/out/hlsl/control-flow.hlsl.config new file mode 100644 index 0000000000..71412ac2d5 --- /dev/null +++ b/tests/out/hlsl/control-flow.hlsl.config @@ -0,0 +1,2 @@ +compute=cs_5_0 +compute_name=main diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 81bab60a58..323c851b5c 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -368,7 +368,7 @@ fn convert_wgsl() { ("access", Targets::SPIRV | Targets::METAL | Targets::WGSL), ( "control-flow", - Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL, + Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ), ( "standard", From 83f34847bdfdcd187325e4009d9d0381c2564fcf Mon Sep 17 00:00:00 2001 From: Gordon-F Date: Wed, 14 Jul 2021 13:30:50 +0300 Subject: [PATCH 2/4] [hlsl-out] Add module documentation --- src/back/hlsl/mod.rs | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/src/back/hlsl/mod.rs b/src/back/hlsl/mod.rs index 27bec76e72..31273f6829 100644 --- a/src/back/hlsl/mod.rs +++ b/src/back/hlsl/mod.rs @@ -1,3 +1,11 @@ +//! HLSL shading language backend +//! +//! # Supported shader model versions: +//! - 5.0 +//! - 5.1 +//! - 6.0 +//! + mod keywords; mod writer; From 05acc4b45946f157e5712977f6831fb1febeafdd Mon Sep 17 00:00:00 2001 From: Gordon-F Date: Wed, 14 Jul 2021 19:21:54 +0300 Subject: [PATCH 3/4] [hlsl-out] Implement all texture functions --- src/back/hlsl/image.rs | 248 ++++++++++++++++++++++++++++ src/back/hlsl/mod.rs | 3 +- src/back/hlsl/writer.rs | 346 +++++++++++++++++++++++++++++++++------- 3 files changed, 538 insertions(+), 59 deletions(-) create mode 100644 src/back/hlsl/image.rs diff --git a/src/back/hlsl/image.rs b/src/back/hlsl/image.rs new file mode 100644 index 0000000000..fe87ac1afe --- /dev/null +++ b/src/back/hlsl/image.rs @@ -0,0 +1,248 @@ +// 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) +// { +// uint4 ret; +// image_1d.GetDimensions(ret.x); +// return ret.x; +// } +// +// int dim_1d = NagaDimensions1D(image_1d); +// ``` + +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 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, + } + } +} + +impl<'a, W: Write> super::Writer<'a, W> { + pub(super) fn write_wrapped_image_query_function_name( + &mut self, + query: WrappedImageQuery, + ) -> super::writer::BackendResult { + let dim_str = super::writer::image_dimension_str(query.dim); + let class_str = match query.class { + crate::ImageClass::Sampled { multi: true, .. } => "MS", + crate::ImageClass::Depth => "Depth", + _ => "", + }; + 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: &crate::back::FunctionCtx<'_>, + ) -> super::writer::BackendResult { + for (handle, _) in func_ctx.expressions.iter() { + if let crate::Expression::ImageQuery { image, query } = func_ctx.expressions[handle] { + let image_ty = func_ctx.info[image].ty.inner_with(&module.types); + match *image_ty { + crate::TypeInner::Image { + dim, + arrayed, + class, + } => { + use crate::back::INDENT; + + let ret_ty = func_ctx.info[handle].ty.inner_with(&module.types); + + let wrapped_image_query = WrappedImageQuery { + dim, + arrayed, + class, + query: query.into(), + }; + + if !self.wrapped_image_queries.contains(&wrapped_image_query) { + // 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 + const MIP_LEVEL_PARAM: &str = "MipLevel"; + if let crate::ImageQuery::Size { level: Some(_) } = query { + write!(self.out, ", uint {}", MIP_LEVEL_PARAM)?; + } + writeln!(self.out, ")")?; + + // Write function body + writeln!(self.out, "{{")?; + const RETURN_VARIABLE_NAME: &str = "ret"; + + use crate::ImageDimension as IDim; + use crate::ImageQuery as Iq; + + 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)?; + } + _ => + // Write zero mipmap level for supported types + { + if let crate::ImageClass::Sampled { multi: true, .. } = class { + } else { + match dim { + 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); + } + } + // Here we work only with image types + _ => unreachable!(), + } + } + } + + Ok(()) + } +} diff --git a/src/back/hlsl/mod.rs b/src/back/hlsl/mod.rs index 31273f6829..dc98bf85dd 100644 --- a/src/back/hlsl/mod.rs +++ b/src/back/hlsl/mod.rs @@ -6,6 +6,7 @@ //! - 6.0 //! +mod image; mod keywords; mod writer; @@ -16,7 +17,7 @@ pub use writer::Writer; /// A HLSL shader model version. #[allow(non_snake_case, non_camel_case_types)] -#[derive(Copy, Clone, Debug, Hash, Eq, PartialEq)] +#[derive(Copy, Clone, Debug, Hash, Eq, PartialEq, PartialOrd)] pub enum ShaderModel { V5_0, V5_1, diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index f61fc3750d..418e735658 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -1,5 +1,3 @@ -//TODO: temp -#![allow(dead_code)] use super::{Error, Options}; use crate::{ back, @@ -11,13 +9,11 @@ use std::fmt::Write; const LOCATION_SEMANTIC: &str = "LOC"; /// Shorthand result used internally by the backend -type BackendResult = Result<(), Error>; +pub(super) type BackendResult = Result<(), Error>; /// Structure contains information required for generating /// wrapped structure of all entry points arguments struct EntryPointBinding { - /// Associated shader stage - stage: ShaderStage, /// Generated structure name name: String, /// Members of generated structure @@ -31,7 +27,7 @@ struct EpStructMember { } pub struct Writer<'a, W> { - out: W, + pub(super) out: W, names: crate::FastHashMap, namer: proc::Namer, /// HLSL backend options @@ -40,6 +36,7 @@ pub struct Writer<'a, W> { ep_inputs: Vec>, /// Set of expressions that have associated temporary variables named_expressions: crate::NamedExpressions, + pub(super) wrapped_image_queries: crate::FastHashSet, } impl<'a, W: Write> Writer<'a, W> { @@ -51,6 +48,7 @@ impl<'a, W: Write> Writer<'a, W> { options, ep_inputs: Vec::new(), named_expressions: crate::NamedExpressions::default(), + wrapped_image_queries: crate::FastHashSet::default(), } } @@ -60,6 +58,7 @@ impl<'a, W: Write> Writer<'a, W> { .reset(module, super::keywords::RESERVED, &[], &mut self.names); self.named_expressions.clear(); self.ep_inputs.clear(); + self.wrapped_image_queries.clear(); } pub fn write( @@ -154,6 +153,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)?; + self.write_function(module, name.as_str(), function, &ctx)?; writeln!(self.out)?; @@ -170,6 +172,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)?; + if ep.stage == ShaderStage::Compute { // HLSL is calling workgroup size, num threads let num_threads = ep.workgroup_size; @@ -266,7 +271,6 @@ impl<'a, W: Write> Writer<'a, W> { writeln!(self.out)?; let ep_input = EntryPointBinding { - stage, name: struct_name, members, }; @@ -290,19 +294,21 @@ impl<'a, W: Write> Writer<'a, W> { let global = &module.global_variables[handle]; let inner = &module.types[global.ty].inner; - if let Some(storage_access) = storage_access(global.storage_access) { - write!(self.out, "{} ", storage_access)?; - } - - let (storage_class, register_ty) = match *inner { - TypeInner::Image { .. } => ("", "t"), + let (storage, register_ty) = match *inner { + TypeInner::Image { .. } => { + if global.storage_access.contains(crate::StorageAccess::STORE) { + ("RW", "u") + } else { + ("", "t") + } + } TypeInner::Sampler { .. } => ("", "s"), TypeInner::Struct { .. } | TypeInner::Vector { .. } => ("static ", ""), // TODO: other register ty https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-variable-register _ => return Err(Error::Unimplemented(format!("register_ty {:?}", inner))), }; - write!(self.out, "{}", storage_class)?; + write!(self.out, "{}", storage)?; self.write_type(module, global.ty)?; if let TypeInner::Array { size, .. } = module.types[global.ty].inner { self.write_array_size(module, size)?; @@ -314,7 +320,11 @@ impl<'a, W: Write> Writer<'a, W> { )?; if let Some(ref binding) = global.binding { - writeln!(self.out, " : register({}{});", register_ty, binding.binding)?; + write!(self.out, " : register({}{}", register_ty, binding.binding)?; + if self.options.shader_model > super::ShaderModel::V5_0 { + write!(self.out, ", space{}", binding.group)?; + } + writeln!(self.out, ");")?; } else { write!(self.out, " = ")?; if let Some(init) = global.init { @@ -507,7 +517,7 @@ impl<'a, W: Write> Writer<'a, W> { /// /// # Notes /// Adds no trailing or leading whitespace - fn write_value_type(&mut self, module: &Module, inner: &TypeInner) -> BackendResult { + 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)?)?; @@ -537,26 +547,34 @@ impl<'a, W: Write> Writer<'a, W> { } TypeInner::Image { dim, - arrayed: _, //TODO: + arrayed, class, } => { + use crate::ImageClass as Ic; + let dim_str = image_dimension_str(dim); - if let crate::ImageClass::Sampled { kind, multi: false } = class { - write!( - self.out, - "Texture{}<{}4>", - dim_str, - scalar_kind_str(kind, 4)? - )? - } else { - return Err(Error::Unimplemented(format!( - "write_value_type {:?}", - inner - ))); + let arrayed_str = if arrayed { "Array" } else { "" }; + write!(self.out, "Texture{}{}", dim_str, arrayed_str)?; + match class { + Ic::Depth => {} + 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)? + } } } - TypeInner::Sampler { comparison: false } => { - write!(self.out, "SamplerState")?; + TypeInner::Sampler { comparison } => { + let sampler = if comparison { + "SamplerComparisonState" + } else { + "SamplerState" + }; + write!(self.out, "{}", sampler)?; } // HLSL arrays are written as `type name[size]` // Current code is written arrays only as `[size]` @@ -893,6 +911,32 @@ impl<'a, W: Write> Writer<'a, W> { )?; } } + Statement::ImageStore { + image, + coordinate, + array_index, + value, + } => { + write!(self.out, "{}", INDENT.repeat(indent))?; + self.write_expr(module, image, func_ctx)?; + + write!(self.out, "[")?; + if let Some(index) = array_index { + // Array index accepted only for texture_storage_2d_array, so we can safety use int3(coordinate, array_index) here + write!(self.out, "int3(")?; + self.write_expr(module, coordinate, func_ctx)?; + write!(self.out, ", ")?; + self.write_expr(module, index, func_ctx)?; + write!(self.out, ")")?; + } else { + self.write_expr(module, coordinate, func_ctx)?; + } + write!(self.out, "]")?; + + write!(self.out, " = ")?; + self.write_expr(module, value, func_ctx)?; + writeln!(self.out, ";")?; + } _ => return Err(Error::Unimplemented(format!("write_stmt {:?}", stmt))), } @@ -903,7 +947,7 @@ impl<'a, W: Write> Writer<'a, W> { /// /// # Notes /// Doesn't add any newlines or leading/trailing spaces - fn write_expr( + pub(super) fn write_expr( &mut self, module: &Module, expr: Handle, @@ -1023,20 +1067,196 @@ impl<'a, W: Write> Writer<'a, W> { } Expression::ImageSample { image, - sampler, // TODO: - coordinate, // TODO: - array_index: _, // TODO: - offset: _, // TODO: - level: _, // TODO: - depth_ref: _, // TODO: + sampler, + coordinate, + array_index, + offset, + level, + depth_ref, } => { + use crate::SampleLevel as Sl; + + let texture_func = match level { + Sl::Auto => { + if depth_ref.is_some() { + "SampleCmp" + } else { + "Sample" + } + } + Sl::Zero => "SampleCmpLevelZero", + Sl::Exact(_) => "SampleLevel", + Sl::Bias(_) => "SampleBias", + Sl::Gradient { .. } => "SampleGrad", + }; + self.write_expr(module, image, func_ctx)?; - write!(self.out, ".Sample(")?; + write!(self.out, ".{}(", texture_func)?; self.write_expr(module, sampler, func_ctx)?; write!(self.out, ", ")?; self.write_expr(module, coordinate, func_ctx)?; + + if let Some(array_index) = array_index { + write!(self.out, ", ")?; + self.write_expr(module, array_index, func_ctx)?; + } + + if let Some(depth_ref) = depth_ref { + write!(self.out, ", ")?; + self.write_expr(module, depth_ref, func_ctx)?; + } + + match level { + Sl::Auto | Sl::Zero => {} + Sl::Exact(expr) => { + write!(self.out, ", ")?; + self.write_expr(module, expr, func_ctx)?; + } + Sl::Bias(expr) => { + write!(self.out, ", ")?; + self.write_expr(module, expr, func_ctx)?; + } + Sl::Gradient { x, y } => { + write!(self.out, ", ")?; + self.write_expr(module, x, func_ctx)?; + write!(self.out, ", ")?; + self.write_expr(module, y, func_ctx)?; + } + } + + if let Some(offset) = offset { + write!(self.out, ", ")?; + self.write_constant(module, offset)?; + } + write!(self.out, ")")?; } + Expression::ImageQuery { image, query } => { + // use wrapped image query function + if let TypeInner::Image { + dim, + arrayed, + class, + } = *func_ctx.info[image].ty.inner_with(&module.types) + { + let wrapped_image_query = super::image::WrappedImageQuery { + dim, + arrayed, + class, + query: query.into(), + }; + + self.write_wrapped_image_query_function_name(wrapped_image_query)?; + write!(self.out, "(")?; + // Image always first param + self.write_expr(module, image, func_ctx)?; + if let crate::ImageQuery::Size { level: Some(level) } = query { + write!(self.out, ", ")?; + self.write_expr(module, level, func_ctx)?; + } + write!(self.out, ")")?; + } + } + Expression::ImageLoad { + image, + coordinate, + array_index, + index, + } => { + // https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-load + let image_ty = func_ctx.info[image].ty.inner_with(&module.types); + if let crate::TypeInner::Image { + dim, + arrayed, + class, + } = *image_ty + { + use crate::ImageDimension as IDim; + + self.write_expr(module, image, func_ctx)?; + write!(self.out, ".Load(")?; + + let ms = if let crate::ImageClass::Sampled { multi: true, .. } = class { + true + } else { + false + }; + + // Input location type based on texture type + let (load_param_ty, load_components_count) = match dim { + IDim::D1 => { + if arrayed { + ("int3", 3) + } else { + ("int2", 2) + } + } + IDim::D2 => match (ms, arrayed) { + (true, true) => ("int3", 3), + (true, false) => ("int2", 2), + (false, true) => ("int4", 4), + (false, false) => ("int3", 3), + }, + IDim::D3 => ("int4", 4), + _ => unreachable!(), + }; + + let coordinate_expr_components_count = + match *func_ctx.info[coordinate].ty.inner_with(&module.types) { + TypeInner::Vector { size, .. } => size as i32, + TypeInner::Scalar { .. } => 1, + // coordinates can be only vector or scalar + _ => unreachable!(), + }; + + let components_count = { + let array_coords = if arrayed { 1 } else { 0 }; + let array_index_param = if array_index.is_some() { 1 } else { 0 }; + let sampler_index_param = if index.is_some() { 1 } else { 0 }; + + coordinate_expr_components_count + + array_coords + + array_index_param + + sampler_index_param + }; + + let is_ty_cast_required = + load_components_count > coordinate_expr_components_count; + + // cast to another type, if required + if is_ty_cast_required { + write!(self.out, "{}(", load_param_ty)?; + } + self.write_expr(module, coordinate, func_ctx)?; + if let Some(array_index) = array_index { + write!(self.out, ", ")?; + self.write_expr(module, array_index, func_ctx)?; + } + if let Some(index) = index { + write!(self.out, ", ")?; + self.write_expr(module, index, func_ctx)?; + } + if index.is_none() && components_count < load_components_count { + // write zero mipmap level if it's not provided but required + write!(self.out, ", 0")?; + }; + + if is_ty_cast_required { + // close bracket for type + write!(self.out, ")")?; + } + + // close bracket for Load function + write!(self.out, ")")?; + + // return x component if return type is scalar + if let TypeInner::Scalar { .. } = + *func_ctx.info[expr].ty.inner_with(&module.types) + { + write!(self.out, ".x")?; + } + } + } // TODO: copy-paste from wgsl-out Expression::GlobalVariable(handle) => { let name = &self.names[&NameKey::GlobalVariable(handle)]; @@ -1378,7 +1598,7 @@ impl<'a, W: Write> Writer<'a, W> { } } -fn image_dimension_str(dim: crate::ImageDimension) -> &'static str { +pub(super) fn image_dimension_str(dim: crate::ImageDimension) -> &'static str { use crate::ImageDimension as IDim; match dim { @@ -1434,24 +1654,6 @@ fn scalar_kind_str(kind: crate::ScalarKind, width: crate::Bytes) -> Result<&'sta } } -fn storage_access(storage_access: crate::StorageAccess) -> Option<&'static str> { - if storage_access == crate::StorageAccess::LOAD { - Some("ByteAddressBuffer") - } else if storage_access.is_all() { - Some("RWByteAddressBuffer") - } else { - None - } -} - -fn number_of_components(vector_size: crate::VectorSize) -> usize { - match vector_size { - crate::VectorSize::Bi => 2, - crate::VectorSize::Tri => 3, - crate::VectorSize::Quad => 4, - } -} - /// 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; @@ -1473,3 +1675,31 @@ fn sampling_str(sampling: crate::Sampling) -> Option<&'static str> { 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", + } +} From 6e715ec985cc899d5d986f645348ef7bea607d6f Mon Sep 17 00:00:00 2001 From: Gordon-F Date: Fri, 16 Jul 2021 13:08:48 +0300 Subject: [PATCH 4/4] [hlsl-out] Enable image snapshot --- tests/out/hlsl/image.hlsl | 209 +++++++++++++++++++++++++++++++ tests/out/hlsl/image.hlsl.config | 8 ++ tests/snapshots.rs | 5 +- 3 files changed, 221 insertions(+), 1 deletion(-) create mode 100644 tests/out/hlsl/image.hlsl create mode 100644 tests/out/hlsl/image.hlsl.config diff --git a/tests/out/hlsl/image.hlsl b/tests/out/hlsl/image.hlsl new file mode 100644 index 0000000000..f5f20b4ba2 --- /dev/null +++ b/tests/out/hlsl/image.hlsl @@ -0,0 +1,209 @@ +Texture2D image_src : register(t1); +RWTexture1D image_dst : register(u2); +Texture1D image_1d : register(t0); +Texture2D image_2d : register(t1); +Texture2DArray image_2d_array : register(t2); +TextureCube image_cube : register(t3); +TextureCubeArray image_cube_array : register(t4); +Texture3D image_3d : register(t5); +Texture2DMS image_aa : register(t6); +SamplerState sampler_reg : register(s0); +SamplerComparisonState sampler_cmp : register(s1); +Texture2D image_2d_depth : register(t2); + +struct ComputeInput_main { + uint3 local_id1 : SV_GroupThreadID; +}; + +int2 NagaDimensions2D(Texture2D) +{ + uint4 ret; + image_src.GetDimensions(0, ret.x, ret.y, ret.z); + return ret.xy; +} + +[numthreads(16, 1, 1)] +void main(ComputeInput_main computeinput_main) +{ + int2 dim = NagaDimensions2D(image_src); + int2 itc = (mul(dim, int2(computeinput_main.local_id1.xy)) % int2(10, 20)); + uint4 value = image_src.Load(int3(itc, 0)); + image_dst[itc.x] = value; + return; +} + +int NagaDimensions1D(Texture1D) +{ + uint4 ret; + image_1d.GetDimensions(ret.x); + return ret.x; +} + +int2 NagaDimensions2D(Texture2D) +{ + uint4 ret; + image_2d.GetDimensions(0, ret.x, ret.y, ret.z); + return ret.xy; +} + +int NagaNumLevels2D(Texture2D) +{ + uint4 ret; + image_2d.GetDimensions(0, ret.x, ret.y, ret.z); + return ret.z; +} + +int2 NagaMipDimensions2D(Texture2D, uint MipLevel) +{ + uint4 ret; + image_2d.GetDimensions(MipLevel, ret.x, ret.y, ret.z); + return ret.xy; +} + +int2 NagaDimensions2DArray(Texture2DArray) +{ + uint4 ret; + image_2d_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.xy; +} + +int NagaNumLevels2DArray(Texture2DArray) +{ + uint4 ret; + image_2d_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.w; +} + +int2 NagaMipDimensions2DArray(Texture2DArray, uint MipLevel) +{ + uint4 ret; + image_2d_array.GetDimensions(MipLevel, ret.x, ret.y, ret.z, ret.w); + return ret.xy; +} + +int NagaNumLayers2DArray(Texture2DArray) +{ + uint4 ret; + image_2d_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.w; +} + +int2 NagaDimensionsCube(TextureCube) +{ + uint4 ret; + image_cube.GetDimensions(0, ret.x, ret.y, ret.z); + return ret.xy; +} + +int NagaNumLevelsCube(TextureCube) +{ + uint4 ret; + image_cube.GetDimensions(0, ret.x, ret.y, ret.z); + return ret.z; +} + +int2 NagaMipDimensionsCube(TextureCube, uint MipLevel) +{ + uint4 ret; + image_cube.GetDimensions(MipLevel, ret.x, ret.y, ret.z); + return ret.xy; +} + +int2 NagaDimensionsCubeArray(TextureCubeArray) +{ + uint4 ret; + image_cube_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.xy; +} + +int NagaNumLevelsCubeArray(TextureCubeArray) +{ + uint4 ret; + image_cube_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.w; +} + +int2 NagaMipDimensionsCubeArray(TextureCubeArray, uint MipLevel) +{ + uint4 ret; + image_cube_array.GetDimensions(MipLevel, ret.x, ret.y, ret.z, ret.w); + return ret.xy; +} + +int NagaNumLayersCubeArray(TextureCubeArray) +{ + uint4 ret; + image_cube_array.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.w; +} + +int3 NagaDimensions3D(Texture3D) +{ + uint4 ret; + image_3d.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.xyz; +} + +int NagaNumLevels3D(Texture3D) +{ + uint4 ret; + image_3d.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.w; +} + +int3 NagaMipDimensions3D(Texture3D, uint MipLevel) +{ + uint4 ret; + image_3d.GetDimensions(MipLevel, ret.x, ret.y, ret.z, ret.w); + return ret.xyz; +} + +int NagaMSNumSamples2D(Texture2DMS) +{ + uint4 ret; + image_aa.GetDimensions(ret.x, ret.y, ret.z); + return ret.z; +} + +float4 queries() : SV_Position +{ + int dim_1d = NagaDimensions1D(image_1d); + int2 dim_2d = NagaDimensions2D(image_2d); + int num_levels_2d = NagaNumLevels2D(image_2d); + int2 dim_2d_lod = NagaMipDimensions2D(image_2d, 1); + int2 dim_2d_array = NagaDimensions2DArray(image_2d_array); + int num_levels_2d_array = NagaNumLevels2DArray(image_2d_array); + int2 dim_2d_array_lod = NagaMipDimensions2DArray(image_2d_array, 1); + int num_layers_2d = NagaNumLayers2DArray(image_2d_array); + int2 dim_cube = NagaDimensionsCube(image_cube); + int num_levels_cube = NagaNumLevelsCube(image_cube); + int2 dim_cube_lod = NagaMipDimensionsCube(image_cube, 1); + int2 dim_cube_array = NagaDimensionsCubeArray(image_cube_array); + int num_levels_cube_array = NagaNumLevelsCubeArray(image_cube_array); + int2 dim_cube_array_lod = NagaMipDimensionsCubeArray(image_cube_array, 1); + int num_layers_cube = NagaNumLayersCubeArray(image_cube_array); + int3 dim_3d = NagaDimensions3D(image_3d); + int num_levels_3d = NagaNumLevels3D(image_3d); + int3 dim_3d_lod = NagaMipDimensions3D(image_3d, 1); + int num_samples_aa = NagaMSNumSamples2D(image_aa); + int sum = ((((((((((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + num_layers_2d) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + num_layers_cube) + dim_3d.z) + dim_3d_lod.z) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array); + return float4(float(sum).xxxx); +} + +float4 sample1() : SV_Target0 +{ + float2 tc = float2(0.5.xx); + float4 s2d = image_2d.Sample(sampler_reg, tc); + float4 s2d_offset = image_2d.Sample(sampler_reg, tc, int2(3, 1)); + float4 s2d_level = image_2d.SampleLevel(sampler_reg, tc, 2.3); + float4 s2d_level_offset = image_2d.SampleLevel(sampler_reg, tc, 2.3, int2(3, 1)); + return (((s2d + s2d_offset) + s2d_level) + s2d_level_offset); +} + +float sample_comparison() : SV_Target0 +{ + float2 tc = float2(0.5.xx); + float s2d_depth = image_2d_depth.SampleCmp(sampler_cmp, tc, 0.5); + float s2d_depth_level = image_2d_depth.SampleCmpLevelZero(sampler_cmp, tc, 0.5); + return (s2d_depth + s2d_depth_level); +} diff --git a/tests/out/hlsl/image.hlsl.config b/tests/out/hlsl/image.hlsl.config new file mode 100644 index 0000000000..3e97dafc35 --- /dev/null +++ b/tests/out/hlsl/image.hlsl.config @@ -0,0 +1,8 @@ +compute=cs_5_0 +compute_name=main +vertex=vs_5_0 +vertex_name=queries +fragment=ps_5_0 +fragment_name=sample1 +fragment=ps_5_0 +fragment_name=sample_comparison diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 323c851b5c..c0b1db87a5 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -355,7 +355,10 @@ fn convert_wgsl() { "shadow", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL, ), - ("image", Targets::SPIRV | Targets::METAL | Targets::WGSL), + ( + "image", + Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL, + ), ("extra", Targets::SPIRV | Targets::METAL | Targets::WGSL), ( "operators",