From 25e4f17a691e22e157abe31df83724d6ea69ab18 Mon Sep 17 00:00:00 2001 From: wicast Date: Wed, 28 Jun 2023 20:30:25 +0800 Subject: [PATCH] add debug info for spv-out (#2379) --- benches/criterion.rs | 4 +- cli/src/bin/naga.rs | 23 +- src/back/spv/block.rs | 49 +- src/back/spv/instructions.rs | 31 +- src/back/spv/mod.rs | 27 +- src/back/spv/writer.rs | 40 +- tests/in/debug-symbol-simple.param.ron | 7 + tests/in/debug-symbol-simple.wgsl | 33 + tests/in/debug-symbol-terrain.param.ron | 7 + tests/in/debug-symbol-terrain.wgsl | 297 ++++ tests/out/spv/access.spvasm | 1 - tests/out/spv/boids.spvasm | 1 - .../spv/bounds-check-image-restrict.spvasm | 1 - tests/out/spv/bounds-check-image-rzsw.spvasm | 1 - tests/out/spv/collatz.spvasm | 1 - tests/out/spv/debug-symbol-simple.spvasm | 217 +++ tests/out/spv/debug-symbol-terrain.spvasm | 1470 +++++++++++++++++ tests/out/spv/image.spvasm | 1 - tests/out/spv/interpolate.spvasm | 1 - tests/out/spv/padding.spvasm | 1 - tests/out/spv/pointers.spvasm | 1 - tests/out/spv/policy-mix.spvasm | 1 - tests/out/spv/quad.spvasm | 1 - tests/out/spv/shadow.spvasm | 1 - tests/out/spv/texture-arg.spvasm | 1 - tests/out/spv/workgroup-var-init.spvasm | 1 - tests/snapshots.rs | 90 +- tests/spirv-capabilities.rs | 4 +- 28 files changed, 2258 insertions(+), 55 deletions(-) create mode 100644 tests/in/debug-symbol-simple.param.ron create mode 100644 tests/in/debug-symbol-simple.wgsl create mode 100644 tests/in/debug-symbol-terrain.param.ron create mode 100644 tests/in/debug-symbol-terrain.wgsl create mode 100644 tests/out/spv/debug-symbol-simple.spvasm create mode 100644 tests/out/spv/debug-symbol-terrain.spvasm diff --git a/benches/criterion.rs b/benches/criterion.rs index 624744555f..9c0dc225fc 100644 --- a/benches/criterion.rs +++ b/benches/criterion.rs @@ -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(); } diff --git a/cli/src/bin/naga.rs b/cli/src/bin/naga.rs index b121dd52db..4904a78cf4 100644 --- a/cli/src/bin/naga.rs +++ b/cli/src/bin/naga.rs @@ -68,6 +68,10 @@ struct Args { #[argh(option)] stdin_file_path: Option, + /// generate debug symbols, only works for spv-out for now + #[argh(option, short = 'g')] + generate_debug_symbols: Option, + /// 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, keep_coordinate_space: bool, spv_block_ctx_dump_prefix: Option, - 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> { { 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> { }; 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, diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index fa95966552..84910e6999 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -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); diff --git a/src/back/spv/instructions.rs b/src/back/spv/instructions.rs index 31ed6e231d..11794fc73b 100644 --- a/src/back/spv/instructions.rs +++ b/src/back/spv/instructions.rs @@ -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, + ) -> 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 // diff --git a/src/back/spv/mod.rs b/src/back/spv/mod.rs index 642c0c7d1b..4e1ab50150 100644 --- a/src/back/spv/mod.rs +++ b/src/back/spv/mod.rs @@ -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, } @@ -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>, } -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, Error> { - let mut words = Vec::new(); + let mut words: Vec = 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) } diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index 4a1bc8ff20..554db6d8a2 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -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, + debug_info: &Option, ) -> Result { 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, ) -> Result { 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, + debug_info: &Option, ) -> 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, words: &mut Vec, ) -> 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); diff --git a/tests/in/debug-symbol-simple.param.ron b/tests/in/debug-symbol-simple.param.ron new file mode 100644 index 0000000000..2869084cd0 --- /dev/null +++ b/tests/in/debug-symbol-simple.param.ron @@ -0,0 +1,7 @@ +( + spv: ( + version: (1, 1), + debug: true, + adjust_coordinate_space: false, + ), +) \ No newline at end of file diff --git a/tests/in/debug-symbol-simple.wgsl b/tests/in/debug-symbol-simple.wgsl new file mode 100644 index 0000000000..86505a1592 --- /dev/null +++ b/tests/in/debug-symbol-simple.wgsl @@ -0,0 +1,33 @@ +struct VertexInput { + @location(0) position: vec3, + @location(1) color: vec3, +}; + +struct VertexOutput { + @builtin(position) clip_position: vec4, + @location(0) color: vec3, +}; + +@vertex +fn vs_main( + model: VertexInput, +) -> VertexOutput { + var out: VertexOutput; + out.color = model.color; + out.clip_position = vec4(model.position, 1.0); + return out; +} + +// Fragment shader + +@fragment +fn fs_main(in: VertexOutput) -> @location(0) vec4 { + 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(color, 1.0); +} \ No newline at end of file diff --git a/tests/in/debug-symbol-terrain.param.ron b/tests/in/debug-symbol-terrain.param.ron new file mode 100644 index 0000000000..2869084cd0 --- /dev/null +++ b/tests/in/debug-symbol-terrain.param.ron @@ -0,0 +1,7 @@ +( + spv: ( + version: (1, 1), + debug: true, + adjust_coordinate_space: false, + ), +) \ No newline at end of file diff --git a/tests/in/debug-symbol-terrain.wgsl b/tests/in/debug-symbol-terrain.wgsl new file mode 100644 index 0000000000..3f3dd58dd3 --- /dev/null +++ b/tests/in/debug-symbol-terrain.wgsl @@ -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) -> vec3 { return (((x * 34.) + 1.) * x) % vec3(289.); } + +fn snoise2(v: vec2) -> f32 { + let C = vec4(0.211324865405187, 0.366025403784439, -0.577350269189626, 0.024390243902439); + var i: vec2 = 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 = select(vec2(1., 0.), vec2(0., 1.), (x0.x < x0.y)); + var x12: vec4 = x0.xyxy + C.xxzz - vec4(i1, 0., 0.); + i = i % vec2(289.); + let p = permute3(permute3(i.y + vec3(0., i1.y, 1.)) + i.x + vec3(0., i1.x, 1.)); + var m: vec3 = max(0.5 - vec3(dot(x0, x0), dot(x12.xy, x12.xy), dot(x12.zw, x12.zw)), vec3(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(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 { + let NUM_OCTAVES: u32 = 5u; + var x = p * 0.01; + var v = 0.0; + var a = 0.5; + let shift = vec2(100.0); + let cs = vec2(cos(0.5), sin(0.5)); + let rot = mat2x2(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, + chunk_corner: vec2, + min_max_height: vec2, +} + +struct Vertex { + @location(0) position: vec3, + @location(1) normal: vec3, +} + +struct VertexBuffer { + data: array, // stride: 32 +} + +struct IndexBuffer { + data: array, +} + +@group(0) @binding(0) var chunk_data: ChunkData; +@group(0) @binding(1) var vertices: VertexBuffer; +@group(0) @binding(2) var indices: IndexBuffer; + +fn terrain_point(p: vec2, min_max_height: vec2) -> vec3 { + return vec3( + p.x, + mix(min_max_height.x, min_max_height.y, fbm(p)), + p.y, + ); +} + +fn terrain_vertex(p: vec2, min_max_height: vec2) -> Vertex { + let v = terrain_point(p, min_max_height); + + let tpx = terrain_point(p + vec2(0.1, 0.0), min_max_height) - v; + let tpz = terrain_point(p + vec2(0.0, 0.1), min_max_height) - v; + let tnx = terrain_point(p + vec2(-0.1, 0.0), min_max_height) - v; + let tnz = terrain_point(p + vec2(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, chunk_corner: vec2) -> vec2 { + return vec2( + f32(vert_index) % f32(chunk_size.x + 1u), + f32(vert_index / (chunk_size.x + 1u)), + ) + vec2(chunk_corner); +} + +@compute @workgroup_size(64) +fn gen_terrain_compute( + @builtin(global_invocation_id) gid: vec3 +) { + // 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, + chunk_corner: vec2, + min_max_height: vec2, + texture_size: u32, + start_index: u32, +} +@group(0) +@binding(0) +var gen_data: GenData; + +struct GenVertexOutput { + @location(0) + index: u32, + @builtin(position) + position: vec4, + @location(1) + uv: vec2, +}; + +@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(u, v); + + let position = vec4(-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(vert_component); + return GenFragmentOutput(ivert_component, index); +} + +// ============================ +// Terrain Rendering +// ============================ + +struct Camera { + view_pos: vec4, + view_proj: mat4x4, +} +@group(0) @binding(0) +var camera: Camera; + +struct Light { + position: vec3, + color: vec3, +} +@group(1) @binding(0) +var light: Light; + +struct VertexOutput { + @builtin(position) clip_position: vec4, + @location(0) normal: vec3, + @location(1) world_pos: vec3, +} + +@vertex +fn vs_main( + vertex: Vertex, +) -> VertexOutput { + let clip_position = camera.view_proj * vec4(vertex.position, 1.); + let normal = vertex.normal; + return VertexOutput(clip_position, normal, vertex.position); +} + +@group(2) @binding(0) +var t_diffuse: texture_2d; +@group(2) @binding(1) +var s_diffuse: sampler; +@group(2) @binding(2) +var t_normal: texture_2d; +@group(2) @binding(3) +var s_normal: sampler; + +fn color23(p: vec2) -> vec3 { + return vec3( + snoise2(p) * 0.5 + 0.5, + snoise2(p + vec2(23., 32.)) * 0.5 + 0.5, + snoise2(p + vec2(-43., 3.)) * 0.5 + 0.5, + ); +} + +@fragment +fn fs_main(in: VertexOutput) -> @location(0) vec4 { + var color = smoothstep(vec3(0.0), vec3(0.1), fract(in.world_pos)); + color = mix(vec3(0.5, 0.1, 0.7), vec3(0.2, 0.2, 0.2), vec3(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(result, 1.0); +} \ No newline at end of file diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index 92e3a2fbed..12d2ceef26 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -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" diff --git a/tests/out/spv/boids.spvasm b/tests/out/spv/boids.spvasm index 002959167f..4b19d9b4fc 100644 --- a/tests/out/spv/boids.spvasm +++ b/tests/out/spv/boids.spvasm @@ -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" diff --git a/tests/out/spv/bounds-check-image-restrict.spvasm b/tests/out/spv/bounds-check-image-restrict.spvasm index c884210046..1a99215ef8 100644 --- a/tests/out/spv/bounds-check-image-restrict.spvasm +++ b/tests/out/spv/bounds-check-image-restrict.spvasm @@ -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" diff --git a/tests/out/spv/bounds-check-image-rzsw.spvasm b/tests/out/spv/bounds-check-image-rzsw.spvasm index e5e12fabd3..53e27a7969 100644 --- a/tests/out/spv/bounds-check-image-rzsw.spvasm +++ b/tests/out/spv/bounds-check-image-rzsw.spvasm @@ -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" diff --git a/tests/out/spv/collatz.spvasm b/tests/out/spv/collatz.spvasm index a08b44670f..a8ed9ab784 100644 --- a/tests/out/spv/collatz.spvasm +++ b/tests/out/spv/collatz.spvasm @@ -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" diff --git a/tests/out/spv/debug-symbol-simple.spvasm b/tests/out/spv/debug-symbol-simple.spvasm new file mode 100644 index 0000000000..67ad8beb29 --- /dev/null +++ b/tests/out/spv/debug-symbol-simple.spvasm @@ -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, + @location(1) color: vec3, +}; + +struct VertexOutput { + @builtin(position) clip_position: vec4, + @location(0) color: vec3, +}; + +@vertex +fn vs_main( + model: VertexInput, +) -> VertexOutput { + var out: VertexOutput; + out.color = model.color; + out.clip_position = vec4(model.position, 1.0); + return out; +} + +// Fragment shader + +@fragment +fn fs_main(in: VertexOutput) -> @location(0) vec4 { + 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(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 \ No newline at end of file diff --git a/tests/out/spv/debug-symbol-terrain.spvasm b/tests/out/spv/debug-symbol-terrain.spvasm new file mode 100644 index 0000000000..9fb3bbff9a --- /dev/null +++ b/tests/out/spv/debug-symbol-terrain.spvasm @@ -0,0 +1,1470 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 653 +OpCapability Shader +OpExtension "SPV_KHR_storage_buffer_storage_class" +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %351 "gen_terrain_compute" %348 +OpEntryPoint Vertex %421 "gen_terrain_vertex" %412 %415 %417 %419 +OpEntryPoint Fragment %475 "gen_terrain_fragment" %465 %467 %470 %473 %474 +OpEntryPoint Vertex %566 "vs_main" %557 %560 %562 %563 %565 +OpEntryPoint Fragment %593 "fs_main" %586 %588 %590 %592 +OpExecutionMode %351 LocalSize 64 1 1 +OpExecutionMode %475 OriginUpperLeft +OpExecutionMode %593 OriginUpperLeft +%3 = OpString "debug-symbol-terrain" +OpSource Unknown 0 %3 "// 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) -> vec3 { return (((x * 34.) + 1.) * x) % vec3(289.); } + +fn snoise2(v: vec2) -> f32 { + let C = vec4(0.211324865405187, 0.366025403784439, -0.577350269189626, 0.024390243902439); + var i: vec2 = 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 = select(vec2(1., 0.), vec2(0., 1.), (x0.x < x0.y)); + var x12: vec4 = x0.xyxy + C.xxzz - vec4(i1, 0., 0.); + i = i % vec2(289.); + let p = permute3(permute3(i.y + vec3(0., i1.y, 1.)) + i.x + vec3(0., i1.x, 1.)); + var m: vec3 = max(0.5 - vec3(dot(x0, x0), dot(x12.xy, x12.xy), dot(x12.zw, x12.zw)), vec3(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(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 { + let NUM_OCTAVES: u32 = 5u; + var x = p * 0.01; + var v = 0.0; + var a = 0.5; + let shift = vec2(100.0); + let cs = vec2(cos(0.5), sin(0.5)); + let rot = mat2x2(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, + chunk_corner: vec2, + min_max_height: vec2, +} + +struct Vertex { + @location(0) position: vec3, + @location(1) normal: vec3, +} + +struct VertexBuffer { + data: array, // stride: 32 +} + +struct IndexBuffer { + data: array, +} + +@group(0) @binding(0) var chunk_data: ChunkData; +@group(0) @binding(1) var vertices: VertexBuffer; +@group(0) @binding(2) var indices: IndexBuffer; + +fn terrain_point(p: vec2, min_max_height: vec2) -> vec3 { + return vec3( + p.x, + mix(min_max_height.x, min_max_height.y, fbm(p)), + p.y, + ); +} + +fn terrain_vertex(p: vec2, min_max_height: vec2) -> Vertex { + let v = terrain_point(p, min_max_height); + + let tpx = terrain_point(p + vec2(0.1, 0.0), min_max_height) - v; + let tpz = terrain_point(p + vec2(0.0, 0.1), min_max_height) - v; + let tnx = terrain_point(p + vec2(-0.1, 0.0), min_max_height) - v; + let tnz = terrain_point(p + vec2(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, chunk_corner: vec2) -> vec2 { + return vec2( + f32(vert_index) % f32(chunk_size.x + 1u), + f32(vert_index / (chunk_size.x + 1u)), + ) + vec2(chunk_corner); +} + +@compute @workgroup_size(64) +fn gen_terrain_compute( + @builtin(global_invocation_id) gid: vec3 +) { + // 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, + chunk_corner: vec2, + min_max_height: vec2, + texture_size: u32, + start_index: u32, +} +@group(0) +@binding(0) +var gen_data: GenData; + +struct GenVertexOutput { + @location(0) + index: u32, + @builtin(position) + position: vec4, + @location(1) + uv: vec2, +}; + +@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(u, v); + + let position = vec4(-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(vert_component); + return GenFragmentOutput(ivert_component, index); +} + +// ============================ +// Terrain Rendering +// ============================ + +struct Camera { + view_pos: vec4, + view_proj: mat4x4, +} +@group(0) @binding(0) +var camera: Camera; + +struct Light { + position: vec3, + color: vec3, +} +@group(1) @binding(0) +var light: Light; + +struct VertexOutput { + @builtin(position) clip_position: vec4, + @location(0) normal: vec3, + @location(1) world_pos: vec3, +} + +@vertex +fn vs_main( + vertex: Vertex, +) -> VertexOutput { + let clip_position = camera.view_proj * vec4(vertex.position, 1.); + let normal = vertex.normal; + return VertexOutput(clip_position, normal, vertex.position); +} + +@group(2) @binding(0) +var t_diffuse: texture_2d; +@group(2) @binding(1) +var s_diffuse: sampler; +@group(2) @binding(2) +var t_normal: texture_2d; +@group(2) @binding(3) +var s_normal: sampler; + +fn color23(p: vec2) -> vec3 { + return vec3( + snoise2(p) * 0.5 + 0.5, + snoise2(p + vec2(23., 32.)) * 0.5 + 0.5, + snoise2(p + vec2(-43., 3.)) * 0.5 + 0.5, + ); +} + +@fragment +fn fs_main(in: VertexOutput) -> @location(0) vec4 { + var color = smoothstep(vec3(0.0), vec3(0.1), fract(in.world_pos)); + color = mix(vec3(0.5, 0.1, 0.7), vec3(0.2, 0.2, 0.2), vec3(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(result, 1.0); +}" +OpMemberName %13 0 "chunk_size" +OpMemberName %13 1 "chunk_corner" +OpMemberName %13 2 "min_max_height" +OpName %13 "ChunkData" +OpMemberName %14 0 "position" +OpMemberName %14 1 "normal" +OpName %14 "Vertex" +OpMemberName %16 0 "data" +OpName %16 "VertexBuffer" +OpMemberName %18 0 "data" +OpName %18 "IndexBuffer" +OpMemberName %20 0 "chunk_size" +OpMemberName %20 1 "chunk_corner" +OpMemberName %20 2 "min_max_height" +OpMemberName %20 3 "texture_size" +OpMemberName %20 4 "start_index" +OpName %20 "GenData" +OpMemberName %21 0 "index" +OpMemberName %21 1 "position" +OpMemberName %21 2 "uv" +OpName %21 "GenVertexOutput" +OpMemberName %22 0 "vert_component" +OpMemberName %22 1 "index" +OpName %22 "GenFragmentOutput" +OpMemberName %24 0 "view_pos" +OpMemberName %24 1 "view_proj" +OpName %24 "Camera" +OpMemberName %25 0 "position" +OpMemberName %25 1 "color" +OpName %25 "Light" +OpMemberName %26 0 "clip_position" +OpMemberName %26 1 "normal" +OpMemberName %26 2 "world_pos" +OpName %26 "VertexOutput" +OpName %29 "chunk_data" +OpName %32 "vertices" +OpName %34 "indices" +OpName %36 "gen_data" +OpName %39 "camera" +OpName %42 "light" +OpName %45 "t_diffuse" +OpName %47 "s_diffuse" +OpName %49 "t_normal" +OpName %50 "s_normal" +OpName %52 "x" +OpName %53 "permute3" +OpName %65 "i" +OpName %68 "i1" +OpName %70 "x12" +OpName %73 "m" +OpName %77 "v" +OpName %78 "snoise2" +OpName %205 "x" +OpName %207 "v" +OpName %210 "a" +OpName %212 "i" +OpName %216 "p" +OpName %217 "fbm" +OpName %261 "p" +OpName %262 "min_max_height" +OpName %263 "terrain_point" +OpName %274 "p" +OpName %275 "min_max_height" +OpName %276 "terrain_vertex" +OpName %306 "vert_index" +OpName %307 "chunk_size" +OpName %308 "chunk_corner" +OpName %309 "index_to_p" +OpName %325 "p" +OpName %326 "color23" +OpName %348 "gid" +OpName %351 "gen_terrain_compute" +OpName %412 "vindex" +OpName %415 "index" +OpName %417 "position" +OpName %419 "uv" +OpName %421 "gen_terrain_vertex" +OpName %459 "vert_component" +OpName %461 "index" +OpName %465 "index" +OpName %467 "position" +OpName %470 "uv" +OpName %473 "vert_component" +OpName %474 "index" +OpName %475 "gen_terrain_fragment" +OpName %557 "position" +OpName %560 "normal" +OpName %562 "clip_position" +OpName %563 "normal" +OpName %565 "world_pos" +OpName %566 "vs_main" +OpName %582 "color" +OpName %586 "clip_position" +OpName %588 "normal" +OpName %590 "world_pos" +OpName %593 "fs_main" +OpMemberDecorate %13 0 Offset 0 +OpMemberDecorate %13 1 Offset 8 +OpMemberDecorate %13 2 Offset 16 +OpMemberDecorate %14 0 Offset 0 +OpMemberDecorate %14 1 Offset 16 +OpDecorate %15 ArrayStride 32 +OpMemberDecorate %16 0 Offset 0 +OpDecorate %17 ArrayStride 4 +OpMemberDecorate %18 0 Offset 0 +OpMemberDecorate %20 0 Offset 0 +OpMemberDecorate %20 1 Offset 8 +OpMemberDecorate %20 2 Offset 16 +OpMemberDecorate %20 3 Offset 24 +OpMemberDecorate %20 4 Offset 28 +OpMemberDecorate %21 0 Offset 0 +OpMemberDecorate %21 1 Offset 16 +OpMemberDecorate %21 2 Offset 32 +OpMemberDecorate %22 0 Offset 0 +OpMemberDecorate %22 1 Offset 4 +OpMemberDecorate %24 0 Offset 0 +OpMemberDecorate %24 1 Offset 16 +OpMemberDecorate %24 1 ColMajor +OpMemberDecorate %24 1 MatrixStride 16 +OpMemberDecorate %25 0 Offset 0 +OpMemberDecorate %25 1 Offset 16 +OpMemberDecorate %26 0 Offset 0 +OpMemberDecorate %26 1 Offset 16 +OpMemberDecorate %26 2 Offset 32 +OpDecorate %29 DescriptorSet 0 +OpDecorate %29 Binding 0 +OpDecorate %30 Block +OpMemberDecorate %30 0 Offset 0 +OpDecorate %32 DescriptorSet 0 +OpDecorate %32 Binding 1 +OpDecorate %16 Block +OpDecorate %34 DescriptorSet 0 +OpDecorate %34 Binding 2 +OpDecorate %18 Block +OpDecorate %36 DescriptorSet 0 +OpDecorate %36 Binding 0 +OpDecorate %37 Block +OpMemberDecorate %37 0 Offset 0 +OpDecorate %39 DescriptorSet 0 +OpDecorate %39 Binding 0 +OpDecorate %40 Block +OpMemberDecorate %40 0 Offset 0 +OpDecorate %42 DescriptorSet 1 +OpDecorate %42 Binding 0 +OpDecorate %43 Block +OpMemberDecorate %43 0 Offset 0 +OpDecorate %45 DescriptorSet 2 +OpDecorate %45 Binding 0 +OpDecorate %47 DescriptorSet 2 +OpDecorate %47 Binding 1 +OpDecorate %49 DescriptorSet 2 +OpDecorate %49 Binding 2 +OpDecorate %50 DescriptorSet 2 +OpDecorate %50 Binding 3 +OpDecorate %348 BuiltIn GlobalInvocationId +OpDecorate %412 BuiltIn VertexIndex +OpDecorate %415 Location 0 +OpDecorate %415 Flat +OpDecorate %417 BuiltIn Position +OpDecorate %419 Location 1 +OpDecorate %465 Location 0 +OpDecorate %465 Flat +OpDecorate %467 BuiltIn FragCoord +OpDecorate %470 Location 1 +OpDecorate %473 Location 0 +OpDecorate %474 Location 1 +OpDecorate %557 Location 0 +OpDecorate %560 Location 1 +OpDecorate %562 BuiltIn Position +OpDecorate %563 Location 0 +OpDecorate %565 Location 1 +OpDecorate %586 BuiltIn FragCoord +OpDecorate %588 Location 0 +OpDecorate %590 Location 1 +OpDecorate %592 Location 0 +%2 = OpTypeVoid +%5 = OpTypeFloat 32 +%4 = OpTypeVector %5 3 +%6 = OpTypeVector %5 2 +%7 = OpTypeVector %5 4 +%8 = OpTypeInt 32 0 +%9 = OpTypeMatrix %6 2 +%10 = OpTypeVector %8 2 +%12 = OpTypeInt 32 1 +%11 = OpTypeVector %12 2 +%13 = OpTypeStruct %10 %11 %6 +%14 = OpTypeStruct %4 %4 +%15 = OpTypeRuntimeArray %14 +%16 = OpTypeStruct %15 +%17 = OpTypeRuntimeArray %8 +%18 = OpTypeStruct %17 +%19 = OpTypeVector %8 3 +%20 = OpTypeStruct %10 %11 %6 %8 %8 +%21 = OpTypeStruct %8 %7 %6 +%22 = OpTypeStruct %8 %8 +%23 = OpTypeMatrix %7 4 +%24 = OpTypeStruct %7 %23 +%25 = OpTypeStruct %4 %4 +%26 = OpTypeStruct %7 %4 %4 +%27 = OpTypeImage %5 2D 0 0 0 1 Unknown +%28 = OpTypeSampler +%30 = OpTypeStruct %13 +%31 = OpTypePointer Uniform %30 +%29 = OpVariable %31 Uniform +%33 = OpTypePointer StorageBuffer %16 +%32 = OpVariable %33 StorageBuffer +%35 = OpTypePointer StorageBuffer %18 +%34 = OpVariable %35 StorageBuffer +%37 = OpTypeStruct %20 +%38 = OpTypePointer Uniform %37 +%36 = OpVariable %38 Uniform +%40 = OpTypeStruct %24 +%41 = OpTypePointer Uniform %40 +%39 = OpVariable %41 Uniform +%43 = OpTypeStruct %25 +%44 = OpTypePointer Uniform %43 +%42 = OpVariable %44 Uniform +%46 = OpTypePointer UniformConstant %27 +%45 = OpVariable %46 UniformConstant +%48 = OpTypePointer UniformConstant %28 +%47 = OpVariable %48 UniformConstant +%49 = OpVariable %46 UniformConstant +%50 = OpVariable %48 UniformConstant +%54 = OpTypeFunction %4 %4 +%55 = OpConstant %5 34.0 +%56 = OpConstant %5 1.0 +%57 = OpConstant %5 289.0 +%66 = OpTypePointer Function %6 +%67 = OpConstantNull %6 +%69 = OpConstantNull %6 +%71 = OpTypePointer Function %7 +%72 = OpConstantNull %7 +%74 = OpTypePointer Function %4 +%75 = OpConstantNull %4 +%79 = OpTypeFunction %5 %6 +%80 = OpConstant %5 0.21132487 +%81 = OpConstant %5 0.36602542 +%82 = OpConstant %5 -0.57735026 +%83 = OpConstant %5 0.024390243 +%84 = OpConstant %5 0.0 +%85 = OpConstant %5 0.5 +%86 = OpConstant %5 2.0 +%87 = OpConstant %5 1.7928429 +%88 = OpConstant %5 0.85373473 +%89 = OpConstant %5 130.0 +%108 = OpTypeBool +%111 = OpTypeVector %108 2 +%122 = OpTypePointer Function %5 +%123 = OpConstant %8 1 +%132 = OpConstant %8 0 +%206 = OpConstantNull %6 +%208 = OpTypePointer Function %5 +%209 = OpConstantNull %5 +%211 = OpConstantNull %5 +%213 = OpTypePointer Function %8 +%214 = OpConstantNull %8 +%218 = OpConstant %8 5 +%219 = OpConstant %5 0.01 +%220 = OpConstant %5 100.0 +%264 = OpTypeFunction %4 %6 %6 +%277 = OpTypeFunction %14 %6 %6 +%278 = OpConstant %5 0.1 +%279 = OpConstant %5 -0.1 +%310 = OpTypeFunction %6 %8 %10 %11 +%327 = OpTypeFunction %4 %6 +%328 = OpConstant %5 23.0 +%329 = OpConstant %5 32.0 +%330 = OpConstant %5 -43.0 +%331 = OpConstant %5 3.0 +%349 = OpTypePointer Input %19 +%348 = OpVariable %349 Input +%352 = OpTypeFunction %2 +%353 = OpTypePointer Uniform %13 +%355 = OpConstant %8 6 +%356 = OpConstant %8 2 +%357 = OpConstant %8 3 +%358 = OpConstant %8 4 +%361 = OpTypePointer Uniform %10 +%364 = OpTypePointer Uniform %11 +%368 = OpTypePointer StorageBuffer %15 +%369 = OpTypePointer StorageBuffer %14 +%370 = OpTypePointer Uniform %6 +%377 = OpTypePointer Uniform %8 +%398 = OpTypePointer StorageBuffer %17 +%399 = OpTypePointer StorageBuffer %8 +%413 = OpTypePointer Input %8 +%412 = OpVariable %413 Input +%416 = OpTypePointer Output %8 +%415 = OpVariable %416 Output +%418 = OpTypePointer Output %7 +%417 = OpVariable %418 Output +%420 = OpTypePointer Output %6 +%419 = OpVariable %420 Output +%422 = OpTypePointer Uniform %20 +%424 = OpConstant %5 -1.0 +%440 = OpTypePointer Uniform %8 +%460 = OpConstantNull %5 +%462 = OpConstantNull %8 +%465 = OpVariable %413 Input +%468 = OpTypePointer Input %7 +%467 = OpVariable %468 Input +%471 = OpTypePointer Input %6 +%470 = OpVariable %471 Input +%473 = OpVariable %416 Output +%474 = OpVariable %416 Output +%477 = OpConstant %5 6.0 +%558 = OpTypePointer Input %4 +%557 = OpVariable %558 Input +%560 = OpVariable %558 Input +%562 = OpVariable %418 Output +%564 = OpTypePointer Output %4 +%563 = OpVariable %564 Output +%565 = OpVariable %564 Output +%567 = OpTypePointer Uniform %24 +%570 = OpTypePointer Uniform %23 +%583 = OpConstantNull %4 +%586 = OpVariable %468 Input +%588 = OpVariable %558 Input +%590 = OpVariable %558 Input +%592 = OpVariable %418 Output +%595 = OpTypePointer Uniform %25 +%597 = OpConstant %5 0.7 +%598 = OpConstant %5 0.2 +%617 = OpTypePointer Uniform %4 +%626 = OpTypePointer Uniform %7 +%53 = OpFunction %4 None %54 +%52 = OpFunctionParameter %4 +%51 = OpLabel +OpBranch %58 +%58 = OpLabel +OpLine %3 10 52 +%59 = OpVectorTimesScalar %4 %52 %55 +OpLine %3 10 50 +%60 = OpCompositeConstruct %4 %56 %56 %56 +%61 = OpFAdd %4 %59 %60 +%62 = OpFMul %4 %61 %52 +OpLine %3 10 49 +%63 = OpCompositeConstruct %4 %57 %57 %57 +%64 = OpFRem %4 %62 %63 +OpReturnValue %64 +OpFunctionEnd +%78 = OpFunction %5 None %79 +%77 = OpFunctionParameter %6 +%76 = OpLabel +%68 = OpVariable %66 Function %69 +%73 = OpVariable %74 Function %75 +%65 = OpVariable %66 Function %67 +%70 = OpVariable %71 Function %72 +OpBranch %90 +%90 = OpLabel +OpLine %3 13 13 +%91 = OpCompositeConstruct %7 %80 %81 %82 %83 +OpLine %3 14 24 +%92 = OpVectorShuffle %6 %91 %91 1 1 +%93 = OpDot %5 %77 %92 +%94 = OpCompositeConstruct %6 %93 %93 +%95 = OpFAdd %6 %77 %94 +%96 = OpExtInst %6 %1 Floor %95 +OpLine %3 14 5 +OpStore %65 %96 +OpLine %3 15 14 +%97 = OpLoad %6 %65 +%98 = OpFSub %6 %77 %97 +%99 = OpLoad %6 %65 +%100 = OpVectorShuffle %6 %91 %91 0 0 +%101 = OpDot %5 %99 %100 +%102 = OpCompositeConstruct %6 %101 %101 +%103 = OpFAdd %6 %98 %102 +OpLine %3 17 32 +%104 = OpCompositeConstruct %6 %56 %84 +OpLine %3 17 25 +%105 = OpCompositeConstruct %6 %84 %56 +%106 = OpCompositeExtract %5 %103 0 +%107 = OpCompositeExtract %5 %103 1 +%109 = OpFOrdLessThan %108 %106 %107 +%112 = OpCompositeConstruct %111 %109 %109 +%110 = OpSelect %6 %112 %105 %104 +OpLine %3 17 5 +OpStore %68 %110 +OpLine %3 18 26 +%113 = OpVectorShuffle %7 %103 %103 0 1 0 1 +%114 = OpVectorShuffle %7 %91 %91 0 0 2 2 +%115 = OpFAdd %7 %113 %114 +%116 = OpLoad %6 %68 +OpLine %3 18 26 +%117 = OpCompositeConstruct %7 %116 %84 %84 +%118 = OpFSub %7 %115 %117 +OpLine %3 18 5 +OpStore %70 %118 +OpLine %3 1 1 +%119 = OpLoad %6 %65 +OpLine %3 19 9 +%120 = OpCompositeConstruct %6 %57 %57 +%121 = OpFRem %6 %119 %120 +OpLine %3 19 5 +OpStore %65 %121 +OpLine %3 20 31 +%124 = OpAccessChain %122 %65 %123 +%125 = OpLoad %5 %124 +OpLine %3 20 51 +%126 = OpAccessChain %122 %68 %123 +%127 = OpLoad %5 %126 +OpLine %3 20 31 +%128 = OpCompositeConstruct %4 %84 %127 %56 +%129 = OpCompositeConstruct %4 %125 %125 %125 +%130 = OpFAdd %4 %129 %128 +OpLine %3 20 22 +%131 = OpFunctionCall %4 %53 %130 +OpLine %3 20 22 +%133 = OpAccessChain %122 %65 %132 +%134 = OpLoad %5 %133 +%135 = OpCompositeConstruct %4 %134 %134 %134 +%136 = OpFAdd %4 %131 %135 +OpLine %3 20 84 +%137 = OpAccessChain %122 %68 %132 +%138 = OpLoad %5 %137 +OpLine %3 20 22 +%139 = OpCompositeConstruct %4 %84 %138 %56 +%140 = OpFAdd %4 %136 %139 +OpLine %3 20 13 +%141 = OpFunctionCall %4 %53 %140 +OpLine %3 21 28 +%142 = OpDot %5 %103 %103 +%143 = OpLoad %7 %70 +%144 = OpVectorShuffle %6 %143 %143 0 1 +%145 = OpLoad %7 %70 +%146 = OpVectorShuffle %6 %145 %145 0 1 +%147 = OpDot %5 %144 %146 +%148 = OpLoad %7 %70 +%149 = OpVectorShuffle %6 %148 %148 2 3 +%150 = OpLoad %7 %70 +%151 = OpVectorShuffle %6 %150 %150 2 3 +%152 = OpDot %5 %149 %151 +%153 = OpCompositeConstruct %4 %142 %147 %152 +%154 = OpCompositeConstruct %4 %85 %85 %85 +%155 = OpFSub %4 %154 %153 +OpLine %3 21 24 +%156 = OpCompositeConstruct %4 %84 %84 %84 +%157 = OpExtInst %4 %1 FMax %155 %156 +OpLine %3 21 5 +OpStore %73 %157 +OpLine %3 22 9 +%158 = OpLoad %4 %73 +%159 = OpLoad %4 %73 +%160 = OpFMul %4 %158 %159 +OpLine %3 22 5 +OpStore %73 %160 +OpLine %3 23 9 +%161 = OpLoad %4 %73 +%162 = OpLoad %4 %73 +%163 = OpFMul %4 %161 %162 +OpLine %3 23 5 +OpStore %73 %163 +OpLine %3 24 13 +%164 = OpVectorShuffle %4 %91 %91 3 3 3 +%165 = OpFMul %4 %141 %164 +%166 = OpExtInst %4 %1 Fract %165 +%167 = OpVectorTimesScalar %4 %166 %86 +OpLine %3 24 13 +%168 = OpCompositeConstruct %4 %56 %56 %56 +%169 = OpFSub %4 %167 %168 +OpLine %3 25 13 +%170 = OpExtInst %4 %1 FAbs %169 +OpLine %3 25 13 +%171 = OpCompositeConstruct %4 %85 %85 %85 +%172 = OpFSub %4 %170 %171 +OpLine %3 26 14 +%173 = OpCompositeConstruct %4 %85 %85 %85 +%174 = OpFAdd %4 %169 %173 +%175 = OpExtInst %4 %1 Floor %174 +OpLine %3 27 14 +%176 = OpFSub %4 %169 %175 +OpLine %3 1 1 +%177 = OpLoad %4 %73 +OpLine %3 28 9 +%178 = OpFMul %4 %176 %176 +%179 = OpFMul %4 %172 %172 +%180 = OpFAdd %4 %178 %179 +%181 = OpVectorTimesScalar %4 %180 %88 +%182 = OpCompositeConstruct %4 %87 %87 %87 +%183 = OpFSub %4 %182 %181 +%184 = OpFMul %4 %177 %183 +OpLine %3 28 5 +OpStore %73 %184 +OpLine %3 29 13 +%185 = OpCompositeExtract %5 %176 0 +%186 = OpCompositeExtract %5 %103 0 +%187 = OpFMul %5 %185 %186 +%188 = OpCompositeExtract %5 %172 0 +%189 = OpCompositeExtract %5 %103 1 +%190 = OpFMul %5 %188 %189 +%191 = OpFAdd %5 %187 %190 +%192 = OpVectorShuffle %6 %176 %176 1 2 +%193 = OpLoad %7 %70 +%194 = OpVectorShuffle %6 %193 %193 0 2 +%195 = OpFMul %6 %192 %194 +%196 = OpVectorShuffle %6 %172 %172 1 2 +%197 = OpLoad %7 %70 +%198 = OpVectorShuffle %6 %197 %197 1 3 +%199 = OpFMul %6 %196 %198 +%200 = OpFAdd %6 %195 %199 +%201 = OpCompositeConstruct %4 %191 %200 +OpLine %3 30 12 +%202 = OpLoad %4 %73 +%203 = OpDot %5 %202 %201 +%204 = OpFMul %5 %89 %203 +OpReturnValue %204 +OpFunctionEnd +%217 = OpFunction %5 None %79 +%216 = OpFunctionParameter %6 +%215 = OpLabel +%207 = OpVariable %208 Function %209 +%212 = OpVariable %213 Function %214 +%205 = OpVariable %66 Function %206 +%210 = OpVariable %208 Function %211 +OpBranch %221 +%221 = OpLabel +OpLine %3 36 13 +%222 = OpVectorTimesScalar %6 %216 %219 +OpLine %3 36 5 +OpStore %205 %222 +OpLine %3 37 5 +OpStore %207 %84 +OpLine %3 38 5 +OpStore %210 %85 +OpLine %3 39 17 +%223 = OpCompositeConstruct %6 %220 %220 +OpLine %3 40 24 +%224 = OpExtInst %5 %1 Cos %85 +OpLine %3 40 14 +%225 = OpExtInst %5 %1 Sin %85 +%226 = OpCompositeConstruct %6 %224 %225 +OpLine %3 41 15 +%227 = OpCompositeExtract %5 %226 0 +%228 = OpCompositeExtract %5 %226 1 +%229 = OpCompositeExtract %5 %226 1 +%230 = OpFNegate %5 %229 +%231 = OpCompositeExtract %5 %226 0 +%232 = OpCompositeConstruct %6 %227 %228 +%233 = OpCompositeConstruct %6 %230 %231 +%234 = OpCompositeConstruct %9 %232 %233 +OpLine %3 43 10 +OpStore %212 %132 +OpBranch %235 +%235 = OpLabel +OpLine %3 43 5 +OpLoopMerge %236 %238 None +OpBranch %237 +%237 = OpLabel +OpLine %3 43 22 +%239 = OpLoad %8 %212 +%240 = OpULessThan %108 %239 %218 +OpLine %3 43 21 +OpSelectionMerge %241 None +OpBranchConditional %240 %241 %242 +%242 = OpLabel +OpBranch %236 +%241 = OpLabel +OpBranch %243 +%243 = OpLabel +OpLine %3 1 1 +%245 = OpLoad %5 %207 +%246 = OpLoad %5 %210 +%247 = OpLoad %6 %205 +OpLine %3 44 21 +%248 = OpFunctionCall %5 %78 %247 +OpLine %3 44 13 +%249 = OpFMul %5 %246 %248 +%250 = OpFAdd %5 %245 %249 +OpLine %3 44 9 +OpStore %207 %250 +OpLine %3 45 13 +%251 = OpLoad %6 %205 +%252 = OpMatrixTimesVector %6 %234 %251 +OpLine %3 45 13 +%253 = OpVectorTimesScalar %6 %252 %86 +%254 = OpFAdd %6 %253 %223 +OpLine %3 45 9 +OpStore %205 %254 +OpLine %3 1 1 +%255 = OpLoad %5 %210 +OpLine %3 46 13 +%256 = OpFMul %5 %255 %85 +OpLine %3 46 9 +OpStore %210 %256 +OpBranch %244 +%244 = OpLabel +OpBranch %238 +%238 = OpLabel +OpLine %3 1 1 +%257 = OpLoad %8 %212 +OpLine %3 43 43 +%258 = OpIAdd %8 %257 %123 +OpLine %3 43 39 +OpStore %212 %258 +OpBranch %235 +%236 = OpLabel +OpLine %3 1 1 +%259 = OpLoad %5 %207 +OpReturnValue %259 +OpFunctionEnd +%263 = OpFunction %4 None %264 +%261 = OpFunctionParameter %6 +%262 = OpFunctionParameter %6 +%260 = OpLabel +OpBranch %265 +%265 = OpLabel +OpLine %3 77 9 +%266 = OpCompositeExtract %5 %261 0 +%267 = OpCompositeExtract %5 %262 0 +%268 = OpCompositeExtract %5 %262 1 +OpLine %3 78 49 +%269 = OpFunctionCall %5 %217 %261 +OpLine %3 76 12 +%270 = OpExtInst %5 %1 FMix %267 %268 %269 +%271 = OpCompositeExtract %5 %261 1 +%272 = OpCompositeConstruct %4 %266 %270 %271 +OpReturnValue %272 +OpFunctionEnd +%276 = OpFunction %14 None %277 +%274 = OpFunctionParameter %6 +%275 = OpFunctionParameter %6 +%273 = OpLabel +OpBranch %280 +%280 = OpLabel +OpLine %3 84 13 +%281 = OpFunctionCall %4 %263 %274 %275 +OpLine %3 86 29 +%282 = OpCompositeConstruct %6 %278 %84 +%283 = OpFAdd %6 %274 %282 +OpLine %3 86 15 +%284 = OpFunctionCall %4 %263 %283 %275 +OpLine %3 86 15 +%285 = OpFSub %4 %284 %281 +OpLine %3 87 29 +%286 = OpCompositeConstruct %6 %84 %278 +%287 = OpFAdd %6 %274 %286 +OpLine %3 87 15 +%288 = OpFunctionCall %4 %263 %287 %275 +OpLine %3 87 15 +%289 = OpFSub %4 %288 %281 +OpLine %3 88 29 +%290 = OpCompositeConstruct %6 %279 %84 +%291 = OpFAdd %6 %274 %290 +OpLine %3 88 15 +%292 = OpFunctionCall %4 %263 %291 %275 +OpLine %3 88 15 +%293 = OpFSub %4 %292 %281 +OpLine %3 89 29 +%294 = OpCompositeConstruct %6 %84 %279 +%295 = OpFAdd %6 %274 %294 +OpLine %3 89 15 +%296 = OpFunctionCall %4 %263 %295 %275 +OpLine %3 89 15 +%297 = OpFSub %4 %296 %281 +OpLine %3 91 14 +%298 = OpExtInst %4 %1 Cross %289 %285 +%299 = OpExtInst %4 %1 Normalize %298 +OpLine %3 92 14 +%300 = OpExtInst %4 %1 Cross %297 %293 +%301 = OpExtInst %4 %1 Normalize %300 +OpLine %3 94 14 +%302 = OpFAdd %4 %299 %301 +OpLine %3 94 13 +%303 = OpVectorTimesScalar %4 %302 %85 +OpLine %3 96 12 +%304 = OpCompositeConstruct %14 %281 %303 +OpReturnValue %304 +OpFunctionEnd +%309 = OpFunction %6 None %310 +%306 = OpFunctionParameter %8 +%307 = OpFunctionParameter %10 +%308 = OpFunctionParameter %11 +%305 = OpLabel +OpBranch %311 +%311 = OpLabel +OpLine %3 101 9 +%312 = OpConvertUToF %5 %306 +%313 = OpCompositeExtract %8 %307 0 +OpLine %3 101 9 +%314 = OpIAdd %8 %313 %123 +%315 = OpConvertUToF %5 %314 +%316 = OpFRem %5 %312 %315 +%317 = OpCompositeExtract %8 %307 0 +OpLine %3 100 12 +%318 = OpIAdd %8 %317 %123 +%319 = OpUDiv %8 %306 %318 +%320 = OpConvertUToF %5 %319 +%321 = OpCompositeConstruct %6 %316 %320 +%322 = OpConvertSToF %6 %308 +%323 = OpFAdd %6 %321 %322 +OpReturnValue %323 +OpFunctionEnd +%326 = OpFunction %4 None %327 +%325 = OpFunctionParameter %6 +%324 = OpLabel +OpBranch %332 +%332 = OpLabel +OpLine %3 270 9 +%333 = OpFunctionCall %5 %78 %325 +OpLine %3 270 9 +%334 = OpFMul %5 %333 %85 +OpLine %3 270 9 +%335 = OpFAdd %5 %334 %85 +OpLine %3 271 17 +%336 = OpCompositeConstruct %6 %328 %329 +%337 = OpFAdd %6 %325 %336 +OpLine %3 271 9 +%338 = OpFunctionCall %5 %78 %337 +OpLine %3 271 9 +%339 = OpFMul %5 %338 %85 +OpLine %3 271 9 +%340 = OpFAdd %5 %339 %85 +OpLine %3 272 17 +%341 = OpCompositeConstruct %6 %330 %331 +%342 = OpFAdd %6 %325 %341 +OpLine %3 272 9 +%343 = OpFunctionCall %5 %78 %342 +OpLine %3 272 9 +%344 = OpFMul %5 %343 %85 +OpLine %3 269 12 +%345 = OpFAdd %5 %344 %85 +%346 = OpCompositeConstruct %4 %335 %340 %345 +OpReturnValue %346 +OpFunctionEnd +%351 = OpFunction %2 None %352 +%347 = OpLabel +%350 = OpLoad %19 %348 +%354 = OpAccessChain %353 %29 %132 +OpBranch %359 +%359 = OpLabel +OpLine %3 111 22 +%360 = OpCompositeExtract %8 %350 0 +OpLine %3 113 36 +%362 = OpAccessChain %361 %354 %132 +%363 = OpLoad %10 %362 +OpLine %3 113 59 +%365 = OpAccessChain %364 %354 %123 +%366 = OpLoad %11 %365 +OpLine %3 113 13 +%367 = OpFunctionCall %6 %309 %360 %363 %366 +OpLine %3 115 5 +OpLine %3 115 51 +%371 = OpAccessChain %370 %354 %356 +%372 = OpLoad %6 %371 +OpLine %3 115 33 +%373 = OpFunctionCall %14 %276 %367 %372 +OpLine %3 115 5 +%374 = OpAccessChain %369 %32 %132 %360 +OpStore %374 %373 +OpLine %3 118 23 +%375 = OpCompositeExtract %8 %350 0 +OpLine %3 118 23 +%376 = OpIMul %8 %375 %355 +OpLine %3 120 25 +%378 = OpAccessChain %377 %354 %132 %132 +%379 = OpLoad %8 %378 +OpLine %3 120 25 +%380 = OpAccessChain %377 %354 %132 %123 +%381 = OpLoad %8 %380 +%382 = OpIMul %8 %379 %381 +OpLine %3 120 9 +%383 = OpIMul %8 %382 %355 +%384 = OpUGreaterThanEqual %108 %376 %383 +OpLine %3 120 5 +OpSelectionMerge %385 None +OpBranchConditional %384 %386 %385 +%386 = OpLabel +OpReturn +%385 = OpLabel +OpLine %3 122 28 +%387 = OpCompositeExtract %8 %350 0 +OpLine %3 122 15 +%388 = OpAccessChain %377 %354 %132 %132 +%389 = OpLoad %8 %388 +%390 = OpUDiv %8 %387 %389 +%391 = OpIAdd %8 %360 %390 +OpLine %3 123 15 +%392 = OpIAdd %8 %391 %123 +OpLine %3 124 15 +%393 = OpAccessChain %377 %354 %132 %132 +%394 = OpLoad %8 %393 +%395 = OpIAdd %8 %391 %394 +OpLine %3 124 15 +%396 = OpIAdd %8 %395 %123 +OpLine %3 125 15 +%397 = OpIAdd %8 %396 %123 +OpLine %3 127 5 +OpLine %3 127 5 +%400 = OpAccessChain %399 %34 %132 %376 +OpStore %400 %391 +OpLine %3 128 5 +OpLine %3 128 5 +%401 = OpIAdd %8 %376 %123 +OpLine %3 128 5 +%402 = OpAccessChain %399 %34 %132 %401 +OpStore %402 %396 +OpLine %3 129 5 +OpLine %3 129 5 +%403 = OpIAdd %8 %376 %356 +OpLine %3 129 5 +%404 = OpAccessChain %399 %34 %132 %403 +OpStore %404 %397 +OpLine %3 130 5 +OpLine %3 130 5 +%405 = OpIAdd %8 %376 %357 +OpLine %3 130 5 +%406 = OpAccessChain %399 %34 %132 %405 +OpStore %406 %391 +OpLine %3 131 5 +OpLine %3 131 5 +%407 = OpIAdd %8 %376 %358 +OpLine %3 131 5 +%408 = OpAccessChain %399 %34 %132 %407 +OpStore %408 %397 +OpLine %3 132 5 +OpLine %3 132 5 +%409 = OpIAdd %8 %376 %218 +OpLine %3 132 5 +%410 = OpAccessChain %399 %34 %132 %409 +OpStore %410 %392 +OpReturn +OpFunctionEnd +%421 = OpFunction %2 None %352 +%411 = OpLabel +%414 = OpLoad %8 %412 +%423 = OpAccessChain %422 %36 %132 +OpBranch %425 +%425 = OpLabel +OpLine %3 161 19 +%426 = OpIAdd %8 %414 %356 +OpLine %3 161 18 +%427 = OpUDiv %8 %426 %357 +OpLine %3 161 13 +%428 = OpUMod %8 %427 %356 +%429 = OpConvertUToF %5 %428 +OpLine %3 162 19 +%430 = OpIAdd %8 %414 %123 +OpLine %3 162 18 +%431 = OpUDiv %8 %430 %357 +OpLine %3 162 13 +%432 = OpUMod %8 %431 %356 +%433 = OpConvertUToF %5 %432 +OpLine %3 163 14 +%434 = OpCompositeConstruct %6 %429 %433 +OpLine %3 165 30 +%435 = OpVectorTimesScalar %6 %434 %86 +%436 = OpCompositeConstruct %6 %424 %424 +%437 = OpFAdd %6 %436 %435 +OpLine %3 165 20 +%438 = OpCompositeConstruct %7 %437 %84 %56 +OpLine %3 168 21 +%439 = OpCompositeExtract %5 %434 0 +OpLine %3 168 21 +%441 = OpAccessChain %440 %423 %357 +%442 = OpLoad %8 %441 +%443 = OpConvertUToF %5 %442 +%444 = OpFMul %5 %439 %443 +%445 = OpCompositeExtract %5 %434 1 +OpLine %3 168 17 +%446 = OpAccessChain %440 %423 %357 +%447 = OpLoad %8 %446 +%448 = OpConvertUToF %5 %447 +%449 = OpFMul %5 %445 %448 +%450 = OpFAdd %5 %444 %449 +%451 = OpConvertFToU %8 %450 +OpLine %3 168 17 +%452 = OpAccessChain %440 %423 %358 +%453 = OpLoad %8 %452 +%454 = OpIAdd %8 %451 %453 +OpLine %3 170 12 +%455 = OpCompositeConstruct %21 %454 %438 %434 +%456 = OpCompositeExtract %8 %455 0 +OpStore %415 %456 +%457 = OpCompositeExtract %7 %455 1 +OpStore %417 %457 +%458 = OpCompositeExtract %6 %455 2 +OpStore %419 %458 +OpReturn +OpFunctionEnd +%475 = OpFunction %2 None %352 +%463 = OpLabel +%459 = OpVariable %208 Function %460 +%461 = OpVariable %213 Function %462 +%466 = OpLoad %8 %465 +%469 = OpLoad %7 %467 +%472 = OpLoad %6 %470 +%464 = OpCompositeConstruct %21 %466 %469 %472 +%476 = OpAccessChain %422 %36 %132 +OpBranch %478 +%478 = OpLabel +OpLine %3 181 17 +%479 = OpCompositeExtract %6 %464 2 +%480 = OpCompositeExtract %5 %479 0 +OpLine %3 181 17 +%481 = OpAccessChain %440 %476 %357 +%482 = OpLoad %8 %481 +%483 = OpConvertUToF %5 %482 +%484 = OpFMul %5 %480 %483 +%485 = OpCompositeExtract %6 %464 2 +%486 = OpCompositeExtract %5 %485 1 +OpLine %3 181 70 +%487 = OpAccessChain %440 %476 %357 +%488 = OpLoad %8 %487 +OpLine %3 181 13 +%489 = OpAccessChain %440 %476 %357 +%490 = OpLoad %8 %489 +%491 = OpIMul %8 %488 %490 +%492 = OpConvertUToF %5 %491 +%493 = OpFMul %5 %486 %492 +%494 = OpFAdd %5 %484 %493 +%495 = OpConvertFToU %8 %494 +OpLine %3 181 13 +%496 = OpAccessChain %440 %476 %358 +%497 = OpLoad %8 %496 +%498 = OpIAdd %8 %495 %497 +OpLine %3 182 32 +%499 = OpConvertUToF %5 %498 +OpLine %3 182 22 +%500 = OpFDiv %5 %499 %477 +%501 = OpExtInst %5 %1 Floor %500 +%502 = OpConvertFToU %8 %501 +OpLine %3 183 22 +%503 = OpUMod %8 %498 %355 +OpLine %3 185 36 +%504 = OpAccessChain %361 %476 %132 +%505 = OpLoad %10 %504 +OpLine %3 185 57 +%506 = OpAccessChain %364 %476 %123 +%507 = OpLoad %11 %506 +OpLine %3 185 13 +%508 = OpFunctionCall %6 %309 %502 %505 %507 +OpLine %3 186 31 +%509 = OpAccessChain %370 %476 %356 +%510 = OpLoad %6 %509 +OpLine %3 186 13 +%511 = OpFunctionCall %14 %276 %508 %510 +OpLine %3 188 5 +OpStore %459 %84 +OpLine %3 190 5 +OpSelectionMerge %512 None +OpSwitch %503 %519 0 %513 1 %514 2 %515 3 %516 4 %517 5 %518 +%513 = OpLabel +OpLine %3 191 37 +%520 = OpCompositeExtract %4 %511 0 +%521 = OpCompositeExtract %5 %520 0 +OpLine %3 191 20 +OpStore %459 %521 +OpBranch %512 +%514 = OpLabel +OpLine %3 192 37 +%522 = OpCompositeExtract %4 %511 0 +%523 = OpCompositeExtract %5 %522 1 +OpLine %3 192 20 +OpStore %459 %523 +OpBranch %512 +%515 = OpLabel +OpLine %3 193 37 +%524 = OpCompositeExtract %4 %511 0 +%525 = OpCompositeExtract %5 %524 2 +OpLine %3 193 20 +OpStore %459 %525 +OpBranch %512 +%516 = OpLabel +OpLine %3 194 37 +%526 = OpCompositeExtract %4 %511 1 +%527 = OpCompositeExtract %5 %526 0 +OpLine %3 194 20 +OpStore %459 %527 +OpBranch %512 +%517 = OpLabel +OpLine %3 195 37 +%528 = OpCompositeExtract %4 %511 1 +%529 = OpCompositeExtract %5 %528 1 +OpLine %3 195 20 +OpStore %459 %529 +OpBranch %512 +%518 = OpLabel +OpLine %3 196 37 +%530 = OpCompositeExtract %4 %511 1 +%531 = OpCompositeExtract %5 %530 2 +OpLine %3 196 20 +OpStore %459 %531 +OpBranch %512 +%519 = OpLabel +OpBranch %512 +%512 = OpLabel +OpLine %3 200 15 +%532 = OpAccessChain %377 %476 %132 %132 +%533 = OpLoad %8 %532 +%534 = OpUDiv %8 %502 %533 +%535 = OpIAdd %8 %502 %534 +OpLine %3 201 15 +%536 = OpIAdd %8 %535 %123 +OpLine %3 202 15 +%537 = OpAccessChain %377 %476 %132 %132 +%538 = OpLoad %8 %537 +%539 = OpIAdd %8 %535 %538 +OpLine %3 202 15 +%540 = OpIAdd %8 %539 %123 +OpLine %3 203 15 +%541 = OpIAdd %8 %540 %123 +OpLine %3 205 5 +OpStore %461 %132 +OpLine %3 206 5 +OpSelectionMerge %542 None +OpSwitch %503 %547 0 %543 3 %543 2 %544 4 %544 1 %545 5 %546 +%543 = OpLabel +OpLine %3 207 24 +OpStore %461 %535 +OpBranch %542 +%544 = OpLabel +OpLine %3 208 24 +OpStore %461 %541 +OpBranch %542 +%545 = OpLabel +OpLine %3 209 20 +OpStore %461 %540 +OpBranch %542 +%546 = OpLabel +OpLine %3 210 20 +OpStore %461 %536 +OpBranch %542 +%547 = OpLabel +OpBranch %542 +%542 = OpLabel +OpLine %3 213 13 +%548 = OpCompositeExtract %8 %464 0 +OpLine %3 213 5 +OpStore %461 %548 +OpLine %3 222 27 +%549 = OpLoad %5 %459 +%550 = OpBitcast %8 %549 +OpLine %3 223 12 +%551 = OpLoad %8 %461 +%552 = OpCompositeConstruct %22 %550 %551 +%553 = OpCompositeExtract %8 %552 0 +OpStore %473 %553 +%554 = OpCompositeExtract %8 %552 1 +OpStore %474 %554 +OpReturn +OpFunctionEnd +%566 = OpFunction %2 None %352 +%555 = OpLabel +%559 = OpLoad %4 %557 +%561 = OpLoad %4 %560 +%556 = OpCompositeConstruct %14 %559 %561 +%568 = OpAccessChain %567 %39 %132 +OpBranch %569 +%569 = OpLabel +OpLine %3 254 25 +%571 = OpAccessChain %570 %568 %123 +%572 = OpLoad %23 %571 +%573 = OpCompositeExtract %4 %556 0 +OpLine %3 254 25 +%574 = OpCompositeConstruct %7 %573 %56 +%575 = OpMatrixTimesVector %7 %572 %574 +OpLine %3 255 18 +%576 = OpCompositeExtract %4 %556 1 +OpLine %3 256 12 +%577 = OpCompositeExtract %4 %556 0 +%578 = OpCompositeConstruct %26 %575 %576 %577 +%579 = OpCompositeExtract %7 %578 0 +OpStore %562 %579 +%580 = OpCompositeExtract %4 %578 1 +OpStore %563 %580 +%581 = OpCompositeExtract %4 %578 2 +OpStore %565 %581 +OpReturn +OpFunctionEnd +%593 = OpFunction %2 None %352 +%584 = OpLabel +%582 = OpVariable %74 Function %583 +%587 = OpLoad %7 %586 +%589 = OpLoad %4 %588 +%591 = OpLoad %4 %590 +%585 = OpCompositeConstruct %26 %587 %589 %591 +%594 = OpAccessChain %567 %39 %132 +%596 = OpAccessChain %595 %42 %132 +OpBranch %599 +%599 = OpLabel +OpLine %3 278 28 +%600 = OpCompositeConstruct %4 %84 %84 %84 +OpLine %3 278 17 +%601 = OpCompositeConstruct %4 %278 %278 %278 +%602 = OpCompositeExtract %4 %585 2 +%603 = OpExtInst %4 %1 Fract %602 +%604 = OpExtInst %4 %1 SmoothStep %600 %601 %603 +OpLine %3 278 5 +OpStore %582 %604 +OpLine %3 279 17 +%605 = OpCompositeConstruct %4 %85 %278 %597 +OpLine %3 279 13 +%606 = OpCompositeConstruct %4 %598 %598 %598 +%607 = OpAccessChain %122 %582 %132 +%608 = OpLoad %5 %607 +%609 = OpAccessChain %122 %582 %123 +%610 = OpLoad %5 %609 +%611 = OpFMul %5 %608 %610 +%612 = OpAccessChain %122 %582 %356 +%613 = OpLoad %5 %612 +%614 = OpFMul %5 %611 %613 +%615 = OpCompositeConstruct %4 %614 %614 %614 +%616 = OpExtInst %4 %1 FMix %605 %606 %615 +OpLine %3 279 5 +OpStore %582 %616 +OpLine %3 282 25 +%618 = OpAccessChain %617 %596 %123 +%619 = OpLoad %4 %618 +%620 = OpVectorTimesScalar %4 %619 %278 +OpLine %3 284 21 +%621 = OpAccessChain %617 %596 %132 +%622 = OpLoad %4 %621 +%623 = OpCompositeExtract %4 %585 2 +%624 = OpFSub %4 %622 %623 +%625 = OpExtInst %4 %1 Normalize %624 +OpLine %3 285 20 +%627 = OpAccessChain %626 %594 %132 +%628 = OpLoad %7 %627 +%629 = OpVectorShuffle %4 %628 %628 0 1 2 +%630 = OpCompositeExtract %4 %585 2 +%631 = OpFSub %4 %629 %630 +%632 = OpExtInst %4 %1 Normalize %631 +OpLine %3 286 20 +%633 = OpFAdd %4 %632 %625 +%634 = OpExtInst %4 %1 Normalize %633 +OpLine %3 288 32 +%635 = OpCompositeExtract %4 %585 1 +%636 = OpDot %5 %635 %625 +OpLine %3 288 28 +%637 = OpExtInst %5 %1 FMax %636 %84 +OpLine %3 289 25 +%638 = OpAccessChain %617 %596 %123 +%639 = OpLoad %4 %638 +%640 = OpVectorTimesScalar %4 %639 %637 +OpLine %3 291 37 +%641 = OpCompositeExtract %4 %585 1 +%642 = OpDot %5 %641 %634 +OpLine %3 291 33 +%643 = OpExtInst %5 %1 FMax %642 %84 +OpLine %3 291 29 +%644 = OpExtInst %5 %1 Pow %643 %329 +OpLine %3 292 26 +%645 = OpAccessChain %617 %596 %123 +%646 = OpLoad %4 %645 +%647 = OpVectorTimesScalar %4 %646 %644 +OpLine %3 294 18 +%648 = OpFAdd %4 %620 %640 +%649 = OpFAdd %4 %648 %647 +%650 = OpLoad %4 %582 +%651 = OpFMul %4 %649 %650 +OpLine %3 296 12 +%652 = OpCompositeConstruct %7 %651 %56 +OpStore %592 %652 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/tests/out/spv/image.spvasm b/tests/out/spv/image.spvasm index 6b94b3e27a..4f9dd442e0 100644 --- a/tests/out/spv/image.spvasm +++ b/tests/out/spv/image.spvasm @@ -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" diff --git a/tests/out/spv/interpolate.spvasm b/tests/out/spv/interpolate.spvasm index b9df1e0593..bcfad15c14 100644 --- a/tests/out/spv/interpolate.spvasm +++ b/tests/out/spv/interpolate.spvasm @@ -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" diff --git a/tests/out/spv/padding.spvasm b/tests/out/spv/padding.spvasm index 485ac93671..0efcef10e4 100644 --- a/tests/out/spv/padding.spvasm +++ b/tests/out/spv/padding.spvasm @@ -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" diff --git a/tests/out/spv/pointers.spvasm b/tests/out/spv/pointers.spvasm index 71a680aa73..3debe27b16 100644 --- a/tests/out/spv/pointers.spvasm +++ b/tests/out/spv/pointers.spvasm @@ -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" diff --git a/tests/out/spv/policy-mix.spvasm b/tests/out/spv/policy-mix.spvasm index d117476050..970da7d103 100644 --- a/tests/out/spv/policy-mix.spvasm +++ b/tests/out/spv/policy-mix.spvasm @@ -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" diff --git a/tests/out/spv/quad.spvasm b/tests/out/spv/quad.spvasm index 9d4ba02268..4c82c92b07 100644 --- a/tests/out/spv/quad.spvasm +++ b/tests/out/spv/quad.spvasm @@ -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" diff --git a/tests/out/spv/shadow.spvasm b/tests/out/spv/shadow.spvasm index d087a25c99..26ee96559b 100644 --- a/tests/out/spv/shadow.spvasm +++ b/tests/out/spv/shadow.spvasm @@ -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" diff --git a/tests/out/spv/texture-arg.spvasm b/tests/out/spv/texture-arg.spvasm index f761edf70d..485b2d7d9a 100644 --- a/tests/out/spv/texture-arg.spvasm +++ b/tests/out/spv/texture-arg.spvasm @@ -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" diff --git a/tests/out/spv/workgroup-var-init.spvasm b/tests/out/spv/workgroup-var-init.spvasm index 122a61283c..e95a11aa90 100644 --- a/tests/out/spv/workgroup-var-init.spvasm +++ b/tests/out/spv/workgroup-var-init.spvasm @@ -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" diff --git a/tests/snapshots.rs b/tests/snapshots.rs index eab1aaae6f..658e5fe7f3 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -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, ¶ms.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, 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")] diff --git a/tests/spirv-capabilities.rs b/tests/spirv-capabilities.rs index f269182859..35f24b7d69 100644 --- a/tests/spirv-capabilities.rs +++ b/tests/spirv-capabilities.rs @@ -23,7 +23,9 @@ fn capabilities_used(source: &str) -> naga::FastIndexSet { 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() }