Implement shader barycentrics (#8320)

2025-12-08 21:26:17 +00:00 · 2025-10-29 13:15:21 -04:00 · 2025-10-29 13:15:21 -04:00 · ddcd89fd69
commit ddcd89fd69
parent 54beef8aca
34 changed files with 422 additions and 4 deletions
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@ -85,6 +85,12 @@ SamplerDescriptor {
 - Using both the wgpu command encoding APIs and `CommandEncoder::as_hal_mut` on the same encoder will now result in a panic.
 - Allow `include_spirv!` and `include_spirv_raw!` macros to be used in constants and statics. By @clarfonthey in [#8250](https://github.com/gfx-rs/wgpu/pull/8250).

+### Added/New Features
+
+## General
+
+- Implement shader triangle barycentric coordinate builtins. By @atlv24 in [#8320](https://github.com/gfx-rs/wgpu/pull/8320).
+
 ### Bug Fixes

 #### naga
--- a/naga/src/back/glsl/features.rs
+++ b/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 {
--- a/naga/src/back/glsl/mod.rs
+++ b/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 {
--- a/naga/src/back/hlsl/conv.rs
+++ b/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
--- a/naga/src/back/msl/mod.rs
+++ b/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
--- a/naga/src/back/spv/writer.rs
+++ b/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",
--- a/naga/src/common/wgsl/to_wgsl.rs
+++ b/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",
--- a/naga/src/front/glsl/variables.rs
+++ b/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,
--- a/naga/src/front/spv/convert.rs
+++ b/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
--- a/naga/src/front/spv/mod.rs
+++ b/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)) =
--- a/naga/src/front/wgsl/parse/conv.rs
+++ b/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
--- a/naga/src/ir/mod.rs
+++ b/naga/src/ir/mod.rs
@ -387,6 +387,7 @@ pub enum BuiltIn {
    PointCoord,
    FrontFacing,
    PrimitiveIndex,
+    Barycentric,
    SampleIndex,
    SampleMask,
    // compute
--- a/naga/src/valid/interface.rs
+++ b/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),
--- a/naga/src/valid/mod.rs
+++ b/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;
    }
 }

--- a/naga/tests/in/wgsl/barycentrics.toml
+++ b/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
--- a/naga/tests/in/wgsl/barycentrics.wgsl
+++ b/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);
+}
--- a/naga/tests/in/wgsl/extra.toml
+++ b/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

--- a/naga/tests/naga/spirv_capabilities.rs
+++ b/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(
--- a/naga/tests/out/glsl/wgsl-barycentrics.fs_main.Fragment.glsl
+++ b/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;
+}
+
--- a/naga/tests/out/hlsl/wgsl-barycentrics.hlsl
+++ b/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);
+}
--- a/naga/tests/out/hlsl/wgsl-barycentrics.ron
+++ b/naga/tests/out/hlsl/wgsl-barycentrics.ron
@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+        (
+            entry_point:"fs_main",
+            target_profile:"ps_6_1",
+        ),
+    ],
+    compute:[
+    ],
+)
--- a/naga/tests/out/msl/wgsl-barycentrics.msl
+++ b/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) };
+}
--- a/naga/tests/out/msl/wgsl-extra.msl
+++ b/naga/tests/out/msl/wgsl-extra.msl
@ -1,4 +1,4 @@
-// language: metal2.2
+// language: metal2.3
 #include <metal_stdlib>
 #include <simd/simd.h>

