mirror of
https://github.com/mozilla/gecko-dev.git
synced 2024-11-24 21:31:04 +00:00
Bug 1641370 - Add optional SPIR-V shader validation to WebGPU, via naga r=jgilbert
This change attempts to parse the incoming SPIR-V shader modules with Naga SPIR-V front-end. It's not complete, but it returns an Error if it's unable to parse, in which case we just continue without the validation (for now). If it succeeds, we extract the reflection information from it, and use it for the pipeline. This is just a start. More states would need to be validated, and SPIR-V front-end needs more work. Differential Revision: https://phabricator.services.mozilla.com/D77170
This commit is contained in:
Dzmitry Malyshau
parent
29b214847d
commit
74b412c826
@ -52,6 +52,11 @@ git = "https://github.com/hsivonen/packed_simd"
|
|||||||
replace-with = "vendored-sources"
|
replace-with = "vendored-sources"
|
||||||
rev = "3541e3818fdc7c2a24f87e3459151a4ce955a67a"
|
rev = "3541e3818fdc7c2a24f87e3459151a4ce955a67a"
|
||||||
|
|
||||||
|
[source."https://github.com/gfx-rs/naga"]
|
||||||
|
git = "https://github.com/gfx-rs/naga"
|
||||||
|
replace-with = "vendored-sources"
|
||||||
|
rev = "bce6358eb1026c13d2f1c6d365af37afe8869a86"
|
||||||
|
|
||||||
[source."https://github.com/djg/cubeb-pulse-rs"]
|
[source."https://github.com/djg/cubeb-pulse-rs"]
|
||||||
git = "https://github.com/djg/cubeb-pulse-rs"
|
git = "https://github.com/djg/cubeb-pulse-rs"
|
||||||
replace-with = "vendored-sources"
|
replace-with = "vendored-sources"
|
||||||
|
|||||||
24
Cargo.lock
generated
24
Cargo.lock
generated
@ -3051,6 +3051,18 @@ version = "0.0.5"
|
|||||||
source = "registry+https://github.com/rust-lang/crates.io-index"
|
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||||
checksum = "a2983372caf4480544083767bf2d27defafe32af49ab4df3a0b7fc90793a3664"
|
checksum = "a2983372caf4480544083767bf2d27defafe32af49ab4df3a0b7fc90793a3664"
|
||||||
|
|
||||||
|
[[package]]
|
||||||
|
name = "naga"
|
||||||
|
version = "0.1.0"
|
||||||
|
source = "git+https://github.com/gfx-rs/naga?rev=bce6358eb1026c13d2f1c6d365af37afe8869a86#bce6358eb1026c13d2f1c6d365af37afe8869a86"
|
||||||
|
dependencies = [
|
||||||
|
"bitflags",
|
||||||
|
"fxhash",
|
||||||
|
"log",
|
||||||
|
"num-traits",
|
||||||
|
"spirv_headers",
|
||||||
|
]
|
||||||
|
|
||||||
[[package]]
|
[[package]]
|
||||||
name = "neqo-common"
|
name = "neqo-common"
|
||||||
version = "0.4.3"
|
version = "0.4.3"
|
||||||
@ -4456,6 +4468,16 @@ dependencies = [
|
|||||||
"spirv-cross-internal",
|
"spirv-cross-internal",
|
||||||
]
|
]
|
||||||
|
|
||||||
|
[[package]]
|
||||||
|
name = "spirv_headers"
|
||||||
|
version = "1.4.2"
|
||||||
|
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||||
|
checksum = "3f1418983d16481227ffa3ab3cf44ef92eebc9a76c092fbcd4c51a64ff032622"
|
||||||
|
dependencies = [
|
||||||
|
"bitflags",
|
||||||
|
"num-traits",
|
||||||
|
]
|
||||||
|
|
||||||
[[package]]
|
[[package]]
|
||||||
name = "sql-support"
|
name = "sql-support"
|
||||||
version = "0.1.0"
|
version = "0.1.0"
|
||||||
@ -5543,11 +5565,13 @@ dependencies = [
|
|||||||
"gfx-hal",
|
"gfx-hal",
|
||||||
"gfx-memory",
|
"gfx-memory",
|
||||||
"log",
|
"log",
|
||||||
|
"naga",
|
||||||
"parking_lot",
|
"parking_lot",
|
||||||
"peek-poke 0.2.0 (registry+https://github.com/rust-lang/crates.io-index)",
|
"peek-poke 0.2.0 (registry+https://github.com/rust-lang/crates.io-index)",
|
||||||
"ron",
|
"ron",
|
||||||
"serde",
|
"serde",
|
||||||
"smallvec",
|
"smallvec",
|
||||||
|
"spirv_headers",
|
||||||
"vec_map",
|
"vec_map",
|
||||||
"wgpu-types",
|
"wgpu-types",
|
||||||
]
|
]
|
||||||
|
|||||||
24
gfx/wgpu/Cargo.lock
generated
24
gfx/wgpu/Cargo.lock
generated
@ -744,6 +744,18 @@ dependencies = [
|
|||||||
"ws2_32-sys",
|
"ws2_32-sys",
|
||||||
]
|
]
|
||||||
|
|
||||||
|
[[package]]
|
||||||
|
name = "naga"
|
||||||
|
version = "0.1.0"
|
||||||
|
source = "git+https://github.com/gfx-rs/naga?rev=bce6358eb1026c13d2f1c6d365af37afe8869a86#bce6358eb1026c13d2f1c6d365af37afe8869a86"
|
||||||
|
dependencies = [
|
||||||
|
"bitflags",
|
||||||
|
"fxhash",
|
||||||
|
"log",
|
||||||
|
"num-traits",
|
||||||
|
"spirv_headers",
|
||||||
|
]
|
||||||
|
|
||||||
[[package]]
|
[[package]]
|
||||||
name = "net2"
|
name = "net2"
|
||||||
version = "0.2.33"
|
version = "0.2.33"
|
||||||
@ -1133,6 +1145,16 @@ dependencies = [
|
|||||||
"wasm-bindgen",
|
"wasm-bindgen",
|
||||||
]
|
]
|
||||||
|
|
||||||
|
[[package]]
|
||||||
|
name = "spirv_headers"
|
||||||
|
version = "1.4.2"
|
||||||
|
source = "registry+https://github.com/rust-lang/crates.io-index"
|
||||||
|
checksum = "3f1418983d16481227ffa3ab3cf44ef92eebc9a76c092fbcd4c51a64ff032622"
|
||||||
|
dependencies = [
|
||||||
|
"bitflags",
|
||||||
|
"num-traits",
|
||||||
|
]
|
||||||
|
|
||||||
[[package]]
|
[[package]]
|
||||||
name = "stb_truetype"
|
name = "stb_truetype"
|
||||||
version = "0.3.1"
|
version = "0.3.1"
|
||||||
@ -1376,12 +1398,14 @@ dependencies = [
|
|||||||
"gfx-memory",
|
"gfx-memory",
|
||||||
"log",
|
"log",
|
||||||
"loom",
|
"loom",
|
||||||
|
"naga",
|
||||||
"parking_lot",
|
"parking_lot",
|
||||||
"peek-poke",
|
"peek-poke",
|
||||||
"raw-window-handle",
|
"raw-window-handle",
|
||||||
"ron",
|
"ron",
|
||||||
"serde",
|
"serde",
|
||||||
"smallvec",
|
"smallvec",
|
||||||
|
"spirv_headers",
|
||||||
"vec_map",
|
"vec_map",
|
||||||
"wgpu-types",
|
"wgpu-types",
|
||||||
]
|
]
|
||||||
|
|||||||
@ -38,8 +38,13 @@ raw-window-handle = { version = "0.3", optional = true }
|
|||||||
ron = { version = "0.5", optional = true }
|
ron = { version = "0.5", optional = true }
|
||||||
serde = { version = "1.0", features = ["serde_derive"], optional = true }
|
serde = { version = "1.0", features = ["serde_derive"], optional = true }
|
||||||
smallvec = "1"
|
smallvec = "1"
|
||||||
|
spirv_headers = { version = "1.4.2" }
|
||||||
vec_map = "0.8.1"
|
vec_map = "0.8.1"
|
||||||
|
|
||||||
|
[dependencies.naga]
|
||||||
|
git = "https://github.com/gfx-rs/naga"
|
||||||
|
rev = "bce6358eb1026c13d2f1c6d365af37afe8869a86"
|
||||||
|
|
||||||
[dependencies.wgt]
|
[dependencies.wgt]
|
||||||
path = "../wgpu-types"
|
path = "../wgpu-types"
|
||||||
package = "wgpu-types"
|
package = "wgpu-types"
|
||||||
|
|||||||
@ -30,6 +30,8 @@ use std::{
|
|||||||
sync::atomic::Ordering,
|
sync::atomic::Ordering,
|
||||||
};
|
};
|
||||||
|
|
||||||
|
use spirv_headers::ExecutionModel;
|
||||||
|
|
||||||
mod life;
|
mod life;
|
||||||
#[cfg(any(feature = "trace", feature = "replay"))]
|
#[cfg(any(feature = "trace", feature = "replay"))]
|
||||||
pub mod trace;
|
pub mod trace;
|
||||||
@ -1511,12 +1513,26 @@ impl<G: GlobalIdentityHandlerFactory> Global<G> {
|
|||||||
|
|
||||||
let spv = unsafe { slice::from_raw_parts(desc.code.bytes, desc.code.length) };
|
let spv = unsafe { slice::from_raw_parts(desc.code.bytes, desc.code.length) };
|
||||||
let raw = unsafe { device.raw.create_shader_module(spv).unwrap() };
|
let raw = unsafe { device.raw.create_shader_module(spv).unwrap() };
|
||||||
|
|
||||||
|
let module = {
|
||||||
|
// Parse the given shader code and store its representation.
|
||||||
|
let spv_iter = spv.into_iter().cloned();
|
||||||
|
let mut parser = naga::front::spirv::Parser::new(spv_iter);
|
||||||
|
parser
|
||||||
|
.parse()
|
||||||
|
.map_err(|err| {
|
||||||
|
log::warn!("Failed to parse shader SPIR-V code: {:?}", err);
|
||||||
|
log::warn!("Shader module will not be validated");
|
||||||
|
})
|
||||||
|
.ok()
|
||||||
|
};
|
||||||
let shader = pipeline::ShaderModule {
|
let shader = pipeline::ShaderModule {
|
||||||
raw,
|
raw,
|
||||||
device_id: Stored {
|
device_id: Stored {
|
||||||
value: device_id,
|
value: device_id,
|
||||||
ref_count: device.life_guard.add_ref(),
|
ref_count: device.life_guard.add_ref(),
|
||||||
},
|
},
|
||||||
|
module,
|
||||||
};
|
};
|
||||||
|
|
||||||
let id = hub
|
let id = hub
|
||||||
@ -2014,23 +2030,55 @@ impl<G: GlobalIdentityHandlerFactory> Global<G> {
|
|||||||
}
|
}
|
||||||
};
|
};
|
||||||
|
|
||||||
let vertex = hal::pso::EntryPoint::<B> {
|
let vertex = {
|
||||||
entry: unsafe { ffi::CStr::from_ptr(desc.vertex_stage.entry_point) }
|
let entry_point_name =
|
||||||
.to_str()
|
unsafe { ffi::CStr::from_ptr(desc.vertex_stage.entry_point) }
|
||||||
.to_owned()
|
|
||||||
.unwrap(), // TODO
|
|
||||||
module: &shader_module_guard[desc.vertex_stage.module].raw,
|
|
||||||
specialization: hal::pso::Specialization::EMPTY,
|
|
||||||
};
|
|
||||||
let fragment =
|
|
||||||
unsafe { desc.fragment_stage.as_ref() }.map(|stage| hal::pso::EntryPoint::<B> {
|
|
||||||
entry: unsafe { ffi::CStr::from_ptr(stage.entry_point) }
|
|
||||||
.to_str()
|
.to_str()
|
||||||
.to_owned()
|
.to_owned()
|
||||||
.unwrap(), // TODO
|
.unwrap();
|
||||||
module: &shader_module_guard[stage.module].raw,
|
|
||||||
|
let shader_module = &shader_module_guard[desc.vertex_stage.module];
|
||||||
|
|
||||||
|
if let Some(ref module) = shader_module.module {
|
||||||
|
if let Err(e) =
|
||||||
|
validate_shader(module, entry_point_name, ExecutionModel::Vertex)
|
||||||
|
{
|
||||||
|
log::error!("Failed validating vertex shader module: {:?}", e);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
hal::pso::EntryPoint::<B> {
|
||||||
|
entry: entry_point_name, // TODO
|
||||||
|
module: &shader_module.raw,
|
||||||
specialization: hal::pso::Specialization::EMPTY,
|
specialization: hal::pso::Specialization::EMPTY,
|
||||||
});
|
}
|
||||||
|
};
|
||||||
|
|
||||||
|
let fragment = {
|
||||||
|
let fragment_stage = unsafe { desc.fragment_stage.as_ref() };
|
||||||
|
fragment_stage.map(|stage| {
|
||||||
|
let entry_point_name = unsafe { ffi::CStr::from_ptr(stage.entry_point) }
|
||||||
|
.to_str()
|
||||||
|
.to_owned()
|
||||||
|
.unwrap();
|
||||||
|
|
||||||
|
let shader_module = &shader_module_guard[stage.module];
|
||||||
|
|
||||||
|
if let Some(ref module) = shader_module.module {
|
||||||
|
if let Err(e) =
|
||||||
|
validate_shader(module, entry_point_name, ExecutionModel::Fragment)
|
||||||
|
{
|
||||||
|
log::error!("Failed validating fragment shader module: {:?}", e);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
hal::pso::EntryPoint::<B> {
|
||||||
|
entry: entry_point_name, // TODO
|
||||||
|
module: &shader_module.raw,
|
||||||
|
specialization: hal::pso::Specialization::EMPTY,
|
||||||
|
}
|
||||||
|
})
|
||||||
|
};
|
||||||
|
|
||||||
let shaders = hal::pso::GraphicsShaderSet {
|
let shaders = hal::pso::GraphicsShaderSet {
|
||||||
vertex,
|
vertex,
|
||||||
@ -2195,12 +2243,23 @@ impl<G: GlobalIdentityHandlerFactory> Global<G> {
|
|||||||
let pipeline_stage = &desc.compute_stage;
|
let pipeline_stage = &desc.compute_stage;
|
||||||
let (shader_module_guard, _) = hub.shader_modules.read(&mut token);
|
let (shader_module_guard, _) = hub.shader_modules.read(&mut token);
|
||||||
|
|
||||||
|
let entry_point_name = unsafe { ffi::CStr::from_ptr(pipeline_stage.entry_point) }
|
||||||
|
.to_str()
|
||||||
|
.to_owned()
|
||||||
|
.unwrap();
|
||||||
|
|
||||||
|
let shader_module = &shader_module_guard[pipeline_stage.module];
|
||||||
|
|
||||||
|
if let Some(ref module) = shader_module.module {
|
||||||
|
if let Err(e) = validate_shader(module, entry_point_name, ExecutionModel::GLCompute)
|
||||||
|
{
|
||||||
|
log::error!("Failed validating compute shader module: {:?}", e);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
let shader = hal::pso::EntryPoint::<B> {
|
let shader = hal::pso::EntryPoint::<B> {
|
||||||
entry: unsafe { ffi::CStr::from_ptr(pipeline_stage.entry_point) }
|
entry: entry_point_name, // TODO
|
||||||
.to_str()
|
module: &shader_module.raw,
|
||||||
.to_owned()
|
|
||||||
.unwrap(), // TODO
|
|
||||||
module: &shader_module_guard[pipeline_stage.module].raw,
|
|
||||||
specialization: hal::pso::Specialization::EMPTY,
|
specialization: hal::pso::Specialization::EMPTY,
|
||||||
};
|
};
|
||||||
|
|
||||||
@ -2575,3 +2634,27 @@ impl<G: GlobalIdentityHandlerFactory> Global<G> {
|
|||||||
buffer.map_state = resource::BufferMapState::Idle;
|
buffer.map_state = resource::BufferMapState::Idle;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
/// Errors produced when validating the shader modules of a pipeline.
|
||||||
|
#[derive(Copy, Clone, Debug, Eq, PartialEq)]
|
||||||
|
enum ShaderValidationError {
|
||||||
|
/// Unable to find an entry point matching the specified execution model.
|
||||||
|
MissingEntryPoint(ExecutionModel),
|
||||||
|
}
|
||||||
|
|
||||||
|
fn validate_shader(
|
||||||
|
module: &naga::Module,
|
||||||
|
entry_point_name: &str,
|
||||||
|
execution_model: ExecutionModel,
|
||||||
|
) -> Result<(), ShaderValidationError> {
|
||||||
|
// Since a shader module can have multiple entry points with the same name,
|
||||||
|
// we need to look for one with the right execution model.
|
||||||
|
let entry_point = module.entry_points.iter().find(|entry_point| {
|
||||||
|
entry_point.name == entry_point_name && entry_point.exec_model == execution_model
|
||||||
|
});
|
||||||
|
|
||||||
|
match entry_point {
|
||||||
|
Some(_) => Ok(()),
|
||||||
|
None => Err(ShaderValidationError::MissingEntryPoint(execution_model)),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|||||||
@ -40,6 +40,7 @@ pub struct ShaderModuleDescriptor {
|
|||||||
pub struct ShaderModule<B: hal::Backend> {
|
pub struct ShaderModule<B: hal::Backend> {
|
||||||
pub(crate) raw: B::ShaderModule,
|
pub(crate) raw: B::ShaderModule,
|
||||||
pub(crate) device_id: Stored<DeviceId>,
|
pub(crate) device_id: Stored<DeviceId>,
|
||||||
|
pub(crate) module: Option<naga::Module>,
|
||||||
}
|
}
|
||||||
|
|
||||||
#[repr(C)]
|
#[repr(C)]
|
||||||
|
|||||||
1
third_party/rust/naga/.cargo-checksum.json
vendored
Normal file
1
third_party/rust/naga/.cargo-checksum.json
vendored
Normal file
@ -0,0 +1 @@
|
|||||||
|
{"files":{".monocodus":"2ed558753daedcd63659a41c82bcd4971f3dffd61c3bbbecbf6aaba433f00c1c",".travis.yml":"d91d0d75087934c2d0503a8c04439fea459a19182021bb0b699644d442f5b6fc","Cargo.toml":"fd51368f651a341628248c0aaafeb5d4d62eea359986dd7ba0e8a368dcacdf94","LICENSE":"c71d239df91726fc519c6eb72d318ec65820627232b2f796219e87dcf35d0ab4","Makefile":"51b23cb293d2409397342e53a8f610dfe6da10e8ece06df614fd6e6d88f6eabd","README.md":"747c69b3edb6fa813975a892922d3f99591667c3aa7da5b6d956e7876bb644f4","examples/convert.rs":"561fffdb3a5b145fb3fb873f59f310310ad0c377450da450a6b86a3bfce9238b","src/arena.rs":"a34c363da8fa308dc70d768e09b6acca22c2d73da0d6151b57122913260dbd37","src/back/mod.rs":"0950e0e938f906292b9eab10313421f2179a943aa2956511a0e2968b41663971","src/back/msl.rs":"65512e81bb33b5ce48ed75c71dee4cf3f48e67884c2b6c0a76d21d4e9bcd1c32","src/front/mod.rs":"2320b45153d52992d7104bd3b60499b589c517c6c69f87990e1c91b69bfb4db0","src/front/spirv.rs":"65e82a007bcd461ed07e975474bb4154c029a12d3aada7bb493751b7feeda2ae","src/front/wgsl.rs":"65fe687124a052359dcec5fbc95111a0bc0ef72e8e96d8c3f38fbe70120e4f60","src/lib.rs":"332ab628226bf602794c8470017ac7a3760fba30cbfd93f6dc617d86867764b5","src/proc/interface.rs":"9d43ee5785ea78e833c4b1ba707608b1c3f726b56d4f4a84fdf9195d022fab02","src/proc/mod.rs":"9bf971cb7e476a4b191a970093306e2a412ad1d6bb961cb690083960277a4684","src/proc/typifier.rs":"d9a99f066077f0650d421d96bbd601d222f659c242d2232bd0ef8af74e01d66f","test-data/boids.ron":"f6adc2d25fbb87694d6c4b38d248935402d93576db59236acb4836884ec80945","test-data/boids.wgsl":"a9616a05d35ecca002be0ec3e07ef092b95fda750a2f8d80b38195146edd0a0a","test-data/quad.wgsl":"9b159d06c543d676c2d055c19029d53b7f72ab66e7592d45b4fdd9ed34377e21","tests/convert.rs":"5407ad831e20ab637340119433ca1605568838668581fc90ee2d546818a24101"},"package":null}
|
||||||
7
third_party/rust/naga/.monocodus
vendored
Normal file
7
third_party/rust/naga/.monocodus
vendored
Normal file
@ -0,0 +1,7 @@
|
|||||||
|
version: 1.0.0
|
||||||
|
|
||||||
|
rust:
|
||||||
|
formatter:
|
||||||
|
name: rustfmt
|
||||||
|
ignore_paths:
|
||||||
|
- ".*"
|
||||||
1
third_party/rust/naga/.travis.yml
vendored
Normal file
1
third_party/rust/naga/.travis.yml
vendored
Normal file
@ -0,0 +1 @@
|
|||||||
|
language: rust
|
||||||
22
third_party/rust/naga/Cargo.toml
vendored
Normal file
22
third_party/rust/naga/Cargo.toml
vendored
Normal file
@ -0,0 +1,22 @@
|
|||||||
|
[package]
|
||||||
|
name = "naga"
|
||||||
|
version = "0.1.0"
|
||||||
|
authors = ["Dzmitry Malyshau <kvarkus@gmail.com>"]
|
||||||
|
edition = "2018"
|
||||||
|
description = "Shader translation infrastructure"
|
||||||
|
homepage = "https://github.com/gfx-rs/naga"
|
||||||
|
repository = "https://github.com/gfx-rs/naga"
|
||||||
|
keywords = ["shader", "SPIR-V"]
|
||||||
|
license = "MPL-2.0"
|
||||||
|
|
||||||
|
[dependencies]
|
||||||
|
bitflags = "1"
|
||||||
|
fxhash = "0.2"
|
||||||
|
log = "0.4"
|
||||||
|
num-traits = "0.2"
|
||||||
|
spirv = { package = "spirv_headers", version = "1" }
|
||||||
|
|
||||||
|
[dev-dependencies]
|
||||||
|
env_logger = "0.6"
|
||||||
|
ron = "0.5"
|
||||||
|
serde = { version = "1", features = ["serde_derive"] }
|
||||||
201
third_party/rust/naga/LICENSE
vendored
Normal file
201
third_party/rust/naga/LICENSE
vendored
Normal file
@ -0,0 +1,201 @@
|
|||||||
|
Apache License
|
||||||
|
Version 2.0, January 2004
|
||||||
|
http://www.apache.org/licenses/
|
||||||
|
|
||||||
|
TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
|
||||||
|
|
||||||
|
1. Definitions.
|
||||||
|
|
||||||
|
"License" shall mean the terms and conditions for use, reproduction,
|
||||||
|
and distribution as defined by Sections 1 through 9 of this document.
|
||||||
|
|
||||||
|
"Licensor" shall mean the copyright owner or entity authorized by
|
||||||
|
the copyright owner that is granting the License.
|
||||||
|
|
||||||
|
"Legal Entity" shall mean the union of the acting entity and all
|
||||||
|
other entities that control, are controlled by, or are under common
|
||||||
|
control with that entity. For the purposes of this definition,
|
||||||
|
"control" means (i) the power, direct or indirect, to cause the
|
||||||
|
direction or management of such entity, whether by contract or
|
||||||
|
otherwise, or (ii) ownership of fifty percent (50%) or more of the
|
||||||
|
outstanding shares, or (iii) beneficial ownership of such entity.
|
||||||
|
|
||||||
|
"You" (or "Your") shall mean an individual or Legal Entity
|
||||||
|
exercising permissions granted by this License.
|
||||||
|
|
||||||
|
"Source" form shall mean the preferred form for making modifications,
|
||||||
|
including but not limited to software source code, documentation
|
||||||
|
source, and configuration files.
|
||||||
|
|
||||||
|
"Object" form shall mean any form resulting from mechanical
|
||||||
|
transformation or translation of a Source form, including but
|
||||||
|
not limited to compiled object code, generated documentation,
|
||||||
|
and conversions to other media types.
|
||||||
|
|
||||||
|
"Work" shall mean the work of authorship, whether in Source or
|
||||||
|
Object form, made available under the License, as indicated by a
|
||||||
|
copyright notice that is included in or attached to the work
|
||||||
|
(an example is provided in the Appendix below).
|
||||||
|
|
||||||
|
"Derivative Works" shall mean any work, whether in Source or Object
|
||||||
|
form, that is based on (or derived from) the Work and for which the
|
||||||
|
editorial revisions, annotations, elaborations, or other modifications
|
||||||
|
represent, as a whole, an original work of authorship. For the purposes
|
||||||
|
of this License, Derivative Works shall not include works that remain
|
||||||
|
separable from, or merely link (or bind by name) to the interfaces of,
|
||||||
|
the Work and Derivative Works thereof.
|
||||||
|
|
||||||
|
"Contribution" shall mean any work of authorship, including
|
||||||
|
the original version of the Work and any modifications or additions
|
||||||
|
to that Work or Derivative Works thereof, that is intentionally
|
||||||
|
submitted to Licensor for inclusion in the Work by the copyright owner
|
||||||
|
or by an individual or Legal Entity authorized to submit on behalf of
|
||||||
|
the copyright owner. For the purposes of this definition, "submitted"
|
||||||
|
means any form of electronic, verbal, or written communication sent
|
||||||
|
to the Licensor or its representatives, including but not limited to
|
||||||
|
communication on electronic mailing lists, source code control systems,
|
||||||
|
and issue tracking systems that are managed by, or on behalf of, the
|
||||||
|
Licensor for the purpose of discussing and improving the Work, but
|
||||||
|
excluding communication that is conspicuously marked or otherwise
|
||||||
|
designated in writing by the copyright owner as "Not a Contribution."
|
||||||
|
|
||||||
|
"Contributor" shall mean Licensor and any individual or Legal Entity
|
||||||
|
on behalf of whom a Contribution has been received by Licensor and
|
||||||
|
subsequently incorporated within the Work.
|
||||||
|
|
||||||
|
2. Grant of Copyright License. Subject to the terms and conditions of
|
||||||
|
this License, each Contributor hereby grants to You a perpetual,
|
||||||
|
worldwide, non-exclusive, no-charge, royalty-free, irrevocable
|
||||||
|
copyright license to reproduce, prepare Derivative Works of,
|
||||||
|
publicly display, publicly perform, sublicense, and distribute the
|
||||||
|
Work and such Derivative Works in Source or Object form.
|
||||||
|
|
||||||
|
3. Grant of Patent License. Subject to the terms and conditions of
|
||||||
|
this License, each Contributor hereby grants to You a perpetual,
|
||||||
|
worldwide, non-exclusive, no-charge, royalty-free, irrevocable
|
||||||
|
(except as stated in this section) patent license to make, have made,
|
||||||
|
use, offer to sell, sell, import, and otherwise transfer the Work,
|
||||||
|
where such license applies only to those patent claims licensable
|
||||||
|
by such Contributor that are necessarily infringed by their
|
||||||
|
Contribution(s) alone or by combination of their Contribution(s)
|
||||||
|
with the Work to which such Contribution(s) was submitted. If You
|
||||||
|
institute patent litigation against any entity (including a
|
||||||
|
cross-claim or counterclaim in a lawsuit) alleging that the Work
|
||||||
|
or a Contribution incorporated within the Work constitutes direct
|
||||||
|
or contributory patent infringement, then any patent licenses
|
||||||
|
granted to You under this License for that Work shall terminate
|
||||||
|
as of the date such litigation is filed.
|
||||||
|
|
||||||
|
4. Redistribution. You may reproduce and distribute copies of the
|
||||||
|
Work or Derivative Works thereof in any medium, with or without
|
||||||
|
modifications, and in Source or Object form, provided that You
|
||||||
|
meet the following conditions:
|
||||||
|
|
||||||
|
(a) You must give any other recipients of the Work or
|
||||||
|
Derivative Works a copy of this License; and
|
||||||
|
|
||||||
|
(b) You must cause any modified files to carry prominent notices
|
||||||
|
stating that You changed the files; and
|
||||||
|
|
||||||
|
(c) You must retain, in the Source form of any Derivative Works
|
||||||
|
that You distribute, all copyright, patent, trademark, and
|
||||||
|
attribution notices from the Source form of the Work,
|
||||||
|
excluding those notices that do not pertain to any part of
|
||||||
|
the Derivative Works; and
|
||||||
|
|
||||||
|
(d) If the Work includes a "NOTICE" text file as part of its
|
||||||
|
distribution, then any Derivative Works that You distribute must
|
||||||
|
include a readable copy of the attribution notices contained
|
||||||
|
within such NOTICE file, excluding those notices that do not
|
||||||
|
pertain to any part of the Derivative Works, in at least one
|
||||||
|
of the following places: within a NOTICE text file distributed
|
||||||
|
as part of the Derivative Works; within the Source form or
|
||||||
|
documentation, if provided along with the Derivative Works; or,
|
||||||
|
within a display generated by the Derivative Works, if and
|
||||||
|
wherever such third-party notices normally appear. The contents
|
||||||
|
of the NOTICE file are for informational purposes only and
|
||||||
|
do not modify the License. You may add Your own attribution
|
||||||
|
notices within Derivative Works that You distribute, alongside
|
||||||
|
or as an addendum to the NOTICE text from the Work, provided
|
||||||
|
that such additional attribution notices cannot be construed
|
||||||
|
as modifying the License.
|
||||||
|
|
||||||
|
You may add Your own copyright statement to Your modifications and
|
||||||
|
may provide additional or different license terms and conditions
|
||||||
|
for use, reproduction, or distribution of Your modifications, or
|
||||||
|
for any such Derivative Works as a whole, provided Your use,
|
||||||
|
reproduction, and distribution of the Work otherwise complies with
|
||||||
|
the conditions stated in this License.
|
||||||
|
|
||||||
|
5. Submission of Contributions. Unless You explicitly state otherwise,
|
||||||
|
any Contribution intentionally submitted for inclusion in the Work
|
||||||
|
by You to the Licensor shall be under the terms and conditions of
|
||||||
|
this License, without any additional terms or conditions.
|
||||||
|
Notwithstanding the above, nothing herein shall supersede or modify
|
||||||
|
the terms of any separate license agreement you may have executed
|
||||||
|
with Licensor regarding such Contributions.
|
||||||
|
|
||||||
|
6. Trademarks. This License does not grant permission to use the trade
|
||||||
|
names, trademarks, service marks, or product names of the Licensor,
|
||||||
|
except as required for reasonable and customary use in describing the
|
||||||
|
origin of the Work and reproducing the content of the NOTICE file.
|
||||||
|
|
||||||
|
7. Disclaimer of Warranty. Unless required by applicable law or
|
||||||
|
agreed to in writing, Licensor provides the Work (and each
|
||||||
|
Contributor provides its Contributions) on an "AS IS" BASIS,
|
||||||
|
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
|
||||||
|
implied, including, without limitation, any warranties or conditions
|
||||||
|
of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
|
||||||
|
PARTICULAR PURPOSE. You are solely responsible for determining the
|
||||||
|
appropriateness of using or redistributing the Work and assume any
|
||||||
|
risks associated with Your exercise of permissions under this License.
|
||||||
|
|
||||||
|
8. Limitation of Liability. In no event and under no legal theory,
|
||||||
|
whether in tort (including negligence), contract, or otherwise,
|
||||||
|
unless required by applicable law (such as deliberate and grossly
|
||||||
|
negligent acts) or agreed to in writing, shall any Contributor be
|
||||||
|
liable to You for damages, including any direct, indirect, special,
|
||||||
|
incidental, or consequential damages of any character arising as a
|
||||||
|
result of this License or out of the use or inability to use the
|
||||||
|
Work (including but not limited to damages for loss of goodwill,
|
||||||
|
work stoppage, computer failure or malfunction, or any and all
|
||||||
|
other commercial damages or losses), even if such Contributor
|
||||||
|
has been advised of the possibility of such damages.
|
||||||
|
|
||||||
|
9. Accepting Warranty or Additional Liability. While redistributing
|
||||||
|
the Work or Derivative Works thereof, You may choose to offer,
|
||||||
|
and charge a fee for, acceptance of support, warranty, indemnity,
|
||||||
|
or other liability obligations and/or rights consistent with this
|
||||||
|
License. However, in accepting such obligations, You may act only
|
||||||
|
on Your own behalf and on Your sole responsibility, not on behalf
|
||||||
|
of any other Contributor, and only if You agree to indemnify,
|
||||||
|
defend, and hold each Contributor harmless for any liability
|
||||||
|
incurred by, or claims asserted against, such Contributor by reason
|
||||||
|
of your accepting any such warranty or additional liability.
|
||||||
|
|
||||||
|
END OF TERMS AND CONDITIONS
|
||||||
|
|
||||||
|
APPENDIX: How to apply the Apache License to your work.
|
||||||
|
|
||||||
|
To apply the Apache License to your work, attach the following
|
||||||
|
boilerplate notice, with the fields enclosed by brackets "[]"
|
||||||
|
replaced with your own identifying information. (Don't include
|
||||||
|
the brackets!) The text should be enclosed in the appropriate
|
||||||
|
comment syntax for the file format. We also recommend that a
|
||||||
|
file or class name and description of purpose be included on the
|
||||||
|
same "printed page" as the copyright notice for easier
|
||||||
|
identification within third-party archives.
|
||||||
|
|
||||||
|
Copyright [yyyy] [name of copyright owner]
|
||||||
|
|
||||||
|
Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
you may not use this file except in compliance with the License.
|
||||||
|
You may obtain a copy of the License at
|
||||||
|
|
||||||
|
http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
|
||||||
|
Unless required by applicable law or agreed to in writing, software
|
||||||
|
distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
See the License for the specific language governing permissions and
|
||||||
|
limitations under the License.
|
||||||
16
third_party/rust/naga/Makefile
vendored
Normal file
16
third_party/rust/naga/Makefile
vendored
Normal file
@ -0,0 +1,16 @@
|
|||||||
|
.PHONY: all clean
|
||||||
|
.SECONDARY: boids.metal quad.metal
|
||||||
|
|
||||||
|
all:
|
||||||
|
|
||||||
|
clean:
|
||||||
|
rm *.metal *.air *.metallib
|
||||||
|
|
||||||
|
%.metal: test-data/%.wgsl $(wildcard src/*.rs src/**/*.rs examples/*.rs)
|
||||||
|
cargo run --example convert -- $< $@
|
||||||
|
|
||||||
|
%.air: %.metal
|
||||||
|
xcrun -sdk macosx metal -c $< -mmacosx-version-min=10.11
|
||||||
|
|
||||||
|
%.metallib: %.air
|
||||||
|
xcrun -sdk macosx metallib $< -o $@
|
||||||
29
third_party/rust/naga/README.md
vendored
Normal file
29
third_party/rust/naga/README.md
vendored
Normal file
@ -0,0 +1,29 @@
|
|||||||
|
# Naga
|
||||||
|
|
||||||
|
[](https://matrix.to/#/#naga:matrix.org)
|
||||||
|
[](https://crates.io/crates/naga)
|
||||||
|
[](https://docs.rs/naga)
|
||||||
|
[](https://travis-ci.org/gfx-rs/naga)
|
||||||
|
|
||||||
|
This is an experimental shader translation library for the needs of gfx-rs project and WebGPU. It's meant to provide a safe and performant way of converting to and from SPIR-V.
|
||||||
|
|
||||||
|
## Supported end-points
|
||||||
|
|
||||||
|
Front-end | Status | Notes |
|
||||||
|
--------------- | ------------------ | ----- |
|
||||||
|
SPIR-V (binary) | :construction: | |
|
||||||
|
WGSL (Tint) | :construction: | |
|
||||||
|
GLSL (Vulkan) | | |
|
||||||
|
Rust | | |
|
||||||
|
|
||||||
|
Back-end | Status | Notes |
|
||||||
|
--------------- | ------------------ | ----- |
|
||||||
|
SPIR-V (binary) | | |
|
||||||
|
WGSL | | |
|
||||||
|
Metal | :construction: | |
|
||||||
|
HLSL | | |
|
||||||
|
GLSL | | |
|
||||||
|
AIR | | |
|
||||||
|
DXIR | | |
|
||||||
|
DXIL | | |
|
||||||
|
DXBC | | |
|
||||||
78
third_party/rust/naga/examples/convert.rs
vendored
Normal file
78
third_party/rust/naga/examples/convert.rs
vendored
Normal file
@ -0,0 +1,78 @@
|
|||||||
|
use serde::{Serialize, Deserialize};
|
||||||
|
use std::{env, fs};
|
||||||
|
|
||||||
|
#[derive(Hash, PartialEq, Eq, Serialize, Deserialize)]
|
||||||
|
struct BindSource {
|
||||||
|
set: u32,
|
||||||
|
binding: u32,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Serialize, Deserialize)]
|
||||||
|
struct BindTarget {
|
||||||
|
buffer: Option<u8>,
|
||||||
|
texture: Option<u8>,
|
||||||
|
sampler: Option<u8>,
|
||||||
|
mutable: bool,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Default, Serialize, Deserialize)]
|
||||||
|
struct Parameters {
|
||||||
|
metal_bindings: naga::FastHashMap<BindSource, BindTarget>,
|
||||||
|
}
|
||||||
|
|
||||||
|
fn main() {
|
||||||
|
env_logger::init();
|
||||||
|
|
||||||
|
let args = env::args().collect::<Vec<_>>();
|
||||||
|
|
||||||
|
let module = if args.len() <= 1 {
|
||||||
|
println!("Call with <input> <output>");
|
||||||
|
return
|
||||||
|
} else if args[1].ends_with(".spv") {
|
||||||
|
let input = fs::read(&args[1]).unwrap();
|
||||||
|
naga::front::spirv::parse_u8_slice(&input).unwrap()
|
||||||
|
} else if args[1].ends_with(".wgsl") {
|
||||||
|
let input = fs::read_to_string(&args[1]).unwrap();
|
||||||
|
naga::front::wgsl::parse_str(&input).unwrap()
|
||||||
|
} else {
|
||||||
|
panic!("Unknown input: {:?}", args[1]);
|
||||||
|
};
|
||||||
|
|
||||||
|
if args.len() <= 2 {
|
||||||
|
println!("{:#?}", module);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
let param_path = std::path::PathBuf::from(&args[1])
|
||||||
|
.with_extension("ron");
|
||||||
|
let params = match fs::read_to_string(param_path) {
|
||||||
|
Ok(string) => ron::de::from_str(&string).unwrap(),
|
||||||
|
Err(_) => Parameters::default(),
|
||||||
|
};
|
||||||
|
|
||||||
|
if args[2].ends_with(".metal") {
|
||||||
|
use naga::back::msl;
|
||||||
|
let mut binding_map = msl::BindingMap::default();
|
||||||
|
for (key, value) in params.metal_bindings {
|
||||||
|
binding_map.insert(
|
||||||
|
msl::BindSource {
|
||||||
|
set: key.set,
|
||||||
|
binding: key.binding,
|
||||||
|
},
|
||||||
|
msl::BindTarget {
|
||||||
|
buffer: value.buffer,
|
||||||
|
texture: value.texture,
|
||||||
|
sampler: value.sampler,
|
||||||
|
mutable: value.mutable,
|
||||||
|
},
|
||||||
|
);
|
||||||
|
}
|
||||||
|
let options = msl::Options {
|
||||||
|
binding_map: &binding_map,
|
||||||
|
};
|
||||||
|
let msl = msl::write_string(&module, options).unwrap();
|
||||||
|
fs::write(&args[2], msl).unwrap();
|
||||||
|
} else {
|
||||||
|
panic!("Unknown output: {:?}", args[2]);
|
||||||
|
}
|
||||||
|
}
|
||||||
168
third_party/rust/naga/src/arena.rs
vendored
Normal file
168
third_party/rust/naga/src/arena.rs
vendored
Normal file
@ -0,0 +1,168 @@
|
|||||||
|
use std::{fmt, hash, marker::PhantomData, num::NonZeroU32};
|
||||||
|
|
||||||
|
/// An unique index in the arena array that a handle points to.
|
||||||
|
///
|
||||||
|
/// This type is independent of `spirv::Word`. `spirv::Word` is used in data
|
||||||
|
/// representation. It holds a SPIR-V and refers to that instruction. In
|
||||||
|
/// structured representation, we use Handle to refer to an SPIR-V instruction.
|
||||||
|
/// `Index` is an implementation detail to `Handle`.
|
||||||
|
type Index = NonZeroU32;
|
||||||
|
|
||||||
|
/// A strongly typed reference to a SPIR-V element.
|
||||||
|
pub struct Handle<T> {
|
||||||
|
index: Index,
|
||||||
|
marker: PhantomData<T>,
|
||||||
|
}
|
||||||
|
|
||||||
|
impl<T> Clone for Handle<T> {
|
||||||
|
fn clone(&self) -> Self {
|
||||||
|
Handle {
|
||||||
|
index: self.index,
|
||||||
|
marker: self.marker,
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
impl<T> Copy for Handle<T> {}
|
||||||
|
impl<T> PartialEq for Handle<T> {
|
||||||
|
fn eq(&self, other: &Self) -> bool {
|
||||||
|
self.index == other.index
|
||||||
|
}
|
||||||
|
}
|
||||||
|
impl<T> Eq for Handle<T> {}
|
||||||
|
impl<T> fmt::Debug for Handle<T> {
|
||||||
|
fn fmt(&self, formatter: &mut fmt::Formatter) -> fmt::Result {
|
||||||
|
write!(formatter, "Handle({})", self.index)
|
||||||
|
}
|
||||||
|
}
|
||||||
|
impl<T> hash::Hash for Handle<T> {
|
||||||
|
fn hash<H: hash::Hasher>(&self, hasher: &mut H) {
|
||||||
|
self.index.hash(hasher)
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
impl<T> Handle<T> {
|
||||||
|
#[cfg(test)]
|
||||||
|
pub const DUMMY: Self = Handle {
|
||||||
|
index: unsafe { NonZeroU32::new_unchecked(!0) },
|
||||||
|
marker: PhantomData,
|
||||||
|
};
|
||||||
|
|
||||||
|
pub(crate) fn new(index: Index) -> Self {
|
||||||
|
Handle {
|
||||||
|
index,
|
||||||
|
marker: PhantomData,
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
/// Returns the zero-based index of this handle.
|
||||||
|
pub fn index(self) -> usize {
|
||||||
|
let index = self.index.get() - 1;
|
||||||
|
index as usize
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
/// An arena holding some kind of component (e.g., type, constant,
|
||||||
|
/// instruction, etc.) that can be referenced.
|
||||||
|
#[derive(Debug)]
|
||||||
|
pub struct Arena<T> {
|
||||||
|
/// Values of this arena.
|
||||||
|
data: Vec<T>,
|
||||||
|
}
|
||||||
|
|
||||||
|
impl<T> Default for Arena<T> {
|
||||||
|
fn default() -> Self {
|
||||||
|
Self::new()
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
impl<T> Arena<T> {
|
||||||
|
pub fn new() -> Self {
|
||||||
|
Arena { data: Vec::new() }
|
||||||
|
}
|
||||||
|
|
||||||
|
pub fn len(&self) -> usize {
|
||||||
|
self.data.len()
|
||||||
|
}
|
||||||
|
|
||||||
|
pub fn iter(&self) -> impl Iterator<Item = (Handle<T>, &T)> {
|
||||||
|
self.data.iter().enumerate().map(|(i, v)| {
|
||||||
|
let position = i + 1;
|
||||||
|
let index = unsafe { Index::new_unchecked(position as u32) };
|
||||||
|
(Handle::new(index), v)
|
||||||
|
})
|
||||||
|
}
|
||||||
|
|
||||||
|
/// Adds a new value to the arena, returning a typed handle.
|
||||||
|
///
|
||||||
|
/// The value is not linked to any SPIR-V module.
|
||||||
|
pub fn append(&mut self, value: T) -> Handle<T> {
|
||||||
|
let position = self.data.len() + 1;
|
||||||
|
let index = unsafe { Index::new_unchecked(position as u32) };
|
||||||
|
self.data.push(value);
|
||||||
|
Handle::new(index)
|
||||||
|
}
|
||||||
|
|
||||||
|
/// Adds a value with a check for uniqueness: returns a handle pointing to
|
||||||
|
/// an existing element if its value matches the given one, or adds a new
|
||||||
|
/// element otherwise.
|
||||||
|
pub fn fetch_or_append(&mut self, value: T) -> Handle<T>
|
||||||
|
where
|
||||||
|
T: PartialEq,
|
||||||
|
{
|
||||||
|
if let Some(index) = self.data.iter().position(|d| d == &value) {
|
||||||
|
let index = unsafe { Index::new_unchecked((index + 1) as u32) };
|
||||||
|
Handle::new(index)
|
||||||
|
} else {
|
||||||
|
self.append(value)
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
impl<T> std::ops::Index<Handle<T>> for Arena<T> {
|
||||||
|
type Output = T;
|
||||||
|
fn index(&self, handle: Handle<T>) -> &T {
|
||||||
|
let index = handle.index.get() - 1;
|
||||||
|
&self.data[index as usize]
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#[cfg(test)]
|
||||||
|
mod tests {
|
||||||
|
use super::*;
|
||||||
|
|
||||||
|
#[test]
|
||||||
|
fn append_non_unique() {
|
||||||
|
let mut arena: Arena<u8> = Arena::new();
|
||||||
|
let t1 = arena.append(0);
|
||||||
|
let t2 = arena.append(0);
|
||||||
|
assert!(t1 != t2);
|
||||||
|
assert!(arena[t1] == arena[t2]);
|
||||||
|
}
|
||||||
|
|
||||||
|
#[test]
|
||||||
|
fn append_unique() {
|
||||||
|
let mut arena: Arena<u8> = Arena::new();
|
||||||
|
let t1 = arena.append(0);
|
||||||
|
let t2 = arena.append(1);
|
||||||
|
assert!(t1 != t2);
|
||||||
|
assert!(arena[t1] != arena[t2]);
|
||||||
|
}
|
||||||
|
|
||||||
|
#[test]
|
||||||
|
fn fetch_or_append_non_unique() {
|
||||||
|
let mut arena: Arena<u8> = Arena::new();
|
||||||
|
let t1 = arena.fetch_or_append(0);
|
||||||
|
let t2 = arena.fetch_or_append(0);
|
||||||
|
assert!(t1 == t2);
|
||||||
|
assert!(arena[t1] == arena[t2])
|
||||||
|
}
|
||||||
|
|
||||||
|
#[test]
|
||||||
|
fn fetch_or_append_unique() {
|
||||||
|
let mut arena: Arena<u8> = Arena::new();
|
||||||
|
let t1 = arena.fetch_or_append(0);
|
||||||
|
let t2 = arena.fetch_or_append(1);
|
||||||
|
assert!(t1 != t2);
|
||||||
|
assert!(arena[t1] != arena[t2]);
|
||||||
|
}
|
||||||
|
}
|
||||||
1
third_party/rust/naga/src/back/mod.rs
vendored
Normal file
1
third_party/rust/naga/src/back/mod.rs
vendored
Normal file
@ -0,0 +1 @@
|
|||||||
|
pub mod msl;
|
||||||
976
third_party/rust/naga/src/back/msl.rs
vendored
Normal file
976
third_party/rust/naga/src/back/msl.rs
vendored
Normal file
@ -0,0 +1,976 @@
|
|||||||
|
/*! Metal Shading Language (MSL) backend
|
||||||
|
|
||||||
|
## Binding model
|
||||||
|
|
||||||
|
Metal's bindings are flat per resource. Since there isn't an obvious mapping
|
||||||
|
from SPIR-V's descriptor sets, we require a separate mapping provided in the options.
|
||||||
|
This mapping may have one or more resource end points for each descriptor set + index
|
||||||
|
pair.
|
||||||
|
|
||||||
|
## Outputs
|
||||||
|
|
||||||
|
In Metal, built-in shader outputs can not be nested into structures within
|
||||||
|
the output struct. If there is a structure in the outputs, and it contains any built-ins,
|
||||||
|
we move them up to the root output structure that we define ourselves.
|
||||||
|
!*/
|
||||||
|
|
||||||
|
use std::{
|
||||||
|
fmt::{
|
||||||
|
Display, Error as FmtError, Formatter, Write,
|
||||||
|
},
|
||||||
|
};
|
||||||
|
|
||||||
|
use crate::{
|
||||||
|
arena::Handle,
|
||||||
|
FastHashMap,
|
||||||
|
};
|
||||||
|
|
||||||
|
/// Expect all the global variables to have a pointer type,
|
||||||
|
/// like in SPIR-V.
|
||||||
|
const GLOBAL_POINTERS: bool = false;
|
||||||
|
|
||||||
|
#[derive(Clone, Debug, Default, PartialEq)]
|
||||||
|
pub struct BindTarget {
|
||||||
|
pub buffer: Option<u8>,
|
||||||
|
pub texture: Option<u8>,
|
||||||
|
pub sampler: Option<u8>,
|
||||||
|
pub mutable: bool,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Clone, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
|
||||||
|
pub struct BindSource {
|
||||||
|
pub set: spirv::Word,
|
||||||
|
pub binding: spirv::Word,
|
||||||
|
}
|
||||||
|
|
||||||
|
pub type BindingMap = FastHashMap<BindSource, BindTarget>;
|
||||||
|
|
||||||
|
enum ResolvedBinding {
|
||||||
|
BuiltIn(spirv::BuiltIn),
|
||||||
|
Attribute(spirv::Word),
|
||||||
|
Color(spirv::Word),
|
||||||
|
User { prefix: &'static str, index: spirv::Word },
|
||||||
|
Resource(BindTarget),
|
||||||
|
}
|
||||||
|
|
||||||
|
struct Level(usize);
|
||||||
|
impl Level {
|
||||||
|
fn next(&self) -> Self {
|
||||||
|
Level(self.0 + 1)
|
||||||
|
}
|
||||||
|
}
|
||||||
|
impl Display for Level {
|
||||||
|
fn fmt(&self, formatter: &mut Formatter<'_>) -> Result<(), FmtError> {
|
||||||
|
(0 .. self.0).map(|_| formatter.write_str("\t")).collect()
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug)]
|
||||||
|
pub enum Error {
|
||||||
|
Format(FmtError),
|
||||||
|
UnsupportedExecutionModel(spirv::ExecutionModel),
|
||||||
|
UnexpectedLocation,
|
||||||
|
MixedExecutionModels(crate::Handle<crate::Function>),
|
||||||
|
MissingBinding(crate::Handle<crate::GlobalVariable>),
|
||||||
|
MissingBindTarget(BindSource),
|
||||||
|
InvalidImageFlags(crate::ImageFlags),
|
||||||
|
MutabilityViolation(crate::Handle<crate::GlobalVariable>),
|
||||||
|
BadName(String),
|
||||||
|
}
|
||||||
|
|
||||||
|
impl From<FmtError> for Error {
|
||||||
|
fn from(e: FmtError) -> Self {
|
||||||
|
Error::Format(e)
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Clone, Copy, Debug)]
|
||||||
|
enum LocationMode {
|
||||||
|
VertexInput,
|
||||||
|
FragmentOutput,
|
||||||
|
Intermediate,
|
||||||
|
Uniform,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug, Clone, Copy)]
|
||||||
|
pub struct Options<'a> {
|
||||||
|
pub binding_map: &'a BindingMap,
|
||||||
|
}
|
||||||
|
|
||||||
|
impl Options<'_> {
|
||||||
|
fn resolve_binding(&self, binding: &crate::Binding, mode: LocationMode) -> Result<ResolvedBinding, Error> {
|
||||||
|
match *binding {
|
||||||
|
crate::Binding::BuiltIn(built_in) => Ok(ResolvedBinding::BuiltIn(built_in)),
|
||||||
|
crate::Binding::Location(index) => match mode {
|
||||||
|
LocationMode::VertexInput => Ok(ResolvedBinding::Attribute(index)),
|
||||||
|
LocationMode::FragmentOutput => Ok(ResolvedBinding::Color(index)),
|
||||||
|
LocationMode::Intermediate => Ok(ResolvedBinding::User {
|
||||||
|
prefix: "loc",
|
||||||
|
index,
|
||||||
|
}),
|
||||||
|
LocationMode::Uniform => Err(Error::UnexpectedLocation),
|
||||||
|
},
|
||||||
|
crate::Binding::Descriptor { set, binding } => {
|
||||||
|
let source = BindSource { set, binding };
|
||||||
|
self.binding_map
|
||||||
|
.get(&source)
|
||||||
|
.cloned()
|
||||||
|
.map(ResolvedBinding::Resource)
|
||||||
|
.ok_or(Error::MissingBindTarget(source))
|
||||||
|
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
trait Indexed {
|
||||||
|
const CLASS: &'static str;
|
||||||
|
const PREFIX: bool = false;
|
||||||
|
fn id(&self) -> usize;
|
||||||
|
}
|
||||||
|
|
||||||
|
impl Indexed for crate::Handle<crate::Type> {
|
||||||
|
const CLASS: &'static str = "Type";
|
||||||
|
fn id(&self) -> usize { self.index() }
|
||||||
|
}
|
||||||
|
impl Indexed for crate::Handle<crate::GlobalVariable> {
|
||||||
|
const CLASS: &'static str = "global";
|
||||||
|
fn id(&self) -> usize { self.index() }
|
||||||
|
}
|
||||||
|
impl Indexed for crate::Handle<crate::LocalVariable> {
|
||||||
|
const CLASS: &'static str = "local";
|
||||||
|
fn id(&self) -> usize { self.index() }
|
||||||
|
}
|
||||||
|
impl Indexed for crate::Handle<crate::Function> {
|
||||||
|
const CLASS: &'static str = "function";
|
||||||
|
fn id(&self) -> usize { self.index() }
|
||||||
|
}
|
||||||
|
|
||||||
|
struct MemberIndex(usize);
|
||||||
|
impl Indexed for MemberIndex {
|
||||||
|
const CLASS: &'static str = "field";
|
||||||
|
fn id(&self) -> usize { self.0 }
|
||||||
|
}
|
||||||
|
struct ParameterIndex(usize);
|
||||||
|
impl Indexed for ParameterIndex {
|
||||||
|
const CLASS: &'static str = "param";
|
||||||
|
fn id(&self) -> usize { self.0 }
|
||||||
|
}
|
||||||
|
struct InputStructIndex(crate::Handle<crate::Function>);
|
||||||
|
impl Indexed for InputStructIndex {
|
||||||
|
const CLASS: &'static str = "Input";
|
||||||
|
const PREFIX: bool = true;
|
||||||
|
fn id(&self) -> usize { self.0.index() }
|
||||||
|
}
|
||||||
|
struct OutputStructIndex(crate::Handle<crate::Function>);
|
||||||
|
impl Indexed for OutputStructIndex {
|
||||||
|
const CLASS: &'static str = "Output";
|
||||||
|
const PREFIX: bool = true;
|
||||||
|
fn id(&self) -> usize { self.0.index() }
|
||||||
|
}
|
||||||
|
|
||||||
|
enum NameSource<'a> {
|
||||||
|
Custom { name: &'a str, prefix: bool },
|
||||||
|
Index(usize),
|
||||||
|
}
|
||||||
|
|
||||||
|
const RESERVED_NAMES: &[&str] = &[
|
||||||
|
"main",
|
||||||
|
];
|
||||||
|
|
||||||
|
struct Name<'a> {
|
||||||
|
class: &'static str,
|
||||||
|
source: NameSource<'a>,
|
||||||
|
}
|
||||||
|
impl Display for Name<'_> {
|
||||||
|
fn fmt(&self, formatter: &mut Formatter<'_>) -> Result<(), FmtError> {
|
||||||
|
match self.source {
|
||||||
|
NameSource::Custom { name, prefix: false } if RESERVED_NAMES.contains(&name) => {
|
||||||
|
write!(formatter, "{}_", name)
|
||||||
|
}
|
||||||
|
NameSource::Custom { name, prefix: false } => formatter.write_str(name),
|
||||||
|
NameSource::Custom { name, prefix: true } => {
|
||||||
|
let (head, tail) = name.split_at(1);
|
||||||
|
write!(formatter, "{}{}{}", self.class, head.to_uppercase(), tail)
|
||||||
|
}
|
||||||
|
NameSource::Index(index) => write!(formatter, "{}{}", self.class, index),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
impl<I: Indexed> From<I> for Name<'_> {
|
||||||
|
fn from(index: I) -> Self {
|
||||||
|
Name {
|
||||||
|
class: I::CLASS,
|
||||||
|
source: NameSource::Index(index.id()),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
trait AsName {
|
||||||
|
fn or_index<I: Indexed>(&self, index: I) -> Name;
|
||||||
|
}
|
||||||
|
impl AsName for Option<String> {
|
||||||
|
fn or_index<I: Indexed>(&self, index: I) -> Name {
|
||||||
|
Name {
|
||||||
|
class: I::CLASS,
|
||||||
|
source: match *self {
|
||||||
|
Some(ref name) if !name.is_empty() => NameSource::Custom { name, prefix: I::PREFIX },
|
||||||
|
_ => NameSource::Index(index.id()),
|
||||||
|
},
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
struct TypedGlobalVariable<'a> {
|
||||||
|
module: &'a crate::Module,
|
||||||
|
handle: crate::Handle<crate::GlobalVariable>,
|
||||||
|
usage: crate::GlobalUse,
|
||||||
|
}
|
||||||
|
impl Display for TypedGlobalVariable<'_> {
|
||||||
|
fn fmt(&self, formatter: &mut Formatter<'_>) -> Result<(), FmtError> {
|
||||||
|
let var = &self.module.global_variables[self.handle];
|
||||||
|
let name = var.name.or_index(self.handle);
|
||||||
|
let (space_qualifier, reference) = match var.class {
|
||||||
|
spirv::StorageClass::Uniform |
|
||||||
|
spirv::StorageClass::UniformConstant |
|
||||||
|
spirv::StorageClass::StorageBuffer => {
|
||||||
|
let space = if self.usage.contains(crate::GlobalUse::STORE) {
|
||||||
|
"device "
|
||||||
|
} else {
|
||||||
|
"constant "
|
||||||
|
};
|
||||||
|
(space, "&")
|
||||||
|
}
|
||||||
|
_ => ("", "")
|
||||||
|
};
|
||||||
|
if GLOBAL_POINTERS {
|
||||||
|
let ty = &self.module.types[var.ty];
|
||||||
|
match ty.inner {
|
||||||
|
crate::TypeInner::Pointer { base, class } => {
|
||||||
|
let ty_handle = match class {
|
||||||
|
spirv::StorageClass::Input |
|
||||||
|
spirv::StorageClass::Output |
|
||||||
|
spirv::StorageClass::Uniform |
|
||||||
|
spirv::StorageClass::UniformConstant => base,
|
||||||
|
_ => var.ty
|
||||||
|
};
|
||||||
|
let ty_name = self.module.types[ty_handle].name.or_index(ty_handle);
|
||||||
|
write!(formatter, "{} {}", ty_name, name)
|
||||||
|
}
|
||||||
|
_ => panic!("Unexpected global type {:?} = {:?}", var.ty, ty),
|
||||||
|
}
|
||||||
|
} else {
|
||||||
|
let ty_name = self.module.types[var.ty].name.or_index(var.ty);
|
||||||
|
write!(formatter, "{}{}{} {}", space_qualifier, ty_name, reference, name)
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
impl Display for ResolvedBinding {
|
||||||
|
fn fmt(&self, formatter: &mut Formatter<'_>) -> Result<(), FmtError> {
|
||||||
|
match *self {
|
||||||
|
ResolvedBinding::BuiltIn(built_in) => {
|
||||||
|
let name = match built_in {
|
||||||
|
spirv::BuiltIn::ClipDistance => "clip_distance",
|
||||||
|
spirv::BuiltIn::GlobalInvocationId => "thread_position_in_grid",
|
||||||
|
spirv::BuiltIn::PointSize => "point_size",
|
||||||
|
spirv::BuiltIn::Position => "position",
|
||||||
|
_ => panic!("Built in {:?} is not implemented", built_in),
|
||||||
|
};
|
||||||
|
formatter.write_str(name)
|
||||||
|
}
|
||||||
|
ResolvedBinding::Attribute(index) => {
|
||||||
|
write!(formatter, "attribute({})", index)
|
||||||
|
}
|
||||||
|
ResolvedBinding::Color(index) => {
|
||||||
|
write!(formatter, "color({})", index)
|
||||||
|
}
|
||||||
|
ResolvedBinding::User { prefix, index } => {
|
||||||
|
write!(formatter, "user({}{})", prefix, index)
|
||||||
|
}
|
||||||
|
ResolvedBinding::Resource(ref target) => {
|
||||||
|
if let Some(id) = target.buffer {
|
||||||
|
write!(formatter, "buffer({})", id)
|
||||||
|
} else if let Some(id) = target.texture {
|
||||||
|
write!(formatter, "texture({})", id)
|
||||||
|
} else if let Some(id) = target.sampler {
|
||||||
|
write!(formatter, "sampler({})", id)
|
||||||
|
} else {
|
||||||
|
unimplemented!()
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
pub struct Writer<W> {
|
||||||
|
out: W,
|
||||||
|
}
|
||||||
|
|
||||||
|
fn scalar_kind_string(kind: crate::ScalarKind) -> &'static str {
|
||||||
|
match kind {
|
||||||
|
crate::ScalarKind::Float => "float",
|
||||||
|
crate::ScalarKind::Sint => "int",
|
||||||
|
crate::ScalarKind::Uint => "uint",
|
||||||
|
crate::ScalarKind::Bool => "bool",
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
fn vector_size_string(size: crate::VectorSize) -> &'static str {
|
||||||
|
match size {
|
||||||
|
crate::VectorSize::Bi => "2",
|
||||||
|
crate::VectorSize::Tri => "3",
|
||||||
|
crate::VectorSize::Quad => "4",
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
const OUTPUT_STRUCT_NAME: &str = "output";
|
||||||
|
const LOCATION_INPUT_STRUCT_NAME: &str = "input";
|
||||||
|
const COMPONENTS: &[char] = &['x', 'y', 'z', 'w'];
|
||||||
|
|
||||||
|
fn separate(is_last: bool) -> &'static str {
|
||||||
|
if is_last { "" } else { "," }
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug)]
|
||||||
|
enum MaybeOwned<'a, T: 'a> {
|
||||||
|
Borrowed(&'a T),
|
||||||
|
Owned(T),
|
||||||
|
}
|
||||||
|
|
||||||
|
impl<T> MaybeOwned<'_, T> {
|
||||||
|
fn borrow(&self) -> &T {
|
||||||
|
match *self {
|
||||||
|
MaybeOwned::Borrowed(inner) => inner,
|
||||||
|
MaybeOwned::Owned(ref inner) => inner,
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
impl crate::Module {
|
||||||
|
fn borrow_type(&self, handle: Handle<crate::Type>) -> MaybeOwned<crate::TypeInner> {
|
||||||
|
MaybeOwned::Borrowed(&self.types[handle].inner)
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
impl<W: Write> Writer<W> {
|
||||||
|
fn put_expression<'a>(
|
||||||
|
&mut self,
|
||||||
|
expr_handle: Handle<crate::Expression>,
|
||||||
|
function: &crate::Function,
|
||||||
|
module: &'a crate::Module,
|
||||||
|
) -> Result<MaybeOwned<'a, crate::TypeInner>, Error> {
|
||||||
|
let expression = &function.expressions[expr_handle];
|
||||||
|
log::trace!("expression {:?} = {:?}", expr_handle, expression);
|
||||||
|
match *expression {
|
||||||
|
crate::Expression::Access { base, index } => {
|
||||||
|
match *self.put_expression(base, function, module)?.borrow() {
|
||||||
|
crate::TypeInner::Array { base, .. } => {
|
||||||
|
//TODO: add size check
|
||||||
|
self.out.write_str("[")?;
|
||||||
|
self.put_expression(index, function, module)?;
|
||||||
|
self.out.write_str("]")?;
|
||||||
|
Ok(module.borrow_type(base))
|
||||||
|
}
|
||||||
|
ref other => panic!("Unexpected indexing of {:?}", other),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
crate::Expression::AccessIndex { base, index } => {
|
||||||
|
match *self.put_expression(base, function, module)?.borrow() {
|
||||||
|
crate::TypeInner::Struct { ref members } => {
|
||||||
|
let member = &members[index as usize];
|
||||||
|
let name = member.name.or_index(MemberIndex(index as usize));
|
||||||
|
write!(self.out, ".{}", name)?;
|
||||||
|
Ok(module.borrow_type(member.ty))
|
||||||
|
}
|
||||||
|
crate::TypeInner::Matrix { rows, kind, width, .. } => {
|
||||||
|
write!(self.out, ".{}", COMPONENTS[index as usize])?;
|
||||||
|
Ok(MaybeOwned::Owned(crate::TypeInner::Vector { size: rows, kind, width }))
|
||||||
|
}
|
||||||
|
crate::TypeInner::Vector { kind, width, .. } => {
|
||||||
|
write!(self.out, ".{}", COMPONENTS[index as usize])?;
|
||||||
|
Ok(MaybeOwned::Owned(crate::TypeInner::Scalar { kind, width }))
|
||||||
|
}
|
||||||
|
crate::TypeInner::Array { base, size } => {
|
||||||
|
if let crate::ArraySize::Static(length) = size {
|
||||||
|
assert!(index < length);
|
||||||
|
}
|
||||||
|
write!(self.out, "[{}]", index)?;
|
||||||
|
Ok(module.borrow_type(base))
|
||||||
|
}
|
||||||
|
ref other => panic!("Unexpected indexing of {:?}", other),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
crate::Expression::Constant(handle) => {
|
||||||
|
self.put_constant(handle, module)
|
||||||
|
}
|
||||||
|
crate::Expression::Compose { ty, ref components } => {
|
||||||
|
let inner = &module.types[ty].inner;
|
||||||
|
match *inner {
|
||||||
|
crate::TypeInner::Vector { size, kind, .. } => {
|
||||||
|
write!(self.out, "{}{}(", scalar_kind_string(kind), vector_size_string(size))?;
|
||||||
|
for (i, &handle) in components.iter().enumerate() {
|
||||||
|
if i != 0 {
|
||||||
|
write!(self.out, ",")?;
|
||||||
|
}
|
||||||
|
self.put_expression(handle, function, module)?;
|
||||||
|
}
|
||||||
|
write!(self.out, ")")?;
|
||||||
|
}
|
||||||
|
_ => panic!("Unsupported compose {:?}", ty),
|
||||||
|
}
|
||||||
|
Ok(MaybeOwned::Borrowed(inner))
|
||||||
|
}
|
||||||
|
crate::Expression::GlobalVariable(handle) => {
|
||||||
|
let var = &module.global_variables[handle];
|
||||||
|
let inner = &module.types[var.ty].inner;
|
||||||
|
match var.class {
|
||||||
|
spirv::StorageClass::Output => {
|
||||||
|
if GLOBAL_POINTERS {
|
||||||
|
if let crate::TypeInner::Pointer { base, .. } = *inner {
|
||||||
|
let base_inner = &module.types[base].inner;
|
||||||
|
if let crate::TypeInner::Struct { .. } = *base_inner {
|
||||||
|
return Ok(MaybeOwned::Borrowed(base_inner));
|
||||||
|
}
|
||||||
|
}
|
||||||
|
} else {
|
||||||
|
if let crate::TypeInner::Struct { .. } = *inner {
|
||||||
|
return Ok(MaybeOwned::Borrowed(inner));
|
||||||
|
}
|
||||||
|
}
|
||||||
|
write!(self.out, "{}.", OUTPUT_STRUCT_NAME)?;
|
||||||
|
}
|
||||||
|
spirv::StorageClass::Input => {
|
||||||
|
if let Some(crate::Binding::Location(_)) = var.binding {
|
||||||
|
write!(self.out, "{}.", LOCATION_INPUT_STRUCT_NAME)?;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
_ => {}
|
||||||
|
}
|
||||||
|
let name = var.name.or_index(handle);
|
||||||
|
write!(self.out, "{}", name)?;
|
||||||
|
Ok(MaybeOwned::Borrowed(inner))
|
||||||
|
}
|
||||||
|
crate::Expression::LocalVariable(handle) => {
|
||||||
|
let var = &function.local_variables[handle];
|
||||||
|
let inner = &module.types[var.ty].inner;
|
||||||
|
let name = var.name.or_index(handle);
|
||||||
|
write!(self.out, "{}", name)?;
|
||||||
|
Ok(MaybeOwned::Borrowed(inner))
|
||||||
|
}
|
||||||
|
crate::Expression::Load { pointer } => {
|
||||||
|
//write!(self.out, "*")?;
|
||||||
|
match *self.put_expression(pointer, function, module)?.borrow() {
|
||||||
|
crate::TypeInner::Pointer { base, .. } => {
|
||||||
|
Ok(module.borrow_type(base))
|
||||||
|
}
|
||||||
|
ref other => panic!("Unexpected load pointer {:?}", other),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
crate::Expression::Unary { op, expr } => {
|
||||||
|
let op_str = match op {
|
||||||
|
crate::UnaryOperator::Negate => "-",
|
||||||
|
crate::UnaryOperator::Not => "!",
|
||||||
|
};
|
||||||
|
write!(self.out, "{}", op_str)?;
|
||||||
|
self.put_expression(expr, function, module)
|
||||||
|
}
|
||||||
|
crate::Expression::Binary { op, left, right } => {
|
||||||
|
let op_str = match op {
|
||||||
|
crate::BinaryOperator::Add => "+",
|
||||||
|
crate::BinaryOperator::Subtract => "-",
|
||||||
|
crate::BinaryOperator::Multiply => "*",
|
||||||
|
crate::BinaryOperator::Divide => "/",
|
||||||
|
crate::BinaryOperator::Modulo => "%",
|
||||||
|
crate::BinaryOperator::Equal => "==",
|
||||||
|
crate::BinaryOperator::NotEqual => "!=",
|
||||||
|
crate::BinaryOperator::Less => "<",
|
||||||
|
crate::BinaryOperator::LessEqual => "<=",
|
||||||
|
crate::BinaryOperator::Greater => "==",
|
||||||
|
crate::BinaryOperator::GreaterEqual => ">=",
|
||||||
|
_ => panic!("Unsupported binary op {:?}", op),
|
||||||
|
};
|
||||||
|
//write!(self.out, "(")?;
|
||||||
|
let ty_left = self.put_expression(left, function, module)?;
|
||||||
|
write!(self.out, " {} ", op_str)?;
|
||||||
|
let ty_right = self.put_expression(right, function, module)?;
|
||||||
|
//write!(self.out, ")")?;
|
||||||
|
|
||||||
|
Ok(if op_str.len() == 1 {
|
||||||
|
match (ty_left.borrow(), ty_right.borrow()) {
|
||||||
|
(&crate::TypeInner::Scalar { kind, width }, &crate::TypeInner::Scalar { .. }) =>
|
||||||
|
MaybeOwned::Owned(crate::TypeInner::Scalar { kind, width }),
|
||||||
|
(&crate::TypeInner::Scalar { .. }, &crate::TypeInner::Vector { size, kind, width }) |
|
||||||
|
(&crate::TypeInner::Vector { size, kind, width }, &crate::TypeInner::Scalar { .. }) |
|
||||||
|
(&crate::TypeInner::Vector { size, kind, width }, &crate::TypeInner::Vector { .. }) =>
|
||||||
|
MaybeOwned::Owned(crate::TypeInner::Vector { size, kind, width }),
|
||||||
|
other => panic!("Unable to infer {:?} for {:?}", op, other),
|
||||||
|
}
|
||||||
|
} else {
|
||||||
|
MaybeOwned::Owned(crate::TypeInner::Scalar { kind: crate::ScalarKind::Bool, width: 1 })
|
||||||
|
})
|
||||||
|
}
|
||||||
|
crate::Expression::ImageSample { image, sampler, coordinate } => {
|
||||||
|
let ty_image = self.put_expression(image, function, module)?;
|
||||||
|
write!(self.out, ".sample(")?;
|
||||||
|
self.put_expression(sampler, function, module)?;
|
||||||
|
write!(self.out, ", ")?;
|
||||||
|
self.put_expression(coordinate, function, module)?;
|
||||||
|
write!(self.out, ")")?;
|
||||||
|
match *ty_image.borrow() {
|
||||||
|
crate::TypeInner::Image { base, .. } => Ok(module.borrow_type(base)),
|
||||||
|
ref other => panic!("Unexpected image type {:?}", other),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
crate::Expression::Call { ref name, ref arguments } => {
|
||||||
|
match name.as_str() {
|
||||||
|
"cos" |
|
||||||
|
"normalize" |
|
||||||
|
"sin" => {
|
||||||
|
write!(self.out, "{}(", name)?;
|
||||||
|
let result = self.put_expression(arguments[0], function, module)?;
|
||||||
|
write!(self.out, ")")?;
|
||||||
|
Ok(result)
|
||||||
|
}
|
||||||
|
"fclamp" => {
|
||||||
|
write!(self.out, "clamp(")?;
|
||||||
|
let result = self.put_expression(arguments[0], function, module)?;
|
||||||
|
write!(self.out, ", ")?;
|
||||||
|
self.put_expression(arguments[1], function, module)?;
|
||||||
|
write!(self.out, ", ")?;
|
||||||
|
self.put_expression(arguments[2], function, module)?;
|
||||||
|
write!(self.out, ")")?;
|
||||||
|
Ok(result)
|
||||||
|
}
|
||||||
|
"atan2" => {
|
||||||
|
write!(self.out, "{}(", name)?;
|
||||||
|
let result = self.put_expression(arguments[0], function, module)?;
|
||||||
|
write!(self.out, ", ")?;
|
||||||
|
self.put_expression(arguments[1], function, module)?;
|
||||||
|
write!(self.out, ")")?;
|
||||||
|
Ok(result)
|
||||||
|
}
|
||||||
|
"distance" => {
|
||||||
|
write!(self.out, "distance(")?;
|
||||||
|
let result = match *self.put_expression(arguments[0], function, module)?.borrow() {
|
||||||
|
crate::TypeInner::Vector { kind, width, .. } => crate::TypeInner::Scalar { kind, width },
|
||||||
|
ref other => panic!("Unexpected distance argument {:?}", other),
|
||||||
|
};
|
||||||
|
write!(self.out, ", ")?;
|
||||||
|
self.put_expression(arguments[1], function, module)?;
|
||||||
|
write!(self.out, ")")?;
|
||||||
|
Ok(MaybeOwned::Owned(result))
|
||||||
|
}
|
||||||
|
"length" => {
|
||||||
|
write!(self.out, "length(")?;
|
||||||
|
let result = match *self.put_expression(arguments[0], function, module)?.borrow() {
|
||||||
|
crate::TypeInner::Vector { kind, width, .. } => crate::TypeInner::Scalar { kind, width },
|
||||||
|
ref other => panic!("Unexpected distance argument {:?}", other),
|
||||||
|
};
|
||||||
|
write!(self.out, ")")?;
|
||||||
|
Ok(MaybeOwned::Owned(result))
|
||||||
|
}
|
||||||
|
|
||||||
|
_ => panic!("Unsupported call to '{}'", name),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
ref other => panic!("Unsupported {:?}", other),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
fn put_constant<'a>(
|
||||||
|
&mut self,
|
||||||
|
handle: Handle<crate::Constant>,
|
||||||
|
module: &'a crate::Module,
|
||||||
|
) -> Result<MaybeOwned<'a, crate::TypeInner>, Error> {
|
||||||
|
let constant = &module.constants[handle];
|
||||||
|
let ty = &module.types[constant.ty];
|
||||||
|
|
||||||
|
match constant.inner {
|
||||||
|
crate::ConstantInner::Sint(value) => {
|
||||||
|
write!(self.out, "{}", value)?;
|
||||||
|
}
|
||||||
|
crate::ConstantInner::Uint(value) => {
|
||||||
|
write!(self.out, "{}", value)?;
|
||||||
|
}
|
||||||
|
crate::ConstantInner::Float(value) => {
|
||||||
|
write!(self.out, "{}", value)?;
|
||||||
|
if value.fract() == 0.0 {
|
||||||
|
self.out.write_str(".0")?;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
crate::ConstantInner::Bool(value) => {
|
||||||
|
write!(self.out, "{}", value)?;
|
||||||
|
}
|
||||||
|
crate::ConstantInner::Composite(ref constituents) => {
|
||||||
|
let ty_name = ty.name.or_index(constant.ty);
|
||||||
|
write!(self.out, "{}(", ty_name)?;
|
||||||
|
for (i, handle) in constituents.iter().enumerate() {
|
||||||
|
if i != 0 {
|
||||||
|
write!(self.out, ", ")?;
|
||||||
|
}
|
||||||
|
self.put_constant(*handle, module)?;
|
||||||
|
}
|
||||||
|
write!(self.out, ")")?;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
Ok(MaybeOwned::Borrowed(&ty.inner))
|
||||||
|
}
|
||||||
|
|
||||||
|
fn put_statement<'a>(
|
||||||
|
&mut self,
|
||||||
|
level: Level,
|
||||||
|
statement: &crate::Statement,
|
||||||
|
function: &crate::Function,
|
||||||
|
has_output: bool,
|
||||||
|
module: &'a crate::Module,
|
||||||
|
) -> Result<(), Error> {
|
||||||
|
log::trace!("statement[{}] {:?}", level.0, statement);
|
||||||
|
match *statement {
|
||||||
|
crate::Statement::Empty => {}
|
||||||
|
crate::Statement::If { condition, ref accept, ref reject } => {
|
||||||
|
write!(self.out, "{}if (", level)?;
|
||||||
|
self.put_expression(condition, function, module)?;
|
||||||
|
writeln!(self.out, ") {{")?;
|
||||||
|
for s in accept {
|
||||||
|
self.put_statement(level.next(), s, function, has_output, module)?;
|
||||||
|
}
|
||||||
|
if !reject.is_empty() {
|
||||||
|
writeln!(self.out, "{}}} else {{", level)?;
|
||||||
|
for s in reject {
|
||||||
|
self.put_statement(level.next(), s, function, has_output, module)?;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
writeln!(self.out, "{}}}", level)?;
|
||||||
|
}
|
||||||
|
crate::Statement::Loop { ref body, ref continuing } => {
|
||||||
|
writeln!(self.out, "{}while(true) {{", level)?;
|
||||||
|
for s in body {
|
||||||
|
self.put_statement(level.next(), s, function, has_output, module)?;
|
||||||
|
}
|
||||||
|
if !continuing.is_empty() {
|
||||||
|
//TODO
|
||||||
|
}
|
||||||
|
writeln!(self.out, "{}}}", level)?;
|
||||||
|
}
|
||||||
|
crate::Statement::Store { pointer, value } => {
|
||||||
|
//write!(self.out, "\t*")?;
|
||||||
|
write!(self.out, "{}", level)?;
|
||||||
|
self.put_expression(pointer, function, module)?;
|
||||||
|
write!(self.out, " = ")?;
|
||||||
|
self.put_expression(value, function, module)?;
|
||||||
|
writeln!(self.out, ";")?;
|
||||||
|
}
|
||||||
|
crate::Statement::Break => {
|
||||||
|
writeln!(self.out, "{}break;", level)?;
|
||||||
|
}
|
||||||
|
crate::Statement::Continue => {
|
||||||
|
writeln!(self.out, "{}continue;", level)?;
|
||||||
|
}
|
||||||
|
crate::Statement::Return { value } => {
|
||||||
|
write!(self.out, "{}return ", level)?;
|
||||||
|
match value {
|
||||||
|
None if has_output => self.out.write_str(OUTPUT_STRUCT_NAME)?,
|
||||||
|
None => {}
|
||||||
|
Some(expr_handle) if has_output => {
|
||||||
|
panic!("Unable to return value {:?} from an entry point!", expr_handle)
|
||||||
|
}
|
||||||
|
Some(expr_handle) => {
|
||||||
|
self.put_expression(expr_handle, function, module)?;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
writeln!(self.out, ";")?;
|
||||||
|
}
|
||||||
|
_ => panic!("Unsupported {:?}", statement),
|
||||||
|
};
|
||||||
|
Ok(())
|
||||||
|
}
|
||||||
|
|
||||||
|
pub fn write(&mut self, module: &crate::Module, options: Options) -> Result<(), Error> {
|
||||||
|
writeln!(self.out, "#include <metal_stdlib>")?;
|
||||||
|
writeln!(self.out, "#include <simd/simd.h>")?;
|
||||||
|
writeln!(self.out, "using namespace metal;")?;
|
||||||
|
|
||||||
|
writeln!(self.out)?;
|
||||||
|
self.write_type_defs(module, options)?;
|
||||||
|
|
||||||
|
writeln!(self.out)?;
|
||||||
|
self.write_functions(module, options)?;
|
||||||
|
|
||||||
|
Ok(())
|
||||||
|
}
|
||||||
|
|
||||||
|
fn write_type_defs(&mut self, module: &crate::Module, options: Options) -> Result<(), Error> {
|
||||||
|
for (handle, ty) in module.types.iter() {
|
||||||
|
let name = ty.name.or_index(handle);
|
||||||
|
match ty.inner {
|
||||||
|
crate::TypeInner::Scalar { kind, .. } => {
|
||||||
|
write!(self.out, "typedef {} {}", scalar_kind_string(kind), name)?;
|
||||||
|
},
|
||||||
|
crate::TypeInner::Vector { size, kind, .. } => {
|
||||||
|
write!(self.out, "typedef {}{} {}", scalar_kind_string(kind), vector_size_string(size), name)?;
|
||||||
|
},
|
||||||
|
crate::TypeInner::Matrix { columns, rows, kind, .. } => {
|
||||||
|
write!(self.out, "typedef {}{}x{} {}", scalar_kind_string(kind), vector_size_string(columns), vector_size_string(rows), name)?;
|
||||||
|
}
|
||||||
|
crate::TypeInner::Pointer { base, class } => {
|
||||||
|
let base_name = module.types[base].name.or_index(base);
|
||||||
|
let class_name = match class {
|
||||||
|
spirv::StorageClass::Input |
|
||||||
|
spirv::StorageClass::Output => continue,
|
||||||
|
spirv::StorageClass::Uniform |
|
||||||
|
spirv::StorageClass::UniformConstant => "constant",
|
||||||
|
other => {
|
||||||
|
log::warn!("Unexpected pointer class {:?}", other);
|
||||||
|
""
|
||||||
|
}
|
||||||
|
};
|
||||||
|
write!(self.out, "typedef {} {} *{}", class_name, base_name, name)?;
|
||||||
|
}
|
||||||
|
crate::TypeInner::Array { base, size } => {
|
||||||
|
let base_name = module.types[base].name.or_index(base);
|
||||||
|
let resolved_size = match size {
|
||||||
|
crate::ArraySize::Static(length) => length,
|
||||||
|
crate::ArraySize::Dynamic => 1,
|
||||||
|
};
|
||||||
|
write!(self.out, "typedef {} {}[{}]", base_name, name, resolved_size)?;
|
||||||
|
}
|
||||||
|
crate::TypeInner::Struct { ref members } => {
|
||||||
|
writeln!(self.out, "struct {} {{", name)?;
|
||||||
|
for (index, member) in members.iter().enumerate() {
|
||||||
|
let name = member.name.or_index(MemberIndex(index));
|
||||||
|
let base_name = module.types[member.ty].name.or_index(member.ty);
|
||||||
|
write!(self.out, "\t{} {}", base_name, name)?;
|
||||||
|
if let Some(ref binding) = member.binding {
|
||||||
|
let resolved = options.resolve_binding(binding, LocationMode::Intermediate)?;
|
||||||
|
write!(self.out, " [[{}]]", resolved)?;
|
||||||
|
}
|
||||||
|
writeln!(self.out, ";")?;
|
||||||
|
}
|
||||||
|
write!(self.out, "}}")?;
|
||||||
|
}
|
||||||
|
crate::TypeInner::Image { base, dim, flags } => {
|
||||||
|
let base_name = module.types[base].name.or_index(base);
|
||||||
|
let dim = match dim {
|
||||||
|
spirv::Dim::Dim1D => "1d",
|
||||||
|
spirv::Dim::Dim2D => "2d",
|
||||||
|
spirv::Dim::Dim3D => "3d",
|
||||||
|
spirv::Dim::DimCube => "Cube",
|
||||||
|
_ => panic!("Unsupported dim {:?}", dim),
|
||||||
|
};
|
||||||
|
let access = if flags.contains(crate::ImageFlags::SAMPLED) {
|
||||||
|
if flags.intersects(crate::ImageFlags::CAN_STORE) {
|
||||||
|
return Err(Error::InvalidImageFlags(flags));
|
||||||
|
}
|
||||||
|
"sample"
|
||||||
|
} else if flags.contains(crate::ImageFlags::CAN_LOAD | crate::ImageFlags::CAN_STORE) {
|
||||||
|
"read_write"
|
||||||
|
} else if flags.contains(crate::ImageFlags::CAN_STORE) {
|
||||||
|
"write"
|
||||||
|
} else if flags.contains(crate::ImageFlags::CAN_LOAD) {
|
||||||
|
"read"
|
||||||
|
} else {
|
||||||
|
return Err(Error::InvalidImageFlags(flags));
|
||||||
|
};
|
||||||
|
write!(self.out, "typedef texture{}<{}, access::{}> {}", dim, base_name, access, name)?;
|
||||||
|
}
|
||||||
|
crate::TypeInner::Sampler => {
|
||||||
|
write!(self.out, "typedef sampler {}", name)?;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
writeln!(self.out, ";")?;
|
||||||
|
}
|
||||||
|
Ok(())
|
||||||
|
}
|
||||||
|
|
||||||
|
fn write_functions(&mut self, module: &crate::Module, options: Options) -> Result<(), Error> {
|
||||||
|
for (fun_handle, fun) in module.functions.iter() {
|
||||||
|
let fun_name = fun.name.or_index(fun_handle);
|
||||||
|
// find the entry point(s) and inputs/outputs
|
||||||
|
let mut exec_model = None;
|
||||||
|
let mut last_used_global = None;
|
||||||
|
for ((handle, var), &usage) in module.global_variables.iter().zip(&fun.global_usage) {
|
||||||
|
match var.class {
|
||||||
|
spirv::StorageClass::Input => {
|
||||||
|
if let Some(crate::Binding::Location(_)) = var.binding {
|
||||||
|
continue
|
||||||
|
}
|
||||||
|
}
|
||||||
|
spirv::StorageClass::Output => continue,
|
||||||
|
_ => {}
|
||||||
|
}
|
||||||
|
if !usage.is_empty() {
|
||||||
|
last_used_global = Some(handle);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
for ep in module.entry_points.iter() {
|
||||||
|
if ep.function == fun_handle {
|
||||||
|
if exec_model.is_some() {
|
||||||
|
if exec_model != Some(ep.exec_model) {
|
||||||
|
return Err(Error::MixedExecutionModels(fun_handle));
|
||||||
|
}
|
||||||
|
} else {
|
||||||
|
exec_model = Some(ep.exec_model);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
let output_name = fun.name.or_index(OutputStructIndex(fun_handle));
|
||||||
|
|
||||||
|
// make dedicated input/output structs
|
||||||
|
if let Some(em) = exec_model {
|
||||||
|
assert_eq!(fun.return_type, None);
|
||||||
|
let (em_str, in_mode, out_mode) = match em {
|
||||||
|
spirv::ExecutionModel::Vertex => ("vertex", LocationMode::VertexInput, LocationMode::Intermediate),
|
||||||
|
spirv::ExecutionModel::Fragment => ("fragment", LocationMode::Intermediate, LocationMode::FragmentOutput),
|
||||||
|
spirv::ExecutionModel::GLCompute => ("kernel", LocationMode::Uniform, LocationMode::Uniform),
|
||||||
|
_ => return Err(Error::UnsupportedExecutionModel(em)),
|
||||||
|
};
|
||||||
|
let location_input_name = fun.name.or_index(InputStructIndex(fun_handle));
|
||||||
|
|
||||||
|
if em != spirv::ExecutionModel::GLCompute {
|
||||||
|
writeln!(self.out, "struct {} {{", location_input_name)?;
|
||||||
|
for ((handle, var), &usage) in module.global_variables.iter().zip(&fun.global_usage) {
|
||||||
|
if var.class != spirv::StorageClass::Input || !usage.contains(crate::GlobalUse::LOAD) {
|
||||||
|
continue
|
||||||
|
}
|
||||||
|
// if it's a struct, lift all the built-in contents up to the root
|
||||||
|
let mut ty_handle = var.ty;
|
||||||
|
if GLOBAL_POINTERS {
|
||||||
|
if let crate::TypeInner::Pointer { base, .. } = module.types[var.ty].inner {
|
||||||
|
ty_handle = base;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
if let crate::TypeInner::Struct { ref members } = module.types[ty_handle].inner {
|
||||||
|
for (index, member) in members.iter().enumerate() {
|
||||||
|
if let Some(ref binding@crate::Binding::Location(_)) = member.binding {
|
||||||
|
let name = member.name.or_index(MemberIndex(index));
|
||||||
|
let ty_name = module.types[member.ty].name.or_index(member.ty);
|
||||||
|
let resolved = options.resolve_binding(binding, in_mode)?;
|
||||||
|
writeln!(self.out, "\t{} {} [[{}]];", ty_name, name, resolved)?;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
} else {
|
||||||
|
if let Some(ref binding@crate::Binding::Location(_)) = var.binding {
|
||||||
|
let tyvar = TypedGlobalVariable { module, handle, usage: crate::GlobalUse::empty() };
|
||||||
|
let resolved = options.resolve_binding(binding, in_mode)?;
|
||||||
|
writeln!(self.out, "\t{} [[{}]];", tyvar, resolved)?;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
writeln!(self.out, "}};")?;
|
||||||
|
writeln!(self.out, "struct {} {{", output_name)?;
|
||||||
|
for ((handle, var), &usage) in module.global_variables.iter().zip(&fun.global_usage) {
|
||||||
|
if var.class != spirv::StorageClass::Output || !usage.contains(crate::GlobalUse::STORE) {
|
||||||
|
continue
|
||||||
|
}
|
||||||
|
// if it's a struct, lift all the built-in contents up to the root
|
||||||
|
let mut ty_handle = var.ty;
|
||||||
|
if GLOBAL_POINTERS {
|
||||||
|
if let crate::TypeInner::Pointer { base, .. } = module.types[var.ty].inner {
|
||||||
|
ty_handle = base;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
if let crate::TypeInner::Struct { ref members } = module.types[ty_handle].inner {
|
||||||
|
for (index, member) in members.iter().enumerate() {
|
||||||
|
let name = member.name.or_index(MemberIndex(index));
|
||||||
|
let ty_name = module.types[member.ty].name.or_index(member.ty);
|
||||||
|
let binding = member.binding
|
||||||
|
.as_ref()
|
||||||
|
.ok_or(Error::MissingBinding(handle))?;
|
||||||
|
let resolved = options.resolve_binding(binding, out_mode)?;
|
||||||
|
writeln!(self.out, "\t{} {} [[{}]];", ty_name, name, resolved)?;
|
||||||
|
}
|
||||||
|
} else {
|
||||||
|
let tyvar = TypedGlobalVariable { module, handle, usage: crate::GlobalUse::empty() };
|
||||||
|
write!(self.out, "\t{}", tyvar)?;
|
||||||
|
if let Some(ref binding) = var.binding {
|
||||||
|
let resolved = options.resolve_binding(binding, out_mode)?;
|
||||||
|
write!(self.out, " [[{}]]", resolved)?;
|
||||||
|
}
|
||||||
|
writeln!(self.out, ";")?;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
writeln!(self.out, "}};")?;
|
||||||
|
writeln!(self.out, "{} {} {}(", em_str, output_name, fun_name)?;
|
||||||
|
let separator = separate(last_used_global.is_none());
|
||||||
|
writeln!(self.out, "\t{} {} [[stage_in]]{}",
|
||||||
|
location_input_name, LOCATION_INPUT_STRUCT_NAME, separator)?;
|
||||||
|
} else {
|
||||||
|
writeln!(self.out, "{} void {}(", em_str, fun_name)?;
|
||||||
|
}
|
||||||
|
|
||||||
|
for ((handle, var), &usage) in module.global_variables.iter().zip(&fun.global_usage) {
|
||||||
|
if usage.is_empty() || var.class == spirv::StorageClass::Output {
|
||||||
|
continue
|
||||||
|
}
|
||||||
|
if var.class == spirv::StorageClass::Input {
|
||||||
|
if let Some(crate::Binding::Location(_)) = var.binding {
|
||||||
|
// location inputs are put into a separate struct
|
||||||
|
continue
|
||||||
|
}
|
||||||
|
}
|
||||||
|
let loc_mode = match (em, var.class) {
|
||||||
|
(spirv::ExecutionModel::Vertex, spirv::StorageClass::Input) => LocationMode::VertexInput,
|
||||||
|
(spirv::ExecutionModel::Vertex, spirv::StorageClass::Output) |
|
||||||
|
(spirv::ExecutionModel::Fragment, spirv::StorageClass::Input) => LocationMode::Intermediate,
|
||||||
|
(spirv::ExecutionModel::Fragment, spirv::StorageClass::Output) => LocationMode::FragmentOutput,
|
||||||
|
_ => LocationMode::Uniform,
|
||||||
|
};
|
||||||
|
let resolved = options.resolve_binding(var.binding.as_ref().unwrap(), loc_mode)?;
|
||||||
|
let tyvar = TypedGlobalVariable { module, handle, usage };
|
||||||
|
let separator = separate(last_used_global == Some(handle));
|
||||||
|
writeln!(self.out, "\t{} [[{}]]{}", tyvar, resolved, separator)?;
|
||||||
|
}
|
||||||
|
} else {
|
||||||
|
let result_type_name = match fun.return_type {
|
||||||
|
Some(type_id) => module.types[type_id].name.or_index(type_id),
|
||||||
|
None => Name {
|
||||||
|
class: "",
|
||||||
|
source: NameSource::Custom { name: "void", prefix: false },
|
||||||
|
},
|
||||||
|
};
|
||||||
|
writeln!(self.out, "{} {}(", result_type_name, fun_name)?;
|
||||||
|
for (index, &ty) in fun.parameter_types.iter().enumerate() {
|
||||||
|
let name = Name::from(ParameterIndex(index));
|
||||||
|
let member_type_name = module.types[ty].name.or_index(ty);
|
||||||
|
let separator = separate(index + 1 == fun.parameter_types.len() && last_used_global.is_none());
|
||||||
|
writeln!(self.out, "\t{} {}{}", member_type_name, name, separator)?;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
writeln!(self.out, ") {{")?;
|
||||||
|
|
||||||
|
// write down function body
|
||||||
|
let has_output = match exec_model {
|
||||||
|
Some(spirv::ExecutionModel::Vertex) |
|
||||||
|
Some(spirv::ExecutionModel::Fragment) => {
|
||||||
|
writeln!(self.out, "\t{} {};", output_name, OUTPUT_STRUCT_NAME)?;
|
||||||
|
true
|
||||||
|
}
|
||||||
|
_ => false
|
||||||
|
};
|
||||||
|
for (local_handle, local) in fun.local_variables.iter() {
|
||||||
|
let ty_name = module.types[local.ty].name.or_index(local.ty);
|
||||||
|
write!(self.out, "\t{} {}", ty_name, local.name.or_index(local_handle))?;
|
||||||
|
if let Some(value) = local.init {
|
||||||
|
write!(self.out, " = ")?;
|
||||||
|
self.put_expression(value, fun, module)?;
|
||||||
|
}
|
||||||
|
writeln!(self.out, ";")?;
|
||||||
|
}
|
||||||
|
for statement in fun.body.iter() {
|
||||||
|
self.put_statement(Level(1), statement, fun, has_output, module)?;
|
||||||
|
}
|
||||||
|
writeln!(self.out, "}}")?;
|
||||||
|
}
|
||||||
|
|
||||||
|
Ok(())
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
pub fn write_string(module: &crate::Module, options: Options) -> Result<String, Error> {
|
||||||
|
let mut w = Writer { out: String::new() };
|
||||||
|
w.write(module, options)?;
|
||||||
|
Ok(w.out)
|
||||||
|
}
|
||||||
26
third_party/rust/naga/src/front/mod.rs
vendored
Normal file
26
third_party/rust/naga/src/front/mod.rs
vendored
Normal file
@ -0,0 +1,26 @@
|
|||||||
|
pub mod spirv;
|
||||||
|
pub mod wgsl;
|
||||||
|
|
||||||
|
use crate::arena::Arena;
|
||||||
|
|
||||||
|
pub const GENERATOR: u32 = 0;
|
||||||
|
|
||||||
|
impl crate::Module {
|
||||||
|
fn from_header(header: crate::Header) -> Self {
|
||||||
|
crate::Module {
|
||||||
|
header,
|
||||||
|
types: Arena::new(),
|
||||||
|
constants: Arena::new(),
|
||||||
|
global_variables: Arena::new(),
|
||||||
|
functions: Arena::new(),
|
||||||
|
entry_points: Vec::new(),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
fn generate_empty() -> Self {
|
||||||
|
Self::from_header(crate::Header {
|
||||||
|
version: (1, 0, 0),
|
||||||
|
generator: GENERATOR,
|
||||||
|
})
|
||||||
|
}
|
||||||
|
}
|
||||||
1515
third_party/rust/naga/src/front/spirv.rs
vendored
Normal file
1515
third_party/rust/naga/src/front/spirv.rs
vendored
Normal file
File diff suppressed because it is too large
Load Diff
1402
third_party/rust/naga/src/front/wgsl.rs
vendored
Normal file
1402
third_party/rust/naga/src/front/wgsl.rs
vendored
Normal file
File diff suppressed because it is too large
Load Diff
292
third_party/rust/naga/src/lib.rs
vendored
Normal file
292
third_party/rust/naga/src/lib.rs
vendored
Normal file
@ -0,0 +1,292 @@
|
|||||||
|
#![allow(clippy::new_without_default)]
|
||||||
|
|
||||||
|
mod arena;
|
||||||
|
pub mod back;
|
||||||
|
pub mod front;
|
||||||
|
pub mod proc;
|
||||||
|
|
||||||
|
use crate::arena::{Arena, Handle};
|
||||||
|
|
||||||
|
use std::{
|
||||||
|
collections::{HashMap, HashSet},
|
||||||
|
hash::BuildHasherDefault,
|
||||||
|
};
|
||||||
|
|
||||||
|
|
||||||
|
pub type FastHashMap<K, T> = HashMap<K, T, BuildHasherDefault<fxhash::FxHasher>>;
|
||||||
|
pub type FastHashSet<K> = HashSet<K, BuildHasherDefault<fxhash::FxHasher>>;
|
||||||
|
|
||||||
|
#[derive(Clone, Debug)]
|
||||||
|
pub struct Header {
|
||||||
|
pub version: (u8, u8, u8),
|
||||||
|
pub generator: u32,
|
||||||
|
}
|
||||||
|
|
||||||
|
pub type Bytes = u8;
|
||||||
|
|
||||||
|
#[repr(u8)]
|
||||||
|
#[derive(Clone, Copy, Debug, PartialEq)]
|
||||||
|
pub enum VectorSize {
|
||||||
|
Bi = 2,
|
||||||
|
Tri = 3,
|
||||||
|
Quad = 4,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[repr(u8)]
|
||||||
|
#[derive(Clone, Copy, Debug, PartialEq)]
|
||||||
|
pub enum ScalarKind {
|
||||||
|
Sint,
|
||||||
|
Uint,
|
||||||
|
Float,
|
||||||
|
Bool,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[repr(u8)]
|
||||||
|
#[derive(Clone, Copy, Debug, PartialEq)]
|
||||||
|
pub enum ArraySize {
|
||||||
|
Static(spirv::Word),
|
||||||
|
Dynamic,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug, PartialEq)]
|
||||||
|
pub struct StructMember {
|
||||||
|
pub name: Option<String>,
|
||||||
|
pub binding: Option<Binding>,
|
||||||
|
pub ty: Handle<Type>,
|
||||||
|
}
|
||||||
|
|
||||||
|
bitflags::bitflags! {
|
||||||
|
pub struct ImageFlags: u32 {
|
||||||
|
const ARRAYED = 0x1;
|
||||||
|
const MULTISAMPLED = 0x2;
|
||||||
|
const SAMPLED = 0x4;
|
||||||
|
const CAN_LOAD = 0x10;
|
||||||
|
const CAN_STORE = 0x20;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug)]
|
||||||
|
pub struct Type {
|
||||||
|
pub name: Option<String>,
|
||||||
|
pub inner: TypeInner,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug, PartialEq)]
|
||||||
|
pub enum TypeInner {
|
||||||
|
Scalar { kind: ScalarKind, width: Bytes },
|
||||||
|
Vector { size: VectorSize, kind: ScalarKind, width: Bytes },
|
||||||
|
Matrix { columns: VectorSize, rows: VectorSize, kind: ScalarKind, width: Bytes },
|
||||||
|
Pointer { base: Handle<Type>, class: spirv::StorageClass },
|
||||||
|
Array { base: Handle<Type>, size: ArraySize },
|
||||||
|
Struct { members: Vec<StructMember> },
|
||||||
|
Image { base: Handle<Type>, dim: spirv::Dim, flags: ImageFlags },
|
||||||
|
Sampler,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug, PartialEq)]
|
||||||
|
pub struct Constant {
|
||||||
|
pub name: Option<String>,
|
||||||
|
pub specialization: Option<spirv::Word>,
|
||||||
|
pub inner: ConstantInner,
|
||||||
|
pub ty: Handle<Type>,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug, PartialEq)]
|
||||||
|
pub enum ConstantInner {
|
||||||
|
Sint(i64),
|
||||||
|
Uint(u64),
|
||||||
|
Float(f64),
|
||||||
|
Bool(bool),
|
||||||
|
Composite(Vec<Handle<Constant>>),
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Clone, Debug, PartialEq)]
|
||||||
|
pub enum Binding {
|
||||||
|
BuiltIn(spirv::BuiltIn),
|
||||||
|
Location(spirv::Word),
|
||||||
|
Descriptor { set: spirv::Word, binding: spirv::Word },
|
||||||
|
}
|
||||||
|
|
||||||
|
bitflags::bitflags! {
|
||||||
|
pub struct GlobalUse: u8 {
|
||||||
|
const LOAD = 0x1;
|
||||||
|
const STORE = 0x2;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Clone, Debug)]
|
||||||
|
pub struct GlobalVariable {
|
||||||
|
pub name: Option<String>,
|
||||||
|
pub class: spirv::StorageClass,
|
||||||
|
pub binding: Option<Binding>,
|
||||||
|
pub ty: Handle<Type>,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Clone, Debug)]
|
||||||
|
pub struct LocalVariable {
|
||||||
|
pub name: Option<String>,
|
||||||
|
pub ty: Handle<Type>,
|
||||||
|
pub init: Option<Handle<Expression>>,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Clone, Copy, Debug, PartialEq)]
|
||||||
|
pub enum UnaryOperator {
|
||||||
|
Negate,
|
||||||
|
Not,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Clone, Copy, Debug, PartialEq)]
|
||||||
|
pub enum BinaryOperator {
|
||||||
|
Add,
|
||||||
|
Subtract,
|
||||||
|
Multiply,
|
||||||
|
Divide,
|
||||||
|
Modulo,
|
||||||
|
Equal,
|
||||||
|
NotEqual,
|
||||||
|
Less,
|
||||||
|
LessEqual,
|
||||||
|
Greater,
|
||||||
|
GreaterEqual,
|
||||||
|
And,
|
||||||
|
ExclusiveOr,
|
||||||
|
InclusiveOr,
|
||||||
|
LogicalAnd,
|
||||||
|
LogicalOr,
|
||||||
|
ShiftLeftLogical,
|
||||||
|
ShiftRightLogical,
|
||||||
|
ShiftRightArithmetic,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Clone, Copy, Debug, PartialEq)]
|
||||||
|
pub enum IntrinsicFunction {
|
||||||
|
Any,
|
||||||
|
All,
|
||||||
|
IsNan,
|
||||||
|
IsInf,
|
||||||
|
IsFinite,
|
||||||
|
IsNormal,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Clone, Copy, Debug, PartialEq)]
|
||||||
|
pub enum DerivativeAxis {
|
||||||
|
X,
|
||||||
|
Y,
|
||||||
|
Width,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Clone, Debug)]
|
||||||
|
pub enum Expression {
|
||||||
|
Access {
|
||||||
|
base: Handle<Expression>,
|
||||||
|
index: Handle<Expression>, //int
|
||||||
|
},
|
||||||
|
AccessIndex {
|
||||||
|
base: Handle<Expression>,
|
||||||
|
index: u32,
|
||||||
|
},
|
||||||
|
Constant(Handle<Constant>),
|
||||||
|
Compose {
|
||||||
|
ty: Handle<Type>,
|
||||||
|
components: Vec<Handle<Expression>>,
|
||||||
|
},
|
||||||
|
FunctionParameter(u32),
|
||||||
|
GlobalVariable(Handle<GlobalVariable>),
|
||||||
|
LocalVariable(Handle<LocalVariable>),
|
||||||
|
Load {
|
||||||
|
pointer: Handle<Expression>,
|
||||||
|
},
|
||||||
|
ImageSample {
|
||||||
|
image: Handle<Expression>,
|
||||||
|
sampler: Handle<Expression>,
|
||||||
|
coordinate: Handle<Expression>,
|
||||||
|
},
|
||||||
|
Unary {
|
||||||
|
op: UnaryOperator,
|
||||||
|
expr: Handle<Expression>,
|
||||||
|
},
|
||||||
|
Binary {
|
||||||
|
op: BinaryOperator,
|
||||||
|
left: Handle<Expression>,
|
||||||
|
right: Handle<Expression>,
|
||||||
|
},
|
||||||
|
Intrinsic {
|
||||||
|
fun: IntrinsicFunction,
|
||||||
|
argument: Handle<Expression>,
|
||||||
|
},
|
||||||
|
DotProduct(Handle<Expression>, Handle<Expression>),
|
||||||
|
CrossProduct(Handle<Expression>, Handle<Expression>),
|
||||||
|
Derivative {
|
||||||
|
axis: DerivativeAxis,
|
||||||
|
//modifier,
|
||||||
|
expr: Handle<Expression>,
|
||||||
|
},
|
||||||
|
Call {
|
||||||
|
name: String,
|
||||||
|
arguments: Vec<Handle<Expression>>,
|
||||||
|
},
|
||||||
|
}
|
||||||
|
|
||||||
|
pub type Block = Vec<Statement>;
|
||||||
|
#[derive(Debug)]
|
||||||
|
pub struct FallThrough;
|
||||||
|
|
||||||
|
#[derive(Debug)]
|
||||||
|
pub enum Statement {
|
||||||
|
Empty,
|
||||||
|
Block(Block),
|
||||||
|
If {
|
||||||
|
condition: Handle<Expression>, //bool
|
||||||
|
accept: Block,
|
||||||
|
reject: Block,
|
||||||
|
},
|
||||||
|
Switch {
|
||||||
|
selector: Handle<Expression>, //int
|
||||||
|
cases: FastHashMap<i32, (Block, Option<FallThrough>)>,
|
||||||
|
default: Block,
|
||||||
|
},
|
||||||
|
Loop {
|
||||||
|
body: Block,
|
||||||
|
continuing: Block,
|
||||||
|
},
|
||||||
|
//TODO: move terminator variations into a separate enum?
|
||||||
|
Break,
|
||||||
|
Continue,
|
||||||
|
Return {
|
||||||
|
value: Option<Handle<Expression>>,
|
||||||
|
},
|
||||||
|
Kill,
|
||||||
|
Store {
|
||||||
|
pointer: Handle<Expression>,
|
||||||
|
value: Handle<Expression>,
|
||||||
|
},
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug)]
|
||||||
|
pub struct Function {
|
||||||
|
pub name: Option<String>,
|
||||||
|
pub control: spirv::FunctionControl,
|
||||||
|
pub parameter_types: Vec<Handle<Type>>,
|
||||||
|
pub return_type: Option<Handle<Type>>,
|
||||||
|
pub global_usage: Vec<GlobalUse>,
|
||||||
|
pub local_variables: Arena<LocalVariable>,
|
||||||
|
pub expressions: Arena<Expression>,
|
||||||
|
pub body: Block,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug)]
|
||||||
|
pub struct EntryPoint {
|
||||||
|
pub exec_model: spirv::ExecutionModel,
|
||||||
|
pub name: String,
|
||||||
|
pub function: Handle<Function>,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug)]
|
||||||
|
pub struct Module {
|
||||||
|
pub header: Header,
|
||||||
|
pub types: Arena<Type>,
|
||||||
|
pub constants: Arena<Constant>,
|
||||||
|
pub global_variables: Arena<GlobalVariable>,
|
||||||
|
pub functions: Arena<Function>,
|
||||||
|
pub entry_points: Vec<EntryPoint>,
|
||||||
|
}
|
||||||
139
third_party/rust/naga/src/proc/interface.rs
vendored
Normal file
139
third_party/rust/naga/src/proc/interface.rs
vendored
Normal file
@ -0,0 +1,139 @@
|
|||||||
|
use crate::{
|
||||||
|
arena::{Arena, Handle},
|
||||||
|
};
|
||||||
|
|
||||||
|
struct Interface<'a> {
|
||||||
|
expressions: &'a Arena<crate::Expression>,
|
||||||
|
uses: Vec<crate::GlobalUse>,
|
||||||
|
}
|
||||||
|
|
||||||
|
impl<'a> Interface<'a> {
|
||||||
|
fn add_inputs(&mut self, handle: Handle<crate::Expression>) {
|
||||||
|
use crate::Expression as E;
|
||||||
|
match self.expressions[handle] {
|
||||||
|
E::Access { base, index } => {
|
||||||
|
self.add_inputs(base);
|
||||||
|
self.add_inputs(index);
|
||||||
|
}
|
||||||
|
E::AccessIndex { base, .. } => {
|
||||||
|
self.add_inputs(base);
|
||||||
|
}
|
||||||
|
E::Constant(_) => {}
|
||||||
|
E::Compose { ref components, .. } => {
|
||||||
|
for &comp in components {
|
||||||
|
self.add_inputs(comp);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
E::FunctionParameter(_) => {},
|
||||||
|
E::GlobalVariable(handle) => {
|
||||||
|
self.uses[handle.index()] |= crate::GlobalUse::LOAD;
|
||||||
|
}
|
||||||
|
E::LocalVariable(_) => {}
|
||||||
|
E::Load { pointer } => {
|
||||||
|
self.add_inputs(pointer);
|
||||||
|
}
|
||||||
|
E::ImageSample { image, sampler, coordinate } => {
|
||||||
|
self.add_inputs(image);
|
||||||
|
self.add_inputs(sampler);
|
||||||
|
self.add_inputs(coordinate);
|
||||||
|
}
|
||||||
|
E::Unary { expr, .. } => {
|
||||||
|
self.add_inputs(expr);
|
||||||
|
}
|
||||||
|
E::Binary { left, right, .. } => {
|
||||||
|
self.add_inputs(left);
|
||||||
|
self.add_inputs(right);
|
||||||
|
}
|
||||||
|
E::Intrinsic { argument, .. } => {
|
||||||
|
self.add_inputs(argument);
|
||||||
|
}
|
||||||
|
E::DotProduct(left, right) => {
|
||||||
|
self.add_inputs(left);
|
||||||
|
self.add_inputs(right);
|
||||||
|
}
|
||||||
|
E::CrossProduct(left, right) => {
|
||||||
|
self.add_inputs(left);
|
||||||
|
self.add_inputs(right);
|
||||||
|
}
|
||||||
|
E::Derivative { expr, .. } => {
|
||||||
|
self.add_inputs(expr);
|
||||||
|
}
|
||||||
|
E::Call { ref arguments, .. } => {
|
||||||
|
for &argument in arguments {
|
||||||
|
self.add_inputs(argument);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
fn collect(&mut self, block: &[crate::Statement]) {
|
||||||
|
for statement in block {
|
||||||
|
use crate::Statement as S;
|
||||||
|
match *statement {
|
||||||
|
S::Empty |
|
||||||
|
S::Break |
|
||||||
|
S::Continue |
|
||||||
|
S::Kill => (),
|
||||||
|
S::Block(ref b) => {
|
||||||
|
self.collect(b);
|
||||||
|
}
|
||||||
|
S::If { condition, ref accept, ref reject } => {
|
||||||
|
self.add_inputs(condition);
|
||||||
|
self.collect(accept);
|
||||||
|
self.collect(reject);
|
||||||
|
}
|
||||||
|
S::Switch { selector, ref cases, ref default } => {
|
||||||
|
self.add_inputs(selector);
|
||||||
|
for &(ref case, _) in cases.values() {
|
||||||
|
self.collect(case);
|
||||||
|
}
|
||||||
|
self.collect(default);
|
||||||
|
}
|
||||||
|
S::Loop { ref body, ref continuing } => {
|
||||||
|
self.collect(body);
|
||||||
|
self.collect(continuing);
|
||||||
|
}
|
||||||
|
S::Return { value } => {
|
||||||
|
if let Some(expr) = value {
|
||||||
|
self.add_inputs(expr);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
S::Store { pointer, value } => {
|
||||||
|
let mut left = pointer;
|
||||||
|
loop {
|
||||||
|
match self.expressions[left] {
|
||||||
|
crate::Expression::Access { base, index } => {
|
||||||
|
self.add_inputs(index);
|
||||||
|
left = base;
|
||||||
|
}
|
||||||
|
crate::Expression::AccessIndex { base, .. } => {
|
||||||
|
left = base;
|
||||||
|
}
|
||||||
|
crate::Expression::GlobalVariable(handle) => {
|
||||||
|
self.uses[handle.index()] |= crate::GlobalUse::STORE;
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
_ => break,
|
||||||
|
}
|
||||||
|
}
|
||||||
|
self.add_inputs(value);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
impl crate::GlobalUse {
|
||||||
|
pub fn scan(
|
||||||
|
expressions: &Arena<crate::Expression>,
|
||||||
|
body: &[crate::Statement],
|
||||||
|
globals: &Arena<crate::GlobalVariable>,
|
||||||
|
) -> Vec<Self> {
|
||||||
|
let mut io = Interface {
|
||||||
|
expressions,
|
||||||
|
uses: vec![crate::GlobalUse::empty(); globals.len()],
|
||||||
|
};
|
||||||
|
io.collect(body);
|
||||||
|
io.uses
|
||||||
|
}
|
||||||
|
}
|
||||||
4
third_party/rust/naga/src/proc/mod.rs
vendored
Normal file
4
third_party/rust/naga/src/proc/mod.rs
vendored
Normal file
@ -0,0 +1,4 @@
|
|||||||
|
mod interface;
|
||||||
|
mod typifier;
|
||||||
|
|
||||||
|
pub use typifier::{ResolveError, Typifier};
|
||||||
155
third_party/rust/naga/src/proc/typifier.rs
vendored
Normal file
155
third_party/rust/naga/src/proc/typifier.rs
vendored
Normal file
@ -0,0 +1,155 @@
|
|||||||
|
use crate::arena::{Arena, Handle};
|
||||||
|
|
||||||
|
pub struct Typifier {
|
||||||
|
types: Vec<Handle<crate::Type>>,
|
||||||
|
}
|
||||||
|
|
||||||
|
#[derive(Debug)]
|
||||||
|
pub enum ResolveError {
|
||||||
|
InvalidAccessIndex,
|
||||||
|
}
|
||||||
|
|
||||||
|
impl Typifier {
|
||||||
|
pub fn new() -> Self {
|
||||||
|
Typifier {
|
||||||
|
types: Vec::new(),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
pub fn resolve(
|
||||||
|
&mut self,
|
||||||
|
expr_handle: Handle<crate::Expression>,
|
||||||
|
expressions: &Arena<crate::Expression>,
|
||||||
|
types: &mut Arena<crate::Type>,
|
||||||
|
constants: &Arena<crate::Constant>,
|
||||||
|
global_vars: &Arena<crate::GlobalVariable>,
|
||||||
|
local_vars: &Arena<crate::LocalVariable>,
|
||||||
|
) -> Result<Handle<crate::Type>, ResolveError> {
|
||||||
|
if self.types.len() <= expr_handle.index() {
|
||||||
|
for (eh, expr) in expressions.iter().skip(self.types.len()) {
|
||||||
|
let ty = match *expr {
|
||||||
|
crate::Expression::Access { base, .. } => {
|
||||||
|
match types[self.types[base.index()]].inner {
|
||||||
|
crate::TypeInner::Array { base, .. } => base,
|
||||||
|
ref other => panic!("Can't access into {:?}", other),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
crate::Expression::AccessIndex { base, index } => {
|
||||||
|
match types[self.types[base.index()]].inner {
|
||||||
|
crate::TypeInner::Vector { size, kind, width } => {
|
||||||
|
if index >= size as u32 {
|
||||||
|
return Err(ResolveError::InvalidAccessIndex)
|
||||||
|
}
|
||||||
|
let inner = crate::TypeInner::Scalar { kind, width };
|
||||||
|
Self::deduce_type_handle(inner, types)
|
||||||
|
}
|
||||||
|
crate::TypeInner::Matrix { columns, rows, kind, width } => {
|
||||||
|
if index >= columns as u32 {
|
||||||
|
return Err(ResolveError::InvalidAccessIndex)
|
||||||
|
}
|
||||||
|
let inner = crate::TypeInner::Vector { size: rows, kind, width };
|
||||||
|
Self::deduce_type_handle(inner, types)
|
||||||
|
}
|
||||||
|
crate::TypeInner::Array { base, .. } => base,
|
||||||
|
crate::TypeInner::Struct { ref members } => {
|
||||||
|
members.get(index as usize)
|
||||||
|
.ok_or(ResolveError::InvalidAccessIndex)?
|
||||||
|
.ty
|
||||||
|
}
|
||||||
|
ref other => panic!("Can't access into {:?}", other),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
crate::Expression::Constant(h) => constants[h].ty,
|
||||||
|
crate::Expression::Compose { ty, .. } => ty,
|
||||||
|
crate::Expression::FunctionParameter(_) => unimplemented!(),
|
||||||
|
crate::Expression::GlobalVariable(h) => global_vars[h].ty,
|
||||||
|
crate::Expression::LocalVariable(h) => local_vars[h].ty,
|
||||||
|
crate::Expression::Load { .. } => unimplemented!(),
|
||||||
|
crate::Expression::ImageSample { .. } => unimplemented!(),
|
||||||
|
crate::Expression::Unary { expr, .. } => self.types[expr.index()],
|
||||||
|
crate::Expression::Binary { op, left, right } => {
|
||||||
|
match op {
|
||||||
|
crate::BinaryOperator::Add |
|
||||||
|
crate::BinaryOperator::Subtract |
|
||||||
|
crate::BinaryOperator::Divide |
|
||||||
|
crate::BinaryOperator::Modulo => {
|
||||||
|
self.types[left.index()]
|
||||||
|
}
|
||||||
|
crate::BinaryOperator::Multiply => {
|
||||||
|
let ty_left = self.types[left.index()];
|
||||||
|
let ty_right = self.types[right.index()];
|
||||||
|
if ty_left == ty_right {
|
||||||
|
ty_left
|
||||||
|
} else if let crate::TypeInner::Scalar { .. } = types[ty_right].inner {
|
||||||
|
ty_left
|
||||||
|
} else if let crate::TypeInner::Matrix { columns, kind, width, .. } = types[ty_left].inner {
|
||||||
|
let inner = crate::TypeInner::Vector { size: columns, kind, width};
|
||||||
|
Self::deduce_type_handle(inner, types)
|
||||||
|
} else {
|
||||||
|
panic!("Incompatible arguments {:?} x {:?}", types[ty_left], types[ty_right]);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
crate::BinaryOperator::Equal |
|
||||||
|
crate::BinaryOperator::NotEqual |
|
||||||
|
crate::BinaryOperator::Less |
|
||||||
|
crate::BinaryOperator::LessEqual |
|
||||||
|
crate::BinaryOperator::Greater |
|
||||||
|
crate::BinaryOperator::GreaterEqual |
|
||||||
|
crate::BinaryOperator::LogicalAnd |
|
||||||
|
crate::BinaryOperator::LogicalOr => {
|
||||||
|
self.types[left.index()]
|
||||||
|
}
|
||||||
|
crate::BinaryOperator::And |
|
||||||
|
crate::BinaryOperator::ExclusiveOr |
|
||||||
|
crate::BinaryOperator::InclusiveOr |
|
||||||
|
crate::BinaryOperator::ShiftLeftLogical |
|
||||||
|
crate::BinaryOperator::ShiftRightLogical |
|
||||||
|
crate::BinaryOperator::ShiftRightArithmetic => {
|
||||||
|
self.types[left.index()]
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
crate::Expression::Intrinsic { .. } => unimplemented!(),
|
||||||
|
crate::Expression::DotProduct(_, _) => unimplemented!(),
|
||||||
|
crate::Expression::CrossProduct(_, _) => unimplemented!(),
|
||||||
|
crate::Expression::Derivative { .. } => unimplemented!(),
|
||||||
|
crate::Expression::Call { ref name, ref arguments } => {
|
||||||
|
match name.as_str() {
|
||||||
|
"distance" | "length" => {
|
||||||
|
let ty_handle = self.types[arguments[0].index()];
|
||||||
|
let inner = match types[ty_handle].inner {
|
||||||
|
crate::TypeInner::Vector { kind, width, .. } => {
|
||||||
|
crate::TypeInner::Scalar { kind, width }
|
||||||
|
}
|
||||||
|
ref other => panic!("Unexpected argument {:?}", other),
|
||||||
|
};
|
||||||
|
Self::deduce_type_handle(inner, types)
|
||||||
|
}
|
||||||
|
"normalize" | "fclamp" => self.types[arguments[0].index()],
|
||||||
|
_ => panic!("Unknown '{}' call", name),
|
||||||
|
}
|
||||||
|
}
|
||||||
|
};
|
||||||
|
log::debug!("Resolving {:?} = {:?} : {:?}", eh, expr, ty);
|
||||||
|
self.types.push(ty);
|
||||||
|
};
|
||||||
|
}
|
||||||
|
Ok(self.types[expr_handle.index()])
|
||||||
|
}
|
||||||
|
|
||||||
|
pub fn deduce_type_handle(
|
||||||
|
inner: crate::TypeInner,
|
||||||
|
arena: &mut Arena<crate::Type>,
|
||||||
|
) -> Handle<crate::Type> {
|
||||||
|
if let Some((token, _)) = arena
|
||||||
|
.iter()
|
||||||
|
.find(|(_, ty)| ty.inner == inner)
|
||||||
|
{
|
||||||
|
return token;
|
||||||
|
}
|
||||||
|
arena.append(crate::Type {
|
||||||
|
name: None,
|
||||||
|
inner,
|
||||||
|
})
|
||||||
|
}
|
||||||
|
}
|
||||||
7
third_party/rust/naga/test-data/boids.ron
vendored
Normal file
7
third_party/rust/naga/test-data/boids.ron
vendored
Normal file
@ -0,0 +1,7 @@
|
|||||||
|
(
|
||||||
|
metal_bindings: {
|
||||||
|
(set: 0, binding: 0): (buffer: Some(0), texture: None, sampler: None, mutable: false),
|
||||||
|
(set: 0, binding: 1): (buffer: Some(1), texture: None, sampler: None, mutable: true),
|
||||||
|
(set: 0, binding: 2): (buffer: Some(2), texture: None, sampler: None, mutable: true),
|
||||||
|
}
|
||||||
|
)
|
||||||
152
third_party/rust/naga/test-data/boids.wgsl
vendored
Normal file
152
third_party/rust/naga/test-data/boids.wgsl
vendored
Normal file
@ -0,0 +1,152 @@
|
|||||||
|
# Copyright 2020 The Tint Authors.
|
||||||
|
#
|
||||||
|
# Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
# you may not use this file except in compliance with the License.
|
||||||
|
# You may obtain a copy of the License at
|
||||||
|
#
|
||||||
|
# http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
#
|
||||||
|
# Unless required by applicable law or agreed to in writing, software
|
||||||
|
# distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
# See the License for the specific language governing permissions and
|
||||||
|
# limitations under the License.
|
||||||
|
|
||||||
|
import "GLSL.std.450" as std;
|
||||||
|
|
||||||
|
# vertex shader
|
||||||
|
|
||||||
|
[[location 0]] var<in> a_particlePos : vec2<f32>;
|
||||||
|
[[location 1]] var<in> a_particleVel : vec2<f32>;
|
||||||
|
[[location 2]] var<in> a_pos : vec2<f32>;
|
||||||
|
[[builtin position]] var gl_Position : vec4<f32>;
|
||||||
|
|
||||||
|
fn vtx_main() -> void {
|
||||||
|
var angle : f32 = -std::atan2(a_particleVel.x, a_particleVel.y);
|
||||||
|
var pos : vec2<f32> = vec2<f32>(
|
||||||
|
(a_pos.x * std::cos(angle)) - (a_pos.y * std::sin(angle)),
|
||||||
|
(a_pos.x * std::sin(angle)) + (a_pos.y * std::cos(angle)));
|
||||||
|
gl_Position = vec4<f32>(pos + a_particlePos, 0, 1);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
entry_point vertex as "main" = vtx_main;
|
||||||
|
|
||||||
|
# fragment shader
|
||||||
|
[[location 0]] var<out> fragColor : vec4<f32>;
|
||||||
|
|
||||||
|
fn frag_main() -> void {
|
||||||
|
fragColor = vec4<f32>(1.0, 1.0, 1.0, 1.0);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
entry_point fragment as "main" = frag_main;
|
||||||
|
|
||||||
|
# compute shader
|
||||||
|
type Particle = struct {
|
||||||
|
[[offset 0]] pos : vec2<f32>;
|
||||||
|
[[offset 8]] vel : vec2<f32>;
|
||||||
|
};
|
||||||
|
|
||||||
|
type SimParams = struct {
|
||||||
|
[[offset 0]] deltaT : f32;
|
||||||
|
[[offset 4]] rule1Distance : f32;
|
||||||
|
[[offset 8]] rule2Distance : f32;
|
||||||
|
[[offset 12]] rule3Distance : f32;
|
||||||
|
[[offset 16]] rule1Scale : f32;
|
||||||
|
[[offset 20]] rule2Scale : f32;
|
||||||
|
[[offset 24]] rule3Scale : f32;
|
||||||
|
};
|
||||||
|
|
||||||
|
type Particles = struct {
|
||||||
|
[[offset 0]] particles : array<Particle, 5>;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[binding 0, set 0]] var<uniform> params : SimParams;
|
||||||
|
[[binding 1, set 0]] var<storage_buffer> particlesA : Particles;
|
||||||
|
[[binding 2, set 0]] var<storage_buffer> particlesB : Particles;
|
||||||
|
|
||||||
|
[[builtin global_invocation_id]] var gl_GlobalInvocationID : vec3<u32>;
|
||||||
|
|
||||||
|
# https://github.com/austinEng/Project6-Vulkan-Flocking/blob/master/data/shaders/computeparticles/particle.comp
|
||||||
|
fn compute_main() -> void {
|
||||||
|
var index : u32 = gl_GlobalInvocationID.x;
|
||||||
|
if (index >= 5) {
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
var vPos : vec2<f32> = particlesA.particles[index].pos;
|
||||||
|
var vVel : vec2<f32> = particlesA.particles[index].vel;
|
||||||
|
|
||||||
|
var cMass : vec2<f32> = vec2<f32>(0, 0);
|
||||||
|
var cVel : vec2<f32> = vec2<f32>(0, 0);
|
||||||
|
var colVel : vec2<f32> = vec2<f32>(0, 0);
|
||||||
|
var cMassCount : i32 = 0;
|
||||||
|
var cVelCount : i32 = 0;
|
||||||
|
|
||||||
|
var pos : vec2<f32>;
|
||||||
|
var vel : vec2<f32>;
|
||||||
|
var i : u32 = 0;
|
||||||
|
loop {
|
||||||
|
if (i >= 5) {
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
if (i == index) {
|
||||||
|
continue;
|
||||||
|
}
|
||||||
|
|
||||||
|
pos = particlesA.particles[i].pos.xy;
|
||||||
|
vel = particlesA.particles[i].vel.xy;
|
||||||
|
|
||||||
|
if (std::distance(pos, vPos) < params.rule1Distance) {
|
||||||
|
cMass = cMass + pos;
|
||||||
|
cMassCount = cMassCount + 1;
|
||||||
|
}
|
||||||
|
if (std::distance(pos, vPos) < params.rule2Distance) {
|
||||||
|
colVel = colVel - (pos - vPos);
|
||||||
|
}
|
||||||
|
if (std::distance(pos, vPos) < params.rule3Distance) {
|
||||||
|
cVel = cVel + vel;
|
||||||
|
cVelCount = cVelCount + 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
continuing {
|
||||||
|
i = i + 1;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
if (cMassCount > 0) {
|
||||||
|
cMass = (cMass / vec2<f32>(cMassCount, cMassCount)) + vPos;
|
||||||
|
}
|
||||||
|
if (cVelCount > 0) {
|
||||||
|
cVel = cVel / vec2<f32>(cVelCount, cVelCount);
|
||||||
|
}
|
||||||
|
|
||||||
|
vVel = vVel + (cMass * params.rule1Scale) + (colVel * params.rule2Scale) +
|
||||||
|
(cVel * params.rule3Scale);
|
||||||
|
|
||||||
|
# clamp velocity for a more pleasing simulation
|
||||||
|
vVel = std::normalize(vVel) * std::fclamp(std::length(vVel), 0.0, 0.1);
|
||||||
|
|
||||||
|
# kinematic update
|
||||||
|
vPos = vPos + (vVel * params.deltaT);
|
||||||
|
|
||||||
|
# Wrap around boundary
|
||||||
|
if (vPos.x < -1.0) {
|
||||||
|
vPos.x = 1.0;
|
||||||
|
}
|
||||||
|
if (vPos.x > 1.0) {
|
||||||
|
vPos.x = -1.0;
|
||||||
|
}
|
||||||
|
if (vPos.y < -1.0) {
|
||||||
|
vPos.y = 1.0;
|
||||||
|
}
|
||||||
|
if (vPos.y > 1.0) {
|
||||||
|
vPos.y = -1.0;
|
||||||
|
}
|
||||||
|
|
||||||
|
# Write back
|
||||||
|
particlesB.particles[index].pos = vPos;
|
||||||
|
particlesB.particles[index].vel = vVel;
|
||||||
|
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
entry_point compute as "main" = compute_main;
|
||||||
|
|
||||||
24
third_party/rust/naga/test-data/quad.wgsl
vendored
Normal file
24
third_party/rust/naga/test-data/quad.wgsl
vendored
Normal file
@ -0,0 +1,24 @@
|
|||||||
|
# vertex
|
||||||
|
const c_scale: f32 = 1.2;
|
||||||
|
[[location 0]] var<in> a_pos : vec2<f32>;
|
||||||
|
[[location 1]] var<in> a_uv : vec2<f32>;
|
||||||
|
[[location 0]] var<out> v_uv : vec2<f32>;
|
||||||
|
[[builtin position]] var<out> o_position : vec4<f32>;
|
||||||
|
|
||||||
|
fn main_vert() -> void {
|
||||||
|
o_position = vec4<f32>(c_scale * a_pos, 0.0, 1.0);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
entry_point vertex as "main" = main_vert;
|
||||||
|
|
||||||
|
# fragment
|
||||||
|
[[location 0]] var<in> a_uv : vec2<f32>;
|
||||||
|
#layout(set = 0, binding = 0) uniform texture2D u_texture;
|
||||||
|
#layout(set = 0, binding = 1) uniform sampler u_sampler;
|
||||||
|
[[location 0]] var<out> o_color : vec4<f32>;
|
||||||
|
|
||||||
|
fn main_frag() -> void {
|
||||||
|
o_color = vec4<f32>(1, 0, 0, 1); #TODO: sample
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
entry_point fragment as "main" = main_frag;
|
||||||
51
third_party/rust/naga/tests/convert.rs
vendored
Normal file
51
third_party/rust/naga/tests/convert.rs
vendored
Normal file
@ -0,0 +1,51 @@
|
|||||||
|
fn load_wgsl(name: &str) -> naga::Module {
|
||||||
|
let path = format!("{}/test-data/{}.wgsl", env!("CARGO_MANIFEST_DIR"), name);
|
||||||
|
let input = std::fs::read_to_string(path).unwrap();
|
||||||
|
naga::front::wgsl::parse_str(&input).unwrap()
|
||||||
|
}
|
||||||
|
|
||||||
|
#[test]
|
||||||
|
fn convert_quad() {
|
||||||
|
let module = load_wgsl("quad");
|
||||||
|
{
|
||||||
|
use naga::back::msl;
|
||||||
|
let mut binding_map = msl::BindingMap::default();
|
||||||
|
binding_map.insert(
|
||||||
|
msl::BindSource { set: 0, binding: 0 },
|
||||||
|
msl::BindTarget { buffer: None, texture: Some(1), sampler: None, mutable: false },
|
||||||
|
);
|
||||||
|
binding_map.insert(
|
||||||
|
msl::BindSource { set: 0, binding: 1 },
|
||||||
|
msl::BindTarget { buffer: None, texture: None, sampler: Some(1), mutable: false },
|
||||||
|
);
|
||||||
|
let options = msl::Options {
|
||||||
|
binding_map: &binding_map,
|
||||||
|
};
|
||||||
|
msl::write_string(&module, options).unwrap();
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
#[test]
|
||||||
|
fn convert_boids() {
|
||||||
|
let module = load_wgsl("boids");
|
||||||
|
{
|
||||||
|
use naga::back::msl;
|
||||||
|
let mut binding_map = msl::BindingMap::default();
|
||||||
|
binding_map.insert(
|
||||||
|
msl::BindSource { set: 0, binding: 0 },
|
||||||
|
msl::BindTarget { buffer: Some(0), texture: None, sampler: None, mutable: false },
|
||||||
|
);
|
||||||
|
binding_map.insert(
|
||||||
|
msl::BindSource { set: 0, binding: 1 },
|
||||||
|
msl::BindTarget { buffer: Some(1), texture: None, sampler: Some(1), mutable: false },
|
||||||
|
);
|
||||||
|
binding_map.insert(
|
||||||
|
msl::BindSource { set: 0, binding: 2 },
|
||||||
|
msl::BindTarget { buffer: Some(2), texture: None, sampler: Some(1), mutable: false },
|
||||||
|
);
|
||||||
|
let options = msl::Options {
|
||||||
|
binding_map: &binding_map,
|
||||||
|
};
|
||||||
|
msl::write_string(&module, options).unwrap();
|
||||||
|
}
|
||||||
|
}
|
||||||
1
third_party/rust/spirv_headers/.cargo-checksum.json
vendored
Normal file
1
third_party/rust/spirv_headers/.cargo-checksum.json
vendored
Normal file
@ -0,0 +1 @@
|
|||||||
|
{"files":{"Cargo.toml":"a7fb396d7d75f52e15338f5fb389c1d10ef51e4f5bc24c1d6ab834248f2b7efd","README.md":"e6b0558fda2d831b1c291a773fd2baee3c2e835f2d76952de927b9e455bf0fc9","autogen_spirv.rs":"eec69221ce0b44a06aa0dc44dbff1af974f1fd7530fee02e02badd2edaf15071","lib.rs":"968a4b6b083e283b530271cc2cc13d6c69793877647694a15636c0fbecd1bf58"},"package":"3f1418983d16481227ffa3ab3cf44ef92eebc9a76c092fbcd4c51a64ff032622"}
|
||||||
31
third_party/rust/spirv_headers/Cargo.toml
vendored
Normal file
31
third_party/rust/spirv_headers/Cargo.toml
vendored
Normal file
@ -0,0 +1,31 @@
|
|||||||
|
# THIS FILE IS AUTOMATICALLY GENERATED BY CARGO
|
||||||
|
#
|
||||||
|
# When uploading crates to the registry Cargo will automatically
|
||||||
|
# "normalize" Cargo.toml files for maximal compatibility
|
||||||
|
# with all versions of Cargo and also rewrite `path` dependencies
|
||||||
|
# to registry (e.g., crates.io) dependencies
|
||||||
|
#
|
||||||
|
# If you believe there's an error in this file please file an
|
||||||
|
# issue against the rust-lang/cargo repository. If you're
|
||||||
|
# editing this file be aware that the upstream Cargo.toml
|
||||||
|
# will likely look very different (and much more reasonable)
|
||||||
|
|
||||||
|
[package]
|
||||||
|
edition = "2018"
|
||||||
|
name = "spirv_headers"
|
||||||
|
version = "1.4.2"
|
||||||
|
authors = ["Lei Zhang <antiagainst@gmail.com>"]
|
||||||
|
description = "Rust definition of SPIR-V structs and enums"
|
||||||
|
documentation = "https://docs.rs/spirv_headers"
|
||||||
|
readme = "README.md"
|
||||||
|
keywords = ["spirv", "definition", "struct", "enum"]
|
||||||
|
license = "Apache-2.0"
|
||||||
|
repository = "https://github.com/gfx-rs/rspirv"
|
||||||
|
|
||||||
|
[lib]
|
||||||
|
path = "lib.rs"
|
||||||
|
[dependencies.bitflags]
|
||||||
|
version = "1"
|
||||||
|
|
||||||
|
[dependencies.num-traits]
|
||||||
|
version = "0.2"
|
||||||
50
third_party/rust/spirv_headers/README.md
vendored
Normal file
50
third_party/rust/spirv_headers/README.md
vendored
Normal file
@ -0,0 +1,50 @@
|
|||||||
|
spirv-headers of the rspirv project
|
||||||
|
===================================
|
||||||
|
|
||||||
|
[![Crate][img-crate-headers]][crate-headers]
|
||||||
|
[![Documentation][img-doc-headers]][doc-headers]
|
||||||
|
|
||||||
|
The headers crate for the rspirv project which provides Rust definitions of
|
||||||
|
SPIR-V structs, enums, and constants.
|
||||||
|
|
||||||
|
Usage
|
||||||
|
-----
|
||||||
|
|
||||||
|
This project uses associated constants, which became available in the stable channel
|
||||||
|
since [1.20][rust-1.20]. So to compile with a compiler from the stable channel,
|
||||||
|
please make sure that the version is >= 1.20.
|
||||||
|
|
||||||
|
First add to your `Cargo.toml`:
|
||||||
|
|
||||||
|
```toml
|
||||||
|
[dependencies]
|
||||||
|
rspirv_headers = "1.4"
|
||||||
|
```
|
||||||
|
|
||||||
|
Then add to your crate root:
|
||||||
|
|
||||||
|
```rust
|
||||||
|
extern crate spirv_headers;
|
||||||
|
```
|
||||||
|
|
||||||
|
Version
|
||||||
|
-------
|
||||||
|
|
||||||
|
Note that the major and minor version of this create is tracking the SPIR-V spec,
|
||||||
|
while the patch number is used for bugfixes for the crate itself. So version
|
||||||
|
`1.4.2` is tracking SPIR-V 1.4 but not necessarily revision 2. Major client APIs
|
||||||
|
like Vulkan/OpenCL pin to a specific major and minor version, regardless of the
|
||||||
|
revision.
|
||||||
|
|
||||||
|
Examples
|
||||||
|
--------
|
||||||
|
|
||||||
|
Please see the [documentation][doc-headers] and project's
|
||||||
|
[README][project-readme] for examples.
|
||||||
|
|
||||||
|
[img-crate-headers]: https://img.shields.io/crates/v/spirv_headers.svg
|
||||||
|
[img-doc-headers]: https://docs.rs/spirv_headers/badge.svg
|
||||||
|
[crate-headers]: https://crates.io/crates/spirv_headers
|
||||||
|
[doc-headers]: https://docs.rs/spirv_headers
|
||||||
|
[project-readme]: https://github.com/gfx-rs/rspirv/blob/master/README.md
|
||||||
|
[rust-1.20]: https://blog.rust-lang.org/2017/08/31/Rust-1.20.html
|
||||||
3645
third_party/rust/spirv_headers/autogen_spirv.rs
vendored
Normal file
3645
third_party/rust/spirv_headers/autogen_spirv.rs
vendored
Normal file
File diff suppressed because it is too large
Load Diff
16
third_party/rust/spirv_headers/lib.rs
vendored
Normal file
16
third_party/rust/spirv_headers/lib.rs
vendored
Normal file
@ -0,0 +1,16 @@
|
|||||||
|
#![doc(html_root_url = "https://docs.rs/spirv_headers/1.4/")]
|
||||||
|
|
||||||
|
//! The SPIR-V header.
|
||||||
|
//!
|
||||||
|
//! This crate contains Rust definitions of all SPIR-V structs, enums,
|
||||||
|
//! and constants.
|
||||||
|
//!
|
||||||
|
//! The version of this crate is the version of SPIR-V it contains.
|
||||||
|
|
||||||
|
#![allow(non_camel_case_types)]
|
||||||
|
#![cfg_attr(rustfmt, rustfmt_skip)]
|
||||||
|
|
||||||
|
use bitflags::bitflags;
|
||||||
|
use num_traits;
|
||||||
|
|
||||||
|
include!("autogen_spirv.rs");
|
||||||
Loading…
Reference in New Issue
Block a user