github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

add debug info for spv-out (#2379)

Browse Source
This commit is contained in:
wicast
2023-06-28 20:30:25 +08:00
committed by GitHub
parent 614da63366
commit 25e4f17a69
28 changed files with 2258 additions and 55 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
benches/criterion.rs
View File

@@ -181,7 +181,7 @@ fn backends(c: &mut Criterion) {
let options = naga::back::spv::Options::default();
for &(ref module, ref info) in inputs.iter() {
let mut writer = naga::back::spv::Writer::new(&options).unwrap();
writer.write(module, info, None, &mut data).unwrap();
writer.write(module, info, None, &None, &mut data).unwrap();
data.clear();
}
});
@@ -199,7 +199,7 @@ fn backends(c: &mut Criterion) {
entry_point: ep.name.clone(),
};
writer
.write(module, info, Some(&pipeline_options), &mut data)
.write(module, info, Some(&pipeline_options), &None, &mut data)
.unwrap();
data.clear();
}

23
cli/src/bin/naga.rs
View File

@@ -68,6 +68,10 @@ struct Args {
#[argh(option)]
stdin_file_path: Option<String>,
/// generate debug symbols, only works for spv-out for now
#[argh(option, short = 'g')]
generate_debug_symbols: Option<bool>,
/// show version
#[argh(switch)]
version: bool,
@@ -144,13 +148,13 @@ impl FromStr for GlslProfileArg {
}
#[derive(Default)]
struct Parameters {
struct Parameters<'a> {
validation_flags: naga::valid::ValidationFlags,
bounds_check_policies: naga::proc::BoundsCheckPolicies,
entry_point: Option<String>,
keep_coordinate_space: bool,
spv_block_ctx_dump_prefix: Option<String>,
spv: naga::back::spv::Options,
spv: naga::back::spv::Options<'a>,
msl: naga::back::msl::Options,
glsl: naga::back::glsl::Options,
hlsl: naga::back::hlsl::Options,
@@ -336,9 +340,9 @@ fn run() -> Result<(), Box<dyn std::error::Error>> {
{
Ok(info) => Some(info),
Err(error) => {
if let Some(input) = input_text {
if let Some(input) = &input_text {
let filename = input_path.file_name().and_then(std::ffi::OsStr::to_str);
emit_annotated_error(&error, filename.unwrap_or("input"), &input);
emit_annotated_error(&error, filename.unwrap_or("input"), input);
}
print_err(&error);
None
@@ -415,6 +419,17 @@ fn run() -> Result<(), Box<dyn std::error::Error>> {
};
params.spv.bounds_check_policies = params.bounds_check_policies;
//Insert Debug infos
let debug_info = args.generate_debug_symbols.and_then(|debug| {
params.spv.flags.set(spv::WriterFlags::DEBUG, debug);
Some(spv::DebugInfo {
source_code: input_text.as_ref()?,
file_name: input_path.file_name().and_then(std::ffi::OsStr::to_str)?,
})
});
params.spv.debug_info = debug_info;
params.spv.flags.set(
spv::WriterFlags::ADJUST_COORDINATE_SPACE,
!params.keep_coordinate_space,

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

@@ -7,7 +7,7 @@ use super::{
Dimension, Error, Instruction, LocalType, LookupType, LoopContext, ResultMember, Writer,
WriterFlags,
};
use crate::{arena::Handle, proc::TypeResolution};
use crate::{arena::Handle, proc::TypeResolution, Statement};
use spirv::Word;
fn get_dimension(type_inner: &crate::TypeInner) -> Dimension {
@@ -60,6 +60,12 @@ pub enum BlockExit {
},
}
#[derive(Debug)]
pub(crate) struct DebugInfoInner<'a> {
pub source_code: &'a str,
pub source_file_id: Word,
}
impl Writer {
// Flip Y coordinate to adjust for coordinate space difference
// between SPIR-V and our IR.
@@ -1699,13 +1705,32 @@ impl<'w> BlockContext<'w> {
pub(super) fn write_block(
&mut self,
label_id: Word,
statements: &[crate::Statement],
naga_block: &crate::Block,
exit: BlockExit,
loop_context: LoopContext,
debug_info: Option<&DebugInfoInner>,
) -> Result<(), Error> {
let mut block = Block::new(label_id);
for statement in statements {
for (statement, span) in naga_block.span_iter() {
if let (Some(debug_info), false) = (
debug_info,
matches!(
statement,
&(Statement::Block(..)
| Statement::Break
| Statement::Continue
| Statement::Kill
| Statement::Return { .. }
| Statement::Loop { .. })
),
) {
let loc: crate::SourceLocation = span.location(debug_info.source_code);
block.body.push(Instruction::line(
debug_info.source_file_id,
loc.line_number,
loc.line_position,
));
};
match *statement {
crate::Statement::Emit(ref range) => {
for handle in range.clone() {
@@ -1722,6 +1747,7 @@ impl<'w> BlockContext<'w> {
block_statements,
BlockExit::Branch { target: merge_id },
loop_context,
debug_info,
)?;
block = Block::new(merge_id);
@@ -1765,6 +1791,7 @@ impl<'w> BlockContext<'w> {
accept,
BlockExit::Branch { target: merge_id },
loop_context,
debug_info,
)?;
}
if let Some(block_id) = reject_id {
@@ -1773,6 +1800,7 @@ impl<'w> BlockContext<'w> {
reject,
BlockExit::Branch { target: merge_id },
loop_context,
debug_info,
)?;
}
@@ -1852,6 +1880,7 @@ impl<'w> BlockContext<'w> {
target: case_finish_id,
},
inner_context,
debug_info,
)?;
}
@@ -1873,6 +1902,16 @@ impl<'w> BlockContext<'w> {
// SPIR-V requires the continuing to the `OpLoopMerge`,
// so we have to start a new block with it.
block = Block::new(preamble_id);
// HACK the loop statement is begin with branch instruction,
// so we need to put `OpLine` debug info before merge instruction
if let Some(debug_info) = debug_info {
let loc: crate::SourceLocation = span.location(debug_info.source_code);
block.body.push(Instruction::line(
debug_info.source_file_id,
loc.line_number,
loc.line_position,
))
}
block.body.push(Instruction::loop_merge(
merge_id,
continuing_id,
@@ -1890,6 +1929,7 @@ impl<'w> BlockContext<'w> {
continuing_id: Some(continuing_id),
break_id: Some(merge_id),
},
debug_info,
)?;
let exit = match break_if {
@@ -1910,6 +1950,7 @@ impl<'w> BlockContext<'w> {
continuing_id: None,
break_id: Some(merge_id),
},
debug_info,
)?;
block = Block::new(merge_id);

31
src/back/spv/instructions.rs
View File

@@ -1,4 +1,4 @@
use super::helpers;
use super::{block::DebugInfoInner, helpers};
use spirv::{Op, Word};
pub(super) enum Signedness {
@@ -21,10 +21,25 @@ impl super::Instruction {
// Debug Instructions
//
pub(super) fn source(source_language: spirv::SourceLanguage, version: u32) -> Self {
pub(super) fn string(name: &str, id: Word) -> Self {
let mut instruction = Self::new(Op::String);
instruction.set_result(id);
instruction.add_operands(helpers::string_to_words(name));
instruction
}
pub(super) fn source(
source_language: spirv::SourceLanguage,
version: u32,
source: &Option<DebugInfoInner>,
) -> Self {
let mut instruction = Self::new(Op::Source);
instruction.add_operand(source_language as u32);
instruction.add_operands(helpers::bytes_to_words(&version.to_le_bytes()));
if let Some(source) = source.as_ref() {
instruction.add_operand(source.source_file_id);
instruction.add_operands(helpers::string_to_words(source.source_code));
}
instruction
}
@@ -43,6 +58,18 @@ impl super::Instruction {
instruction
}
pub(super) fn line(file: Word, line: Word, column: Word) -> Self {
let mut instruction = Self::new(Op::Line);
instruction.add_operand(file);
instruction.add_operand(line);
instruction.add_operand(column);
instruction
}
pub(super) const fn no_line() -> Self {
Self::new(Op::NoLine)
}
//
// Annotation Instructions
//

27
src/back/spv/mod.rs
View File

@@ -82,6 +82,12 @@ impl IdGenerator {
}
}
#[derive(Debug, Clone)]
pub struct DebugInfo<'a> {
pub source_code: &'a str,
pub file_name: &'a str,
}
/// A SPIR-V block to which we are still adding instructions.
///
/// A `Block` represents a SPIR-V block that does not yet have a termination
@@ -620,6 +626,7 @@ pub struct Writer {
saved_cached: CachedExpressions,
gl450_ext_inst_id: Word,
// Just a temporary list of SPIR-V ids
temp_list: Vec<Word>,
}
@@ -664,7 +671,7 @@ pub enum ZeroInitializeWorkgroupMemoryMode {
}
#[derive(Debug, Clone)]
pub struct Options {
pub struct Options<'a> {
/// (Major, Minor) target version of the SPIR-V.
pub lang_version: (u8, u8),
@@ -686,9 +693,11 @@ pub struct Options {
/// Dictates the way workgroup variables should be zero initialized
pub zero_initialize_workgroup_memory: ZeroInitializeWorkgroupMemoryMode,
pub debug_info: Option<DebugInfo<'a>>,
}
impl Default for Options {
impl<'a> Default for Options<'a> {
fn default() -> Self {
let mut flags = WriterFlags::ADJUST_COORDINATE_SPACE
| WriterFlags::LABEL_VARYINGS
@@ -703,6 +712,7 @@ impl Default for Options {
capabilities: None,
bounds_check_policies: crate::proc::BoundsCheckPolicies::default(),
zero_initialize_workgroup_memory: ZeroInitializeWorkgroupMemoryMode::Polyfill,
debug_info: None,
}
}
}
@@ -726,8 +736,17 @@ pub fn write_vec(
options: &Options,
pipeline_options: Option<&PipelineOptions>,
) -> Result<Vec<u32>, Error> {
let mut words = Vec::new();
let mut words: Vec<u32> = Vec::new();
let mut w = Writer::new(options)?;
w.write(module, info, pipeline_options, &mut words)?;
if options.flags.contains(WriterFlags::DEBUG) {}
w.write(
module,
info,
pipeline_options,
&options.debug_info,
&mut words,
)?;
Ok(words)
}

40
src/back/spv/writer.rs
View File

@@ -1,9 +1,10 @@
use super::{
block::DebugInfoInner,
helpers::{contains_builtin, global_needs_wrapper, map_storage_class},
make_local, Block, BlockContext, CachedConstant, CachedExpressions, EntryPointContext, Error,
Function, FunctionArgument, GlobalVariable, IdGenerator, Instruction, LocalType, LocalVariable,
LogicalLayout, LookupFunctionType, LookupType, LoopContext, Options, PhysicalLayout,
PipelineOptions, ResultMember, Writer, WriterFlags, BITS_PER_BYTE,
make_local, Block, BlockContext, CachedConstant, CachedExpressions, DebugInfo,
EntryPointContext, Error, Function, FunctionArgument, GlobalVariable, IdGenerator, Instruction,
LocalType, LocalVariable, LogicalLayout, LookupFunctionType, LookupType, LoopContext, Options,
PhysicalLayout, PipelineOptions, ResultMember, Writer, WriterFlags, BITS_PER_BYTE,
};
use crate::{
arena::{Handle, UniqueArena},
@@ -329,6 +330,7 @@ impl Writer {
info: &FunctionInfo,
ir_module: &crate::Module,
mut interface: Option<FunctionInterface>,
debug_info: &Option<DebugInfoInner>,
) -> Result<Word, Error> {
let mut function = Function::default();
@@ -693,6 +695,7 @@ impl Writer {
&ir_function.body,
super::block::BlockExit::Return,
LoopContext::default(),
debug_info.as_ref(),
)?;
// Consume the `BlockContext`, ending its borrows and letting the
@@ -726,6 +729,7 @@ impl Writer {
entry_point: &crate::EntryPoint,
info: &FunctionInfo,
ir_module: &crate::Module,
debug_info: &Option<DebugInfoInner>,
) -> Result<Instruction, Error> {
let mut interface_ids = Vec::new();
let function_id = self.write_function(
@@ -736,6 +740,7 @@ impl Writer {
varying_ids: &mut interface_ids,
stage: entry_point.stage,
}),
debug_info,
)?;
let exec_model = match entry_point.stage {
@@ -1724,6 +1729,7 @@ impl Writer {
ir_module: &crate::Module,
mod_info: &ModuleInfo,
ep_index: Option<usize>,
debug_info: &Option<DebugInfo>,
) -> Result<(), Error> {
fn has_view_index_check(
ir_module: &crate::Module,
@@ -1771,9 +1777,23 @@ impl Writer {
Instruction::ext_inst_import(self.gl450_ext_inst_id, "GLSL.std.450")
.to_words(&mut self.logical_layout.ext_inst_imports);
let mut debug_info_inner = None;
if self.flags.contains(WriterFlags::DEBUG) {
self.debugs
.push(Instruction::source(spirv::SourceLanguage::GLSL, 450));
if let Some(debug_info) = debug_info.as_ref() {
let source_file_id = self.id_gen.next();
self.debugs
.push(Instruction::string(debug_info.file_name, source_file_id));
debug_info_inner = Some(DebugInfoInner {
source_code: debug_info.source_code,
source_file_id,
});
self.debugs.push(Instruction::source(
spirv::SourceLanguage::Unknown,
0,
&debug_info_inner,
));
}
}
self.constant_ids.resize(ir_module.constants.len(), 0);
@@ -1858,7 +1878,7 @@ impl Writer {
continue;
}
}
let id = self.write_function(ir_function, info, ir_module, None)?;
let id = self.write_function(ir_function, info, ir_module, None, &debug_info_inner)?;
self.lookup_function.insert(handle, id);
}
@@ -1868,7 +1888,8 @@ impl Writer {
continue;
}
let info = mod_info.get_entry_point(index);
let ep_instruction = self.write_entry_point(ir_ep, info, ir_module)?;
let ep_instruction =
self.write_entry_point(ir_ep, info, ir_module, &debug_info_inner)?;
ep_instruction.to_words(&mut self.logical_layout.entry_points);
}
@@ -1910,6 +1931,7 @@ impl Writer {
ir_module: &crate::Module,
info: &ModuleInfo,
pipeline_options: Option<&PipelineOptions>,
debug_info: &Option<DebugInfo>,
words: &mut Vec<Word>,
) -> Result<(), Error> {
self.reset();
@@ -1927,7 +1949,7 @@ impl Writer {
None => None,
};
self.write_logical_layout(ir_module, info, ep_index)?;
self.write_logical_layout(ir_module, info, ep_index, debug_info)?;
self.write_physical_layout();
self.physical_layout.in_words(words);

7
tests/in/debug-symbol-simple.param.ron Normal file
View File

@@ -0,0 +1,7 @@
(
spv: (
version: (1, 1),
debug: true,
adjust_coordinate_space: false,
),
)

33
tests/in/debug-symbol-simple.wgsl Normal file
View File

@@ -0,0 +1,33 @@
struct VertexInput {
@location(0) position: vec3<f32>,
@location(1) color: vec3<f32>,
};
struct VertexOutput {
@builtin(position) clip_position: vec4<f32>,
@location(0) color: vec3<f32>,
};
@vertex
fn vs_main(
model: VertexInput,
) -> VertexOutput {
var out: VertexOutput;
out.color = model.color;
out.clip_position = vec4<f32>(model.position, 1.0);
return out;
}
// Fragment shader
@fragment
fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
var color = in.color;
for (var i = 0; i < 10; i += 1) {
var ii = f32(i);
color.x += ii*0.001;
color.y += ii*0.002;
}
return vec4<f32>(color, 1.0);
}

7
tests/in/debug-symbol-terrain.param.ron Normal file
View File

@@ -0,0 +1,7 @@
(
spv: (
version: (1, 1),
debug: true,
adjust_coordinate_space: false,
),
)

297
tests/in/debug-symbol-terrain.wgsl Normal file
View File

@@ -0,0 +1,297 @@
// Taken from https://github.com/sotrh/learn-wgpu/blob/11820796f5e1dbce42fb1119f04ddeb4b167d2a0/code/intermediate/tutorial13-terrain/src/terrain.wgsl
// ============================
// Terrain Generation
// ============================
// https://gist.github.com/munrocket/236ed5ba7e409b8bdf1ff6eca5dcdc39
// MIT License. © Ian McEwan, Stefan Gustavson, Munrocket
// - Less condensed glsl implementation with comments can be found at https://weber.itn.liu.se/~stegu/jgt2012/article.pdf
fn permute3(x: vec3<f32>) -> vec3<f32> { return (((x * 34.) + 1.) * x) % vec3<f32>(289.); }
fn snoise2(v: vec2<f32>) -> f32 {
let C = vec4<f32>(0.211324865405187, 0.366025403784439, -0.577350269189626, 0.024390243902439);
var i: vec2<f32> = floor(v + dot(v, C.yy));
let x0 = v - i + dot(i, C.xx);
// I flipped the condition here from > to < as it fixed some artifacting I was observing
var i1: vec2<f32> = select(vec2<f32>(1., 0.), vec2<f32>(0., 1.), (x0.x < x0.y));
var x12: vec4<f32> = x0.xyxy + C.xxzz - vec4<f32>(i1, 0., 0.);
i = i % vec2<f32>(289.);
let p = permute3(permute3(i.y + vec3<f32>(0., i1.y, 1.)) + i.x + vec3<f32>(0., i1.x, 1.));
var m: vec3<f32> = max(0.5 - vec3<f32>(dot(x0, x0), dot(x12.xy, x12.xy), dot(x12.zw, x12.zw)), vec3<f32>(0.));
m = m * m;
m = m * m;
let x = 2. * fract(p * C.www) - 1.;
let h = abs(x) - 0.5;
let ox = floor(x + 0.5);
let a0 = x - ox;
m = m * (1.79284291400159 - 0.85373472095314 * (a0 * a0 + h * h));
let g = vec3<f32>(a0.x * x0.x + h.x * x0.y, a0.yz * x12.xz + h.yz * x12.yw);
return 130. * dot(m, g);
}
fn fbm(p: vec2<f32>) -> f32 {
let NUM_OCTAVES: u32 = 5u;
var x = p * 0.01;
var v = 0.0;
var a = 0.5;
let shift = vec2<f32>(100.0);
let cs = vec2<f32>(cos(0.5), sin(0.5));
let rot = mat2x2<f32>(cs.x, cs.y, -cs.y, cs.x);
for (var i = 0u; i < NUM_OCTAVES; i = i + 1u) {
v = v + a * snoise2(x);
x = rot * x * 2.0 + shift;
a = a * 0.5;
}
return v;
}
struct ChunkData {
chunk_size: vec2<u32>,
chunk_corner: vec2<i32>,
min_max_height: vec2<f32>,
}
struct Vertex {
@location(0) position: vec3<f32>,
@location(1) normal: vec3<f32>,
}
struct VertexBuffer {
data: array<Vertex>, // stride: 32
}
struct IndexBuffer {
data: array<u32>,
}
@group(0) @binding(0) var<uniform> chunk_data: ChunkData;
@group(0) @binding(1) var<storage, read_write> vertices: VertexBuffer;
@group(0) @binding(2) var<storage, read_write> indices: IndexBuffer;
fn terrain_point(p: vec2<f32>, min_max_height: vec2<f32>) -> vec3<f32> {
return vec3<f32>(
p.x,
mix(min_max_height.x, min_max_height.y, fbm(p)),
p.y,
);
}
fn terrain_vertex(p: vec2<f32>, min_max_height: vec2<f32>) -> Vertex {
let v = terrain_point(p, min_max_height);
let tpx = terrain_point(p + vec2<f32>(0.1, 0.0), min_max_height) - v;
let tpz = terrain_point(p + vec2<f32>(0.0, 0.1), min_max_height) - v;
let tnx = terrain_point(p + vec2<f32>(-0.1, 0.0), min_max_height) - v;
let tnz = terrain_point(p + vec2<f32>(0.0, -0.1), min_max_height) - v;
let pn = normalize(cross(tpz, tpx));
let nn = normalize(cross(tnz, tnx));
let n = (pn + nn) * 0.5;
return Vertex(v, n);
}
fn index_to_p(vert_index: u32, chunk_size: vec2<u32>, chunk_corner: vec2<i32>) -> vec2<f32> {
return vec2(
f32(vert_index) % f32(chunk_size.x + 1u),
f32(vert_index / (chunk_size.x + 1u)),
) + vec2<f32>(chunk_corner);
}
@compute @workgroup_size(64)
fn gen_terrain_compute(
@builtin(global_invocation_id) gid: vec3<u32>
) {
// Create vert_component
let vert_index = gid.x;
let p = index_to_p(vert_index, chunk_data.chunk_size, chunk_data.chunk_corner);
vertices.data[vert_index] = terrain_vertex(p, chunk_data.min_max_height);
// Create indices
let start_index = gid.x * 6u; // using TriangleList
if (start_index >= (chunk_data.chunk_size.x * chunk_data.chunk_size.y * 6u)) { return; }
let v00 = vert_index + gid.x / chunk_data.chunk_size.x;
let v10 = v00 + 1u;
let v01 = v00 + chunk_data.chunk_size.x + 1u;
let v11 = v01 + 1u;
indices.data[start_index] = v00;
indices.data[start_index + 1u] = v01;
indices.data[start_index + 2u] = v11;
indices.data[start_index + 3u] = v00;
indices.data[start_index + 4u] = v11;
indices.data[start_index + 5u] = v10;
}
// ============================
// Terrain Gen (Fragment Shader)
// ============================
struct GenData {
chunk_size: vec2<u32>,
chunk_corner: vec2<i32>,
min_max_height: vec2<f32>,
texture_size: u32,
start_index: u32,
}
@group(0)
@binding(0)
var<uniform> gen_data: GenData;
struct GenVertexOutput {
@location(0)
index: u32,
@builtin(position)
position: vec4<f32>,
@location(1)
uv: vec2<f32>,
};
@vertex
fn gen_terrain_vertex(@builtin(vertex_index) vindex: u32) -> GenVertexOutput {
let u = f32(((vindex + 2u) / 3u) % 2u);
let v = f32(((vindex + 1u) / 3u) % 2u);
let uv = vec2<f32>(u, v);
let position = vec4<f32>(-1.0 + uv * 2.0, 0.0, 1.0);
// TODO: maybe replace this with u32(dot(uv, vec2(f32(gen_data.texture_dim.x))))
let index = u32(uv.x * f32(gen_data.texture_size) + uv.y * f32(gen_data.texture_size)) + gen_data.start_index;
return GenVertexOutput(index, position, uv);
}
struct GenFragmentOutput {
@location(0) vert_component: u32,
@location(1) index: u32,
}
@fragment
fn gen_terrain_fragment(in: GenVertexOutput) -> GenFragmentOutput {
let i = u32(in.uv.x * f32(gen_data.texture_size) + in.uv.y * f32(gen_data.texture_size * gen_data.texture_size)) + gen_data.start_index;
let vert_index = u32(floor(f32(i) / 6.));
let comp_index = i % 6u;
let p = index_to_p(vert_index, gen_data.chunk_size, gen_data.chunk_corner);
let v = terrain_vertex(p, gen_data.min_max_height);
var vert_component: f32 = 0.;
switch comp_index {
case 0u: { vert_component = v.position.x; }
case 1u: { vert_component = v.position.y; }
case 2u: { vert_component = v.position.z; }
case 3u: { vert_component = v.normal.x; }
case 4u: { vert_component = v.normal.y; }
case 5u: { vert_component = v.normal.z; }
default: {}
}
let v00 = vert_index + vert_index / gen_data.chunk_size.x;
let v10 = v00 + 1u;
let v01 = v00 + gen_data.chunk_size.x + 1u;
let v11 = v01 + 1u;
var index = 0u;
switch comp_index {
case 0u, 3u: { index = v00; }
case 2u, 4u: { index = v11; }
case 1u: { index = v01; }
case 5u: { index = v10; }
default: {}
}
index = in.index;
// index = gen_data.start_index;
// indices.data[start_index] = v00;
// indices.data[start_index + 1u] = v01;
// indices.data[start_index + 2u] = v11;
// indices.data[start_index + 3u] = v00;
// indices.data[start_index + 4u] = v11;
// indices.data[start_index + 5u] = v10;
let ivert_component = bitcast<u32>(vert_component);
return GenFragmentOutput(ivert_component, index);
}
// ============================
// Terrain Rendering
// ============================
struct Camera {
view_pos: vec4<f32>,
view_proj: mat4x4<f32>,
}
@group(0) @binding(0)
var<uniform> camera: Camera;
struct Light {
position: vec3<f32>,
color: vec3<f32>,
}
@group(1) @binding(0)
var<uniform> light: Light;
struct VertexOutput {
@builtin(position) clip_position: vec4<f32>,
@location(0) normal: vec3<f32>,
@location(1) world_pos: vec3<f32>,
}
@vertex
fn vs_main(
vertex: Vertex,
) -> VertexOutput {
let clip_position = camera.view_proj * vec4<f32>(vertex.position, 1.);
let normal = vertex.normal;
return VertexOutput(clip_position, normal, vertex.position);
}
@group(2) @binding(0)
var t_diffuse: texture_2d<f32>;
@group(2) @binding(1)
var s_diffuse: sampler;
@group(2) @binding(2)
var t_normal: texture_2d<f32>;
@group(2) @binding(3)
var s_normal: sampler;
fn color23(p: vec2<f32>) -> vec3<f32> {
return vec3<f32>(
snoise2(p) * 0.5 + 0.5,
snoise2(p + vec2<f32>(23., 32.)) * 0.5 + 0.5,
snoise2(p + vec2<f32>(-43., 3.)) * 0.5 + 0.5,
);
}
@fragment
fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
var color = smoothstep(vec3<f32>(0.0), vec3<f32>(0.1), fract(in.world_pos));
color = mix(vec3<f32>(0.5, 0.1, 0.7), vec3<f32>(0.2, 0.2, 0.2), vec3<f32>(color.x * color.y * color.z));
let ambient_strength = 0.1;
let ambient_color = light.color * ambient_strength;
let light_dir = normalize(light.position - in.world_pos);
let view_dir = normalize(camera.view_pos.xyz - in.world_pos);
let half_dir = normalize(view_dir + light_dir);
let diffuse_strength = max(dot(in.normal, light_dir), 0.0);
let diffuse_color = diffuse_strength * light.color;
let specular_strength = pow(max(dot(in.normal, half_dir), 0.0), 32.0);
let specular_color = specular_strength * light.color;
let result = (ambient_color + diffuse_color + specular_color) * color;
return vec4<f32>(result, 1.0);
}

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

@@ -11,7 +11,6 @@ OpEntryPoint Fragment %286 "foo_frag" %285
OpEntryPoint GLCompute %306 "assign_through_ptr" %309
OpExecutionMode %286 OriginUpperLeft
OpExecutionMode %306 LocalSize 1 1 1
OpSource GLSL 450
OpMemberName %12 0 "a"
OpMemberName %12 1 "b"
OpMemberName %12 2 "c"

1
tests/out/spv/boids.spvasm
View File

@@ -8,7 +8,6 @@ OpExtension "SPV_KHR_storage_buffer_storage_class"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %46 "main" %43
OpExecutionMode %46 LocalSize 64 1 1
OpSource GLSL 450
OpName %3 "NUM_PARTICLES"
OpMemberName %7 0 "pos"
OpMemberName %7 1 "vel"

1
tests/out/spv/bounds-check-image-restrict.spvasm
View File

@@ -10,7 +10,6 @@ OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %269 "fragment_shader" %267
OpExecutionMode %269 OriginUpperLeft
OpSource GLSL 450
OpName %21 "image_1d"
OpName %23 "image_2d"
OpName %25 "image_2d_array"

1
tests/out/spv/bounds-check-image-rzsw.spvasm
View File

@@ -10,7 +10,6 @@ OpCapability ImageQuery
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %306 "fragment_shader" %304
OpExecutionMode %306 OriginUpperLeft
OpSource GLSL 450
OpName %21 "image_1d"
OpName %23 "image_2d"
OpName %25 "image_2d_array"

1
tests/out/spv/collatz.spvasm
View File

@@ -8,7 +8,6 @@ OpExtension "SPV_KHR_storage_buffer_storage_class"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %52 "main" %49
OpExecutionMode %52 LocalSize 1 1 1
OpSource GLSL 450
OpMemberName %5 0 "data"
OpName %5 "PrimeIndices"
OpName %7 "v_indices"

217
tests/out/spv/debug-symbol-simple.spvasm Normal file
View File

@@ -0,0 +1,217 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 95
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %24 "vs_main" %15 %18 %20 %22
OpEntryPoint Fragment %57 "fs_main" %51 %54 %56
OpExecutionMode %57 OriginUpperLeft
%3 = OpString "debug-symbol-simple"
OpSource Unknown 0 %3 "struct VertexInput {
@location(0) position: vec3<f32>,
@location(1) color: vec3<f32>,
};
struct VertexOutput {
@builtin(position) clip_position: vec4<f32>,
@location(0) color: vec3<f32>,
};
@vertex
fn vs_main(
model: VertexInput,
) -> VertexOutput {
var out: VertexOutput;
out.color = model.color;
out.clip_position = vec4<f32>(model.position, 1.0);
return out;
}
// Fragment shader
@fragment
fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
var color = in.color;
for (var i = 0; i < 10; i += 1) {
var ii = f32(i);
color.x += ii*0.001;
color.y += ii*0.002;
}
return vec4<f32>(color, 1.0);
}"
OpMemberName %6 0 "position"
OpMemberName %6 1 "color"
OpName %6 "VertexInput"
OpMemberName %8 0 "clip_position"
OpMemberName %8 1 "color"
OpName %8 "VertexOutput"
OpName %10 "out"
OpName %15 "position"
OpName %18 "color"
OpName %20 "clip_position"
OpName %22 "color"
OpName %24 "vs_main"
OpName %41 "color"
OpName %43 "i"
OpName %46 "ii"
OpName %51 "clip_position"
OpName %54 "color"
OpName %57 "fs_main"
OpMemberDecorate %6 0 Offset 0
OpMemberDecorate %6 1 Offset 16
OpMemberDecorate %8 0 Offset 0
OpMemberDecorate %8 1 Offset 16
OpDecorate %15 Location 0
OpDecorate %18 Location 1
OpDecorate %20 BuiltIn Position
OpDecorate %22 Location 0
OpDecorate %51 BuiltIn FragCoord
OpDecorate %54 Location 0
OpDecorate %56 Location 0
%2 = OpTypeVoid
%5 = OpTypeFloat 32
%4 = OpTypeVector %5 3
%6 = OpTypeStruct %4 %4
%7 = OpTypeVector %5 4
%8 = OpTypeStruct %7 %4
%9 = OpTypeInt 32 1
%11 = OpTypePointer Function %8
%12 = OpConstantNull %8
%16 = OpTypePointer Input %4
%15 = OpVariable %16 Input
%18 = OpVariable %16 Input
%21 = OpTypePointer Output %7
%20 = OpVariable %21 Output
%23 = OpTypePointer Output %4
%22 = OpVariable %23 Output
%25 = OpTypeFunction %2
%26 = OpConstant %5 1.0
%28 = OpTypePointer Function %4
%31 = OpTypeInt 32 0
%30 = OpConstant %31 1
%33 = OpTypePointer Function %7
%36 = OpConstant %31 0
%42 = OpConstantNull %4
%44 = OpTypePointer Function %9
%45 = OpConstantNull %9
%47 = OpTypePointer Function %5
%48 = OpConstantNull %5
%52 = OpTypePointer Input %7
%51 = OpVariable %52 Input
%54 = OpVariable %16 Input
%56 = OpVariable %21 Output
%58 = OpConstant %9 0
%59 = OpConstant %9 10
%60 = OpConstant %5 0.001
%61 = OpConstant %5 0.002
%62 = OpConstant %9 1
%70 = OpTypeBool
%78 = OpTypePointer Function %5
%24 = OpFunction %2 None %25
%13 = OpLabel
%10 = OpVariable %11 Function %12
%17 = OpLoad %4 %15
%19 = OpLoad %4 %18
%14 = OpCompositeConstruct %6 %17 %19
OpBranch %27
%27 = OpLabel
OpLine %3 16 5
%29 = OpCompositeExtract %4 %14 1
OpLine %3 16 5
%32 = OpAccessChain %28 %10 %30
OpStore %32 %29
OpLine %3 17 5
%34 = OpCompositeExtract %4 %14 0
OpLine %3 17 25
%35 = OpCompositeConstruct %7 %34 %26
OpLine %3 17 5
%37 = OpAccessChain %33 %10 %36
OpStore %37 %35
OpLine %3 1 1
%38 = OpLoad %8 %10
%39 = OpCompositeExtract %7 %38 0
OpStore %20 %39
%40 = OpCompositeExtract %4 %38 1
OpStore %22 %40
OpReturn
OpFunctionEnd
%57 = OpFunction %2 None %25
%49 = OpLabel
%41 = OpVariable %28 Function %42
%43 = OpVariable %44 Function %45
%46 = OpVariable %47 Function %48
%53 = OpLoad %7 %51
%55 = OpLoad %4 %54
%50 = OpCompositeConstruct %8 %53 %55
OpBranch %63
%63 = OpLabel
OpLine %3 25 17
%64 = OpCompositeExtract %4 %50 1
OpLine %3 25 5
OpStore %41 %64
OpLine %3 26 10
OpStore %43 %58
OpBranch %65
%65 = OpLabel
OpLine %3 26 5
OpLoopMerge %66 %68 None
OpBranch %67
%67 = OpLabel
OpLine %3 1 1
%69 = OpLoad %9 %43
OpLine %3 26 21
%71 = OpSLessThan %70 %69 %59
OpLine %3 26 20
OpSelectionMerge %72 None
OpBranchConditional %71 %72 %73
%73 = OpLabel
OpBranch %66
%72 = OpLabel
OpBranch %74
%74 = OpLabel
OpLine %3 27 18
%76 = OpLoad %9 %43
%77 = OpConvertSToF %5 %76
OpLine %3 27 9
OpStore %46 %77
OpLine %3 28 9
%79 = OpLoad %5 %46
OpLine %3 28 9
%80 = OpFMul %5 %79 %60
%81 = OpAccessChain %78 %41 %36
%82 = OpLoad %5 %81
%83 = OpFAdd %5 %82 %80
OpLine %3 28 9
%84 = OpAccessChain %78 %41 %36
OpStore %84 %83
OpLine %3 29 9
%85 = OpLoad %5 %46
OpLine %3 29 9
%86 = OpFMul %5 %85 %61
%87 = OpAccessChain %78 %41 %30
%88 = OpLoad %5 %87
%89 = OpFAdd %5 %88 %86
OpLine %3 29 9
%90 = OpAccessChain %78 %41 %30
OpStore %90 %89
OpBranch %75
%75 = OpLabel
OpBranch %68
%68 = OpLabel
OpLine %3 26 29
%91 = OpLoad %9 %43
%92 = OpIAdd %9 %91 %62
OpLine %3 26 29
OpStore %43 %92
OpBranch %65
%66 = OpLabel
OpLine %3 1 1
%93 = OpLoad %4 %41
OpLine %3 32 12
%94 = OpCompositeConstruct %7 %93 %26
OpStore %56 %94
OpReturn
OpFunctionEnd

1470
tests/out/spv/debug-symbol-terrain.spvasm Normal file
View File

File diff suppressed because it is too large Load Diff

1
tests/out/spv/image.spvasm
View File

@@ -23,7 +23,6 @@ OpExecutionMode %299 OriginUpperLeft
OpExecutionMode %444 OriginUpperLeft
OpExecutionMode %499 OriginUpperLeft
OpExecutionMode %534 OriginUpperLeft
OpSource GLSL 450
OpName %35 "image_mipmapped_src"
OpName %37 "image_multisampled_src"
OpName %39 "image_depth_multisampled_src"

1
tests/out/spv/interpolate.spvasm
View File

@@ -9,7 +9,6 @@ OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %29 "vert_main" %13 %15 %17 %19 %21 %23 %24 %25 %26
OpEntryPoint Fragment %109 "frag_main" %88 %91 %94 %97 %100 %103 %105 %107
OpExecutionMode %109 OriginUpperLeft
OpSource GLSL 450
OpMemberName %8 0 "position"
OpMemberName %8 1 "_flat"
OpMemberName %8 2 "_linear"

1
tests/out/spv/padding.spvasm
View File

@@ -6,7 +6,6 @@ OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %28 "vertex" %26
OpSource GLSL 450
OpMemberName %7 0 "a"
OpName %7 "S"
OpMemberName %8 0 "a"

1
tests/out/spv/pointers.spvasm
View File

@@ -7,7 +7,6 @@ OpCapability Linkage
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpSource GLSL 450
OpMemberName %7 0 "arr"
OpName %7 "DynamicArray"
OpName %10 "dynamic_array"

1
tests/out/spv/policy-mix.spvasm
View File

@@ -8,7 +8,6 @@ OpCapability Linkage
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpSource GLSL 450
OpMemberName %14 0 "a"
OpName %14 "InStorage"
OpMemberName %17 0 "a"

1
tests/out/spv/quad.spvasm
View File

@@ -10,7 +10,6 @@ OpEntryPoint Fragment %44 "frag_main" %41 %43
OpEntryPoint Fragment %60 "fs_extra" %59
OpExecutionMode %44 OriginUpperLeft
OpExecutionMode %60 OriginUpperLeft
OpSource GLSL 450
OpName %3 "c_scale"
OpMemberName %7 0 "uv"
OpMemberName %7 1 "position"

1
tests/out/spv/shadow.spvasm
View File

@@ -11,7 +11,6 @@ OpEntryPoint Fragment %149 "fs_main" %140 %143 %146 %148
OpEntryPoint Fragment %218 "fs_main_without_storage" %211 %213 %215 %217
OpExecutionMode %149 OriginUpperLeft
OpExecutionMode %218 OriginUpperLeft
OpSource GLSL 450
OpName %7 "c_max_lights"
OpMemberName %13 0 "view_proj"
OpMemberName %13 1 "num_lights"

1
tests/out/spv/texture-arg.spvasm
View File

@@ -7,7 +7,6 @@ OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %28 "main" %26
OpExecutionMode %28 OriginUpperLeft
OpSource GLSL 450
OpName %8 "Texture"
OpName %10 "Sampler"
OpName %13 "Passed_Texture"

1
tests/out/spv/workgroup-var-init.spvasm
View File

@@ -8,7 +8,6 @@ OpExtension "SPV_KHR_storage_buffer_storage_class"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %19 "main" %27
OpExecutionMode %19 LocalSize 1 1 1
OpSource GLSL 450
OpMemberName %12 0 "arr"
OpMemberName %12 1 "atom"
OpMemberName %12 2 "atom_arr"

90
tests/snapshots.rs
View File

@@ -89,7 +89,7 @@ struct Parameters {
}
#[allow(unused_variables)]
fn check_targets(module: &naga::Module, name: &str, targets: Targets) {
fn check_targets(module: &naga::Module, name: &str, targets: Targets, source_code: Option<&str>) {
let root = env!("CARGO_MANIFEST_DIR");
let filepath = format!("{root}/{BASE_DIR_IN}/{name}.param.ron");
let params = match fs::read_to_string(&filepath) {
@@ -131,12 +131,22 @@ fn check_targets(module: &naga::Module, name: &str, targets: Targets) {
#[cfg(all(feature = "deserialize", feature = "spv-out"))]
{
let debug_info = if cfg!(feature = "span") {
source_code.map(|code| naga::back::spv::DebugInfo {
source_code: code,
file_name: name,
})
} else {
None
};
if targets.contains(Targets::SPIRV) {
write_output_spv(
module,
&info,
&dest,
name,
debug_info,
&params.spv,
params.bounds_check_policies,
);
@@ -204,6 +214,7 @@ fn write_output_spv(
info: &naga::valid::ModuleInfo,
destination: &Path,
file_name: &str,
debug_info: Option<naga::back::spv::DebugInfo>,
params: &SpirvOutParameters,
bounds_check_policies: naga::proc::BoundsCheckPolicies,
) {
@@ -232,6 +243,7 @@ fn write_output_spv(
bounds_check_policies,
binding_map: params.binding_map.clone(),
zero_initialize_workgroup_memory: spv::ZeroInitializeWorkgroupMemoryMode::Polyfill,
debug_info,
};
if params.separate_entry_points {
@@ -240,22 +252,52 @@ fn write_output_spv(
entry_point: ep.name.clone(),
shader_stage: ep.stage,
};
let spv = spv::write_vec(module, info, &options, Some(&pipeline_options)).unwrap();
let dis = rspirv::dr::load_words(spv)
.expect("Produced invalid SPIR-V")
.disassemble();
let path = format!("spv/{}.{}.spvasm", file_name, ep.name);
fs::write(destination.join(path), dis).unwrap();
write_output_spv_inner(
module,
info,
&options,
Some(&pipeline_options),
destination,
format!("spv/{}.{}.spvasm", file_name, ep.name),
);
}
} else {
let spv = spv::write_vec(module, info, &options, None).unwrap();
let dis = rspirv::dr::load_words(spv)
.expect("Produced invalid SPIR-V")
.disassemble();
fs::write(destination.join(format!("spv/{file_name}.spvasm")), dis).unwrap();
write_output_spv_inner(
module,
info,
&options,
None,
destination,
format!("spv/{file_name}.spvasm"),
);
}
}
#[cfg(feature = "spv-out")]
fn write_output_spv_inner(
module: &naga::Module,
info: &naga::valid::ModuleInfo,
options: &naga::back::spv::Options<'_>,
pipeline_options: Option<&naga::back::spv::PipelineOptions>,
destination: &Path,
path: String,
) {
use naga::back::spv;
use rspirv::binary::Disassemble;
let spv = spv::write_vec(module, info, options, pipeline_options).unwrap();
let dis = rspirv::dr::load_words(spv)
.expect("Produced invalid SPIR-V")
.disassemble();
// HACK escape CR/LF if source code is in side.
let dis = if options.debug_info.is_some() {
let dis = dis.replace("\\r", "\r");
dis.replace("\\n", "\n")
} else {
dis
};
fs::write(destination.join(path), dis).unwrap();
}
#[cfg(feature = "msl-out")]
fn write_output_msl(
module: &naga::Module,
@@ -573,10 +615,28 @@ fn convert_wgsl() {
let file = fs::read_to_string(format!("{root}/{BASE_DIR_IN}/{name}.wgsl"))
.expect("Couldn't find wgsl file");
match naga::front::wgsl::parse_str(&file) {
Ok(module) => check_targets(&module, name, targets),
Ok(module) => check_targets(&module, name, targets, None),
Err(e) => panic!("{}", e.emit_to_string(&file)),
}
}
#[cfg(feature = "span")]
{
let inputs = [
("debug-symbol-simple", Targets::SPIRV),
("debug-symbol-terrain", Targets::SPIRV),
];
for &(name, targets) in inputs.iter() {
println!("Processing '{name}'");
// WGSL shaders lives in root dir as a privileged.
let file = fs::read_to_string(format!("{root}/{BASE_DIR_IN}/{name}.wgsl"))
.expect("Couldn't find wgsl file");
match naga::front::wgsl::parse_str(&file) {
Ok(module) => check_targets(&module, name, targets, Some(&file)),
Err(e) => panic!("{}", e.emit_to_string(&file)),
}
}
}
}
#[cfg(feature = "spv-in")]
@@ -594,7 +654,7 @@ fn convert_spv(name: &str, adjust_coordinate_space: bool, targets: Targets) {
},
)
.unwrap();
check_targets(&module, name, targets);
check_targets(&module, name, targets, None);
naga::valid::Validator::new(
naga::valid::ValidationFlags::all(),
naga::valid::Capabilities::default(),
@@ -648,7 +708,7 @@ fn convert_glsl_variations_check() {
&file,
)
.unwrap();
check_targets(&module, "variations-glsl", Targets::GLSL);
check_targets(&module, "variations-glsl", Targets::GLSL, None);
}
#[cfg(feature = "glsl-in")]

4
tests/spirv-capabilities.rs
View File

@@ -23,7 +23,9 @@ fn capabilities_used(source: &str) -> naga::FastIndexSet<Ca> {
let mut words = vec![];
let mut writer = spv::Writer::new(&spv::Options::default()).unwrap();
writer.write(&module, &info, None, &mut words).unwrap();
writer
.write(&module, &info, None, &None, &mut words)
.unwrap();
writer.get_capabilities_used().clone()
}