From 30539d72df83949bcb0c8b1987d0ff7b6378c764 Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Mon, 2 Oct 2023 14:51:27 +0200 Subject: [PATCH 1/3] add test --- tests/in/separate-entry-points.param.ron | 6 +++ tests/in/separate-entry-points.wgsl | 23 +++++++++ ...separate-entry-points.compute.Compute.glsl | 27 +++++++++++ ...parate-entry-points.fragment.Fragment.glsl | 27 +++++++++++ .../spv/separate-entry-points.compute.spvasm | 43 +++++++++++++++++ .../spv/separate-entry-points.fragment.spvasm | 48 +++++++++++++++++++ tests/snapshots.rs | 1 + 7 files changed, 175 insertions(+) create mode 100644 tests/in/separate-entry-points.param.ron create mode 100644 tests/in/separate-entry-points.wgsl create mode 100644 tests/out/glsl/separate-entry-points.compute.Compute.glsl create mode 100644 tests/out/glsl/separate-entry-points.fragment.Fragment.glsl create mode 100644 tests/out/spv/separate-entry-points.compute.spvasm create mode 100644 tests/out/spv/separate-entry-points.fragment.spvasm diff --git a/tests/in/separate-entry-points.param.ron b/tests/in/separate-entry-points.param.ron new file mode 100644 index 0000000000..af0931c111 --- /dev/null +++ b/tests/in/separate-entry-points.param.ron @@ -0,0 +1,6 @@ +( + spv: ( + version: (1, 0), + separate_entry_points: true, + ), +) diff --git a/tests/in/separate-entry-points.wgsl b/tests/in/separate-entry-points.wgsl new file mode 100644 index 0000000000..a7ec3b083a --- /dev/null +++ b/tests/in/separate-entry-points.wgsl @@ -0,0 +1,23 @@ +// only available in the fragment stage +fn derivatives() { + let x = dpdx(0.0); + let y = dpdy(0.0); + let width = fwidth(0.0); +} + +// only available in the compute stage +fn barriers() { + storageBarrier(); + workgroupBarrier(); +} + +@fragment +fn fragment() -> @location(0) vec4 { + derivatives(); + return vec4(); +} + +@compute @workgroup_size(1) +fn compute() { + barriers(); +} \ No newline at end of file diff --git a/tests/out/glsl/separate-entry-points.compute.Compute.glsl b/tests/out/glsl/separate-entry-points.compute.Compute.glsl new file mode 100644 index 0000000000..2a439cdb41 --- /dev/null +++ b/tests/out/glsl/separate-entry-points.compute.Compute.glsl @@ -0,0 +1,27 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + + +void derivatives() { + float x = dFdx(0.0); + float y = dFdy(0.0); + float width = fwidth(0.0); +} + +void barriers() { + memoryBarrierBuffer(); + barrier(); + memoryBarrierShared(); + barrier(); + return; +} + +void main() { + barriers(); + return; +} + diff --git a/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl b/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl new file mode 100644 index 0000000000..71cef76693 --- /dev/null +++ b/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl @@ -0,0 +1,27 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(location = 0) out vec4 _fs2p_location0; + +void derivatives() { + float x = dFdx(0.0); + float y = dFdy(0.0); + float width = fwidth(0.0); +} + +void barriers() { + memoryBarrierBuffer(); + barrier(); + memoryBarrierShared(); + barrier(); + return; +} + +void main() { + derivatives(); + _fs2p_location0 = vec4(0.0); + return; +} + diff --git a/tests/out/spv/separate-entry-points.compute.spvasm b/tests/out/spv/separate-entry-points.compute.spvasm new file mode 100644 index 0000000000..1aeafda00c --- /dev/null +++ b/tests/out/spv/separate-entry-points.compute.spvasm @@ -0,0 +1,43 @@ +; SPIR-V +; Version: 1.0 +; Generator: rspirv +; Bound: 25 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %22 "compute" +OpExecutionMode %22 LocalSize 1 1 1 +%2 = OpTypeVoid +%4 = OpTypeFloat 32 +%3 = OpTypeVector %4 4 +%7 = OpTypeFunction %2 +%8 = OpConstant %4 0.0 +%17 = OpTypeInt 32 0 +%16 = OpConstant %17 2 +%18 = OpConstant %17 1 +%19 = OpConstant %17 72 +%20 = OpConstant %17 264 +%6 = OpFunction %2 None %7 +%5 = OpLabel +OpBranch %9 +%9 = OpLabel +%10 = OpDPdx %4 %8 +%11 = OpDPdy %4 %8 +%12 = OpFwidth %4 %8 +OpReturn +OpFunctionEnd +%14 = OpFunction %2 None %7 +%13 = OpLabel +OpBranch %15 +%15 = OpLabel +OpControlBarrier %16 %18 %19 +OpControlBarrier %16 %16 %20 +OpReturn +OpFunctionEnd +%22 = OpFunction %2 None %7 +%21 = OpLabel +OpBranch %23 +%23 = OpLabel +%24 = OpFunctionCall %2 %14 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/tests/out/spv/separate-entry-points.fragment.spvasm b/tests/out/spv/separate-entry-points.fragment.spvasm new file mode 100644 index 0000000000..215f9b21ca --- /dev/null +++ b/tests/out/spv/separate-entry-points.fragment.spvasm @@ -0,0 +1,48 @@ +; SPIR-V +; Version: 1.0 +; Generator: rspirv +; Bound: 28 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint Fragment %24 "fragment" %22 +OpExecutionMode %24 OriginUpperLeft +OpDecorate %22 Location 0 +%2 = OpTypeVoid +%4 = OpTypeFloat 32 +%3 = OpTypeVector %4 4 +%7 = OpTypeFunction %2 +%8 = OpConstant %4 0.0 +%17 = OpTypeInt 32 0 +%16 = OpConstant %17 2 +%18 = OpConstant %17 1 +%19 = OpConstant %17 72 +%20 = OpConstant %17 264 +%23 = OpTypePointer Output %3 +%22 = OpVariable %23 Output +%25 = OpConstantNull %3 +%6 = OpFunction %2 None %7 +%5 = OpLabel +OpBranch %9 +%9 = OpLabel +%10 = OpDPdx %4 %8 +%11 = OpDPdy %4 %8 +%12 = OpFwidth %4 %8 +OpReturn +OpFunctionEnd +%14 = OpFunction %2 None %7 +%13 = OpLabel +OpBranch %15 +%15 = OpLabel +OpControlBarrier %16 %18 %19 +OpControlBarrier %16 %16 %20 +OpReturn +OpFunctionEnd +%24 = OpFunction %2 None %7 +%21 = OpLabel +OpBranch %26 +%26 = OpLabel +%27 = OpFunctionCall %2 %6 +OpStore %22 %25 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 2bc7f45444..c3455dd864 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -781,6 +781,7 @@ fn convert_wgsl() { "const-exprs", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ), + ("separate-entry-points", Targets::SPIRV | Targets::GLSL), ]; for &(name, targets) in inputs.iter() { From bc732798291318aa377214891bf14be8d165e315 Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Tue, 3 Oct 2023 12:43:55 +0200 Subject: [PATCH 2/3] Skip functions that that are not compatible with entry point's stage --- src/back/glsl/mod.rs | 5 +++ src/back/spv/writer.rs | 5 +++ ...separate-entry-points.compute.Compute.glsl | 6 --- ...parate-entry-points.fragment.Fragment.glsl | 8 ---- .../spv/separate-entry-points.compute.spvasm | 44 +++++++------------ .../spv/separate-entry-points.fragment.spvasm | 37 +++++----------- 6 files changed, 39 insertions(+), 66 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 77e206d3c6..db3b88f15c 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -827,6 +827,11 @@ impl<'a, W: Write> Writer<'a, W> { let fun_info = &self.info[handle]; + // Skip functions that that are not compatible with this entry point's stage + if !fun_info.available_stages.contains(ep_info.available_stages) { + continue; + } + // Write the function self.write_function(back::FunctionType::Function(handle), function, fun_info)?; diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index 2e133985af..8192b6d219 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -1952,6 +1952,11 @@ impl Writer { log::info!("Skip function {:?}", ir_function.name); continue; } + + // Skip functions that that are not compatible with this entry point's stage + if !info.available_stages.contains(ep_info.available_stages) { + continue; + } } let id = self.write_function(ir_function, info, ir_module, None, &debug_info_inner)?; self.lookup_function.insert(handle, id); diff --git a/tests/out/glsl/separate-entry-points.compute.Compute.glsl b/tests/out/glsl/separate-entry-points.compute.Compute.glsl index 2a439cdb41..869b7ca418 100644 --- a/tests/out/glsl/separate-entry-points.compute.Compute.glsl +++ b/tests/out/glsl/separate-entry-points.compute.Compute.glsl @@ -6,12 +6,6 @@ precision highp int; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; -void derivatives() { - float x = dFdx(0.0); - float y = dFdy(0.0); - float width = fwidth(0.0); -} - void barriers() { memoryBarrierBuffer(); barrier(); diff --git a/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl b/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl index 71cef76693..9ea32684cd 100644 --- a/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl +++ b/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl @@ -11,14 +11,6 @@ void derivatives() { float width = fwidth(0.0); } -void barriers() { - memoryBarrierBuffer(); - barrier(); - memoryBarrierShared(); - barrier(); - return; -} - void main() { derivatives(); _fs2p_location0 = vec4(0.0); diff --git a/tests/out/spv/separate-entry-points.compute.spvasm b/tests/out/spv/separate-entry-points.compute.spvasm index 1aeafda00c..38b7ea417e 100644 --- a/tests/out/spv/separate-entry-points.compute.spvasm +++ b/tests/out/spv/separate-entry-points.compute.spvasm @@ -1,43 +1,33 @@ ; SPIR-V ; Version: 1.0 ; Generator: rspirv -; Bound: 25 +; Bound: 18 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %22 "compute" -OpExecutionMode %22 LocalSize 1 1 1 +OpEntryPoint GLCompute %15 "compute" +OpExecutionMode %15 LocalSize 1 1 1 %2 = OpTypeVoid %4 = OpTypeFloat 32 %3 = OpTypeVector %4 4 %7 = OpTypeFunction %2 -%8 = OpConstant %4 0.0 -%17 = OpTypeInt 32 0 -%16 = OpConstant %17 2 -%18 = OpConstant %17 1 -%19 = OpConstant %17 72 -%20 = OpConstant %17 264 +%10 = OpTypeInt 32 0 +%9 = OpConstant %10 2 +%11 = OpConstant %10 1 +%12 = OpConstant %10 72 +%13 = OpConstant %10 264 %6 = OpFunction %2 None %7 %5 = OpLabel -OpBranch %9 -%9 = OpLabel -%10 = OpDPdx %4 %8 -%11 = OpDPdy %4 %8 -%12 = OpFwidth %4 %8 +OpBranch %8 +%8 = OpLabel +OpControlBarrier %9 %11 %12 +OpControlBarrier %9 %9 %13 OpReturn OpFunctionEnd -%14 = OpFunction %2 None %7 -%13 = OpLabel -OpBranch %15 -%15 = OpLabel -OpControlBarrier %16 %18 %19 -OpControlBarrier %16 %16 %20 -OpReturn -OpFunctionEnd -%22 = OpFunction %2 None %7 -%21 = OpLabel -OpBranch %23 -%23 = OpLabel -%24 = OpFunctionCall %2 %14 +%15 = OpFunction %2 None %7 +%14 = OpLabel +OpBranch %16 +%16 = OpLabel +%17 = OpFunctionCall %2 %6 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/spv/separate-entry-points.fragment.spvasm b/tests/out/spv/separate-entry-points.fragment.spvasm index 215f9b21ca..e29ce8f15d 100644 --- a/tests/out/spv/separate-entry-points.fragment.spvasm +++ b/tests/out/spv/separate-entry-points.fragment.spvasm @@ -1,26 +1,21 @@ ; SPIR-V ; Version: 1.0 ; Generator: rspirv -; Bound: 28 +; Bound: 20 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Fragment %24 "fragment" %22 -OpExecutionMode %24 OriginUpperLeft -OpDecorate %22 Location 0 +OpEntryPoint Fragment %16 "fragment" %14 +OpExecutionMode %16 OriginUpperLeft +OpDecorate %14 Location 0 %2 = OpTypeVoid %4 = OpTypeFloat 32 %3 = OpTypeVector %4 4 %7 = OpTypeFunction %2 %8 = OpConstant %4 0.0 -%17 = OpTypeInt 32 0 -%16 = OpConstant %17 2 -%18 = OpConstant %17 1 -%19 = OpConstant %17 72 -%20 = OpConstant %17 264 -%23 = OpTypePointer Output %3 -%22 = OpVariable %23 Output -%25 = OpConstantNull %3 +%15 = OpTypePointer Output %3 +%14 = OpVariable %15 Output +%17 = OpConstantNull %3 %6 = OpFunction %2 None %7 %5 = OpLabel OpBranch %9 @@ -30,19 +25,11 @@ OpBranch %9 %12 = OpFwidth %4 %8 OpReturn OpFunctionEnd -%14 = OpFunction %2 None %7 +%16 = OpFunction %2 None %7 %13 = OpLabel -OpBranch %15 -%15 = OpLabel -OpControlBarrier %16 %18 %19 -OpControlBarrier %16 %16 %20 -OpReturn -OpFunctionEnd -%24 = OpFunction %2 None %7 -%21 = OpLabel -OpBranch %26 -%26 = OpLabel -%27 = OpFunctionCall %2 %6 -OpStore %22 %25 +OpBranch %18 +%18 = OpLabel +%19 = OpFunctionCall %2 %6 +OpStore %14 %17 OpReturn OpFunctionEnd \ No newline at end of file From 3a19bb5b1810fe911905cf37e79520dbfedd87d1 Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Wed, 18 Oct 2023 14:45:21 +0200 Subject: [PATCH 3/3] add more docs --- src/back/glsl/mod.rs | 10 +++++++++- src/back/spv/writer.rs | 10 +++++++++- 2 files changed, 18 insertions(+), 2 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index db3b88f15c..8527ba650e 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -827,7 +827,15 @@ impl<'a, W: Write> Writer<'a, W> { let fun_info = &self.info[handle]; - // Skip functions that that are not compatible with this entry point's stage + // Skip functions that that are not compatible with this entry point's stage. + // + // When validation is enabled, it rejects modules whose entry points try to call + // incompatible functions, so if we got this far, then any functions incompatible + // with our selected entry point must not be used. + // + // When validation is disabled, `fun_info.available_stages` is always just + // `ShaderStages::all()`, so this will write all functions in the module, and + // the downstream GLSL compiler will catch any problems. if !fun_info.available_stages.contains(ep_info.available_stages) { continue; } diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index 8192b6d219..24cb14a161 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -1953,7 +1953,15 @@ impl Writer { continue; } - // Skip functions that that are not compatible with this entry point's stage + // Skip functions that that are not compatible with this entry point's stage. + // + // When validation is enabled, it rejects modules whose entry points try to call + // incompatible functions, so if we got this far, then any functions incompatible + // with our selected entry point must not be used. + // + // When validation is disabled, `fun_info.available_stages` is always just + // `ShaderStages::all()`, so this will write all functions in the module, and + // the downstream GLSL compiler will catch any problems. if !info.available_stages.contains(ep_info.available_stages) { continue; }