--- a/naga/tests/out/spv/wgsl-barycentrics.spvasm
+++ b/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
--- a/naga/tests/out/wgsl/wgsl-barycentrics.wgsl
+++ b/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);
+}
--- a/tests/tests/wgpu-gpu/main.rs
+++ b/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);
--- a/tests/tests/wgpu-gpu/shader_barycentric/barycentric.wgsl
+++ b/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);
+}
--- a/tests/tests/wgpu-gpu/shader_barycentric/mod.rs
+++ b/tests/tests/wgpu-gpu/shader_barycentric/mod.rs
@ -0,0 +1,164 @@
+use wgpu::util::DeviceExt;
+use wgpu_test::{gpu_test, GpuTestConfiguration, TestParameters, TestingContext};
+
+pub fn all_tests(vec: &mut Vec<wgpu_test::GpuTestInitializer>) {
+    vec.push(BARYCENTRIC);
+}
+
+//
+// This test renders one triangle to a 2x2 render target. The triangle
+// covers the bottom-left, bottom-right, and the top-left pixel.
+// XY layout of the render target, with the triangle:
+//
+//     (-1,1)  (0,1)  (1,1)
+//        +------+------+
+//        |      |      |
+//        |   o  |      |
+//        |   |\ |      |
+//        |   | \|      |
+// (-1,0) +---|--\------+ (1,0)
+//        |   |  |\     |
+//        |   |  | \    |
+//        |   o--+--o   |
+//        |      |      |
+//        +------+------+
+//     (-1,-1) (0,-1) (1,-1)
+//
+// The fragment shader outputs color based on builtin(barycentric):
+//
+//     return vec4<f32>(bary * 1.1 - 0.05, 1.0);
+//
+
+#[gpu_test]
+static BARYCENTRIC: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(
+        TestParameters::default()
+            .test_features_limits()
+            .features(wgpu::Features::SHADER_BARYCENTRICS),
+    )
+    .run_async(barycentric);
+
+async fn barycentric(ctx: TestingContext) {
+    let shader = ctx
+        .device
+        .create_shader_module(wgpu::include_wgsl!("barycentric.wgsl"));
+
+    let n = -0.505;
+    let p = 0.51;
+    let triangle_xy: [f32; 6] = [n, n, p, n, n, p];
+    let vertex_buffer = ctx
+        .device
+        .create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: None,
+            contents: bytemuck::cast_slice(&triangle_xy),
+            usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::COPY_DST,
+        });
+
+    let indices = [0u32, 1, 2];
+    let index_buffer = ctx
+        .device
+        .create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: None,
+            contents: bytemuck::cast_slice(&indices),
+            usage: wgpu::BufferUsages::INDEX | wgpu::BufferUsages::COPY_DST,
+        });
+
+    let pipeline = ctx
+        .device
+        .create_render_pipeline(&wgpu::RenderPipelineDescriptor {
+            label: None,
+            layout: None,
+            vertex: wgpu::VertexState {
+                module: &shader,
+                entry_point: Some("vs_main"),
+                compilation_options: Default::default(),
+                buffers: &[wgpu::VertexBufferLayout {
+                    array_stride: 8,
+                    step_mode: wgpu::VertexStepMode::Vertex,
+                    attributes: &[wgpu::VertexAttribute {
+                        format: wgpu::VertexFormat::Float32x2,
+                        offset: 0,
+                        shader_location: 0,
+                    }],
+                }],
+            },
+            primitive: wgpu::PrimitiveState::default(),
+            depth_stencil: None,
+            multisample: wgpu::MultisampleState::default(),
+            fragment: Some(wgpu::FragmentState {
+                module: &shader,
+                entry_point: Some("fs_main"),
+                compilation_options: Default::default(),
+                targets: &[Some(wgpu::ColorTargetState {
+                    format: wgpu::TextureFormat::Rgba8Unorm,
+                    blend: None,
+                    write_mask: wgpu::ColorWrites::ALL,
+                })],
+            }),
+            multiview: None,
+            cache: None,
+        });
+
+    let width = 2;
+    let height = 2;
+    let texture_size = wgpu::Extent3d {
+        width,
+        height,
+        depth_or_array_layers: 1,
+    };
+    let color_texture = ctx.device.create_texture(&wgpu::TextureDescriptor {
+        label: None,
+        size: texture_size,
+        mip_level_count: 1,
+        sample_count: 1,
+        dimension: wgpu::TextureDimension::D2,
+        format: wgpu::TextureFormat::Rgba8Unorm,
+        usage: wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::COPY_SRC,
+        view_formats: &[],
+    });
+    let color_view = color_texture.create_view(&wgpu::TextureViewDescriptor::default());
+
+    let readback_buffer = wgpu_test::image::ReadbackBuffers::new(&ctx.device, &color_texture);
+
+    let mut encoder = ctx
+        .device
+        .create_command_encoder(&wgpu::CommandEncoderDescriptor::default());
+    {
+        let mut rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+            label: None,
+            color_attachments: &[Some(wgpu::RenderPassColorAttachment {
+                ops: wgpu::Operations {
+                    load: wgpu::LoadOp::Clear(wgpu::Color::WHITE),
+                    store: wgpu::StoreOp::Store,
+                },
+                resolve_target: None,
+                view: &color_view,
+                depth_slice: None,
+            })],
+            depth_stencil_attachment: None,
+            timestamp_writes: None,
+            occlusion_query_set: None,
+        });
+
+        rpass.set_pipeline(&pipeline);
+        rpass.set_index_buffer(index_buffer.slice(..), wgpu::IndexFormat::Uint32);
+        rpass.set_vertex_buffer(0, vertex_buffer.slice(..));
+        rpass.draw(0..3, 0..1);
+    }
+    readback_buffer.copy_from(&ctx.device, &mut encoder, &color_texture);
+    ctx.queue.submit(Some(encoder.finish()));
+
+    //
+    //   +-----+-----+
+    //   |blue |white|
+    //   +-----+-----+
+    //   | red |green|
+    //   +-----+-----+
+    //
+    let expected = [
+        0, 0, 255, 255, 255, 255, 255, 255, 255, 0, 0, 255, 0, 255, 0, 255,
+    ];
+    readback_buffer
+        .assert_buffer_contents(&ctx, &expected)
+        .await;
+}
--- a/wgpu-core/src/device/mod.rs
+++ b/wgpu-core/src/device/mod.rs
@ -510,6 +510,10 @@ pub fn create_validator(
        Caps::TEXTURE_EXTERNAL,
        features.intersects(wgt::Features::EXTERNAL_TEXTURE),
    );
