[glsl/spv-out] Cull functions that should not be available for a given stage (#2531)

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)?;
 

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);

tests/in/separate-entry-points.param.ron

@@ -0,0 +1,6 @@
+(
+	spv: (
+        version: (1, 0),
+		separate_entry_points: true,
+	),
+)

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<f32> {
+    derivatives();
+    return vec4<f32>();
+}
+
+@compute @workgroup_size(1)
+fn compute() {
+    barriers();
+}

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;
+}
+

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;
+}
+

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

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

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() {