diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 37d3303c7..c6f211183 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -827,6 +827,19 @@ 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. + // + // 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; + } + // 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 2e133985a..24cb14a16 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -1952,6 +1952,19 @@ impl Writer { log::info!("Skip function {:?}", ir_function.name); continue; } + + // 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; + } } let id = self.write_function(ir_function, info, ir_module, None, &debug_info_inner)?; self.lookup_function.insert(handle, id); diff --git a/tests/in/separate-entry-points.param.ron b/tests/in/separate-entry-points.param.ron new file mode 100644 index 000000000..af0931c11 --- /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 000000000..a7ec3b083 --- /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 000000000..869b7ca41 --- /dev/null +++ b/tests/out/glsl/separate-entry-points.compute.Compute.glsl @@ -0,0 +1,21 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + + +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 000000000..9ea32684c --- /dev/null +++ b/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl @@ -0,0 +1,19 @@ +#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 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 000000000..38b7ea417 --- /dev/null +++ b/tests/out/spv/separate-entry-points.compute.spvasm @@ -0,0 +1,33 @@ +; SPIR-V +; Version: 1.0 +; Generator: rspirv +; Bound: 18 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %15 "compute" +OpExecutionMode %15 LocalSize 1 1 1 +%2 = OpTypeVoid +%4 = OpTypeFloat 32 +%3 = OpTypeVector %4 4 +%7 = OpTypeFunction %2 +%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 %8 +%8 = OpLabel +OpControlBarrier %9 %11 %12 +OpControlBarrier %9 %9 %13 +OpReturn +OpFunctionEnd +%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 new file mode 100644 index 000000000..e29ce8f15 --- /dev/null +++ b/tests/out/spv/separate-entry-points.fragment.spvasm @@ -0,0 +1,35 @@ +; SPIR-V +; Version: 1.0 +; Generator: rspirv +; Bound: 20 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +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 +%15 = OpTypePointer Output %3 +%14 = OpVariable %15 Output +%17 = 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 +%16 = OpFunction %2 None %7 +%13 = OpLabel +OpBranch %18 +%18 = OpLabel +%19 = OpFunctionCall %2 %6 +OpStore %14 %17 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 2bc7f4544..c3455dd86 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() {