+    caps.set(
+        Caps::SHADER_BARYCENTRICS,
+        features.intersects(wgt::Features::SHADER_BARYCENTRICS),
+    );

    naga::valid::Validator::new(flags, caps)
 }
--- a/wgpu-hal/src/dx12/adapter.rs
+++ b/wgpu-hal/src/dx12/adapter.rs
@ -568,6 +568,23 @@ impl super::Adapter {
            wgt::Features::EXPERIMENTAL_MESH_SHADER,
            mesh_shader_supported,
        );
+        let shader_barycentrics_supported = {
+            let mut features3 = Direct3D12::D3D12_FEATURE_DATA_D3D12_OPTIONS3::default();
+            unsafe {
+                device.CheckFeatureSupport(
+                    Direct3D12::D3D12_FEATURE_D3D12_OPTIONS3,
+                    <*mut _>::cast(&mut features3),
+                    size_of_val(&features3) as u32,
+                )
+            }
+            .is_ok()
+                && features3.BarycentricsSupported.as_bool()
+                && shader_model >= naga::back::hlsl::ShaderModel::V6_1
+        };
+        features.set(
+            wgt::Features::SHADER_BARYCENTRICS,
+            shader_barycentrics_supported,
+        );

        // TODO: Determine if IPresentationManager is supported
        let presentation_timer = auxil::dxgi::time::PresentationTimer::new_dxgi();
