Implement shader barycentrics

CHANGELOG.md

@@ -78,6 +78,12 @@ SamplerDescriptor {
 
 - Texture now has `from_custom`. By @R-Cramer4 in [#8315](https://github.com/gfx-rs/wgpu/pull/8315).
 
+### Added/New Features
+
+## General
+
+- Implement shader triangle barycentric coordinate builtins. By @atlv24 in [#8320](https://github.com/gfx-rs/wgpu/pull/8320).
+
 ### Bug Fixes
 
 - Fixed bug where mapping sub-ranges of a buffer on web would fail with `OperationError: GPUBuffer.getMappedRange: GetMappedRange range extends beyond buffer's mapped range`. By @ryankaplan in [#8349](https://github.com/gfx-rs/wgpu/pull/8349)

naga/src/back/glsl/features.rs

@@ -55,6 +55,8 @@ bitflags::bitflags! {
         const SUBGROUP_OPERATIONS = 1 << 24;
         /// Image atomics
         const TEXTURE_ATOMICS = 1 << 25;
+        /// Image atomics
+        const SHADER_BARYCENTRICS = 1 << 26;
     }
 }
 
@@ -288,6 +290,14 @@ impl FeaturesManager {
             writeln!(out, "#extension GL_OES_shader_image_atomic : require")?;
         }
 
+        if self.0.contains(Features::SHADER_BARYCENTRICS) {
+            // https://github.com/KhronosGroup/GLSL/blob/main/extensions/ext/GLSL_EXT_fragment_shader_barycentric.txt
+            writeln!(
+                out,
+                "#extension GL_EXT_fragment_shader_barycentric : require"
+            )?;
+        }
+
         Ok(())
     }
 }
@@ -603,6 +613,9 @@ impl<W> Writer<'_, W> {
                     crate::BuiltIn::InstanceIndex | crate::BuiltIn::DrawID => {
                         self.features.request(Features::INSTANCE_INDEX)
                     }
+                    crate::BuiltIn::Barycentric => {
+                        self.features.request(Features::SHADER_BARYCENTRICS)
+                    }
                     _ => {}
                 },
                 Binding::Location {

naga/src/back/glsl/mod.rs

@@ -5227,6 +5227,7 @@ const fn glsl_built_in(built_in: crate::BuiltIn, options: VaryingOptions) -> &'s
         Bi::PointCoord => "gl_PointCoord",
         Bi::FrontFacing => "gl_FrontFacing",
         Bi::PrimitiveIndex => "uint(gl_PrimitiveID)",
+        Bi::Barycentric => "gl_BaryCoordEXT",
         Bi::SampleIndex => "gl_SampleID",
         Bi::SampleMask => {
             if options.output {

naga/src/back/hlsl/conv.rs

@@ -161,6 +161,7 @@ impl crate::BuiltIn {
             Self::FragDepth => "SV_Depth",
             Self::FrontFacing => "SV_IsFrontFace",
             Self::PrimitiveIndex => "SV_PrimitiveID",
+            Self::Barycentric => "SV_Barycentrics",
             Self::SampleIndex => "SV_SampleIndex",
             Self::SampleMask => "SV_Coverage",
             // compute

naga/src/back/msl/mod.rs

@@ -526,10 +526,15 @@ impl Options {
                         return Err(Error::UnsupportedAttribute("instance_id".to_string()));
                     }
                     // macOS: Since Metal 2.2
-                    // iOS: Since Metal 2.3 (check depends on https://github.com/gfx-rs/naga/issues/2164)
-                    crate::BuiltIn::PrimitiveIndex if self.lang_version < (2, 2) => {
+                    // iOS: Since Metal 2.3 (check depends on https://github.com/gfx-rs/wgpu/issues/4414)
+                    crate::BuiltIn::PrimitiveIndex if self.lang_version < (2, 3) => {
                         return Err(Error::UnsupportedAttribute("primitive_id".to_string()));
                     }
+                    // macOS: Since Metal 2.2
+                    // iOS: Since Metal 2.3 (check depends on https://github.com/gfx-rs/wgpu/issues/4414)
+                    crate::BuiltIn::Barycentric if self.lang_version < (2, 3) => {
+                        return Err(Error::UnsupportedAttribute("barycentric_coord".to_string()));
+                    }
                     _ => {}
                 }
 
@@ -680,6 +685,7 @@ impl ResolvedBinding {
                     Bi::PointCoord => "point_coord",
                     Bi::FrontFacing => "front_facing",
                     Bi::PrimitiveIndex => "primitive_id",
+                    Bi::Barycentric => "barycentric_coord",
                     Bi::SampleIndex => "sample_id",
                     Bi::SampleMask => "sample_mask",
                     // compute

naga/src/back/spv/writer.rs

@@ -2089,6 +2089,14 @@ impl Writer {
                         )?;
                         BuiltIn::PrimitiveId
                     }
+                    Bi::Barycentric => {
+                        self.require_any(
+                            "`barycentric` built-in",
+                            &[spirv::Capability::FragmentBarycentricKHR],
+                        )?;
+                        self.use_extension("SPV_KHR_fragment_shader_barycentric");
+                        BuiltIn::BaryCoordKHR
+                    }
                     Bi::SampleIndex => {
                         self.require_any(
                             "`sample_index` built-in",

naga/src/common/wgsl/to_wgsl.rs

@@ -169,6 +169,7 @@ impl TryToWgsl for crate::BuiltIn {
             Bi::FragDepth => "frag_depth",
             Bi::FrontFacing => "front_facing",
             Bi::PrimitiveIndex => "primitive_index",
+            Bi::Barycentric => "barycentric",
             Bi::SampleIndex => "sample_index",
             Bi::SampleMask => "sample_mask",
             Bi::GlobalInvocationId => "global_invocation_id",

naga/src/front/glsl/variables.rs

@@ -200,6 +200,7 @@ impl Frontend {
                     "gl_BaseVertex" => BuiltIn::BaseVertex,
                     "gl_BaseInstance" => BuiltIn::BaseInstance,
                     "gl_PrimitiveID" => BuiltIn::PrimitiveIndex,
+                    "gl_BaryCoordEXT" => BuiltIn::Barycentric,
                     "gl_InstanceIndex" => BuiltIn::InstanceIndex,
                     "gl_VertexIndex" => BuiltIn::VertexIndex,
                     "gl_SampleID" => BuiltIn::SampleIndex,

naga/src/front/spv/convert.rs

@@ -147,6 +147,7 @@ pub(super) fn map_builtin(word: spirv::Word, invariant: bool) -> Result<crate::B
         Some(Bi::PointCoord) => crate::BuiltIn::PointCoord,
         Some(Bi::FrontFacing) => crate::BuiltIn::FrontFacing,
         Some(Bi::PrimitiveId) => crate::BuiltIn::PrimitiveIndex,
+        Some(Bi::BaryCoordKHR) => crate::BuiltIn::Barycentric,
         Some(Bi::SampleId) => crate::BuiltIn::SampleIndex,
         Some(Bi::SampleMask) => crate::BuiltIn::SampleMask,
         // compute

naga/src/front/spv/mod.rs

@@ -83,6 +83,7 @@ pub const SUPPORTED_CAPABILITIES: &[spirv::Capability] = &[
     spirv::Capability::GroupNonUniformShuffleRelative,
     spirv::Capability::RuntimeDescriptorArray,
     spirv::Capability::StorageImageMultisample,
+    spirv::Capability::FragmentBarycentricKHR,
     // tricky ones
     spirv::Capability::UniformBufferArrayDynamicIndexing,
     spirv::Capability::StorageBufferArrayDynamicIndexing,
@@ -6038,6 +6039,10 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                             size: crate::VectorSize::Tri,
                             scalar: crate::Scalar::U32,
                         }),
+                        crate::BuiltIn::Barycentric => Some(crate::TypeInner::Vector {
+                            size: crate::VectorSize::Tri,
+                            scalar: crate::Scalar::F32,
+                        }),
                         _ => None,
                     };
                     if let (Some(inner), Some(crate::ScalarKind::Sint)) =

naga/src/front/wgsl/parse/conv.rs

@@ -36,6 +36,7 @@ pub fn map_built_in(
         "front_facing" => crate::BuiltIn::FrontFacing,
         "frag_depth" => crate::BuiltIn::FragDepth,
         "primitive_index" => crate::BuiltIn::PrimitiveIndex,
+        "barycentric" => crate::BuiltIn::Barycentric,
         "sample_index" => crate::BuiltIn::SampleIndex,
         "sample_mask" => crate::BuiltIn::SampleMask,
         // compute

naga/src/ir/mod.rs

@@ -387,6 +387,7 @@ pub enum BuiltIn {
     PointCoord,
     FrontFacing,
     PrimitiveIndex,
+    Barycentric,
     SampleIndex,
     SampleMask,
     // compute

naga/src/valid/interface.rs

@@ -180,6 +180,7 @@ impl VaryingContext<'_> {
                     Bi::ClipDistance => Capabilities::CLIP_DISTANCE,
                     Bi::CullDistance => Capabilities::CULL_DISTANCE,
                     Bi::PrimitiveIndex => Capabilities::PRIMITIVE_INDEX,
+                    Bi::Barycentric => Capabilities::SHADER_BARYCENTRICS,
                     Bi::ViewIndex => Capabilities::MULTIVIEW,
                     Bi::SampleIndex => Capabilities::MULTISAMPLED_SHADING,
                     Bi::NumSubgroups
@@ -267,6 +268,14 @@ impl VaryingContext<'_> {
                         self.stage == St::Fragment && !self.output,
                         *ty_inner == Ti::Scalar(crate::Scalar::U32),
                     ),
+                    Bi::Barycentric => (
+                        self.stage == St::Fragment && !self.output,
+                        *ty_inner
+                            == Ti::Vector {
+                                size: Vs::Tri,
+                                scalar: crate::Scalar::F32,
+                            },
+                    ),
                     Bi::SampleIndex => (
                         self.stage == St::Fragment && !self.output,
                         *ty_inner == Ti::Scalar(crate::Scalar::U32),

naga/src/valid/mod.rs

@@ -186,6 +186,8 @@ bitflags::bitflags! {
         /// Support for `quantizeToF16`, `pack2x16float`, and `unpack2x16float`, which store
         /// `f16`-precision values in `f32`s.
         const SHADER_FLOAT16_IN_FLOAT32 = 1 << 28;
+        /// Support for fragment shader barycentric coordinates.
+        const SHADER_BARYCENTRICS = 1 << 29;
     }
 }
 

naga/tests/in/wgsl/barycentrics.toml

@@ -0,0 +1,10 @@
+god_mode = true
+
+[msl]
+lang_version = [2, 3]
+
+[hlsl]
+shader_model = "V6_1"
+
+[glsl]
+version.Desktop = 450

naga/tests/in/wgsl/barycentrics.wgsl

@@ -0,0 +1,4 @@
+@fragment
+fn fs_main(@builtin(barycentric) bary: vec3<f32>) -> @location(0) vec4<f32> {
+    return vec4(bary, 1.0);
+}

naga/tests/in/wgsl/extra.toml

@@ -3,7 +3,7 @@ targets = "SPIRV | METAL | WGSL"
 
 [msl]
 fake_missing_bindings = false
-lang_version = [2, 2]
+lang_version = [2, 3]
 spirv_cross_compatibility = false
 zero_initialize_workgroup_memory = true
 

naga/tests/naga/spirv_capabilities.rs

@@ -148,6 +148,17 @@ fn sample_rate_shading() {
     );
 }
 
+#[test]
+fn barycentrics() {
+    require(
+        &[Ca::FragmentBarycentricKHR],
+        r#"
+        @fragment
+        fn f(@builtin(barycentric) x: vec3<f32>) { }
+    "#,
+    );
+}
+
 #[test]
 fn geometry() {
     require(

naga/tests/out/glsl/wgsl-barycentrics.fs_main.Fragment.glsl

@@ -0,0 +1,10 @@
+#version 450 core
+#extension GL_EXT_fragment_shader_barycentric : require
+layout(location = 0) out vec4 _fs2p_location0;
+
+void main() {
+    vec3 bary = gl_BaryCoordEXT;
+    _fs2p_location0 = vec4(bary, 1.0);
+    return;
+}
+

naga/tests/out/hlsl/wgsl-barycentrics.hlsl

@@ -0,0 +1,9 @@
+struct FragmentInput_fs_main {
+    float3 bary_1 : SV_Barycentrics;
+};
+
+float4 fs_main(FragmentInput_fs_main fragmentinput_fs_main) : SV_Target0
+{
+    float3 bary = fragmentinput_fs_main.bary_1;
+    return float4(bary, 1.0);
+}

naga/tests/out/hlsl/wgsl-barycentrics.ron

@@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+        (
+            entry_point:"fs_main",
+            target_profile:"ps_6_1",
+        ),
+    ],
+    compute:[
+    ],
+)

naga/tests/out/msl/wgsl-barycentrics.msl

@@ -0,0 +1,17 @@
+// language: metal2.3
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+
+struct fs_mainInput {
+};
+struct fs_mainOutput {
+    metal::float4 member [[color(0)]];
+};
+fragment fs_mainOutput fs_main(
+  metal::float3 bary [[barycentric_coord]]
+) {
+    return fs_mainOutput { metal::float4(bary, 1.0) };
+}

naga/tests/out/msl/wgsl-extra.msl

@@ -1,4 +1,4 @@
-// language: metal2.2
+// language: metal2.3
 #include <metal_stdlib>
 #include <simd/simd.h>
 

naga/tests/out/spv/wgsl-barycentrics.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 17
+OpCapability Shader
+OpCapability FragmentBarycentricKHR
+OpExtension "SPV_KHR_fragment_shader_barycentric"
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %12 "fs_main" %7 %10
+OpExecutionMode %12 OriginUpperLeft
+OpDecorate %7 BuiltIn BaryCoordKHR
+OpDecorate %10 Location 0
+%2 = OpTypeVoid
+%4 = OpTypeFloat 32
+%3 = OpTypeVector %4 3
+%5 = OpTypeVector %4 4
+%8 = OpTypePointer Input %3
+%7 = OpVariable  %8  Input
+%11 = OpTypePointer Output %5
+%10 = OpVariable  %11  Output
+%13 = OpTypeFunction %2
+%14 = OpConstant  %4  1
+%12 = OpFunction  %2  None %13
+%6 = OpLabel
+%9 = OpLoad  %3  %7
+OpBranch %15
+%15 = OpLabel
+%16 = OpCompositeConstruct  %5  %9 %14
+OpStore %10 %16
+OpReturn
+OpFunctionEnd

naga/tests/out/wgsl/wgsl-barycentrics.wgsl

@@ -0,0 +1,4 @@
+@fragment 
+fn fs_main(@builtin(barycentric) bary: vec3<f32>) -> @location(0) vec4<f32> {
+    return vec4<f32>(bary, 1f);
+}

tests/tests/wgpu-gpu/main.rs

@@ -54,6 +54,7 @@ mod resource_error;
 mod samplers;
 mod scissor_tests;
 mod shader;
+mod shader_barycentric;
 mod shader_primitive_index;
 mod shader_view_format;
 mod subgroup_operations;
@@ -126,6 +127,7 @@ fn all_tests() -> Vec<wgpu_test::GpuTestInitializer> {
     samplers::all_tests(&mut tests);
     scissor_tests::all_tests(&mut tests);
     shader_primitive_index::all_tests(&mut tests);
+    shader_barycentric::all_tests(&mut tests);
     shader_view_format::all_tests(&mut tests);
     shader::all_tests(&mut tests);
     subgroup_operations::all_tests(&mut tests);

tests/tests/wgpu-gpu/shader_barycentric/barycentric.wgsl

@@ -0,0 +1,9 @@
+@vertex
+fn vs_main(@location(0) xy: vec2<f32>) -> @builtin(position) vec4<f32> {
+    return vec4<f32>(xy, 0.0, 1.0);
+}
+
+@fragment
+fn fs_main(@builtin(barycentric) bary: vec3<f32>) -> @location(0) vec4<f32> {
+    return vec4<f32>(bary * 1.1 - 0.05, 1.0);
+}