--- a/wgpu-hal/src/metal/adapter.rs
+++ b/wgpu-hal/src/metal/adapter.rs
@ -902,6 +902,7 @@ impl super::PrivateCapabilities {
                && (device.supports_family(MTLGPUFamily::Apple7)
                    || device.supports_family(MTLGPUFamily::Mac2)),
            supports_shared_event: version.at_least((10, 14), (12, 0), os_is_mac),
+            shader_barycentrics: device.supports_shader_barycentric_coordinates(),
            // https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf#page=3
            supports_memoryless_storage: if family_check {
                device.supports_family(MTLGPUFamily::Apple2)
@ -1003,6 +1004,11 @@ impl super::PrivateCapabilities {

        features.set(F::RG11B10UFLOAT_RENDERABLE, self.format_rg11b10_all);

+        features.set(
+            F::SHADER_BARYCENTRICS,
+            self.shader_barycentrics && self.msl_version >= MTLLanguageVersion::V2_2,
+        );
+
        if self.supports_simd_scoped_operations {
            features.insert(F::SUBGROUP | F::SUBGROUP_BARRIER);
        }
--- a/wgpu-hal/src/metal/mod.rs
+++ b/wgpu-hal/src/metal/mod.rs
@ -302,6 +302,7 @@ struct PrivateCapabilities {
    int64_atomics: bool,
    float_atomics: bool,
    supports_shared_event: bool,
+    shader_barycentrics: bool,
    supports_memoryless_storage: bool,
 }

--- a/wgpu-hal/src/vulkan/adapter.rs
+++ b/wgpu-hal/src/vulkan/adapter.rs
@ -128,6 +128,9 @@ pub struct PhysicalDeviceFeatures {
    /// Features provided by `VK_KHR_shader_integer_dot_product`, promoted to Vulkan 1.3.
    shader_integer_dot_product:
        Option<vk::PhysicalDeviceShaderIntegerDotProductFeaturesKHR<'static>>,
+
+    /// Features provided by `VK_KHR_fragment_shader_barycentric`
+    shader_barycentrics: Option<vk::PhysicalDeviceFragmentShaderBarycentricFeaturesKHR<'static>>,
 }

 impl PhysicalDeviceFeatures {
@ -201,6 +204,9 @@ impl PhysicalDeviceFeatures {
        if let Some(ref mut feature) = self.shader_integer_dot_product {
            info = info.push_next(feature);
        }
+        if let Some(ref mut feature) = self.shader_barycentrics {
+            info = info.push_next(feature);
+        }
        info
    }

@ -535,6 +541,17 @@ impl PhysicalDeviceFeatures {
            } else {
                None
            },
+            shader_barycentrics: if enabled_extensions
+                .contains(&khr::fragment_shader_barycentric::NAME)
+            {
+                let needed = requested_features.intersects(wgt::Features::SHADER_BARYCENTRICS);
+                Some(
+                    vk::PhysicalDeviceFragmentShaderBarycentricFeaturesKHR::default()
+                        .fragment_shader_barycentric(needed),
+                )
+            } else {
+                None
+            },
        }
    }

@ -669,6 +686,13 @@ impl PhysicalDeviceFeatures {
            );
        }

+        if let Some(ref shader_barycentrics) = self.shader_barycentrics {
+            features.set(
+                F::SHADER_BARYCENTRICS,
+                shader_barycentrics.fragment_shader_barycentric != 0,
+            );
+        }
+
        //if caps.supports_extension(khr::sampler_mirror_clamp_to_edge::NAME) {
        //if caps.supports_extension(ext::sampler_filter_minmax::NAME) {
        features.set(
@ -1184,6 +1208,11 @@ impl PhysicalDeviceProperties {
            extensions.push(ext::mesh_shader::NAME);
        }

+        // Require `VK_KHR_fragment_shader_barycentric` if the associated feature was requested
+        if requested_features.intersects(wgt::Features::SHADER_BARYCENTRICS) {
+            extensions.push(khr::fragment_shader_barycentric::NAME);
+        }
+
        extensions
    }

@ -1638,6 +1667,13 @@ impl super::InstanceShared {
                features2 = features2.push_next(next);
            }

+            if capabilities.supports_extension(khr::fragment_shader_barycentric::NAME) {
+                let next = features
+                    .shader_barycentrics
+                    .insert(vk::PhysicalDeviceFragmentShaderBarycentricFeaturesKHR::default());
+                features2 = features2.push_next(next);
+            }
+
            unsafe { get_device_properties.get_physical_device_features2(phd, &mut features2) };
            features2.features
        } else {
@ -2135,6 +2171,10 @@ impl super::Adapter {
                capabilities.push(spv::Capability::ClipDistance);
            }

+            if features.intersects(wgt::Features::SHADER_BARYCENTRICS) {
+                capabilities.push(spv::Capability::FragmentBarycentricKHR);
+            }
+
            let mut flags = spv::WriterFlags::empty();
            flags.set(
                spv::WriterFlags::DEBUG,
--- a/wgpu-types/src/features.rs
+++ b/wgpu-types/src/features.rs
@ -1231,6 +1231,16 @@ bitflags_array! {
        ///
        /// [`Device::create_shader_module_passthrough`]: https://docs.rs/wgpu/latest/wgpu/struct.Device.html#method.create_shader_module_passthrough
        const EXPERIMENTAL_PASSTHROUGH_SHADERS = 1 << 52;
+
+        /// Enables shader barycentric coordinates.
+        ///
+        /// Supported platforms:
+        /// - Vulkan (with VK_KHR_fragment_shader_barycentric)
+        /// - DX12 (with SM 6.1+)
+        /// - Metal (with MSL 2.2+)
+        ///
+        /// This is a native only feature.
+        const SHADER_BARYCENTRICS = 1 << 53;
    }

    /// Features that are not guaranteed to be supported.