resource.rs

use alloc::{
    borrow::Cow,
    boxed::Box,
    string::{String, ToString as _},
    sync::{Arc, Weak},
    vec::Vec,
};
use core::{
    fmt,
    mem::{self, ManuallyDrop},
    num::NonZeroU32,
    sync::atomic::{AtomicBool, Ordering},
};
use hal::ShouldBeNonZeroExt;

use arrayvec::ArrayVec;
use bitflags::Flags;
use smallvec::SmallVec;
use wgt::{
    math::align_to, ColorWrites, DeviceLostReason, TextureFormat, TextureSampleType,
    TextureViewDimension,
};

#[cfg(feature = "trace")]
use crate::device::trace::{self, IntoTrace as _};
use crate::{
    api_log,
    binding_model::{
        self, BindGroup, BindGroupLateBufferBindingInfo, BindGroupLayout,
        BindGroupLayoutEntryError, BindGroupLayoutState, BindGroupState, CreateBindGroupError,
        CreateBindGroupLayoutError,
    },
    command, conv,
    device::{
        bgl, create_validator, life::WaitIdleError, map_buffer, AttachmentData,
        BufferMapPendingClosure, DeviceLostInvocation, HostMap, MissingDownlevelFlags,
        MissingFeatures, RenderPassContext,
    },
    hal_label,
    init_tracker::{
        BufferInitTracker, BufferInitTrackerAction, MemoryInitKind, TextureInitRange,
        TextureInitTrackerAction,
    },
    instance::{Adapter, RequestDeviceError},
    lock::{rank, Mutex, RwLock},
    pipeline::{self, ColorStateError},
    pool::ResourcePool,
    resource::{
        self, Buffer, BufferState, ExternalTexture, ExternalTextureState, Labeled, ParentDevice,
        QuerySet, QuerySetState, RawResourceAccess, ResourceState, Sampler, StagingBuffer, Texture,
        TextureView, Tlas, TrackingData,
    },
    resource_log,
    snatch::{SnatchGuard, SnatchLock, Snatchable},
    timestamp_normalization::TIMESTAMP_NORMALIZATION_BUFFER_USES,
    track::{BindGroupStates, DeviceTracker, TrackerIndexAllocators, UsageScope, UsageScopePool},
    validation::{self, check_color_attachment_count, PassthroughInterface, ShaderMetaData},
    weak_vec::WeakVec,
    FastHashMap, LabelHelpers, OnceCellOrLock,
};

use super::{
    queue::Queue, DeviceDescriptor, DeviceError, DeviceLostClosure, UserClosures,
    ENTRYPOINT_FAILURE_ERROR, ZERO_BUFFER_SIZE,
};

#[cfg(supports_64bit_atomics)]
use core::sync::atomic::AtomicU64;
#[cfg(not(supports_64bit_atomics))]
use portable_atomic::AtomicU64;

pub(crate) struct CommandIndices {
    /// The index of the last command submission that was attempted.
    ///
    /// Note that `fence` may never be signalled with this value, if the command
    /// submission failed. If you need to wait for everything running on a
    /// `Queue` to complete, wait for [`last_successful_submission_index`].
    ///
    /// [`last_successful_submission_index`]: Device::last_successful_submission_index
    pub(crate) active_submission_index: hal::FenceValue,
    pub(crate) next_acceleration_structure_build_command_index: u64,
}

/// Parameters provided to shaders via a uniform buffer of the type
/// [`NagaExternalTextureParams`], describing an [`ExternalTexture`] resource
/// binding.
///
/// [`NagaExternalTextureParams`]: naga::SpecialTypes::external_texture_params
/// [`ExternalTexture`]: binding_model::BindingResource::ExternalTexture
#[repr(C)]
#[derive(Copy, Clone, bytemuck::Zeroable, bytemuck::Pod)]
pub struct ExternalTextureParams {
    /// 4x4 column-major matrix with which to convert sampled YCbCr values
    /// to RGBA.
    ///
    /// This is ignored when `num_planes` is 1.
    pub yuv_conversion_matrix: [f32; 16],

    /// 3x3 column-major matrix to transform linear RGB values in the source
    /// color space to linear RGB values in the destination color space. In
    /// combination with [`Self::src_transfer_function`] and
    /// [`Self::dst_transfer_function`] this can be used to ensure that
    /// [`ImageSample`] and [`ImageLoad`] operations return values in the
    /// desired destination color space rather than the source color space of
    /// the underlying planes.
    ///
    /// Includes a padding element after each column.
    ///
    /// [`ImageSample`]: naga::ir::Expression::ImageSample
    /// [`ImageLoad`]: naga::ir::Expression::ImageLoad
    pub gamut_conversion_matrix: [f32; 12],

    /// Transfer function for the source color space. The *inverse* of this
    /// will be applied to decode non-linear RGB to linear RGB in the source
    /// color space.
    pub src_transfer_function: wgt::ExternalTextureTransferFunction,

    /// Transfer function for the destination color space. This will be applied
    /// to encode linear RGB to non-linear RGB in the destination color space.
    pub dst_transfer_function: wgt::ExternalTextureTransferFunction,

    /// Transform to apply to [`ImageSample`] coordinates.
    ///
    /// This is a 3x2 column-major matrix representing an affine transform from
    /// normalized texture coordinates to the normalized coordinates that should
    /// be sampled from the external texture's underlying plane(s).
    ///
    /// This transform may scale, translate, flip, and rotate in 90-degree
    /// increments, but the result of transforming the rectangle (0,0)..(1,1)
    /// must be an axis-aligned rectangle that falls within the bounds of
    /// (0,0)..(1,1).
    ///
    /// [`ImageSample`]: naga::ir::Expression::ImageSample
    pub sample_transform: [f32; 6],

    /// Transform to apply to [`ImageLoad`] coordinates.
    ///
    /// This is a 3x2 column-major matrix representing an affine transform from
    /// non-normalized texel coordinates to the non-normalized coordinates of
    /// the texel that should be loaded from the external texture's underlying
    /// plane 0. For planes 1 and 2, if present, plane 0's coordinates are
    /// scaled according to the textures' relative sizes.
    ///
    /// This transform may scale, translate, flip, and rotate in 90-degree
    /// increments, but the result of transforming the rectangle (0,0)..[`size`]
    /// must be an axis-aligned rectangle that falls within the bounds of
    /// (0,0)..[`size`].
    ///
    /// [`ImageLoad`]: naga::ir::Expression::ImageLoad
    /// [`size`]: Self::size
    pub load_transform: [f32; 6],

    /// Size of the external texture.
    ///
    /// This is the value that should be returned by size queries in shader
    /// code; it does not necessarily match the dimensions of the underlying
    /// texture(s). As a special case, if this is `[0, 0]`, the actual size of
    /// plane 0 should be used instead.
    ///
    /// This must be consistent with [`sample_transform`]: it should be the size
    /// in texels of the rectangle covered by the square (0,0)..(1,1) after
    /// [`sample_transform`] has been applied to it.
    ///
    /// [`sample_transform`]: Self::sample_transform
    pub size: [u32; 2],

    /// Number of planes. 1 indicates a single RGBA plane. 2 indicates a Y
    /// plane and an interleaved CbCr plane. 3 indicates separate Y, Cb, and Cr
    /// planes.
    pub num_planes: u32,
    // Ensure the size of this struct matches the type generated by Naga.
    pub _padding: [u8; 4],
}

impl ExternalTextureParams {
    pub fn from_desc<L>(desc: &wgt::ExternalTextureDescriptor<L>) -> Self {
        let gamut_conversion_matrix = [
            desc.gamut_conversion_matrix[0],
            desc.gamut_conversion_matrix[1],
            desc.gamut_conversion_matrix[2],
            0.0, // padding
            desc.gamut_conversion_matrix[3],
            desc.gamut_conversion_matrix[4],
            desc.gamut_conversion_matrix[5],
            0.0, // padding
            desc.gamut_conversion_matrix[6],
            desc.gamut_conversion_matrix[7],
            desc.gamut_conversion_matrix[8],
            0.0, // padding
        ];

        Self {
            yuv_conversion_matrix: desc.yuv_conversion_matrix,
            gamut_conversion_matrix,
            src_transfer_function: desc.src_transfer_function,
            dst_transfer_function: desc.dst_transfer_function,
            size: [desc.width, desc.height],
            sample_transform: desc.sample_transform,
            load_transform: desc.load_transform,
            num_planes: desc.num_planes() as u32,
            _padding: Default::default(),
        }
    }
}

/// Because all operations are push/swap (no longlived lock),
/// we can have mutex without lock rank
pub(crate) struct DeferredBufferMapPendingClosures(
    parking_lot::Mutex<Vec<BufferMapPendingClosure>>,
);

impl DeferredBufferMapPendingClosures {
    pub(crate) fn new() -> Self {
        Self(parking_lot::Mutex::new(Vec::new()))
    }

    pub(crate) fn push(&self, closure: BufferMapPendingClosure) {
        self.0.lock().push(closure);
    }

    pub(crate) fn swap(&self, other: &mut Vec<BufferMapPendingClosure>) {
        mem::swap(&mut *self.0.lock(), other)
    }
}

/// Resources associated with a device.
///
/// This struct exists so that resources can be cleaned up properly on error returns
/// from [`Device::new`].
///
/// [`Device::timestamp_normalizer`] is late-initialized after [`Device::new`], so it is not
/// included here.
struct DeviceResources<'a> {
    raw: &'a dyn hal::DynDevice,
    zero_buffer: Option<Box<dyn hal::DynBuffer>>,
    empty_bgl: Option<Box<dyn hal::DynBindGroupLayout>>,
    default_external_texture_params_buffer: Option<Box<dyn hal::DynBuffer>>,
    fence: Option<Box<dyn hal::DynFence>>,
    indirect_validation: Option<crate::indirect_validation::IndirectValidation>,
}

/// Structure describing a logical device. Some members are internally mutable,
/// stored behind mutexes.
pub struct Device {
    raw: Box<dyn hal::DynDevice>,
    pub(crate) adapter: Arc<Adapter>,
    pub(crate) queue: OnceCellOrLock<Weak<Queue>>,
    pub(crate) zero_buffer: ManuallyDrop<Box<dyn hal::DynBuffer>>,
    pub(crate) empty_bgl: ManuallyDrop<Box<dyn hal::DynBindGroupLayout>>,
    /// The `label` from the descriptor used to create the resource.
    label: String,

    pub(crate) command_allocator: command::CommandAllocator,

    pub(crate) command_indices: RwLock<CommandIndices>,

    /// The index of the last successful submission to this device's
    /// [`hal::Queue`].
    ///
    /// Unlike [`active_submission_index`], which is incremented each time
    /// submission is attempted, this is updated only when submission succeeds,
    /// so waiting for this value won't hang waiting for work that was never
    /// submitted.
    ///
    /// [`active_submission_index`]: CommandIndices::active_submission_index
    pub(crate) last_successful_submission_index: hal::AtomicFenceValue,

    pub(crate) fence: ManuallyDrop<Box<dyn hal::DynFence>>,
    pub(crate) snatchable_lock: SnatchLock,

    /// Is this device valid? Valid is closely associated with "lose the device",
    /// which can be triggered by various methods, including at the end of device
    /// destroy, and by any GPU errors that cause us to no longer trust the state
    /// of the device. Ideally we would like to fold valid into the storage of
    /// the device itself (for example as an Error enum), but unfortunately we
    /// need to continue to be able to retrieve the device in poll_devices to
    /// determine if it can be dropped. If our internal accesses of devices were
    /// done through ref-counted references and external accesses checked for
    /// Error enums, we wouldn't need this. For now, we need it. All the call
    /// sites where we check it are areas that should be revisited if we start
    /// using ref-counted references for internal access.
    pub(crate) valid: AtomicBool,

    /// Closure to be called on "lose the device". This is invoked directly by
    /// device.lose or by the UserCallbacks returned from maintain when the device
    /// has been destroyed and its queues are empty.
    pub(crate) device_lost_closure: Mutex<Option<DeviceLostClosure>>,

    /// Stores the state of buffers and textures.
    pub(crate) trackers: Mutex<DeviceTracker>,
    pub(crate) tracker_indices: TrackerIndexAllocators,
    /// Pool of bind group layouts, allowing deduplication.
    pub(crate) bgl_pool: ResourcePool<bgl::EntryMap, BindGroupLayout>,
    pub(crate) alignments: hal::Alignments,
    pub(crate) limits: wgt::Limits,
    pub(crate) features: wgt::Features,
    pub(crate) downlevel: wgt::DownlevelCapabilities,
    /// Buffer uses listed here, are expected to be ordered by the underlying hardware.
    /// If a usage is ordered, then if the buffer state doesn't change between draw calls,
    /// there are no barriers needed for synchronization.
    /// See the implementations of [`hal::Adapter::get_ordered_buffer_usages`] for hardware specific info
    pub(crate) ordered_buffer_usages: wgt::BufferUses,
    /// Texture uses listed here, are expected to be ordered by the underlying hardware.
    /// If a usage is ordered, then if the buffer state doesn't change between draw calls,
    /// there are no barriers needed for synchronization.
    /// See the implementations of [`hal::Adapter::get_ordered_texture_usages`] for hardware specific info
    pub(crate) ordered_texture_usages: wgt::TextureUses,
    pub(crate) instance_flags: wgt::InstanceFlags,
    pub(crate) deferred_destroy: Mutex<Vec<DeferredDestroy>>,
    /// This closures were created in [`Buffer::drop`] where we do not run them to prevent locking problems.
    pub(crate) deferred_buffer_map_pending_closures: DeferredBufferMapPendingClosures,
    pub(crate) usage_scopes: UsageScopePool,
    pub(crate) indirect_validation: Option<crate::indirect_validation::IndirectValidation>,
    // Optional so that we can late-initialize this after the queue is created.
    pub(crate) timestamp_normalizer:
        OnceCellOrLock<crate::timestamp_normalization::TimestampNormalizer>,
    /// Uniform buffer containing [`ExternalTextureParams`] with values such
    /// that a [`TextureView`] bound to a [`wgt::BindingType::ExternalTexture`]
    /// binding point will be rendered correctly. Intended to be used as the
    /// [`hal::ExternalTextureBinding::params`] field.
    pub(crate) default_external_texture_params_buffer: ManuallyDrop<Box<dyn hal::DynBuffer>>,
    // needs to be dropped last
    #[cfg(feature = "trace")]
    pub(crate) trace: Mutex<Option<Box<dyn trace::Trace + Send + Sync + 'static>>>,
}

pub(crate) enum DeferredDestroy {
    TextureViews(WeakVec<TextureView>),
    BindGroups(WeakVec<BindGroup>),
}

impl fmt::Debug for Device {
    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
        f.debug_struct("Device")
            .field("label", &self.label())
            .field("limits", &self.limits)
            .field("features", &self.features)
            .field("downlevel", &self.downlevel)
            .finish()
    }
}

impl Drop for DeviceResources<'_> {
    fn drop(&mut self) {
        if let Some(indirect_validation) = self.indirect_validation.take() {
            indirect_validation.dispose(self.raw);
        }
        unsafe {
            if let Some(zero_buffer) = self.zero_buffer.take() {
                self.raw.destroy_buffer(zero_buffer);
            }
            if let Some(empty_bgl) = self.empty_bgl.take() {
                self.raw.destroy_bind_group_layout(empty_bgl);
            }
            if let Some(default_external_texture_params_buffer) =
                self.default_external_texture_params_buffer.take()
            {
                self.raw
                    .destroy_buffer(default_external_texture_params_buffer);
            }
            if let Some(fence) = self.fence.take() {
                self.raw.destroy_fence(fence);
            }
        }
    }
}

impl Drop for Device {
    #[allow(trivial_casts)]
    fn drop(&mut self) {
        profiling::scope!("Device::drop");
        api_log!("Device::drop {:?}", self as *const _);
        resource_log!("Drop {}", self.error_ident());

        // The timestamp normalizer is late-initialized, so it is not included in `DeviceResources`.
        if let Some(timestamp_normalizer) = self.timestamp_normalizer.take() {
            timestamp_normalizer.dispose(self.raw.as_ref());
        }

        // Transfer the rest of the resources back to `DeviceResources`, which cleans them
        // up for us.

        // SAFETY: We are in the Drop impl and we don't use self.zero_buffer anymore after this
        // point.
        let zero_buffer = unsafe { ManuallyDrop::take(&mut self.zero_buffer) };
        // SAFETY: We are in the Drop impl and we don't use self.empty_bgl anymore after this point.
        let empty_bgl = unsafe { ManuallyDrop::take(&mut self.empty_bgl) };
        // SAFETY: We are in the Drop impl and we don't use
        // self.default_external_texture_params_buffer anymore after this point.
        let default_external_texture_params_buffer =
            unsafe { ManuallyDrop::take(&mut self.default_external_texture_params_buffer) };
        // SAFETY: We are in the Drop impl and we don't use self.fence anymore after this point.
        let fence = unsafe { ManuallyDrop::take(&mut self.fence) };

        drop(DeviceResources {
            raw: self.raw.as_ref(),
            zero_buffer: Some(zero_buffer),
            empty_bgl: Some(empty_bgl),
            default_external_texture_params_buffer: Some(default_external_texture_params_buffer),
            fence: Some(fence),
            indirect_validation: self.indirect_validation.take(),
        });
    }
}

impl Device {
    pub fn features(&self) -> &wgt::Features {
        &self.features
    }

    pub fn limits(&self) -> &wgt::Limits {
        &self.limits
    }

    pub fn downlevel(&self) -> &wgt::DownlevelCapabilities {
        &self.downlevel
    }
}

impl Device {
    pub(crate) fn raw(&self) -> &dyn hal::DynDevice {
        self.raw.as_ref()
    }
    pub(crate) fn require_features(&self, feature: wgt::Features) -> Result<(), MissingFeatures> {
        if self.features.contains(feature) {
            Ok(())
        } else {
            Err(MissingFeatures(feature))
        }
    }

    pub(crate) fn require_downlevel_flags(
        &self,
        flags: wgt::DownlevelFlags,
    ) -> Result<(), MissingDownlevelFlags> {
        if self.downlevel.flags.contains(flags) {
            Ok(())
        } else {
            Err(MissingDownlevelFlags(flags))
        }
    }

    /// # Safety
    ///
    /// - See [wgpu::Device::start_graphics_debugger_capture][api] for details the safety.
    ///
    /// [api]: ../../wgpu/struct.Device.html#method.start_graphics_debugger_capture
    pub unsafe fn start_graphics_debugger_capture(&self) {
        api_log!("Device::start_graphics_debugger_capture");

        if !self.is_valid() {
            return;
        }
        unsafe { self.raw().start_graphics_debugger_capture() };
    }

    /// # Safety
    ///
    /// - See [wgpu::Device::stop_graphics_debugger_capture][api] for details the safety.
    ///
    /// [api]: ../../wgpu/struct.Device.html#method.stop_graphics_debugger_capture
    pub unsafe fn stop_graphics_debugger_capture(&self) {
        api_log!("Device::stop_graphics_debugger_capture");

        if !self.is_valid() {
            return;
        }
        unsafe { self.raw().stop_graphics_debugger_capture() };
    }
}

impl Device {
    pub(crate) fn new(
        raw_device: Box<dyn hal::DynDevice>,
        adapter: &Arc<Adapter>,
        desc: &DeviceDescriptor,
        instance_flags: wgt::InstanceFlags,
    ) -> Result<Self, DeviceError> {
        #[cfg(not(feature = "trace"))]
        match &desc.trace {
            wgt::Trace::Off => {}
            _ => {
                log::error!("wgpu-core feature 'trace' is not enabled");
            }
        };
        #[cfg(feature = "trace")]
        let trace: Option<Box<dyn trace::Trace + Send + Sync + 'static>> = match &desc.trace {
            wgt::Trace::Off => None,
            wgt::Trace::Directory(dir) => match trace::DiskTrace::new(dir.clone()) {
                Ok(mut trace) => {
                    trace::Trace::add(
                        &mut trace,
                        trace::Action::Init {
                            desc: wgt::DeviceDescriptor {
                                trace: wgt::Trace::Off,
                                ..desc.clone()
                            },
                            backend: adapter.backend(),
                        },
                    );
                    Some(Box::new(trace))
                }
                Err(e) => {
                    log::error!("Unable to start a trace in '{dir:?}': {e}");
                    None
                }
            },
            wgt::Trace::Memory => {
                let mut trace = trace::MemoryTrace::new();
                trace::Trace::add(
                    &mut trace,
                    trace::Action::Init {
                        desc: wgt::DeviceDescriptor {
                            trace: wgt::Trace::Off,
                            ..desc.clone()
                        },
                        backend: adapter.backend(),
                    },
                );
                Some(Box::new(trace))
            }
            // The enum is non_exhaustive, so we must have a fallback arm (that should be
            // unreachable in practice).
            t => {
                log::error!("unimplemented wgpu_types::Trace variant {t:?}");
                None
            }
        };

        let ordered_buffer_usages = adapter.raw.adapter.get_ordered_buffer_usages();
        let ordered_texture_usages = adapter.raw.adapter.get_ordered_texture_usages();

        let mut resources = DeviceResources {
            raw: raw_device.as_ref(),
            zero_buffer: None,
            empty_bgl: None,
            default_external_texture_params_buffer: None,
            fence: None,
            indirect_validation: None,
        };

        resources.fence =
            Some(unsafe { raw_device.create_fence() }.map_err(DeviceError::from_hal)?);

        let command_allocator = command::CommandAllocator::new();

        let rt_uses = if desc
            .required_features
            .intersects(wgt::Features::EXPERIMENTAL_RAY_QUERY)
        {
            wgt::BufferUses::TOP_LEVEL_ACCELERATION_STRUCTURE_INPUT
        } else {
            wgt::BufferUses::empty()
        };

        // Create zeroed buffer used for texture clears (and raytracing if required).
        resources.zero_buffer = Some(
            unsafe {
                raw_device.create_buffer(&hal::BufferDescriptor {
                    label: hal_label(Some("(wgpu internal) zero init buffer"), instance_flags),
                    size: ZERO_BUFFER_SIZE,
                    usage: wgt::BufferUses::COPY_SRC | wgt::BufferUses::COPY_DST | rt_uses,
                    memory_flags: hal::MemoryFlags::empty(),
                })
            }
            .map_err(DeviceError::from_hal)?,
        );

        resources.empty_bgl = Some(
            unsafe {
                raw_device.create_bind_group_layout(&hal::BindGroupLayoutDescriptor {
                    label: None,
                    flags: hal::BindGroupLayoutFlags::empty(),
                    entries: &[],
                })
            }
            .map_err(DeviceError::from_hal)?,
        );

        resources.default_external_texture_params_buffer = Some(
            unsafe {
                raw_device.create_buffer(&hal::BufferDescriptor {
                    label: hal_label(
                        Some("(wgpu internal) default external texture params buffer"),
                        instance_flags,
                    ),
                    size: size_of::<ExternalTextureParams>() as _,
                    usage: wgt::BufferUses::COPY_DST | wgt::BufferUses::UNIFORM,
                    memory_flags: hal::MemoryFlags::empty(),
                })
            }
            .map_err(DeviceError::from_hal)?,
        );

        // Cloned as we need them below anyway.
        let alignments = adapter.raw.capabilities.alignments.clone();
        let downlevel = adapter.raw.capabilities.downlevel.clone();
        let limits = &adapter.raw.capabilities.limits;

        let enable_indirect_validation = instance_flags
            .contains(wgt::InstanceFlags::VALIDATION_INDIRECT_CALL)
            && downlevel.flags.contains(
                wgt::DownlevelFlags::INDIRECT_EXECUTION | wgt::DownlevelFlags::COMPUTE_SHADERS,
            )
            && limits.max_storage_buffers_per_shader_stage >= 2;

        if enable_indirect_validation {
            resources.indirect_validation =
                Some(crate::indirect_validation::IndirectValidation::new(
                    raw_device.as_ref(),
                    &desc.required_limits,
                    &desc.required_features,
                    instance_flags,
                    adapter.backend(),
                )?);
        }

        // Error returns after this point could bypass resource cleanup.
        #[deny(clippy::question_mark_used)]
        {
            let zero_buffer = resources.zero_buffer.take().unwrap();
            let empty_bgl = resources.empty_bgl.take().unwrap();
            let default_external_texture_params_buffer = resources
                .default_external_texture_params_buffer
                .take()
                .unwrap();
            let fence = resources.fence.take().unwrap();
            let indirect_validation = resources.indirect_validation.take();
            drop(resources);

            Ok(Self {
                raw: raw_device,
                adapter: adapter.clone(),
                queue: OnceCellOrLock::new(),
                zero_buffer: ManuallyDrop::new(zero_buffer),
                empty_bgl: ManuallyDrop::new(empty_bgl),
                default_external_texture_params_buffer: ManuallyDrop::new(
                    default_external_texture_params_buffer,
                ),
                label: desc.label.to_string(),
                command_allocator,
                command_indices: RwLock::new(
                    rank::DEVICE_COMMAND_INDICES,
                    CommandIndices {
                        active_submission_index: 0,
                        // By starting at one, we can put the result in a NonZeroU64.
                        next_acceleration_structure_build_command_index: 1,
                    },
                ),
                last_successful_submission_index: AtomicU64::new(0),
                fence: ManuallyDrop::new(fence),
                snatchable_lock: unsafe { SnatchLock::new(rank::DEVICE_SNATCHABLE_LOCK) },
                valid: AtomicBool::new(true),
                device_lost_closure: Mutex::new(rank::DEVICE_LOST_CLOSURE, None),
                trackers: Mutex::new(
                    rank::DEVICE_TRACKERS,
                    DeviceTracker::new(ordered_buffer_usages, ordered_texture_usages),
                ),
                tracker_indices: TrackerIndexAllocators::new(),
                bgl_pool: ResourcePool::new(),
                #[cfg(feature = "trace")]
                trace: Mutex::new(rank::DEVICE_TRACE, trace),
                alignments,
                limits: desc.required_limits.clone(),
                features: desc.required_features,
                downlevel,
                ordered_buffer_usages,
                ordered_texture_usages,
                instance_flags,
                deferred_destroy: Mutex::new(rank::DEVICE_DEFERRED_DESTROY, Vec::new()),
                usage_scopes: Mutex::new(rank::DEVICE_USAGE_SCOPES, Default::default()),
                timestamp_normalizer: OnceCellOrLock::new(),
                indirect_validation,
                deferred_buffer_map_pending_closures: DeferredBufferMapPendingClosures::new(),
            })
        }
    }

    /// Initializes [`Device::default_external_texture_params_buffer`] with
    /// required values such that a [`TextureView`] bound to a
    /// [`wgt::BindingType::ExternalTexture`] binding point will be rendered
    /// correctly.
    fn init_default_external_texture_params_buffer(self: &Arc<Self>) -> Result<(), DeviceError> {
        let data = ExternalTextureParams {
            #[rustfmt::skip]
            yuv_conversion_matrix: [
                1.0, 0.0, 0.0, 0.0,
                0.0, 1.0, 0.0, 0.0,
                0.0, 0.0, 1.0, 0.0,
                0.0, 0.0, 0.0, 1.0,
            ],
            #[rustfmt::skip]
            gamut_conversion_matrix: [
                1.0, 0.0, 0.0, /* padding */ 0.0,
                0.0, 1.0, 0.0, /* padding */ 0.0,
                0.0, 0.0, 1.0, /* padding */ 0.0,
            ],
            src_transfer_function: Default::default(),
            dst_transfer_function: Default::default(),
            size: [0, 0],
            #[rustfmt::skip]
            sample_transform: [
                1.0, 0.0,
                0.0, 1.0,
                0.0, 0.0
            ],
            #[rustfmt::skip]
            load_transform: [
                1.0, 0.0,
                0.0, 1.0,
                0.0, 0.0
            ],
            num_planes: 1,
            _padding: Default::default(),
        };
        let mut staging_buffer =
            StagingBuffer::new(self, wgt::BufferSize::new(size_of_val(&data) as _).unwrap())?;
        staging_buffer.write(bytemuck::bytes_of(&data));
        let staging_buffer = staging_buffer.flush();

        let params_buffer = self.default_external_texture_params_buffer.as_ref();
        let queue = self.get_queue().unwrap();
        let mut pending_writes = queue.pending_writes.lock();

        unsafe {
            pending_writes
                .command_encoder
                .transition_buffers(&[hal::BufferBarrier {
                    buffer: params_buffer,
                    usage: hal::StateTransition {
                        from: wgt::BufferUses::MAP_WRITE,
                        to: wgt::BufferUses::COPY_DST,
                    },
                }]);
            pending_writes.command_encoder.copy_buffer_to_buffer(
                staging_buffer.raw(),
                params_buffer,
                &[hal::BufferCopy {
                    src_offset: 0,
                    dst_offset: 0,
                    size: staging_buffer.size,
                }],
            );
            pending_writes.consume(staging_buffer);
            pending_writes
                .command_encoder
                .transition_buffers(&[hal::BufferBarrier {
                    buffer: params_buffer,
                    usage: hal::StateTransition {
                        from: wgt::BufferUses::COPY_DST,
                        to: wgt::BufferUses::UNIFORM,
                    },
                }]);
        }

        Ok(())
    }

    pub fn late_init_resources_with_queue(self: &Arc<Self>) -> Result<(), RequestDeviceError> {
        let queue = self.get_queue().unwrap();

        let timestamp_normalizer = crate::timestamp_normalization::TimestampNormalizer::new(
            self,
            queue.get_timestamp_period(),
        )?;

        self.timestamp_normalizer
            .set(timestamp_normalizer)
            .unwrap_or_else(|_| panic!("Called late_init_resources_with_queue twice"));

        self.init_default_external_texture_params_buffer()?;

        Ok(())
    }

    /// Returns the backend this device is using.
    pub fn backend(&self) -> wgt::Backend {
        self.adapter.backend()
    }

    pub fn is_valid(&self) -> bool {
        self.valid.load(Ordering::Acquire)
    }

    pub fn check_is_valid(&self) -> Result<(), DeviceError> {
        if self.is_valid() {
            Ok(())
        } else {
            Err(DeviceError::Lost)
        }
    }

    /// Stop tracing and return the trace object.
    ///
    /// This is mostly useful for in-memory traces.
    #[cfg(feature = "trace")]
    pub fn take_trace(&self) -> Option<Box<dyn trace::Trace + Send + Sync + 'static>> {
        self.trace.lock().take()
    }

    /// Checks that we are operating within the memory budget reported by the native APIs.
    ///
    /// If we are not, the device gets invalidated.
    ///
    /// The budget might fluctuate over the lifetime of the application, so it should be checked
    /// somewhat frequently.
    pub fn lose_if_oom(&self) {
        let _ = self
            .raw()
            .check_if_oom()
            .map_err(|e| self.handle_hal_error(e));
    }

    pub fn handle_hal_error(&self, error: hal::DeviceError) -> DeviceError {
        match error {
            hal::DeviceError::OutOfMemory
            | hal::DeviceError::Lost
            | hal::DeviceError::Unexpected => {
                self.lose(&error.to_string());
            }
        }
        DeviceError::from_hal(error)
    }

    pub fn handle_hal_error_with_nonfatal_oom(&self, error: hal::DeviceError) -> DeviceError {
        match error {
            hal::DeviceError::OutOfMemory => DeviceError::from_hal(error),
            error => self.handle_hal_error(error),
        }
    }

    /// Run some destroy operations that were deferred.
    ///
    /// Destroying the resources requires taking a write lock on the device's snatch lock,
    /// so a good reason for deferring resource destruction is when we don't know for sure
    /// how risky it is to take the lock (typically, it shouldn't be taken from the drop
    /// implementation of a reference-counted structure).
    /// The snatch lock must not be held while this function is called.
    pub(crate) fn deferred_resource_destruction(&self) {
        // Note that the deferred_destroy list may contain duplicate entries.
        let deferred_destroy = mem::take(&mut *self.deferred_destroy.lock());
        for item in deferred_destroy {
            match item {
                DeferredDestroy::TextureViews(views) => {
                    for view in views {
                        let Some(view) = view.upgrade() else {
                            continue;
                        };
                        let Ok(view_state) = view.state() else {
                            continue;
                        };
                        let Some(raw_view) =
                            view_state.raw.snatch(&mut self.snatchable_lock.write())
                        else {
                            continue;
                        };

                        resource_log!("Destroy raw {}", view.error_ident());

                        unsafe {
                            self.raw().destroy_texture_view(raw_view);
                        }
                    }
                }
                DeferredDestroy::BindGroups(bind_groups) => {
                    for bind_group in bind_groups {
                        let Some(bind_group) = bind_group.upgrade() else {
                            continue;
                        };
                        let Ok(bind_group_state) = bind_group.state() else {
                            continue;
                        };
                        let Some(raw_bind_group) = bind_group_state
                            .raw
                            .snatch(&mut self.snatchable_lock.write())
                        else {
                            continue;
                        };

                        resource_log!("Destroy raw {}", bind_group.error_ident());

                        unsafe {
                            self.raw().destroy_bind_group(raw_bind_group);
                        }
                    }
                }
            }
        }
    }

    pub fn get_queue(&self) -> Option<Arc<Queue>> {
        self.queue.get().as_ref()?.upgrade()
    }

    pub fn set_queue(&self, queue: &Arc<Queue>) {
        assert!(self.queue.set(Arc::downgrade(queue)).is_ok());
    }

    /// Check device for freeable resources and completed buffer mappings.
    ///
    /// Return `queue_empty` indicating whether there are more queue submissions still in flight.
    pub fn poll(
        &self,
        poll_type: wgt::PollType<crate::SubmissionIndex>,
    ) -> Result<wgt::PollStatus, WaitIdleError> {
        api_log!("Device::poll {poll_type:?}");
        let (user_closures, result) = self.poll_and_return_closures(poll_type);
        user_closures.fire();
        result
    }

    /// Poll the device, returning any `UserClosures` that need to be executed.
    ///
    /// The caller must invoke the `UserClosures` even if this function returns
    /// an error. This is an internal helper, used by [`Device::poll`] and
    /// [`Instance::poll_all_devices`], so that `poll_all_devices` can invoke
    /// closures once after all devices have been polled.
    ///
    /// [`Instance::poll_all_devices`]: crate::instance::Instance::poll_all_devices
    pub(crate) fn poll_and_return_closures(
        &self,
        poll_type: wgt::PollType<crate::SubmissionIndex>,
    ) -> (UserClosures, Result<wgt::PollStatus, WaitIdleError>) {
        let snatch_guard = self.snatchable_lock.read();
        let maintain_result = self.maintain(poll_type, snatch_guard);

        self.lose_if_oom();

        // Some deferred destroys are scheduled in maintain so run this right after
        // to avoid holding on to them until the next device poll.
        self.deferred_resource_destruction();

        maintain_result
    }

    /// Check the current status of the GPU and process any submissions that have
    /// finished.
    ///
    /// The `poll_type` argument tells if this function should wait for a particular
    /// submission index to complete, or if it should just poll the current status.
    ///
    /// This will process _all_ completed submissions, even if the caller only asked
    /// us to poll to a given submission index.
    ///
    /// Return a pair `(closures, result)`, where:
    ///
    /// - `closures` is a list of callbacks that need to be invoked informing the user
    ///   about various things occurring. These happen and should be handled even if
    ///   this function returns an error, hence they are outside of the result.
    ///
    /// - `results` is a boolean indicating the result of the wait operation, including
    ///   if there was a timeout or a validation error.
    pub(crate) fn maintain<'this>(
        &'this self,
        poll_type: wgt::PollType<crate::SubmissionIndex>,
        snatch_guard: SnatchGuard,
    ) -> (UserClosures, Result<wgt::PollStatus, WaitIdleError>) {
        profiling::scope!("Device::maintain");

        let mut user_closures = UserClosures::default();

        self.deferred_buffer_map_pending_closures
            .swap(&mut user_closures.mappings);

        // If a wait was requested, determine which submission index to wait for.
        let wait_submission_index = match poll_type {
            wgt::PollType::Wait {
                submission_index: Some(submission_index),
                ..
            } => {
                let last_successful_submission_index = self
                    .last_successful_submission_index
                    .load(Ordering::Acquire);

                if submission_index > last_successful_submission_index {
                    let result = Err(WaitIdleError::WrongSubmissionIndex(
                        submission_index,
                        last_successful_submission_index,
                    ));

                    return (user_closures, result);
                }

                Some(submission_index)
            }
            wgt::PollType::Wait {
                submission_index: None,
                ..
            } => Some(
                self.last_successful_submission_index
                    .load(Ordering::Acquire),
            ),
            wgt::PollType::Poll => None,
        };

        // Wait for the submission index if requested.
        if let Some(target_submission_index) = wait_submission_index {
            log::trace!("Device::maintain: waiting for submission index {target_submission_index}");

            let wait_timeout = match poll_type {
                wgt::PollType::Wait { timeout, .. } => timeout,
                wgt::PollType::Poll => unreachable!(
                    "`wait_submission_index` index for poll type `Poll` should be None"
                ),
            };

            let wait_result = unsafe {
                self.raw()
                    .wait(self.fence.as_ref(), target_submission_index, wait_timeout)
            };

            // This error match is only about `DeviceErrors`. At this stage we do not care if
            // the wait succeeded or not, and the `Ok(bool)`` variant is ignored.
            if let Err(e) = wait_result {
                let hal_error: WaitIdleError = self.handle_hal_error(e).into();
                return (user_closures, Err(hal_error));
            }
        }

        // Get the currently finished submission index. This may be higher than the requested
        // wait, or it may be less than the requested wait if the wait failed.
        let fence_value_result = unsafe { self.raw().get_fence_value(self.fence.as_ref()) };
        let current_finished_submission = match fence_value_result {
            Ok(fence_value) => fence_value,
            Err(e) => {
                let hal_error: WaitIdleError = self.handle_hal_error(e).into();
                return (user_closures, Err(hal_error));
            }
        };

        // Prevent new commands from being submitted as we want to act on `queue_empty`.
        let command_indices = self.command_indices.read();
        // Check that the device is valid. This is combined with queue empty to decide whether
        // to destroy all resources. Queue.submit blocks on command indices being writable
        // and rejects if invalid so if the device in now invalid, and all submissions are
        // finished, there will be no more submissions.
        let device_valid = self.is_valid();
        drop(command_indices);

        // Maintain all finished submissions on the queue, updating the relevant user closures and
        // collecting if the queue is empty.
        //
        // We don't use the result of the wait here, as we want to progress forward as far as
        // possible and the wait could have been for submissions that finished long ago.
        let mut queue_empty = false;
        if let Some(queue) = self.get_queue() {
            let queue_result = queue.maintain(current_finished_submission, &snatch_guard);
            (
                user_closures.submissions,
                user_closures.mappings,
                user_closures.blas_compact_ready,
                queue_empty,
            ) = queue_result;
            // DEADLOCK PREVENTION: We must drop `snatch_guard` before `queue` goes out of scope.
            //
            // `Queue::drop` acquires the snatch guard. If we still hold it when `queue` is dropped
            // at the end of this block, we would deadlock. This can happen in the following
            // scenario:
            //
            // - Thread A calls `Device::maintain` while Thread B holds the last strong ref to the
            //   queue.
            // - Thread A calls `self.get_queue()`, obtaining a new strong ref, and enters this
            //   branch.
            // - Thread B drops its strong ref, making Thread A's ref the last one.
            // - When `queue` goes out of scope here, `Queue::drop` runs and tries to acquire the
            //   snatch guard — but Thread A (this thread) still holds it, causing a deadlock.
            drop(snatch_guard);
        } else {
            drop(snatch_guard);
        };

        // Based on the queue empty status, and the current finished submission index, determine
        // the result of the poll.
        let result = if queue_empty {
            if let Some(wait_submission_index) = wait_submission_index {
                // Assert to ensure that if we received a queue empty status, the fence shows the
                // correct value. This is defensive, as this should never be hit.
                assert!(
                    current_finished_submission >= wait_submission_index,
                    concat!(
                        "If the queue is empty, the current submission index ",
                        "({}) should be at least the wait submission index ({})",
                    ),
                    current_finished_submission,
                    wait_submission_index,
                );
            }

            Ok(wgt::PollStatus::QueueEmpty)
        } else if let Some(wait_submission_index) = wait_submission_index {
            // This is theoretically possible to succeed more than checking on the poll result
            // as submissions could have finished in the time between the timeout resolving,
            // the thread getting scheduled again, and us checking the fence value.
            if current_finished_submission >= wait_submission_index {
                Ok(wgt::PollStatus::WaitSucceeded)
            } else {
                Err(WaitIdleError::Timeout)
            }
        } else {
            Ok(wgt::PollStatus::Poll)
        };

        // Detect if we have been destroyed and now need to lose the device.
        //
        // If we are invalid (set at start of destroy) and our queue is empty,
        // and we have a DeviceLostClosure, return the closure to be called by
        // our caller. This will complete the steps for both destroy and for
        // "lose the device".
        let mut should_release_gpu_resource = false;
        if !device_valid && queue_empty {
            // We can release gpu resources associated with this device (but not
            // while holding the life_tracker lock).
            should_release_gpu_resource = true;

            // If we have a DeviceLostClosure, build an invocation with the
            // reason DeviceLostReason::Destroyed and no message.
            if let Some(device_lost_closure) = self.device_lost_closure.lock().take() {
                user_closures
                    .device_lost_invocations
                    .push(DeviceLostInvocation {
                        closure: device_lost_closure,
                        reason: DeviceLostReason::Destroyed,
                        message: String::new(),
                    });
            }
        }

        if should_release_gpu_resource {
            self.release_gpu_resources();
        }

        (user_closures, result)
    }

    pub fn create_buffer_inner(
        self: &Arc<Self>,
        desc: &resource::BufferDescriptor,
    ) -> Result<Arc<Buffer>, resource::CreateBufferError> {
        self.check_is_valid()?;

        if desc.size > self.limits.max_buffer_size {
            return Err(resource::CreateBufferError::MaxBufferSize {
                requested: desc.size,
                maximum: self.limits.max_buffer_size,
            });
        }

        if desc
            .usage
            .intersects(wgt::BufferUsages::BLAS_INPUT | wgt::BufferUsages::TLAS_INPUT)
        {
            self.require_features(wgt::Features::EXPERIMENTAL_RAY_QUERY)?;
        }

        if desc.usage.contains(wgt::BufferUsages::INDEX)
            && desc.usage.contains(
                wgt::BufferUsages::VERTEX
                    | wgt::BufferUsages::UNIFORM
                    | wgt::BufferUsages::INDIRECT
                    | wgt::BufferUsages::STORAGE,
            )
        {
            self.require_downlevel_flags(wgt::DownlevelFlags::UNRESTRICTED_INDEX_BUFFER)?;
        }

        if desc.usage.is_empty() || desc.usage.contains_unknown_bits() {
            return Err(resource::CreateBufferError::InvalidUsage(desc.usage));
        }

        if !self
            .features
            .contains(wgt::Features::MAPPABLE_PRIMARY_BUFFERS)
        {
            use wgt::BufferUsages as Bu;
            let write_mismatch = desc.usage.contains(Bu::MAP_WRITE)
                && !(Bu::MAP_WRITE | Bu::COPY_SRC).contains(desc.usage);
            let read_mismatch = desc.usage.contains(Bu::MAP_READ)
                && !(Bu::MAP_READ | Bu::COPY_DST).contains(desc.usage);
            if write_mismatch || read_mismatch {
                return Err(resource::CreateBufferError::UsageMismatch(desc.usage));
            }
        }

        let mut usage = conv::map_buffer_usage(desc.usage);

        if desc.usage.contains(wgt::BufferUsages::INDIRECT) {
            self.require_downlevel_flags(wgt::DownlevelFlags::INDIRECT_EXECUTION)?;
            // We are going to be reading from it, internally;
            // when validating the content of the buffer
            usage |= wgt::BufferUses::STORAGE_READ_ONLY | wgt::BufferUses::STORAGE_READ_WRITE;
        }

        if desc.usage.contains(wgt::BufferUsages::QUERY_RESOLVE) {
            usage |= TIMESTAMP_NORMALIZATION_BUFFER_USES;
        }

        if desc.mapped_at_creation {
            if !desc.size.is_multiple_of(wgt::COPY_BUFFER_ALIGNMENT) {
                return Err(resource::CreateBufferError::UnalignedSize);
            }
            if !desc.usage.contains(wgt::BufferUsages::MAP_WRITE) {
                // we are going to be copying into it, internally
                usage |= wgt::BufferUses::COPY_DST;
            }
        } else {
            // We are required to zero out (initialize) all memory. This is done
            // on demand using clear_buffer which requires write transfer usage!
            usage |= wgt::BufferUses::COPY_DST;
        }

        // The great thing about buffer sizes is there are so many to choose from!
        //  - The application may request an arbitrary byte-aligned size.
        //  - We add an extra byte if it's a vertex buffer so that we can simulate binding an
        //    empty range at the end of the buffer, which Vulkan does not natively allow.
        //  - The overall buffer size must be a non-zero multiple of 4.
        // Because initialization operates at multiples of 4 bytes, and initialization
        // tracking will never determine that it is necessary to initialize a region outside
        // the application-visible range, we eagerly zero-initialize anything beyond that.
        let actual_size = if desc.size == 0 {
            wgt::COPY_BUFFER_ALIGNMENT
        } else if desc.usage.contains(wgt::BufferUsages::VERTEX) {
            desc.size + 1
        } else {
            desc.size
        };
        let actual_size = align_to(actual_size, wgt::COPY_BUFFER_ALIGNMENT);
        let tail_start = desc.size & !(wgt::COPY_BUFFER_ALIGNMENT - 1);
        debug_assert!(actual_size - 4 <= tail_start);

        let hal_desc = hal::BufferDescriptor {
            label: desc.label.to_hal(self.instance_flags),
            size: actual_size,
            usage,
            memory_flags: hal::MemoryFlags::empty(),
        };
        let buffer = unsafe { self.raw().create_buffer(&hal_desc) }
            .map_err(|e| self.handle_hal_error_with_nonfatal_oom(e))?;

        let timestamp_normalization_bind_group = Snatchable::new(unsafe {
            // SAFETY: The size passed here must not overflow the buffer.
            self.timestamp_normalizer
                .get()
                .unwrap()
                .create_normalization_bind_group(
                    self,
                    &*buffer,
                    desc.label.as_deref(),
                    wgt::BufferSize::new(hal_desc.size).unwrap(),
                    desc.usage,
                )
        }?);

        let indirect_validation_bind_groups =
            self.create_indirect_validation_bind_groups(buffer.as_ref(), desc.size, desc.usage)?;

        let buffer = Buffer {
            state: ResourceState::Valid(BufferState {
                raw: Snatchable::new(buffer),
            }),
            device: self.clone(),
            usage: desc.usage,
            size: desc.size,
            initialization_status: RwLock::new(
                rank::BUFFER_INITIALIZATION_STATUS,
                BufferInitTracker::new(tail_start),
            ),
            map_state: Mutex::new(rank::BUFFER_MAP_STATE, resource::BufferMapState::Idle),
            label: desc.label.to_string(),
            tracking_data: TrackingData::new(self.tracker_indices.buffers.clone()),
            bind_groups: Mutex::new(rank::BUFFER_BIND_GROUPS, WeakVec::new()),
            timestamp_normalization_bind_group,
            indirect_validation_bind_groups,
        };

        let buffer = Arc::new(buffer);

        let init_buffer_tail = |buffer, snatch_guard| {
            let mapping = map_buffer(
                buffer,
                tail_start,
                actual_size - tail_start,
                HostMap::Write,
                snatch_guard,
            )?;
            let raw = buffer
                .raw(snatch_guard)
                .expect("newly-created buffer cannot be destroyed");
            unsafe {
                // SAFETY: The buffer tail is valid and aligned.
                mapping.ptr.cast::<u32>().as_ptr().write(0);
                if !mapping.is_coherent {
                    #[allow(clippy::single_range_in_vec_init)]
                    self.raw()
                        .flush_mapped_ranges(raw, &[tail_start..actual_size]);
                }
                self.raw().unmap_buffer(raw);
            }
            Ok::<_, resource::CreateBufferError>(())
        };

        let buffer_use = if !desc.mapped_at_creation {
            if tail_start == actual_size {
                // No tail init necessary
            } else if desc.usage.contains(wgt::BufferUsages::MAP_WRITE) {
                // Tail init by mapping
                let snatch_guard = self.snatchable_lock.read();
                init_buffer_tail(&buffer, &snatch_guard)?;
            } else if let Some(queue) = self.get_queue() {
                // Schedule tail init in pending writes
                let snatch_guard = self.snatchable_lock.read();
                let mut pending_writes = queue.pending_writes.lock();
                self.trackers
                    .lock()
                    .buffers
                    .insert_single(&buffer, wgt::BufferUses::COPY_DST);
                pending_writes.clear_buffer(
                    self,
                    &buffer,
                    tail_start..actual_size,
                    &snatch_guard,
                )?;
                return Ok(buffer);
            } else {
                // Queue is gone, buffer will not be used.
            }
            wgt::BufferUses::empty()
        } else if desc.usage.contains(wgt::BufferUsages::MAP_WRITE) {
            // Mapped-at-creation with MAP_WRITE. Tail init by mapping, then map the
            // application-visible portion of the buffer. Don't do both in the same
            // mapping, because the application shouldn't see the tail.
            let snatch_guard = self.snatchable_lock.read();
            if tail_start != actual_size {
                init_buffer_tail(&buffer, &snatch_guard)?;
            }
            let map_size = buffer.size;
            let mapping = if map_size == 0 {
                hal::BufferMapping {
                    ptr: core::ptr::NonNull::dangling(),
                    is_coherent: true,
                }
            } else {
                map_buffer(&buffer, 0, map_size, HostMap::Write, &snatch_guard)?
            };
            drop(snatch_guard);
            *buffer.map_state.lock() = resource::BufferMapState::Active {
                mapping,
                range: 0..map_size,
                host: HostMap::Write,
            };
            wgt::BufferUses::MAP_WRITE
        } else {
            // Mapped-at-creation without MAP_WRITE. Create and zero-initialize a staging
            // buffer for the entire buffer (including tail). It is okay to use a single
            // mapping, `get_mapped_range` will still only expose the application-visible
            // range. The application could corrupt the tail with an out-of-bounds write,
            // which may cause problems for its shaders if they read out-of-bounds, but
            // won't cause problems in wgpu.
            let mut staging_buffer =
                StagingBuffer::new(self, wgt::BufferSize::new(actual_size).unwrap())?;

            // Zero initialize memory and then mark the buffer as initialized
            // (it's guaranteed that this is the case by the time the buffer is usable)
            staging_buffer.write_zeros();
            buffer.initialization_status.write().drain(0..actual_size);

            *buffer.map_state.lock() = resource::BufferMapState::Init { staging_buffer };
            wgt::BufferUses::COPY_DST
        };

        // If we didn't add the buffer to the tracker already for initialization via
        // PendingWrites, do so now.

        self.trackers
            .lock()
            .buffers
            .insert_single(&buffer, buffer_use);

        Ok(buffer)
    }

    pub fn create_buffer(
        self: &Arc<Self>,
        desc: &resource::BufferDescriptor,
    ) -> (Arc<Buffer>, Option<resource::CreateBufferError>) {
        profiling::scope!("Device::create_buffer");

        let (buffer, error) = match self.create_buffer_inner(desc) {
            Ok(buffer) => (buffer, None),
            Err(e) => (Buffer::invalid(Arc::clone(self), desc), Some(e)),
        };
        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use trace::IntoTrace;
            let mut desc = desc.clone();
            let mapped_at_creation = mem::replace(&mut desc.mapped_at_creation, false);
            if mapped_at_creation && !desc.usage.contains(wgt::BufferUsages::MAP_WRITE) {
                desc.usage |= wgt::BufferUsages::COPY_DST;
            }
            trace.add(trace::Action::CreateBuffer(buffer.to_trace(), desc));
        }
        api_log!(
            "Device::create_buffer({:?}{}) -> {:?}",
            desc.label.as_deref().unwrap_or(""),
            if desc.mapped_at_creation {
                ", mapped_at_creation"
            } else {
                ""
            },
            Arc::as_ptr(&buffer)
        );
        (buffer, error)
    }

    #[cfg(feature = "replay")]
    pub fn set_buffer_data(
        self: &Arc<Self>,
        buffer: &Arc<Buffer>,
        offset: wgt::BufferAddress,
        data: &[u8],
    ) -> resource::BufferAccessResult {
        use crate::resource::RawResourceAccess;

        let device = &buffer.device;

        device.check_is_valid()?;
        buffer.check_usage(wgt::BufferUsages::MAP_WRITE)?;

        let last_submission = device
            .get_queue()
            .and_then(|queue| queue.lock_life().get_buffer_latest_submission_index(buffer));

        if let Some(last_submission) = last_submission {
            device.wait_for_submit(last_submission)?;
        }

        let snatch_guard = device.snatchable_lock.read();
        let raw_buf = buffer.try_raw(&snatch_guard)?;

        if offset > buffer.size {
            return Err(resource::BufferAccessError::OutOfBoundsStartOffsetOverrun {
                index: offset,
                max: buffer.size,
            });
        } else if buffer.size - offset < u64::try_from(data.len()).unwrap() {
            return Err(resource::BufferAccessError::OutOfBoundsEndOffsetOverrun {
                index: offset,
                size: u64::try_from(data.len()).unwrap(),
                max: buffer.size,
            });
        }

        let mapping = unsafe {
            device
                .raw()
                .map_buffer(raw_buf, offset..offset + u64::try_from(data.len()).unwrap())
        }
        .map_err(|e| device.handle_hal_error(e))?;

        unsafe { core::ptr::copy_nonoverlapping(data.as_ptr(), mapping.ptr.as_ptr(), data.len()) };

        if !mapping.is_coherent {
            #[allow(clippy::single_range_in_vec_init)]
            unsafe {
                device
                    .raw()
                    .flush_mapped_ranges(raw_buf, &[offset..offset + data.len() as u64])
            };
        }

        unsafe { device.raw().unmap_buffer(raw_buf) };

        Ok(())
    }

    /// # Safety
    ///
    /// - `hal_texture` must be created from `device_id` corresponding raw handle.
    /// - `hal_texture` must be created respecting `desc`
    /// - `hal_texture` must be initialized
    /// - The `initial_state` must match the actual driver-side state of
    ///   the wrapped resource at the moment of wrap.
    pub unsafe fn create_texture_from_hal(
        self: &Arc<Self>,
        hal_texture: Box<dyn hal::DynTexture>,
        desc: &resource::TextureDescriptor,
        initial_state: wgt::TextureUses,
    ) -> (Arc<Texture>, Option<resource::CreateTextureError>) {
        profiling::scope!("Device::create_texture_from_hal");

        let (texture, error) =
            match self.create_texture_from_hal_inner(hal_texture, desc, initial_state) {
                Ok(texture) => (texture, None),
                Err(e) => (Texture::invalid(self, desc), Some(e)),
            };

        // NB: Any change done through the raw texture handle will not be
        // recorded in the replay
        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            trace.add(trace::Action::CreateTexture(
                texture.to_trace(),
                desc.clone(),
            ));
        }

        api_log!(
            "Device::create_texture({desc:?}) -> {:?}",
            Arc::as_ptr(&texture)
        );

        (texture, error)
    }

    pub(crate) fn create_texture_from_hal_inner(
        self: &Arc<Self>,
        hal_texture: Box<dyn hal::DynTexture>,
        desc: &resource::TextureDescriptor,
        initial_state: wgt::TextureUses,
    ) -> Result<Arc<Texture>, resource::CreateTextureError> {
        self.check_is_valid()?;

        let format_features = self
            .describe_format_features(desc.format)
            .map_err(|error| resource::CreateTextureError::MissingFeatures(desc.format, error))?;

        unsafe { self.raw().add_raw_texture(&*hal_texture) };

        let texture = Texture::new(
            self,
            resource::TextureInner::Native { raw: hal_texture },
            conv::map_texture_usage(desc.usage, desc.format.into(), format_features.flags),
            desc,
            format_features,
            resource::TextureClearMode::None,
            false,
        );

        let texture = Arc::new(texture);

        self.trackers
            .lock()
            .textures
            .insert_single(&texture, initial_state);

        Ok(texture)
    }

    /// # Safety
    ///
    /// - `hal_buffer` must have been created on this device.
    /// - `hal_buffer` must have been created respecting `desc` (in particular, the size).
    /// - `hal_buffer` must be initialized.
    /// - `hal_buffer` must not have zero size.
    pub unsafe fn create_buffer_from_hal(
        self: &Arc<Self>,
        hal_buffer: Box<dyn hal::DynBuffer>,
        desc: &resource::BufferDescriptor,
    ) -> (Arc<Buffer>, Option<resource::CreateBufferError>) {
        profiling::scope!("Device::create_buffer");
        let (buffer, error) = match unsafe { self.create_buffer_from_hal_inner(hal_buffer, desc) } {
            Ok(buffer) => (buffer, None),
            Err(e) => (Buffer::invalid(Arc::clone(self), desc), Some(e)),
        };

        // NB: Any change done through the raw buffer handle will not be
        // recorded in the replay
        #[cfg(feature = "trace")]
        if let Some(trace) = self.trace.lock().as_mut() {
            use trace::IntoTrace;
            trace.add(trace::Action::CreateBuffer(buffer.to_trace(), desc.clone()));
        }
        api_log!("Device::create_buffer -> {:?}", Arc::as_ptr(&buffer));
        (buffer, error)
    }

    /// # Safety
    ///
    /// - `hal_buffer` must have been created on this device.
    /// - `hal_buffer` must have been created respecting `desc` (in particular, the size).
    /// - `hal_buffer` must be initialized.
    /// - `hal_buffer` must not have zero size.
    pub(crate) unsafe fn create_buffer_from_hal_inner(
        self: &Arc<Self>,
        hal_buffer: Box<dyn hal::DynBuffer>,
        desc: &resource::BufferDescriptor,
    ) -> Result<Arc<Buffer>, resource::CreateBufferError> {
        let timestamp_normalization_bind_group = Snatchable::new(unsafe {
            self.timestamp_normalizer
                .get()
                .unwrap()
                .create_normalization_bind_group(
                    self,
                    &*hal_buffer,
                    desc.label.as_deref(),
                    wgt::BufferSize::new(desc.size).unwrap(),
                    desc.usage,
                )?
        });

        let indirect_validation_bind_groups = self.create_indirect_validation_bind_groups(
            hal_buffer.as_ref(),
            desc.size,
            desc.usage,
        )?;

        unsafe { self.raw().add_raw_buffer(&*hal_buffer) };

        let buffer = Buffer {
            state: ResourceState::Valid(BufferState {
                raw: Snatchable::new(hal_buffer),
            }),
            device: self.clone(),
            usage: desc.usage,
            size: desc.size,
            initialization_status: RwLock::new(
                rank::BUFFER_INITIALIZATION_STATUS,
                BufferInitTracker::new(0),
            ),
            map_state: Mutex::new(rank::BUFFER_MAP_STATE, resource::BufferMapState::Idle),
            label: desc.label.to_string(),
            tracking_data: TrackingData::new(self.tracker_indices.buffers.clone()),
            bind_groups: Mutex::new(rank::BUFFER_BIND_GROUPS, WeakVec::new()),
            timestamp_normalization_bind_group,
            indirect_validation_bind_groups,
        };

        let buffer = Arc::new(buffer);

        self.trackers
            .lock()
            .buffers
            .insert_single(&buffer, wgt::BufferUses::empty());

        Ok(buffer)
    }

    fn create_indirect_validation_bind_groups(
        &self,
        raw_buffer: &dyn hal::DynBuffer,
        buffer_size: u64,
        usage: wgt::BufferUsages,
    ) -> Result<Snatchable<crate::indirect_validation::BindGroups>, resource::CreateBufferError>
    {
        if !usage.contains(wgt::BufferUsages::INDIRECT) {
            return Ok(Snatchable::empty());
        }

        let Some(ref indirect_validation) = self.indirect_validation else {
            return Ok(Snatchable::empty());
        };

        let bind_groups = crate::indirect_validation::BindGroups::new(
            indirect_validation,
            self,
            buffer_size,
            raw_buffer,
        )
        .map_err(resource::CreateBufferError::IndirectValidationBindGroup)?;

        if let Some(bind_groups) = bind_groups {
            Ok(Snatchable::new(bind_groups))
        } else {
            Ok(Snatchable::empty())
        }
    }

    /// Validate a texture descriptor.
    ///
    /// This applies the same validation as [`Self::create_texture`], without
    /// actually creating a texture.
    pub fn validate_texture_descriptor(
        self: &Arc<Self>,
        desc: &resource::TextureDescriptor,
    ) -> Result<(), resource::CreateTextureError> {
        self.validate_texture_descriptor_inner(desc)?;
        Ok(())
    }

    /// Validate a texture descriptor.
    ///
    /// Implements the [validating GPUTextureDescriptor] algorithm, with some
    /// `wgpu` extensions.
    ///
    /// [validating GPUTextureDescriptor]: https://www.w3.org/TR/webgpu/#abstract-opdef-validating-gputexturedescriptor
    fn validate_texture_descriptor_inner(
        self: &Arc<Self>,
        desc: &resource::TextureDescriptor,
    ) -> Result<(wgt::TextureFormatFeatures, Vec<TextureFormat>), resource::CreateTextureError>
    {
        use resource::{CreateTextureError, TextureDimensionError};

        self.check_is_valid()?;

        if desc.usage.is_empty() || desc.usage.contains_unknown_bits() {
            return Err(CreateTextureError::InvalidUsage(desc.usage));
        }

        conv::check_texture_dimension_size(
            desc.dimension,
            desc.size,
            desc.sample_count,
            &self.limits,
        )?;

        if desc.dimension != wgt::TextureDimension::D2 {
            // Depth textures can only be 2D
            if desc.format.is_depth_stencil_format() {
                return Err(CreateTextureError::InvalidDepthDimension(
                    desc.dimension,
                    desc.format,
                ));
            }
            // Transient textures can only be 2D
            if desc
                .usage
                .contains(wgt::TextureUsages::TRANSIENT_ATTACHMENT)
            {
                return Err(CreateTextureError::InvalidDimensionUsages(
                    wgt::TextureUsages::TRANSIENT_ATTACHMENT,
                    desc.dimension,
                ));
            }
        }

        if desc.dimension != wgt::TextureDimension::D2
            && desc.dimension != wgt::TextureDimension::D3
        {
            // Compressed textures can only be 2D or 3D
            if desc.format.is_compressed() {
                return Err(CreateTextureError::InvalidCompressedDimension(
                    desc.dimension,
                    desc.format,
                ));
            }

            // Renderable textures can only be 2D or 3D
            if desc.usage.contains(wgt::TextureUsages::RENDER_ATTACHMENT) {
                return Err(CreateTextureError::InvalidDimensionUsages(
                    wgt::TextureUsages::RENDER_ATTACHMENT,
                    desc.dimension,
                ));
            }
        }

        if desc.format.is_compressed() {
            let (block_width, block_height) = desc.format.block_dimensions();

            if !desc.size.width.is_multiple_of(block_width) {
                return Err(CreateTextureError::InvalidDimension(
                    TextureDimensionError::NotMultipleOfBlockWidth {
                        width: desc.size.width,
                        block_width,
                        format: desc.format,
                    },
                ));
            }

            if !desc.size.height.is_multiple_of(block_height) {
                return Err(CreateTextureError::InvalidDimension(
                    TextureDimensionError::NotMultipleOfBlockHeight {
                        height: desc.size.height,
                        block_height,
                        format: desc.format,
                    },
                ));
            }

            if desc.dimension == wgt::TextureDimension::D3 {
                // Only BCn formats with Sliced 3D feature can be used for 3D textures
                if desc.format.is_bcn() {
                    self.require_features(wgt::Features::TEXTURE_COMPRESSION_BC_SLICED_3D)
                        .map_err(|error| CreateTextureError::MissingFeatures(desc.format, error))?;
                } else if desc.format.is_astc() {
                    self.require_features(wgt::Features::TEXTURE_COMPRESSION_ASTC_SLICED_3D)
                        .map_err(|error| CreateTextureError::MissingFeatures(desc.format, error))?;
                } else {
                    return Err(CreateTextureError::InvalidCompressedDimension(
                        desc.dimension,
                        desc.format,
                    ));
                }
            }
        }

        let mips = desc.mip_level_count;
        let max_levels_allowed = desc.size.max_mips(desc.dimension).min(hal::MAX_MIP_LEVELS);
        if mips == 0 || mips > max_levels_allowed {
            return Err(CreateTextureError::InvalidMipLevelCount {
                requested: mips,
                maximum: max_levels_allowed,
            });
        }

        {
            let (mut width_multiple, mut height_multiple) = desc.format.size_multiple_requirement();

            if desc.format.is_multi_planar_format() {
                // TODO(https://github.com/gfx-rs/wgpu/issues/8491): fix
                // `mip_level_size` calculation for these formats and relax this
                // restriction.
                width_multiple <<= desc.mip_level_count.saturating_sub(1);
                height_multiple <<= desc.mip_level_count.saturating_sub(1);
            }

            if !desc.size.width.is_multiple_of(width_multiple) {
                return Err(CreateTextureError::InvalidDimension(
                    TextureDimensionError::WidthNotMultipleOf {
                        width: desc.size.width,
                        multiple: width_multiple,
                        format: desc.format,
                    },
                ));
            }

            if !desc.size.height.is_multiple_of(height_multiple) {
                return Err(CreateTextureError::InvalidDimension(
                    TextureDimensionError::HeightNotMultipleOf {
                        height: desc.size.height,
                        multiple: height_multiple,
                        format: desc.format,
                    },
                ));
            }
        }

        if desc
            .usage
            .contains(wgt::TextureUsages::TRANSIENT_ATTACHMENT)
        {
            if desc.usage
                != (wgt::TextureUsages::TRANSIENT_ATTACHMENT
                    | wgt::TextureUsages::RENDER_ATTACHMENT)
            {
                return Err(CreateTextureError::InvalidTransientTextureUsage(desc.usage));
            }

            if desc.mip_level_count != 1 {
                return Err(CreateTextureError::InvalidTransientTextureMipLevelCount(
                    desc.mip_level_count,
                ));
            }

            if desc.size.depth_or_array_layers != 1 {
                return Err(CreateTextureError::InvalidTransientTextureLayerCount(
                    desc.size.depth_or_array_layers,
                ));
            }

            if !desc.view_formats.is_empty() {
                return Err(CreateTextureError::InvalidTransientTextureViewFormats);
            }
        }

        let format_features = self
            .describe_format_features(desc.format)
            .map_err(|error| CreateTextureError::MissingFeatures(desc.format, error))?;

        if desc.sample_count > 1 {
            // <https://www.w3.org/TR/2025/CRD-webgpu-20251120/#:~:text=If%20descriptor%2EsampleCount%20%3E%201>
            //
            // Note that there are also some checks related to the sample count
            // in [`conv::check_texture_dimension_size`].

            if desc.mip_level_count != 1 {
                return Err(CreateTextureError::InvalidMipLevelCount {
                    requested: desc.mip_level_count,
                    maximum: 1,
                });
            }

            if desc.size.depth_or_array_layers != 1
                && !self.features.contains(wgt::Features::MULTISAMPLE_ARRAY)
            {
                return Err(CreateTextureError::InvalidDimension(
                    TextureDimensionError::MultisampledDepthOrArrayLayer(
                        desc.size.depth_or_array_layers,
                    ),
                ));
            }

            if desc.usage.contains(wgt::TextureUsages::STORAGE_BINDING) {
                return Err(CreateTextureError::InvalidMultisampledStorageBinding);
            }

            if !desc.usage.contains(wgt::TextureUsages::RENDER_ATTACHMENT) {
                return Err(CreateTextureError::MultisampledNotRenderAttachment);
            }

            if !format_features.flags.intersects(
                wgt::TextureFormatFeatureFlags::MULTISAMPLE_X4
                    | wgt::TextureFormatFeatureFlags::MULTISAMPLE_X2
                    | wgt::TextureFormatFeatureFlags::MULTISAMPLE_X8
                    | wgt::TextureFormatFeatureFlags::MULTISAMPLE_X16,
            ) {
                return Err(CreateTextureError::InvalidMultisampledFormat(desc.format));
            }

            if !format_features
                .flags
                .sample_count_supported(desc.sample_count)
            {
                return Err(CreateTextureError::InvalidSampleCount(
                    desc.sample_count,
                    desc.format,
                    desc.format
                        .guaranteed_format_features(self.features)
                        .flags
                        .supported_sample_counts(),
                    self.adapter
                        .get_texture_format_features(desc.format)
                        .flags
                        .supported_sample_counts(),
                ));
            };
        }

        let missing_allowed_usages = match desc.format.planes() {
            Some(planes) => {
                let mut planes_usages = wgt::TextureUsages::all();
                for plane in 0..planes {
                    let aspect = wgt::TextureAspect::from_plane(plane).unwrap();
                    let format = desc.format.aspect_specific_format(aspect).unwrap();
                    let format_features = self
                        .describe_format_features(format)
                        .map_err(|error| CreateTextureError::MissingFeatures(desc.format, error))?;

                    planes_usages &= format_features.allowed_usages;
                }

                desc.usage - planes_usages
            }
            None => desc.usage - format_features.allowed_usages,
        };

        if !missing_allowed_usages.is_empty() {
            // detect downlevel incompatibilities
            let wgpu_allowed_usages = desc
                .format
                .guaranteed_format_features(self.features)
                .allowed_usages;
            let wgpu_missing_usages = desc.usage - wgpu_allowed_usages;
            return Err(CreateTextureError::InvalidFormatUsages(
                missing_allowed_usages,
                desc.format,
                wgpu_missing_usages.is_empty(),
            ));
        }

        let mut hal_view_formats = Vec::new();
        for format in desc.view_formats.iter() {
            if desc.format == *format {
                continue;
            }
            if desc.format.remove_srgb_suffix() != format.remove_srgb_suffix() {
                return Err(CreateTextureError::InvalidViewFormat(*format, desc.format));
            }
            hal_view_formats.push(*format);
        }
        if !hal_view_formats.is_empty() {
            self.require_downlevel_flags(wgt::DownlevelFlags::VIEW_FORMATS)?;
        }

        Ok((format_features, hal_view_formats))
    }

    fn create_texture_inner(
        self: &Arc<Self>,
        desc: &resource::TextureDescriptor,
    ) -> Result<Arc<Texture>, resource::CreateTextureError> {
        let (format_features, hal_view_formats) = self.validate_texture_descriptor_inner(desc)?;

        let hal_usage = conv::map_texture_usage_for_texture(desc, &format_features);

        let hal_desc = hal::TextureDescriptor {
            label: desc.label.to_hal(self.instance_flags),
            size: desc.size,
            mip_level_count: desc.mip_level_count,
            sample_count: desc.sample_count,
            dimension: desc.dimension,
            format: desc.format,
            usage: hal_usage,
            memory_flags: hal::MemoryFlags::empty(),
            view_formats: hal_view_formats,
        };

        let raw_texture = unsafe { self.raw().create_texture(&hal_desc) }
            .map_err(|e| self.handle_hal_error_with_nonfatal_oom(e))?;

        let clear_mode = if hal_usage
            .intersects(wgt::TextureUses::DEPTH_STENCIL_WRITE | wgt::TextureUses::COLOR_TARGET)
            && desc.dimension == wgt::TextureDimension::D2
        {
            let (is_color, usage) = if desc.format.is_depth_stencil_format() {
                (false, wgt::TextureUses::DEPTH_STENCIL_WRITE)
            } else {
                (true, wgt::TextureUses::COLOR_TARGET)
            };

            let clear_label = hal_label(
                Some("(wgpu internal) clear texture view"),
                self.instance_flags,
            );

            let mut clear_views = SmallVec::new();
            for mip_level in 0..desc.mip_level_count {
                for array_layer in 0..desc.size.depth_or_array_layers {
                    macro_rules! push_clear_view {
                        ($format:expr, $aspect:expr) => {
                            let desc = hal::TextureViewDescriptor {
                                label: clear_label,
                                format: $format,
                                dimension: TextureViewDimension::D2,
                                usage,
                                range: wgt::ImageSubresourceRange {
                                    aspect: $aspect,
                                    base_mip_level: mip_level,
                                    mip_level_count: Some(1),
                                    base_array_layer: array_layer,
                                    array_layer_count: Some(1),
                                },
                            };
                            clear_views.push(ManuallyDrop::new(
                                unsafe {
                                    self.raw().create_texture_view(raw_texture.as_ref(), &desc)
                                }
                                .map_err(|e| self.handle_hal_error(e))?,
                            ));
                        };
                    }

                    if let Some(planes) = desc.format.planes() {
                        for plane in 0..planes {
                            let aspect = wgt::TextureAspect::from_plane(plane).unwrap();
                            let format = desc.format.aspect_specific_format(aspect).unwrap();
                            push_clear_view!(format, aspect);
                        }
                    } else {
                        push_clear_view!(desc.format, wgt::TextureAspect::All);
                    }
                }
            }
            resource::TextureClearMode::RenderPass {
                clear_views,
                is_color,
            }
        } else {
            resource::TextureClearMode::BufferCopy
        };

        let texture = Texture::new(
            self,
            resource::TextureInner::Native { raw: raw_texture },
            hal_usage,
            desc,
            format_features,
            clear_mode,
            true,
        );

        let texture = Arc::new(texture);

        self.trackers
            .lock()
            .textures
            .insert_single(&texture, wgt::TextureUses::UNINITIALIZED);

        Ok(texture)
    }

    pub fn create_texture(
        self: &Arc<Self>,
        desc: &resource::TextureDescriptor,
    ) -> (Arc<Texture>, Option<resource::CreateTextureError>) {
        profiling::scope!("Device::create_texture");
        let (texture, error) = match self.create_texture_inner(desc) {
            Ok(texture) => (texture, None),
            Err(e) => {
                let texture = Texture::invalid(self, desc);
                (texture, Some(e))
            }
        };
        api_log!(
            "Device::create_texture({desc:?}) -> {:?}",
            Arc::as_ptr(&texture)
        );

        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use crate::device::trace::IntoTrace as _;

            trace.add(trace::Action::CreateTexture(
                texture.to_trace(),
                desc.clone(),
            ));
        }
        (texture, error)
    }

    /// Creates a texture that is guaranteed to be invalid
    pub fn create_texture_error(
        self: &Arc<Self>,
        desc: &resource::TextureDescriptor,
    ) -> Arc<Texture> {
        let texture = Texture::invalid(self, desc);
        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use crate::device::trace::IntoTrace as _;

            trace.add(trace::Action::CreateTextureError(
                texture.to_trace(),
                desc.clone(),
            ));
        }
        texture
    }

    pub fn create_external_texture(
        self: &Arc<Self>,
        desc: &resource::ExternalTextureDescriptor,
        planes: &[Arc<TextureView>],
    ) -> (
        Arc<ExternalTexture>,
        Option<resource::CreateExternalTextureError>,
    ) {
        profiling::scope!("Device::create_external_texture");

        let (external_texture, error) = match self.create_external_texture_inner(desc, planes) {
            Ok(external_texture) => (external_texture, None),
            Err(e) => (ExternalTexture::invalid(Arc::clone(self), desc), Some(e)),
        };

        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use crate::device::trace;
            use trace::IntoTrace as _;

            let planes = Box::from(
                planes
                    .iter()
                    .map(|plane| plane.to_trace())
                    .collect::<Vec<_>>(),
            );
            trace.add(trace::Action::CreateExternalTexture {
                id: external_texture.to_trace(),
                desc: desc.clone(),
                planes,
            });
        }

        api_log!(
            "Device::create_external_texture({desc:?}) -> {:?}",
            Arc::as_ptr(&external_texture)
        );

        (external_texture, error)
    }

    pub(crate) fn create_external_texture_inner(
        self: &Arc<Self>,
        desc: &resource::ExternalTextureDescriptor,
        planes: &[Arc<TextureView>],
    ) -> Result<Arc<ExternalTexture>, resource::CreateExternalTextureError> {
        use resource::CreateExternalTextureError;
        self.require_features(wgt::Features::EXTERNAL_TEXTURE)?;
        self.check_is_valid()?;

        if desc.num_planes() != planes.len() {
            return Err(CreateExternalTextureError::IncorrectPlaneCount {
                format: desc.format,
                expected: desc.num_planes(),
                provided: planes.len(),
            });
        }

        let planes = planes
            .iter()
            .enumerate()
            .map(|(i, plane)| {
                if plane.samples != 1 {
                    return Err(CreateExternalTextureError::InvalidPlaneMultisample(
                        plane.samples,
                    ));
                }

                let sample_type = plane
                    .desc
                    .format
                    .sample_type(Some(plane.desc.range.aspect), Some(self.features))
                    .unwrap();
                if !matches!(sample_type, TextureSampleType::Float { filterable: true }) {
                    return Err(CreateExternalTextureError::InvalidPlaneSampleType {
                        format: plane.desc.format,
                        sample_type,
                    });
                }

                if plane.desc.dimension != TextureViewDimension::D2 {
                    return Err(CreateExternalTextureError::InvalidPlaneDimension(
                        plane.desc.dimension,
                    ));
                }

                let expected_components = match desc.format {
                    wgt::ExternalTextureFormat::Rgba => 4,
                    wgt::ExternalTextureFormat::Nv12 => match i {
                        0 => 1,
                        1 => 2,
                        _ => unreachable!(),
                    },
                    wgt::ExternalTextureFormat::Yu12 => 1,
                };
                if plane.desc.format.components() != expected_components {
                    return Err(CreateExternalTextureError::InvalidPlaneFormat {
                        format: desc.format,
                        plane: i,
                        expected: expected_components,
                        provided: plane.desc.format,
                    });
                }

                plane.check_usage(wgt::TextureUsages::TEXTURE_BINDING)?;
                Ok(plane.clone())
            })
            .collect::<Result<_, _>>()?;

        let params_data = ExternalTextureParams::from_desc(desc);
        let label = desc.label.as_ref().map(|l| alloc::format!("{l} params"));
        let params_desc = resource::BufferDescriptor {
            label: label.map(Cow::Owned),
            size: size_of_val(&params_data) as wgt::BufferAddress,
            usage: wgt::BufferUsages::UNIFORM | wgt::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        };
        let params = self.create_buffer_inner(&params_desc)?;
        self.get_queue().unwrap().write_buffer(
            params.clone(),
            0,
            bytemuck::bytes_of(&params_data),
        )?;

        let external_texture = ExternalTexture {
            state: ResourceState::Valid(ExternalTextureState { params }),
            device: self.clone(),
            planes,
            label: desc.label.to_string(),
            tracking_data: TrackingData::new(self.tracker_indices.external_textures.clone()),
        };
        let external_texture = Arc::new(external_texture);

        Ok(external_texture)
    }

    pub fn create_sampler(
        self: &Arc<Self>,
        desc: &resource::SamplerDescriptor,
    ) -> (Arc<Sampler>, Option<resource::CreateSamplerError>) {
        profiling::scope!("Device::create_sampler");

        let (sampler, error) = match self.create_sampler_inner(desc) {
            Ok(sampler) => (sampler, None),
            Err(e) => (Sampler::invalid(Arc::clone(self), desc), Some(e)),
        };

        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use crate::device::trace::{Action, IntoTrace as _};
            trace.add(Action::CreateSampler(sampler.to_trace(), desc.clone()));
        }

        api_log!("Device::create_sampler -> {:?}", Arc::as_ptr(&sampler));

        (sampler, error)
    }

    pub(crate) fn create_sampler_inner(
        self: &Arc<Self>,
        desc: &resource::SamplerDescriptor,
    ) -> Result<Arc<Sampler>, resource::CreateSamplerError> {
        self.check_is_valid()?;

        if desc
            .address_modes
            .iter()
            .any(|am| am == &wgt::AddressMode::ClampToBorder)
        {
            self.require_features(wgt::Features::ADDRESS_MODE_CLAMP_TO_BORDER)?;
        }

        if desc.border_color == Some(wgt::SamplerBorderColor::Zero) {
            self.require_features(wgt::Features::ADDRESS_MODE_CLAMP_TO_ZERO)?;
        }

        if desc.lod_min_clamp < 0.0 {
            return Err(resource::CreateSamplerError::InvalidLodMinClamp(
                desc.lod_min_clamp,
            ));
        }
        if desc.lod_max_clamp < desc.lod_min_clamp {
            return Err(resource::CreateSamplerError::InvalidLodMaxClamp {
                lod_min_clamp: desc.lod_min_clamp,
                lod_max_clamp: desc.lod_max_clamp,
            });
        }

        if desc.anisotropy_clamp < 1 {
            return Err(resource::CreateSamplerError::InvalidAnisotropy(
                desc.anisotropy_clamp,
            ));
        }

        if desc.anisotropy_clamp != 1 {
            if !matches!(desc.min_filter, wgt::FilterMode::Linear) {
                return Err(
                    resource::CreateSamplerError::InvalidFilterModeWithAnisotropy {
                        filter_type: resource::SamplerFilterErrorType::MinFilter,
                        filter_mode: desc.min_filter,
                        anisotropic_clamp: desc.anisotropy_clamp,
                    },
                );
            }
            if !matches!(desc.mag_filter, wgt::FilterMode::Linear) {
                return Err(
                    resource::CreateSamplerError::InvalidFilterModeWithAnisotropy {
                        filter_type: resource::SamplerFilterErrorType::MagFilter,
                        filter_mode: desc.mag_filter,
                        anisotropic_clamp: desc.anisotropy_clamp,
                    },
                );
            }
            if !matches!(desc.mipmap_filter, wgt::MipmapFilterMode::Linear) {
                return Err(
                    resource::CreateSamplerError::InvalidMipmapFilterModeWithAnisotropy {
                        filter_type: resource::SamplerFilterErrorType::MipmapFilter,
                        filter_mode: desc.mipmap_filter,
                        anisotropic_clamp: desc.anisotropy_clamp,
                    },
                );
            }
        }

        let anisotropy_clamp = if self
            .downlevel
            .flags
            .contains(wgt::DownlevelFlags::ANISOTROPIC_FILTERING)
        {
            // Clamp anisotropy clamp to [1, 16] per the wgpu-hal interface
            desc.anisotropy_clamp.min(16)
        } else {
            // If it isn't supported, set this unconditionally to 1
            1
        };

        //TODO: check for wgt::DownlevelFlags::COMPARISON_SAMPLERS

        let hal_desc = hal::SamplerDescriptor {
            label: desc.label.to_hal(self.instance_flags),
            address_modes: desc.address_modes,
            mag_filter: desc.mag_filter,
            min_filter: desc.min_filter,
            mipmap_filter: desc.mipmap_filter,
            lod_clamp: desc.lod_min_clamp..desc.lod_max_clamp,
            compare: desc.compare,
            anisotropy_clamp,
            border_color: desc.border_color,
        };

        let raw = unsafe { self.raw().create_sampler(&hal_desc) }
            .map_err(|e| self.handle_hal_error_with_nonfatal_oom(e))?;

        let sampler = Sampler {
            raw: ResourceState::Valid(raw),
            device: self.clone(),
            label: desc.label.to_string(),
            tracking_data: TrackingData::new(self.tracker_indices.samplers.clone()),
            comparison: desc.compare.is_some(),
            filtering: desc.min_filter == wgt::FilterMode::Linear
                || desc.mag_filter == wgt::FilterMode::Linear
                || desc.mipmap_filter == wgt::MipmapFilterMode::Linear,
        };

        let sampler = Arc::new(sampler);

        Ok(sampler)
    }

    pub fn create_shader_module<'a>(
        self: &Arc<Self>,
        desc: &pipeline::ShaderModuleDescriptor<'a>,
        source: pipeline::ShaderModuleSource<'a>,
    ) -> (
        Arc<pipeline::ShaderModule>,
        Option<pipeline::CreateShaderModuleError>,
    ) {
        profiling::scope!("Device::create_shader_module");
        #[cfg(feature = "trace")]
        let data = self.trace.lock().as_mut().map(|trace| {
            use crate::device::trace::DataKind;

            match source {
                #[cfg(feature = "wgsl")]
                pipeline::ShaderModuleSource::Wgsl(ref code) => {
                    trace.make_binary(DataKind::Wgsl, code.as_bytes())
                }
                #[cfg(feature = "glsl")]
                pipeline::ShaderModuleSource::Glsl(ref code, _) => {
                    trace.make_binary(DataKind::Glsl, code.as_bytes())
                }
                #[cfg(feature = "spirv")]
                pipeline::ShaderModuleSource::SpirV(ref code, _) => {
                    trace.make_binary(DataKind::Spv, bytemuck::cast_slice::<u32, u8>(code))
                }
                pipeline::ShaderModuleSource::Naga(ref module) => {
                    let string =
                        ron::ser::to_string_pretty(module, ron::ser::PrettyConfig::default())
                            .unwrap();
                    trace.make_binary(DataKind::Ron, string.as_bytes())
                }
                pipeline::ShaderModuleSource::Dummy(_) => {
                    panic!("found `ShaderModuleSource::Dummy`")
                }
            }
        });
        let (shader, error) = match self.create_shader_module_inner(desc, source) {
            Ok(shader) => (shader, None),
            Err(e) => {
                let shader =
                    pipeline::ShaderModule::invalid(Arc::clone(self), desc.label.to_string());
                (shader, Some(e))
            }
        };
        api_log!("Device::create_shader_module -> {:?}", Arc::as_ptr(&shader));

        #[cfg(feature = "trace")]
        if let Some(data) = data {
            // We don't need these two operations with the trace to be atomic.

            use crate::device::trace::IntoTrace as _;
            self.trace
                .lock()
                .as_mut()
                .expect("trace went away during create_shader_module?")
                .add(trace::Action::CreateShaderModule {
                    id: shader.to_trace(),
                    desc: desc.clone(),
                    data,
                });
        };
        (shader, error)
    }

    pub(crate) fn create_shader_module_inner<'a>(
        self: &Arc<Self>,
        desc: &pipeline::ShaderModuleDescriptor<'a>,
        source: pipeline::ShaderModuleSource<'a>,
    ) -> Result<Arc<pipeline::ShaderModule>, pipeline::CreateShaderModuleError> {
        self.check_is_valid()?;

        let (module, source) = match source {
            #[cfg(feature = "wgsl")]
            pipeline::ShaderModuleSource::Wgsl(code) => {
                profiling::scope!("naga::front::wgsl::parse");
                let capabilities = crate::device::features_to_naga_capabilities(
                    self.features,
                    self.downlevel.flags,
                );
                let mut options = naga::front::wgsl::Options::new();
                options.capabilities = capabilities;
                let mut frontend = naga::front::wgsl::Frontend::new_with_options(options);
                let module = frontend.parse(&code).map_err(|inner| {
                    pipeline::CreateShaderModuleError::Parsing(naga::error::ShaderError {
                        source: code.to_string(),
                        label: desc.label.as_ref().map(|l| l.to_string()),
                        inner: Box::new(inner),
                    })
                })?;
                (Cow::Owned(module), code.into_owned())
            }
            #[cfg(feature = "spirv")]
            pipeline::ShaderModuleSource::SpirV(spv, options) => {
                let parser = naga::front::spv::Frontend::new(spv.iter().cloned(), &options);
                profiling::scope!("naga::front::spv::Frontend");
                let module = parser.parse().map_err(|inner| {
                    pipeline::CreateShaderModuleError::ParsingSpirV(naga::error::ShaderError {
                        source: String::new(),
                        label: desc.label.as_ref().map(|l| l.to_string()),
                        inner: Box::new(inner),
                    })
                })?;
                (Cow::Owned(module), String::new())
            }
            #[cfg(feature = "glsl")]
            pipeline::ShaderModuleSource::Glsl(code, options) => {
                let mut parser = naga::front::glsl::Frontend::default();
                profiling::scope!("naga::front::glsl::Frontend.parse");
                let module = parser.parse(&options, &code).map_err(|inner| {
                    pipeline::CreateShaderModuleError::ParsingGlsl(naga::error::ShaderError {
                        source: code.to_string(),
                        label: desc.label.as_ref().map(|l| l.to_string()),
                        inner: Box::new(inner),
                    })
                })?;
                (Cow::Owned(module), code.into_owned())
            }
            pipeline::ShaderModuleSource::Naga(module) => (module, String::new()),
            pipeline::ShaderModuleSource::Dummy(_) => panic!("found `ShaderModuleSource::Dummy`"),
        };
        for (_, var) in module.global_variables.iter() {
            match var.binding {
                Some(br) if br.group >= self.limits.max_bind_groups => {
                    return Err(pipeline::CreateShaderModuleError::InvalidGroupIndex {
                        bind: br,
                        group: br.group,
                        limit: self.limits.max_bind_groups,
                    });
                }
                _ => continue,
            };
        }

        profiling::scope!("naga::validate");
        let debug_source =
            if self.instance_flags.contains(wgt::InstanceFlags::DEBUG) && !source.is_empty() {
                Some(hal::DebugSource {
                    file_name: Cow::Owned(
                        desc.label
                            .as_ref()
                            .map_or("shader".to_string(), |l| l.to_string()),
                    ),
                    source_code: Cow::Owned(source.clone()),
                })
            } else {
                None
            };

        let info = create_validator(
            self.features,
            self.downlevel.flags,
            naga::valid::ValidationFlags::all(),
        )
        .validate(&module)
        .map_err(|inner| {
            pipeline::CreateShaderModuleError::Validation(naga::error::ShaderError {
                source,
                label: desc.label.as_ref().map(|l| l.to_string()),
                inner,
            })
        })?;

        let interface = validation::Interface::new(&module, &info, self.limits.clone());
        let hal_shader = hal::ShaderInput::Naga(hal::NagaShader {
            module,
            info,
            debug_source,
        });
        let hal_desc = hal::ShaderModuleDescriptor {
            label: desc.label.to_hal(self.instance_flags),
            runtime_checks: desc.runtime_checks,
        };
        let raw = match unsafe { self.raw().create_shader_module(&hal_desc, hal_shader) } {
            Ok(raw) => raw,
            Err(error) => {
                return Err(match error {
                    hal::ShaderError::Device(error) => {
                        pipeline::CreateShaderModuleError::Device(self.handle_hal_error(error))
                    }
                    hal::ShaderError::Compilation(ref msg) => {
                        log::error!("Shader error: {msg}");
                        pipeline::CreateShaderModuleError::Generation
                    }
                })
            }
        };

        let module = pipeline::ShaderModule {
            state: ResourceState::Valid(pipeline::ShaderModuleState {
                raw,
                interface: ShaderMetaData::Interface(interface),
            }),
            device: self.clone(),
            label: desc.label.to_string(),
        };

        let module = Arc::new(module);

        Ok(module)
    }

    /// # Safety
    ///
    /// This function passes source code or binary to the backend as-is and can potentially result in a
    /// driver crash.
    pub unsafe fn create_shader_module_passthrough<'a>(
        self: &Arc<Self>,
        desc: &pipeline::ShaderModuleDescriptorPassthrough<'a>,
    ) -> (
        Arc<pipeline::ShaderModule>,
        Option<pipeline::CreateShaderModuleError>,
    ) {
        profiling::scope!("Device::create_shader_module_passthrough");

        let (shader, error) = match unsafe { self.create_shader_module_passthrough_inner(desc) } {
            Ok(shader) => (shader, None),
            Err(e) => {
                let shader =
                    pipeline::ShaderModule::invalid(Arc::clone(self), desc.label.to_string());
                (shader, Some(e))
            }
        };
        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use crate::device::trace::{DataKind, IntoTrace as _};

            let mut file_names = Vec::new();
            for (data, kind) in [
                (
                    desc.spirv.as_ref().map(|a| bytemuck::cast_slice(a)),
                    DataKind::Spv,
                ),
                (desc.dxil.as_deref(), DataKind::Dxil),
                (desc.hlsl.as_ref().map(|a| a.as_bytes()), DataKind::Hlsl),
                (desc.metallib.as_deref(), DataKind::MetalLib),
                (desc.msl.as_ref().map(|a| a.as_bytes()), DataKind::Msl),
                (desc.glsl.as_ref().map(|a| a.as_bytes()), DataKind::Glsl),
                (desc.wgsl.as_ref().map(|a| a.as_bytes()), DataKind::Wgsl),
            ] {
                if let Some(data) = data {
                    file_names.push(trace.make_binary(kind, data));
                }
            }
            trace.add(trace::Action::CreateShaderModulePassthrough {
                id: shader.to_trace(),
                data: file_names,
                label: desc.label.clone(),
                entry_points: desc.entry_points.clone(),
            });
        };
        api_log!(
            "Device::create_shader_module_spirv -> {:?}",
            Arc::as_ptr(&shader)
        );
        (shader, error)
    }

    pub(crate) unsafe fn create_shader_module_passthrough_inner<'a>(
        self: &Arc<Self>,
        descriptor: &pipeline::ShaderModuleDescriptorPassthrough<'a>,
    ) -> Result<Arc<pipeline::ShaderModule>, pipeline::CreateShaderModuleError> {
        self.check_is_valid()?;
        self.require_features(wgt::Features::PASSTHROUGH_SHADERS)?;

        // Mainly important for GLSL or SPIR-V or DXIL, which each take exactly 1 entry point.
        if (descriptor.dxil.is_some() || descriptor.glsl.is_some())
            && descriptor.entry_points.len() != 1
        {
            return Err(pipeline::CreateShaderModuleError::IncorrectPassthroughEntryPointCount);
        }

        let entry_point_hashmap = || {
            descriptor
                .entry_points
                .iter()
                .map(|e| (e.name.to_string(), e.workgroup_size))
                .collect()
        };

        let hal_shader = match self.backend() {
            wgt::Backend::Vulkan => hal::ShaderInput::SpirV(
                descriptor
                    .spirv
                    .as_ref()
                    .ok_or(pipeline::CreateShaderModuleError::NotCompiledForBackend)?,
            ),
            wgt::Backend::Dx12 => {
                if let Some(dxil) = &descriptor.dxil {
                    hal::ShaderInput::Dxil { shader: dxil }
                } else if let Some(hlsl) = &descriptor.hlsl {
                    hal::ShaderInput::Hlsl { shader: hlsl }
                } else {
                    return Err(pipeline::CreateShaderModuleError::NotCompiledForBackend);
                }
            }
            wgt::Backend::Metal => {
                if let Some(metallib) = &descriptor.metallib {
                    hal::ShaderInput::MetalLib {
                        file: metallib,
                        num_workgroups: entry_point_hashmap(),
                    }
                } else if let Some(msl) = &descriptor.msl {
                    hal::ShaderInput::Msl {
                        shader: msl,
                        num_workgroups: entry_point_hashmap(),
                    }
                } else {
                    return Err(pipeline::CreateShaderModuleError::NotCompiledForBackend);
                }
            }
            wgt::Backend::Gl => hal::ShaderInput::Glsl {
                shader: descriptor
                    .glsl
                    .as_ref()
                    .ok_or(pipeline::CreateShaderModuleError::NotCompiledForBackend)?,
            },
            wgt::Backend::Noop => {
                return Err(pipeline::CreateShaderModuleError::NotCompiledForBackend)
            }
            wgt::Backend::BrowserWebGpu => unreachable!(),
        };

        let hal_desc = hal::ShaderModuleDescriptor {
            label: descriptor.label.to_hal(self.instance_flags),
            runtime_checks: wgt::ShaderRuntimeChecks::unchecked(),
        };

        let raw = match unsafe { self.raw().create_shader_module(&hal_desc, hal_shader) } {
            Ok(raw) => raw,
            Err(error) => {
                return Err(match error {
                    hal::ShaderError::Device(error) => {
                        pipeline::CreateShaderModuleError::Device(self.handle_hal_error(error))
                    }
                    hal::ShaderError::Compilation(ref msg) => {
                        log::error!("Shader error: {msg}");
                        pipeline::CreateShaderModuleError::Generation
                    }
                })
            }
        };

        let module = pipeline::ShaderModule {
            state: ResourceState::Valid(pipeline::ShaderModuleState {
                raw,
                interface: ShaderMetaData::Passthrough(PassthroughInterface {
                    entry_point_names: descriptor
                        .entry_points
                        .iter()
                        .map(|e| e.name.to_string())
                        .collect(),
                }),
            }),
            device: self.clone(),
            label: descriptor.label.to_string(),
        };

        Ok(Arc::new(module))
    }

    pub fn create_command_encoder(
        self: &Arc<Self>,
        desc: &wgt::CommandEncoderDescriptor<crate::Label>,
    ) -> (Arc<command::CommandEncoder>, Option<DeviceError>) {
        profiling::scope!("Device::create_command_encoder");

        let (cmd_enc, error) = match self.create_command_encoder_inner(&desc.label) {
            Ok(cmd_enc) => (cmd_enc, None),
            Err(e) => (
                command::CommandEncoder::new_invalid(self, &desc.label, e.clone().into()),
                Some(e),
            ),
        };

        api_log!(
            "Device::create_command_encoder -> {:?}",
            Arc::as_ptr(&cmd_enc)
        );

        (cmd_enc, error)
    }

    pub(crate) fn create_command_encoder_inner(
        self: &Arc<Self>,
        label: &crate::Label,
    ) -> Result<Arc<command::CommandEncoder>, DeviceError> {
        self.check_is_valid()?;

        let queue = self.get_queue().unwrap();

        let encoder = self
            .command_allocator
            .acquire_encoder(self.raw(), queue.raw())
            .map_err(|e| self.handle_hal_error(e))?;

        let cmd_enc = command::CommandEncoder::new(encoder, self, label);

        let cmd_enc = Arc::new(cmd_enc);

        Ok(cmd_enc)
    }

    pub fn create_render_bundle_encoder(
        self: &Arc<Self>,
        desc: &command::RenderBundleEncoderDescriptor,
    ) -> (
        Box<command::RenderBundleEncoder>,
        Option<command::CreateRenderBundleError>,
    ) {
        profiling::scope!("Device::create_render_bundle_encoder");
        api_log!("Device::create_render_bundle_encoder");
        let (encoder, error) = match command::RenderBundleEncoder::new(self, desc) {
            Ok(encoder) => (encoder, None),
            Err(e) => (command::RenderBundleEncoder::dummy(self), Some(e)),
        };
        (Box::new(encoder), error)
    }

    /// Generate information about late-validated buffer bindings for pipelines.
    //TODO: should this be combined with `get_introspection_bind_group_layouts` in some way?
    fn make_late_sized_buffer_groups(
        shader_binding_sizes: &FastHashMap<naga::ResourceBinding, wgt::BufferSize>,
        layout: &binding_model::PipelineLayout,
    ) -> ArrayVec<pipeline::LateSizedBufferGroup, { hal::MAX_BIND_GROUPS }> {
        // Given the shader-required binding sizes and the pipeline layout,
        // return the filtered list of them in the layout order,
        // removing those with given `min_binding_size`.
        layout
            .bind_group_layouts
            .iter()
            .enumerate()
            .map(|(group_index, bgl)| {
                let Some(bgl) = bgl else {
                    return pipeline::LateSizedBufferGroup::default();
                };

                let shader_sizes = bgl
                    .entries
                    .values()
                    .filter_map(|entry| match entry.ty {
                        wgt::BindingType::Buffer {
                            min_binding_size: None,
                            ..
                        } => {
                            let rb = naga::ResourceBinding {
                                group: group_index as u32,
                                binding: entry.binding,
                            };
                            let shader_size =
                                shader_binding_sizes.get(&rb).map_or(0, |nz| nz.get());
                            Some(shader_size)
                        }
                        _ => None,
                    })
                    .collect();
                pipeline::LateSizedBufferGroup { shader_sizes }
            })
            .collect()
    }

    pub fn create_bind_group_layout(
        self: &Arc<Self>,
        desc: &binding_model::BindGroupLayoutDescriptor,
    ) -> (Arc<BindGroupLayout>, Option<CreateBindGroupLayoutError>) {
        profiling::scope!("Device::create_bind_group_layout");

        let (bgl, error) = match self.create_bind_group_layout_inner(desc) {
            Ok(layout) => (layout, None),
            Err(e) => (
                BindGroupLayout::invalid(self, desc.label.to_string()),
                Some(e),
            ),
        };
        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use crate::device::trace::IntoTrace;

            trace.add(trace::Action::CreateBindGroupLayout(
                bgl.to_trace(),
                desc.clone(),
            ));
        }
        api_log!(
            "Device::create_bind_group_layout -> {:?}",
            Arc::as_ptr(&bgl)
        );
        (bgl, error)
    }

    fn create_bind_group_layout_inner(
        self: &Arc<Device>,
        desc: &binding_model::BindGroupLayoutDescriptor,
    ) -> Result<Arc<BindGroupLayout>, CreateBindGroupLayoutError> {
        self.check_is_valid()?;

        let entry_map = bgl::EntryMap::from_entries(&desc.entries)?;

        let bgl_result = self.bgl_pool.get_or_init(entry_map, |entry_map| {
            let bgl =
                self.create_bind_group_layout_impl(&desc.label, entry_map, bgl::Origin::Pool)?;
            bgl.exclusive_pipeline
                .set(binding_model::ExclusivePipeline::None)
                .unwrap();
            Ok(bgl)
        });

        match bgl_result {
            Ok(layout) => Ok(layout),
            Err(e) => Err(e),
        }
    }

    fn create_bind_group_layout_impl(
        self: &Arc<Self>,
        label: &crate::Label,
        entry_map: bgl::EntryMap,
        origin: bgl::Origin,
    ) -> Result<Arc<BindGroupLayout>, CreateBindGroupLayoutError> {
        #[derive(PartialEq)]
        enum WritableStorage {
            Yes,
            No,
        }

        for entry in entry_map.values() {
            if entry.binding >= self.limits.max_bindings_per_bind_group {
                return Err(CreateBindGroupLayoutError::InvalidBindingIndex {
                    binding: entry.binding,
                    maximum: self.limits.max_bindings_per_bind_group,
                });
            }

            use wgt::BindingType as Bt;

            let mut required_features = wgt::Features::empty();
            let mut required_downlevel_flags = wgt::DownlevelFlags::empty();
            let (array_feature, writable_storage) = match entry.ty {
                Bt::Buffer {
                    ty: wgt::BufferBindingType::Uniform,
                    has_dynamic_offset: false,
                    min_binding_size: _,
                } => (
                    Some(wgt::Features::BUFFER_BINDING_ARRAY),
                    WritableStorage::No,
                ),
                Bt::Buffer {
                    ty: wgt::BufferBindingType::Uniform,
                    has_dynamic_offset: true,
                    min_binding_size: _,
                } => (
                    Some(wgt::Features::BUFFER_BINDING_ARRAY),
                    WritableStorage::No,
                ),
                Bt::Buffer {
                    ty: wgt::BufferBindingType::Storage { read_only },
                    ..
                } => (
                    Some(
                        wgt::Features::BUFFER_BINDING_ARRAY
                            | wgt::Features::STORAGE_RESOURCE_BINDING_ARRAY,
                    ),
                    match read_only {
                        true => WritableStorage::No,
                        false => WritableStorage::Yes,
                    },
                ),
                Bt::Sampler { .. } => (
                    Some(wgt::Features::TEXTURE_BINDING_ARRAY),
                    WritableStorage::No,
                ),
                Bt::Texture {
                    multisampled: true,
                    sample_type: TextureSampleType::Float { filterable: true },
                    ..
                } => {
                    return Err(CreateBindGroupLayoutError::Entry {
                        binding: entry.binding,
                        error:
                            BindGroupLayoutEntryError::SampleTypeFloatFilterableBindingMultisampled,
                    });
                }
                Bt::Texture {
                    multisampled,
                    view_dimension,
                    ..
                } => {
                    if multisampled && view_dimension != TextureViewDimension::D2 {
                        return Err(CreateBindGroupLayoutError::Entry {
                            binding: entry.binding,
                            error: BindGroupLayoutEntryError::Non2DMultisampled(view_dimension),
                        });
                    }

                    (
                        Some(wgt::Features::TEXTURE_BINDING_ARRAY),
                        WritableStorage::No,
                    )
                }
                Bt::StorageTexture {
                    access,
                    view_dimension,
                    format,
                } => {
                    use wgt::{StorageTextureAccess as Access, TextureFormatFeatureFlags as Flags};

                    match view_dimension {
                        TextureViewDimension::Cube | TextureViewDimension::CubeArray => {
                            return Err(CreateBindGroupLayoutError::Entry {
                                binding: entry.binding,
                                error: BindGroupLayoutEntryError::StorageTextureCube,
                            })
                        }
                        _ => (),
                    }
                    match access {
                        wgt::StorageTextureAccess::Atomic
                            if !self.features.contains(wgt::Features::TEXTURE_ATOMIC) =>
                        {
                            return Err(CreateBindGroupLayoutError::Entry {
                                binding: entry.binding,
                                error: BindGroupLayoutEntryError::StorageTextureAtomic,
                            });
                        }
                        _ => (),
                    }

                    let format_features =
                        self.describe_format_features(format).map_err(|error| {
                            CreateBindGroupLayoutError::Entry {
                                binding: entry.binding,
                                error: BindGroupLayoutEntryError::MissingFeatures(error),
                            }
                        })?;

                    let required_feature_flag = match access {
                        Access::WriteOnly => Flags::STORAGE_WRITE_ONLY,
                        Access::ReadOnly => Flags::STORAGE_READ_ONLY,
                        Access::ReadWrite => Flags::STORAGE_READ_WRITE,
                        Access::Atomic => Flags::STORAGE_ATOMIC,
                    };

                    if !format_features.flags.contains(required_feature_flag) {
                        return Err(
                            CreateBindGroupLayoutError::UnsupportedStorageTextureAccess {
                                binding: entry.binding,
                                access,
                                format,
                            },
                        );
                    }

                    (
                        Some(
                            wgt::Features::TEXTURE_BINDING_ARRAY
                                | wgt::Features::STORAGE_RESOURCE_BINDING_ARRAY,
                        ),
                        match access {
                            wgt::StorageTextureAccess::WriteOnly => WritableStorage::Yes,
                            wgt::StorageTextureAccess::ReadOnly => WritableStorage::No,
                            wgt::StorageTextureAccess::ReadWrite => WritableStorage::Yes,
                            wgt::StorageTextureAccess::Atomic => {
                                required_features |= wgt::Features::TEXTURE_ATOMIC;
                                WritableStorage::Yes
                            }
                        },
                    )
                }
                Bt::AccelerationStructure { vertex_return } => {
                    self.require_features(wgt::Features::EXPERIMENTAL_RAY_QUERY)
                        .map_err(|e| CreateBindGroupLayoutError::Entry {
                            binding: entry.binding,
                            error: e.into(),
                        })?;
                    if vertex_return {
                        self.require_features(wgt::Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN)
                            .map_err(|e| CreateBindGroupLayoutError::Entry {
                                binding: entry.binding,
                                error: e.into(),
                            })?;
                    }
                    (
                        Some(wgt::Features::ACCELERATION_STRUCTURE_BINDING_ARRAY),
                        WritableStorage::No,
                    )
                }
                Bt::ExternalTexture => {
                    self.require_features(wgt::Features::EXTERNAL_TEXTURE)
                        .map_err(|e| CreateBindGroupLayoutError::Entry {
                            binding: entry.binding,
                            error: e.into(),
                        })?;
                    (None, WritableStorage::No)
                }
            };

            // Validate the count parameter
            if entry.count.is_some() {
                required_features |= array_feature
                    .ok_or(BindGroupLayoutEntryError::ArrayUnsupported)
                    .map_err(|error| CreateBindGroupLayoutError::Entry {
                        binding: entry.binding,
                        error,
                    })?;
            }

            if entry.visibility.contains_unknown_bits() {
                return Err(CreateBindGroupLayoutError::InvalidVisibility(
                    entry.visibility,
                ));
            }

            if entry.visibility.contains(wgt::ShaderStages::VERTEX) {
                if writable_storage == WritableStorage::Yes {
                    required_features |= wgt::Features::VERTEX_WRITABLE_STORAGE;
                }
                if let Bt::Buffer {
                    ty: wgt::BufferBindingType::Storage { .. },
                    ..
                } = entry.ty
                {
                    required_downlevel_flags |= wgt::DownlevelFlags::VERTEX_STORAGE;
                }
            }
            if writable_storage == WritableStorage::Yes
                && entry.visibility.contains(wgt::ShaderStages::FRAGMENT)
            {
                required_downlevel_flags |= wgt::DownlevelFlags::FRAGMENT_WRITABLE_STORAGE;
            }

            self.require_features(required_features)
                .map_err(BindGroupLayoutEntryError::MissingFeatures)
                .map_err(|error| CreateBindGroupLayoutError::Entry {
                    binding: entry.binding,
                    error,
                })?;
            self.require_downlevel_flags(required_downlevel_flags)
                .map_err(BindGroupLayoutEntryError::MissingDownlevelFlags)
                .map_err(|error| CreateBindGroupLayoutError::Entry {
                    binding: entry.binding,
                    error,
                })?;
        }

        let bgl_flags = conv::bind_group_layout_flags(self.features);

        let hal_bindings = entry_map.values().copied().collect::<Vec<_>>();
        let hal_desc = hal::BindGroupLayoutDescriptor {
            label: label.to_hal(self.instance_flags),
            flags: bgl_flags,
            entries: &hal_bindings,
        };

        let mut count_validator = binding_model::BindingTypeMaxCountValidator::default();
        for entry in entry_map.values() {
            count_validator.add_binding(entry);
        }
        // If a single bind group layout violates limits, the pipeline layout is
        // definitely going to violate limits too, lets catch it now.
        count_validator
            .validate(&self.limits, self.instance_flags)
            .map_err(CreateBindGroupLayoutError::TooManyBindings)?;

        // Validate that binding arrays don't conflict with dynamic offsets.
        count_validator.validate_binding_arrays()?;

        let raw = unsafe { self.raw().create_bind_group_layout(&hal_desc) }
            .map_err(|e| self.handle_hal_error(e))?;

        let bgl = BindGroupLayout {
            state: ResourceState::Valid(BindGroupLayoutState {
                raw: binding_model::RawBindGroupLayout::Owning(ManuallyDrop::new(raw)),
                origin,
                binding_count_validator: count_validator,
            }),
            device: self.clone(),
            entries: entry_map,
            exclusive_pipeline: OnceCellOrLock::new(),
            label: label.to_string(),
        };

        let bgl = Arc::new(bgl);

        Ok(bgl)
    }

    fn create_buffer_binding<'a>(
        &self,
        bb: &'a binding_model::ResolvedBufferBinding,
        binding: u32,
        decl: &wgt::BindGroupLayoutEntry,
        buffer_init_actions: &mut Vec<BufferInitTrackerAction>,
        dynamic_binding_info: &mut Vec<binding_model::BindGroupDynamicBindingData>,
        late_buffer_binding_sizes: &mut FastHashMap<u32, wgt::BufferSize>,
        used: &mut BindGroupStates,
        snatch_guard: &'a SnatchGuard<'a>,
    ) -> Result<hal::BufferBinding<'a, dyn hal::DynBuffer>, CreateBindGroupError> {
        use crate::binding_model::CreateBindGroupError as Error;

        let (binding_ty, dynamic, min_size) = match decl.ty {
            wgt::BindingType::Buffer {
                ty,
                has_dynamic_offset,
                min_binding_size,
            } => (ty, has_dynamic_offset, min_binding_size),
            _ => {
                return Err(Error::WrongBindingType {
                    binding,
                    actual: decl.ty,
                    expected: "UniformBuffer, StorageBuffer or ReadonlyStorageBuffer",
                })
            }
        };

        let (pub_usage, internal_use, range_limit) = match binding_ty {
            wgt::BufferBindingType::Uniform => (
                wgt::BufferUsages::UNIFORM,
                wgt::BufferUses::UNIFORM,
                self.limits.max_uniform_buffer_binding_size,
            ),
            wgt::BufferBindingType::Storage { read_only } => (
                wgt::BufferUsages::STORAGE,
                if read_only {
                    wgt::BufferUses::STORAGE_READ_ONLY
                } else {
                    wgt::BufferUses::STORAGE_READ_WRITE
                },
                self.limits.max_storage_buffer_binding_size,
            ),
        };

        let (align, align_limit_name) =
            binding_model::buffer_binding_type_alignment(&self.limits, binding_ty);
        if !bb.offset.is_multiple_of(align as u64) {
            return Err(Error::UnalignedBufferOffset(
                bb.offset,
                align_limit_name,
                align,
            ));
        }

        let buffer = &bb.buffer;

        used.buffers.insert_single(buffer.clone(), internal_use);

        buffer.check_is_valid()?;
        buffer.same_device(self)?;

        buffer.check_usage(pub_usage)?;

        let req_size = match bb.size.map(wgt::BufferSize::new) {
            // Requested a non-zero size
            Some(non_zero @ Some(_)) => non_zero,
            // Requested size not specified
            None => None,
            // Requested zero size
            Some(None) => return Err(CreateBindGroupError::BindingZeroSize(buffer.error_ident())),
        };
        let (bb, bind_size) = buffer.binding(bb.offset, req_size, snatch_guard)?;

        if matches!(binding_ty, wgt::BufferBindingType::Storage { .. })
            && bind_size % u64::from(wgt::STORAGE_BINDING_SIZE_ALIGNMENT) != 0
        {
            return Err(Error::UnalignedEffectiveBufferBindingSizeForStorage {
                alignment: wgt::STORAGE_BINDING_SIZE_ALIGNMENT,
                size: bind_size,
            });
        }

        let bind_end = bb.offset + bind_size;

        if bind_size > range_limit {
            return Err(Error::BufferRangeTooLarge {
                binding,
                given: bind_size,
                limit: range_limit,
            });
        }

        // Record binding info for validating dynamic offsets
        if dynamic {
            dynamic_binding_info.push(binding_model::BindGroupDynamicBindingData {
                binding_idx: binding,
                buffer_size: buffer.size,
                binding_range: bb.offset..bind_end,
                maximum_dynamic_offset: buffer.size - bind_end,
                binding_type: binding_ty,
            });
        }

        if let Some(non_zero) = min_size {
            let min_size = non_zero.get();
            if min_size > bind_size {
                return Err(Error::BindingSizeTooSmall {
                    buffer: buffer.error_ident(),
                    actual: bind_size,
                    min: min_size,
                });
            }
        } else {
            let late_size = wgt::BufferSize::new(bind_size)
                .ok_or_else(|| Error::BindingZeroSize(buffer.error_ident()))?;
            late_buffer_binding_sizes.insert(binding, late_size);
        }

        // This was checked against the device's alignment requirements above,
        // which should always be a multiple of `COPY_BUFFER_ALIGNMENT`.
        assert_eq!(bb.offset % wgt::COPY_BUFFER_ALIGNMENT, 0);

        let init_range = if dynamic {
            // We don't know what part of the buffer will be bound, so require that it
            // is fully initialized.
            0..buffer.size
        } else {
            // `wgpu_hal` only restricts shader access to bound buffer regions with
            // a certain resolution. For the sake of lazy initialization, round up
            // the size of the bound range to reflect how much of the buffer is
            // actually going to be visible to the shader.
            let bounds_check_alignment = binding_model::buffer_binding_type_bounds_check_alignment(
                &self.alignments,
                binding_ty,
            );
            let visible_size = align_to(bind_size, bounds_check_alignment);

            bb.offset..bb.offset + visible_size
        };

        buffer_init_actions.extend(buffer.initialization_status.read().create_action(
            buffer,
            init_range,
            MemoryInitKind::NeedsInitializedMemory,
        ));

        Ok(bb)
    }

    fn create_sampler_binding<'a>(
        &self,
        used: &mut BindGroupStates,
        binding: u32,
        decl: &wgt::BindGroupLayoutEntry,
        sampler: &'a Arc<Sampler>,
    ) -> Result<&'a dyn hal::DynSampler, CreateBindGroupError> {
        use crate::binding_model::CreateBindGroupError as Error;

        used.samplers.insert_single(sampler.clone());

        sampler.same_device(self)?;

        match decl.ty {
            wgt::BindingType::Sampler(ty) => {
                let (allowed_filtering, allowed_comparison) = match ty {
                    wgt::SamplerBindingType::Filtering => (None, false),
                    wgt::SamplerBindingType::NonFiltering => (Some(false), false),
                    wgt::SamplerBindingType::Comparison => (None, true),
                };
                if let Some(allowed_filtering) = allowed_filtering {
                    if allowed_filtering != sampler.filtering {
                        return Err(Error::WrongSamplerFiltering {
                            binding,
                            layout_flt: allowed_filtering,
                            sampler_flt: sampler.filtering,
                        });
                    }
                }
                if allowed_comparison != sampler.comparison {
                    return Err(Error::WrongSamplerComparison {
                        binding,
                        layout_cmp: allowed_comparison,
                        sampler_cmp: sampler.comparison,
                    });
                }
            }
            _ => {
                return Err(Error::WrongBindingType {
                    binding,
                    actual: decl.ty,
                    expected: "Sampler",
                })
            }
        }

        Ok(sampler.raw()?)
    }

    fn create_texture_binding<'a>(
        &self,
        binding: u32,
        decl: &wgt::BindGroupLayoutEntry,
        view: &'a Arc<TextureView>,
        used: &mut BindGroupStates,
        texture_init_actions: &mut Vec<TextureInitTrackerAction>,
        snatch_guard: &'a SnatchGuard<'a>,
    ) -> Result<hal::TextureBinding<'a, dyn hal::DynTextureView>, CreateBindGroupError> {
        view.check_valid()?;
        view.same_device(self)?;

        let internal_use = self.texture_use_parameters(
            binding,
            decl,
            view,
            "SampledTexture, ReadonlyStorageTexture or WriteonlyStorageTexture",
        )?;

        used.views.insert_single(view.clone(), internal_use);

        let texture = &view.parent;

        texture_init_actions.push(TextureInitTrackerAction {
            texture: texture.clone(),
            range: TextureInitRange {
                mip_range: view.desc.range.mip_range(texture.desc.mip_level_count),
                layer_range: view
                    .desc
                    .range
                    .layer_range(texture.desc.array_layer_count()),
            },
            kind: MemoryInitKind::NeedsInitializedMemory,
        });

        Ok(hal::TextureBinding {
            view: view.try_raw(snatch_guard)?,
            usage: internal_use,
        })
    }

    fn create_tlas_binding<'a>(
        self: &Arc<Self>,
        used: &mut BindGroupStates,
        binding: u32,
        decl: &wgt::BindGroupLayoutEntry,
        tlas: &'a Arc<Tlas>,
        snatch_guard: &'a SnatchGuard<'a>,
    ) -> Result<&'a dyn hal::DynAccelerationStructure, CreateBindGroupError> {
        use crate::binding_model::CreateBindGroupError as Error;

        used.acceleration_structures.insert_single(tlas.clone());

        tlas.check_is_valid()?;
        tlas.same_device(self)?;

        match decl.ty {
            wgt::BindingType::AccelerationStructure { vertex_return } => {
                if vertex_return
                    && !tlas.flags.contains(
                        wgpu_types::AccelerationStructureFlags::ALLOW_RAY_HIT_VERTEX_RETURN,
                    )
                {
                    return Err(Error::MissingTLASVertexReturn { binding });
                }
            }
            _ => {
                return Err(Error::WrongBindingType {
                    binding,
                    actual: decl.ty,
                    expected: "Tlas",
                });
            }
        }

        Ok(tlas.try_raw(snatch_guard)?)
    }

    fn create_external_texture_binding<'a>(
        &'a self,
        binding: u32,
        decl: &wgt::BindGroupLayoutEntry,
        external_texture: &'a Arc<ExternalTexture>,
        used: &mut BindGroupStates,
        snatch_guard: &'a SnatchGuard,
    ) -> Result<
        hal::ExternalTextureBinding<'a, dyn hal::DynBuffer, dyn hal::DynTextureView>,
        CreateBindGroupError,
    > {
        use crate::binding_model::CreateBindGroupError as Error;

        let external_texture_state = external_texture.state()?;
        external_texture.same_device(self)?;

        used.external_textures
            .insert_single(external_texture.clone());

        match decl.ty {
            wgt::BindingType::ExternalTexture => {}
            _ => {
                return Err(Error::WrongBindingType {
                    binding,
                    actual: decl.ty,
                    expected: "ExternalTexture",
                });
            }
        }

        let planes = (0..3)
            .map(|i| {
                // We always need 3 bindings. If we have fewer than 3 planes
                // just bind plane 0 multiple times. The shader will only
                // sample from valid planes anyway.
                let plane = external_texture
                    .planes
                    .get(i)
                    .unwrap_or(&external_texture.planes[0]);
                let internal_use = wgt::TextureUses::RESOURCE;
                used.views.insert_single(plane.clone(), internal_use);
                let view = plane.try_raw(snatch_guard)?;
                Ok(hal::TextureBinding {
                    view,
                    usage: internal_use,
                })
            })
            // We can remove this intermediate Vec by using
            // array::try_from_fn() above, once it stabilizes.
            .collect::<Result<Vec<_>, Error>>()?;
        let planes = planes.try_into().unwrap();

        used.buffers.insert_single(
            external_texture_state.params.clone(),
            wgt::BufferUses::UNIFORM,
        );
        let params = external_texture_state
            .params
            .binding(0, None, snatch_guard)?
            .0;

        Ok(hal::ExternalTextureBinding { planes, params })
    }

    fn create_external_texture_binding_from_view<'a>(
        &'a self,
        binding: u32,
        decl: &wgt::BindGroupLayoutEntry,
        view: &'a Arc<TextureView>,
        used: &mut BindGroupStates,
        snatch_guard: &'a SnatchGuard,
    ) -> Result<
        hal::ExternalTextureBinding<'a, dyn hal::DynBuffer, dyn hal::DynTextureView>,
        CreateBindGroupError,
    > {
        use crate::binding_model::CreateBindGroupError as Error;

        view.same_device(self)?;
        view.check_valid()?;

        let internal_use = self.texture_use_parameters(binding, decl, view, "SampledTexture")?;
        used.views.insert_single(view.clone(), internal_use);

        match decl.ty {
            wgt::BindingType::ExternalTexture => {}
            _ => {
                return Err(Error::WrongBindingType {
                    binding,
                    actual: decl.ty,
                    expected: "ExternalTexture",
                });
            }
        }

        // We need 3 bindings, so just repeat the same texture view 3 times.
        let planes = [
            hal::TextureBinding {
                view: view.try_raw(snatch_guard)?,
                usage: internal_use,
            },
            hal::TextureBinding {
                view: view.try_raw(snatch_guard)?,
                usage: internal_use,
            },
            hal::TextureBinding {
                view: view.try_raw(snatch_guard)?,
                usage: internal_use,
            },
        ];
        let params = hal::BufferBinding::new_unchecked(
            self.default_external_texture_params_buffer.as_ref(),
            0,
            None,
        );

        Ok(hal::ExternalTextureBinding { planes, params })
    }

    pub fn create_bind_group(
        self: &Arc<Self>,
        desc: &binding_model::ResolvedBindGroupDescriptor,
    ) -> (Arc<BindGroup>, Option<CreateBindGroupError>) {
        profiling::scope!("Device::create_bind_group");
        #[cfg(feature = "trace")]
        let trace_desc = (&desc).to_trace();

        let (bind_group, error) = match self.create_bind_group_inner(desc) {
            Ok(bind_group) => (bind_group, None),
            Err(e) => (
                BindGroup::invalid(self.clone(), desc.label.to_string(), desc.layout.clone()),
                Some(e),
            ),
        };

        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            trace.add(trace::Action::CreateBindGroup(
                bind_group.to_trace(),
                trace_desc,
            ));
        }

        api_log!(
            "Device::create_bind_group -> {:?}",
            Arc::as_ptr(&bind_group)
        );

        (bind_group, error)
    }

    // This function expects the provided bind group layout to be resolved
    // (not passing a duplicate) beforehand.
    pub fn create_bind_group_inner(
        self: &Arc<Self>,
        desc: &binding_model::ResolvedBindGroupDescriptor,
    ) -> Result<Arc<BindGroup>, CreateBindGroupError> {
        use crate::binding_model::{CreateBindGroupError as Error, ResolvedBindingResource as Br};

        self.check_is_valid()?;

        let layout = desc.layout.clone();

        layout.same_device(self)?;
        layout.check_is_valid()?;

        {
            // Check that the number of entries in the descriptor matches
            // the number of entries in the layout.
            let actual = desc.entries.len();
            let expected = layout.entries.len();
            if actual != expected {
                return Err(Error::BindingsNumMismatch { expected, actual });
            }
        }

        // TODO: arrayvec/smallvec, or re-use allocations
        // Record binding info for dynamic offset validation
        let mut dynamic_binding_info = Vec::new();
        // Map of binding -> shader reflected size
        //Note: we can't collect into a vector right away because
        // it needs to be in BGL iteration order, not BG entry order.
        let mut late_buffer_binding_sizes = FastHashMap::default();
        // fill out the descriptors
        let mut used = BindGroupStates::new();

        let mut buffer_init_actions = Vec::new();
        let mut texture_init_actions = Vec::new();
        let mut hal_entries = Vec::with_capacity(desc.entries.len());
        let mut hal_buffers = Vec::new();
        let mut hal_samplers = Vec::new();
        let mut hal_textures = Vec::new();
        let mut hal_tlas_s = Vec::new();
        let mut hal_external_textures = Vec::new();
        let snatch_guard = self.snatchable_lock.read();
        for entry in desc.entries.iter() {
            let binding = entry.binding;
            // Find the corresponding declaration in the layout
            let decl = layout
                .entries
                .get(binding)
                .ok_or(Error::MissingBindingDeclaration(binding))?;
            let (res_index, count) = match entry.resource {
                Br::Buffer(ref bb) => {
                    let bb = self.create_buffer_binding(
                        bb,
                        binding,
                        decl,
                        &mut buffer_init_actions,
                        &mut dynamic_binding_info,
                        &mut late_buffer_binding_sizes,
                        &mut used,
                        &snatch_guard,
                    )?;

                    let res_index = hal_buffers.len();
                    hal_buffers.push(bb);
                    (res_index, 1)
                }
                Br::BufferArray(ref bindings_array) => {
                    let num_bindings = bindings_array.len();
                    Self::check_array_binding(self.features, decl.count, num_bindings)?;

                    let res_index = hal_buffers.len();
                    for bb in bindings_array.iter() {
                        let bb = self.create_buffer_binding(
                            bb,
                            binding,
                            decl,
                            &mut buffer_init_actions,
                            &mut dynamic_binding_info,
                            &mut late_buffer_binding_sizes,
                            &mut used,
                            &snatch_guard,
                        )?;
                        hal_buffers.push(bb);
                    }
                    (res_index, num_bindings)
                }
                Br::Sampler(ref sampler) => {
                    let sampler = self.create_sampler_binding(&mut used, binding, decl, sampler)?;

                    let res_index = hal_samplers.len();
                    hal_samplers.push(sampler);
                    (res_index, 1)
                }
                Br::SamplerArray(ref samplers) => {
                    let num_bindings = samplers.len();
                    Self::check_array_binding(self.features, decl.count, num_bindings)?;

                    let res_index = hal_samplers.len();
                    for sampler in samplers.iter() {
                        let sampler =
                            self.create_sampler_binding(&mut used, binding, decl, sampler)?;

                        hal_samplers.push(sampler);
                    }

                    (res_index, num_bindings)
                }
                Br::TextureView(ref view) => match decl.ty {
                    wgt::BindingType::ExternalTexture => {
                        let et = self.create_external_texture_binding_from_view(
                            binding,
                            decl,
                            view,
                            &mut used,
                            &snatch_guard,
                        )?;
                        let res_index = hal_external_textures.len();
                        hal_external_textures.push(et);
                        (res_index, 1)
                    }
                    _ => {
                        let tb = self.create_texture_binding(
                            binding,
                            decl,
                            view,
                            &mut used,
                            &mut texture_init_actions,
                            &snatch_guard,
                        )?;
                        let res_index = hal_textures.len();
                        hal_textures.push(tb);
                        (res_index, 1)
                    }
                },
                Br::TextureViewArray(ref views) => {
                    let num_bindings = views.len();
                    Self::check_array_binding(self.features, decl.count, num_bindings)?;

                    let res_index = hal_textures.len();
                    for view in views.iter() {
                        let tb = self.create_texture_binding(
                            binding,
                            decl,
                            view,
                            &mut used,
                            &mut texture_init_actions,
                            &snatch_guard,
                        )?;

                        hal_textures.push(tb);
                    }

                    (res_index, num_bindings)
                }
                Br::AccelerationStructure(ref tlas) => {
                    let tlas =
                        self.create_tlas_binding(&mut used, binding, decl, tlas, &snatch_guard)?;
                    let res_index = hal_tlas_s.len();
                    hal_tlas_s.push(tlas);
                    (res_index, 1)
                }
                Br::AccelerationStructureArray(ref tlas_array) => {
                    // Feature validation for TLAS binding arrays happens at bind group layout
                    // creation time (mirroring other binding-array resource types). By the time we
                    // get here, `decl.count` has already been validated against device features.
                    let num_bindings = tlas_array.len();
                    Self::check_array_binding(self.features, decl.count, num_bindings)?;

                    let res_index = hal_tlas_s.len();
                    for tlas in tlas_array.iter() {
                        let tlas = self.create_tlas_binding(
                            &mut used,
                            binding,
                            decl,
                            tlas,
                            &snatch_guard,
                        )?;
                        hal_tlas_s.push(tlas);
                    }
                    (res_index, num_bindings)
                }
                Br::ExternalTexture(ref et) => {
                    let et = self.create_external_texture_binding(
                        binding,
                        decl,
                        et,
                        &mut used,
                        &snatch_guard,
                    )?;
                    let res_index = hal_external_textures.len();
                    hal_external_textures.push(et);
                    (res_index, 1)
                }
            };

            hal_entries.push(hal::BindGroupEntry {
                binding,
                resource_index: res_index as u32,
                count: count as u32,
            });
        }

        used.optimize();

        hal_entries.sort_by_key(|entry| entry.binding);
        for (a, b) in hal_entries.iter().zip(hal_entries.iter().skip(1)) {
            if a.binding == b.binding {
                return Err(Error::DuplicateBinding(a.binding));
            }
        }

        dynamic_binding_info.sort_by_key(|i| i.binding_idx);

        let hal_desc = hal::BindGroupDescriptor {
            label: desc.label.to_hal(self.instance_flags),
            layout: layout.try_raw()?,
            entries: &hal_entries,
            buffers: &hal_buffers,
            samplers: &hal_samplers,
            textures: &hal_textures,
            acceleration_structures: &hal_tlas_s,
            external_textures: &hal_external_textures,
        };
        let raw = unsafe { self.raw().create_bind_group(&hal_desc) }
            .map_err(|e| self.handle_hal_error(e))?;

        // collect in the order of BGL iteration
        let late_buffer_binding_infos = layout
            .entries
            .indices()
            .flat_map(|binding| {
                let size = late_buffer_binding_sizes.get(&binding).cloned()?;
                Some(BindGroupLateBufferBindingInfo {
                    binding_index: binding,
                    size,
                })
            })
            .collect();

        let bind_group = BindGroup {
            state: ResourceState::Valid(BindGroupState {
                raw: Snatchable::new(raw),
            }),
            device: self.clone(),
            layout,
            label: desc.label.to_string(),
            tracking_data: TrackingData::new(self.tracker_indices.bind_groups.clone()),
            used,
            buffer_init_actions,
            texture_init_actions,
            dynamic_binding_info,
            late_buffer_binding_infos,
        };

        let bind_group = Arc::new(bind_group);

        let weak_ref = Arc::downgrade(&bind_group);
        for texture in bind_group.used.views.used_textures() {
            let mut bind_groups = texture.bind_groups.lock();
            bind_groups.push(weak_ref.clone());
        }
        for buffer in bind_group.used.buffers.used_resources() {
            let mut bind_groups = buffer.bind_groups.lock();
            bind_groups.push(weak_ref.clone());
        }

        Ok(bind_group)
    }

    fn check_array_binding(
        features: wgt::Features,
        count: Option<NonZeroU32>,
        num_bindings: usize,
    ) -> Result<(), CreateBindGroupError> {
        use super::binding_model::CreateBindGroupError as Error;

        if let Some(count) = count {
            let count = count.get() as usize;
            if count < num_bindings {
                return Err(Error::BindingArrayPartialLengthMismatch {
                    actual: num_bindings,
                    expected: count,
                });
            }
            if count != num_bindings
                && !features.contains(wgt::Features::PARTIALLY_BOUND_BINDING_ARRAY)
            {
                return Err(Error::BindingArrayLengthMismatch {
                    actual: num_bindings,
                    expected: count,
                });
            }
            if num_bindings == 0 {
                return Err(Error::BindingArrayZeroLength);
            }
        } else {
            return Err(Error::SingleBindingExpected);
        };

        Ok(())
    }

    fn texture_use_parameters(
        &self,
        binding: u32,
        decl: &wgt::BindGroupLayoutEntry,
        view: &TextureView,
        expected: &'static str,
    ) -> Result<wgt::TextureUses, CreateBindGroupError> {
        use crate::binding_model::CreateBindGroupError as Error;
        if view
            .desc
            .aspects()
            .contains(hal::FormatAspects::DEPTH | hal::FormatAspects::STENCIL)
        {
            return Err(Error::DepthStencilAspect);
        }
        match decl.ty {
            wgt::BindingType::Texture {
                sample_type,
                view_dimension,
                multisampled,
            } => {
                use wgt::TextureSampleType as Tst;
                if multisampled != (view.samples != 1) {
                    return Err(Error::InvalidTextureMultisample {
                        binding,
                        layout_multisampled: multisampled,
                        view_samples: view.samples,
                    });
                }
                let compat_sample_type = view
                    .desc
                    .format
                    .sample_type(Some(view.desc.range.aspect), Some(self.features))
                    .unwrap();
                match (sample_type, compat_sample_type) {
                    (Tst::Uint, Tst::Uint) |
                        (Tst::Sint, Tst::Sint) |
                        (Tst::Depth, Tst::Depth) |
                        // if we expect non-filterable, accept anything float
                        (Tst::Float { filterable: false }, Tst::Float { .. }) |
                        // if we expect filterable, require it
                        (Tst::Float { filterable: true }, Tst::Float { filterable: true }) |
                        // if we expect non-filterable, also accept depth
                        (Tst::Float { filterable: false }, Tst::Depth) => {}
                    // if we expect filterable, also accept Float that is defined as
                    // unfilterable if filterable feature is explicitly enabled (only hit
                    // if wgt::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES is
                    // enabled)
                    (Tst::Float { filterable: true }, Tst::Float { .. })
                        if view.format_features.flags
                            .contains(wgt::TextureFormatFeatureFlags::FILTERABLE) => {}
                    _ => {
                        return Err(Error::InvalidTextureSampleType {
                            binding,
                            layout_sample_type: sample_type,
                            view_format: view.desc.format,
                            view_sample_type: compat_sample_type,
                        })
                    }
                }
                if view_dimension != view.desc.dimension {
                    return Err(Error::InvalidTextureDimension {
                        binding,
                        layout_dimension: view_dimension,
                        view_dimension: view.desc.dimension,
                    });
                }
                view.check_usage(wgt::TextureUsages::TEXTURE_BINDING)?;
                Ok(wgt::TextureUses::RESOURCE)
            }
            wgt::BindingType::StorageTexture {
                access,
                format,
                view_dimension,
            } => {
                if format != view.desc.format {
                    return Err(Error::InvalidStorageTextureFormat {
                        binding,
                        layout_format: format,
                        view_format: view.desc.format,
                    });
                }
                if view_dimension != view.desc.dimension {
                    return Err(Error::InvalidTextureDimension {
                        binding,
                        layout_dimension: view_dimension,
                        view_dimension: view.desc.dimension,
                    });
                }

                let mip_level_count = view.selector.mips.end - view.selector.mips.start;
                if mip_level_count != 1 {
                    return Err(Error::InvalidStorageTextureMipLevelCount {
                        binding,
                        mip_level_count,
                    });
                }

                view.check_usage(wgt::TextureUsages::STORAGE_BINDING)?;

                Ok(match access {
                    wgt::StorageTextureAccess::ReadOnly => wgt::TextureUses::STORAGE_READ_ONLY,
                    wgt::StorageTextureAccess::WriteOnly => wgt::TextureUses::STORAGE_WRITE_ONLY,
                    wgt::StorageTextureAccess::ReadWrite => wgt::TextureUses::STORAGE_READ_WRITE,
                    wgt::StorageTextureAccess::Atomic => wgt::TextureUses::STORAGE_ATOMIC,
                })
            }
            wgt::BindingType::ExternalTexture => {
                if view.desc.dimension != TextureViewDimension::D2 {
                    return Err(Error::InvalidTextureDimension {
                        binding,
                        layout_dimension: TextureViewDimension::D2,
                        view_dimension: view.desc.dimension,
                    });
                }
                let mip_level_count = view.selector.mips.end - view.selector.mips.start;
                if mip_level_count != 1 {
                    return Err(Error::InvalidExternalTextureMipLevelCount {
                        binding,
                        mip_level_count,
                    });
                }
                if view.desc.format != TextureFormat::Rgba8Unorm
                    && view.desc.format != TextureFormat::Bgra8Unorm
                    && view.desc.format != TextureFormat::Rgba16Float
                {
                    return Err(Error::InvalidExternalTextureFormat {
                        binding,
                        format: view.desc.format,
                    });
                }
                if view.samples != 1 {
                    return Err(Error::InvalidTextureMultisample {
                        binding,
                        layout_multisampled: false,
                        view_samples: view.samples,
                    });
                }

                view.check_usage(wgt::TextureUsages::TEXTURE_BINDING)?;
                Ok(wgt::TextureUses::RESOURCE)
            }
            _ => Err(Error::WrongBindingType {
                binding,
                actual: decl.ty,
                expected,
            }),
        }
    }

    pub fn create_pipeline_layout(
        self: &Arc<Self>,
        desc: &binding_model::ResolvedPipelineLayoutDescriptor,
    ) -> (
        Arc<binding_model::PipelineLayout>,
        Option<binding_model::CreatePipelineLayoutError>,
    ) {
        profiling::scope!("Device::create_pipeline_layout");
        let (layout, error) = match self.create_pipeline_layout_impl(desc, false) {
            Ok(layout) => (layout, None),
            Err(e) => (
                binding_model::PipelineLayout::invalid(Arc::clone(self), desc.label.to_string()),
                Some(e),
            ),
        };
        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use crate::device::trace::IntoTrace;
            trace.add(trace::Action::CreatePipelineLayout(
                layout.to_trace(),
                desc.to_trace(),
            ));
        }
        api_log!(
            "Device::create_pipeline_layout -> {:?}",
            Arc::as_ptr(&layout)
        );
        (layout, error)
    }

    fn create_pipeline_layout_impl(
        self: &Arc<Self>,
        desc: &binding_model::ResolvedPipelineLayoutDescriptor,
        ignore_exclusive_pipeline_check: bool,
    ) -> Result<Arc<binding_model::PipelineLayout>, binding_model::CreatePipelineLayoutError> {
        use crate::binding_model::CreatePipelineLayoutError as Error;

        self.check_is_valid()?;

        let bind_group_layouts_count = desc.bind_group_layouts.len();
        let device_max_bind_groups = self.limits.max_bind_groups as usize;
        if bind_group_layouts_count > device_max_bind_groups {
            return Err(Error::TooManyGroups {
                actual: bind_group_layouts_count,
                max: device_max_bind_groups,
            });
        }

        if desc.immediate_size != 0 {
            self.require_features(wgt::Features::IMMEDIATES)?;
        }
        if self.limits.max_immediate_size < desc.immediate_size {
            return Err(Error::ImmediateRangeTooLarge {
                size: desc.immediate_size,
                max: self.limits.max_immediate_size,
            });
        }
        if !desc
            .immediate_size
            .is_multiple_of(wgt::IMMEDIATE_DATA_ALIGNMENT)
        {
            return Err(Error::MisalignedImmediateSize {
                size: desc.immediate_size,
            });
        }

        let mut count_validator = binding_model::BindingTypeMaxCountValidator::default();

        for (index, bgl) in desc.bind_group_layouts.iter().enumerate() {
            let Some(bgl) = bgl else {
                continue;
            };

            bgl.same_device(self)?;

            if !ignore_exclusive_pipeline_check {
                let exclusive_pipeline = bgl.exclusive_pipeline.get().unwrap();
                if !matches!(exclusive_pipeline, binding_model::ExclusivePipeline::None) {
                    return Err(Error::BglHasExclusivePipeline {
                        index,
                        pipeline: alloc::format!("{exclusive_pipeline}"),
                    });
                }
            }

            count_validator.merge(&bgl.state()?.binding_count_validator);
        }

        count_validator
            .validate(&self.limits, self.instance_flags)
            .map_err(Error::TooManyBindings)?;

        let buffers_and_acceleration_structures_in_vertex_stage =
            count_validator.buffers_and_acceleration_structures_in_vertex_stage();

        let get_bgl_iter = || {
            desc.bind_group_layouts
                .iter()
                .map(|bgl| bgl.as_ref().filter(|bgl| !bgl.entries.is_empty()))
        };

        let bind_group_layouts = get_bgl_iter()
            .map(|bgl| bgl.cloned())
            .collect::<ArrayVec<_, { hal::MAX_BIND_GROUPS }>>();

        let raw_bind_group_layouts = get_bgl_iter()
            .map(|bgl| bgl.map(|bgl| bgl.try_raw()).transpose())
            .collect::<Result<ArrayVec<_, { hal::MAX_BIND_GROUPS }>, _>>()?;

        let additional_flags = if self.indirect_validation.is_some() {
            hal::PipelineLayoutFlags::INDIRECT_BUILTIN_UPDATE
        } else {
            hal::PipelineLayoutFlags::empty()
        };

        let hal_desc = hal::PipelineLayoutDescriptor {
            label: desc.label.to_hal(self.instance_flags),
            flags: hal::PipelineLayoutFlags::FIRST_VERTEX_INSTANCE
                | hal::PipelineLayoutFlags::NUM_WORK_GROUPS
                | additional_flags,
            bind_group_layouts: &raw_bind_group_layouts,
            immediate_size: desc.immediate_size,
        };

        let raw = unsafe { self.raw().create_pipeline_layout(&hal_desc) }
            .map_err(|e| self.handle_hal_error(e))?;

        drop(raw_bind_group_layouts);

        let layout = binding_model::PipelineLayout {
            raw: ResourceState::Valid(raw),
            device: self.clone(),
            label: desc.label.to_string(),
            bind_group_layouts,
            immediate_size: desc.immediate_size,
            buffers_and_acceleration_structures_in_vertex_stage,
        };

        let layout = Arc::new(layout);

        Ok(layout)
    }

    fn create_derived_pipeline_layout(
        self: &Arc<Self>,
        mut derived_group_layouts: Box<ArrayVec<bgl::EntryMap, { hal::MAX_BIND_GROUPS }>>,
        immediate_size: u32,
    ) -> Result<Arc<binding_model::PipelineLayout>, pipeline::ImplicitLayoutError> {
        // <https://gpuweb.github.io/gpuweb/#abstract-opdef-default-pipeline-layout>
        // Round up the immediate size for pipeline layout as it is required to be a multiple of 4
        let immediate_size = align_to(immediate_size, wgt::IMMEDIATE_DATA_ALIGNMENT);

        while derived_group_layouts
            .last()
            .is_some_and(|map| map.is_empty())
        {
            derived_group_layouts.pop();
        }

        let mut unique_bind_group_layouts = FastHashMap::default();

        let bind_group_layouts = derived_group_layouts
            .into_iter()
            .map(|mut bgl_entry_map| {
                if bgl_entry_map.is_empty() {
                    return Ok(None);
                }

                bgl_entry_map.sort();
                match unique_bind_group_layouts.entry(bgl_entry_map) {
                    hashbrown::hash_map::Entry::Occupied(v) => Ok(Some(Arc::clone(v.get()))),
                    hashbrown::hash_map::Entry::Vacant(e) => {
                        match self.create_bind_group_layout_impl(
                            &None,
                            e.key().clone(),
                            bgl::Origin::Derived,
                        ) {
                            Ok(bgl) => {
                                e.insert(bgl.clone());
                                Ok(Some(bgl))
                            }
                            Err(e) => Err(e),
                        }
                    }
                }
            })
            .collect::<Result<Vec<_>, _>>()?;

        let layout_desc = binding_model::ResolvedPipelineLayoutDescriptor {
            label: None,
            bind_group_layouts: Cow::Owned(bind_group_layouts),
            immediate_size,
        };

        let layout = self.create_pipeline_layout_impl(&layout_desc, true)?;
        Ok(layout)
    }

    pub fn create_compute_pipeline(
        self: &Arc<Self>,
        desc: pipeline::ResolvedComputePipelineDescriptor,
    ) -> (
        Arc<pipeline::ComputePipeline>,
        Option<pipeline::CreateComputePipelineError>,
    ) {
        profiling::scope!("Device::create_compute_pipeline");
        let (compute_pipeline, error) = match self.create_compute_pipeline_inner(desc.clone()) {
            Ok(compute_pipeline) => (compute_pipeline, None),
            Err(error) => (
                pipeline::ComputePipeline::invalid(self.clone(), desc.label.to_string()),
                Some(error),
            ),
        };
        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use crate::device::trace;
            use crate::device::trace::IntoTrace;
            trace.add(trace::Action::CreateComputePipeline {
                id: compute_pipeline.to_trace(),
                desc: desc.to_trace(),
            });
        }
        api_log!(
            "Device::create_compute_pipeline -> {:?}",
            Arc::as_ptr(&compute_pipeline)
        );
        (compute_pipeline, error)
    }

    pub fn create_compute_pipeline_inner(
        self: &Arc<Self>,
        desc: pipeline::ResolvedComputePipelineDescriptor,
    ) -> Result<Arc<pipeline::ComputePipeline>, pipeline::CreateComputePipelineError> {
        self.check_is_valid()?;

        self.require_downlevel_flags(wgt::DownlevelFlags::COMPUTE_SHADERS)?;

        let shader_module = desc.stage.module;

        let shader_module_state = shader_module.state()?;
        shader_module.same_device(self)?;

        let is_auto_layout = desc.layout.is_none();

        // Get the pipeline layout from the desc if it is provided.
        let pipeline_layout = match desc.layout {
            Some(pipeline_layout) => {
                pipeline_layout.same_device(self)?;
                pipeline_layout.check_valid()?;
                Some(pipeline_layout)
            }
            None => None,
        };

        if shader_module_state.interface.interface().is_none() && pipeline_layout.is_none() {
            return Err(pipeline::CreateComputePipelineError::Implicit(
                pipeline::ImplicitLayoutError::Passthrough(wgt::ShaderStages::COMPUTE),
            ));
        }

        let mut binding_layout_source = match pipeline_layout {
            Some(pipeline_layout) => validation::BindingLayoutSource::Provided(pipeline_layout),
            None => validation::BindingLayoutSource::new_derived(&self.limits),
        };
        let mut minimum_binding_sizes = FastHashMap::default();
        let mut io = validation::StageIo::default();

        let final_entry_point_name;

        {
            let stage = validation::ShaderStageForValidation::Compute;

            final_entry_point_name = shader_module.finalize_entry_point_name(
                stage.to_naga(),
                desc.stage.entry_point.as_ref().map(|ep| ep.as_ref()),
            )?;

            if let Some(interface) = shader_module_state.interface.interface() {
                io = interface.check_stage(
                    &mut binding_layout_source,
                    &mut minimum_binding_sizes,
                    &final_entry_point_name,
                    stage,
                    io,
                    None,
                )?;
            }
        }

        let pipeline_layout = match binding_layout_source {
            validation::BindingLayoutSource::Provided(pipeline_layout) => pipeline_layout,
            validation::BindingLayoutSource::Derived(entries) => {
                self.create_derived_pipeline_layout(entries, io.immediates.size())?
            }
        };

        let naga::valid::ImmediateUsage::Valid {
            slots: immediate_slots_required,
            size: _,
        } = io.immediates
        else {
            unreachable!("Immediates exceeding maxImmediateSize should have been rejected");
        };

        let late_sized_buffer_groups =
            Device::make_late_sized_buffer_groups(&minimum_binding_sizes, &pipeline_layout);

        let cache = match desc.cache {
            Some(cache) => {
                cache.check_is_valid()?;
                cache.same_device(self)?;
                Some(cache)
            }
            None => None,
        };

        let pipeline_desc = hal::ComputePipelineDescriptor {
            label: desc.label.to_hal(self.instance_flags),
            layout: pipeline_layout.raw()?,
            stage: hal::ProgrammableStage {
                module: shader_module_state.raw.as_ref(),
                entry_point: final_entry_point_name.as_ref(),
                constants: &desc.stage.constants,
                zero_initialize_workgroup_memory: desc.stage.zero_initialize_workgroup_memory,
            },
            cache: cache.as_ref().map(|it| it.raw()).transpose()?,
        };

        let raw =
            unsafe { self.raw().create_compute_pipeline(&pipeline_desc) }.map_err(
                |err| match err {
                    hal::PipelineError::Device(error) => {
                        pipeline::CreateComputePipelineError::Device(self.handle_hal_error(error))
                    }
                    hal::PipelineError::Linkage(_stages, msg) => {
                        pipeline::CreateComputePipelineError::Internal(msg)
                    }
                    hal::PipelineError::EntryPoint(_stage) => {
                        pipeline::CreateComputePipelineError::Internal(
                            ENTRYPOINT_FAILURE_ERROR.to_string(),
                        )
                    }
                    hal::PipelineError::PipelineConstants(_stages, msg) => {
                        pipeline::CreateComputePipelineError::PipelineConstants(msg)
                    }
                },
            )?;

        let pipeline = pipeline::ComputePipeline {
            state: ResourceState::Valid(pipeline::ComputePipelineState {
                raw: ManuallyDrop::new(raw),
                layout: pipeline_layout.clone(),
                _shader_module: shader_module,
            }),
            device: self.clone(),
            late_sized_buffer_groups,
            immediate_slots_required,
            label: desc.label.to_string(),
            tracking_data: TrackingData::new(self.tracker_indices.compute_pipelines.clone()),
        };

        let pipeline = Arc::new(pipeline);

        if is_auto_layout {
            for bgl in pipeline_layout.bind_group_layouts.iter() {
                let Some(bgl) = bgl else {
                    continue;
                };

                // `bind_group_layouts` might contain duplicate entries, so we need to ignore the
                // result.
                let _ = bgl.exclusive_pipeline.set((&pipeline).into());
            }
        }

        Ok(pipeline)
    }

    pub fn create_render_pipeline(
        self: &Arc<Self>,
        desc: pipeline::ResolvedGeneralRenderPipelineDescriptor,
    ) -> (
        Arc<pipeline::RenderPipeline>,
        Option<pipeline::CreateRenderPipelineError>,
    ) {
        profiling::scope!("Device::create_render_pipeline");
        let (render_pipeline, error) = match self.create_render_pipeline_inner(desc.clone()) {
            Ok(pipeline) => (pipeline, None),
            Err(e) => (
                pipeline::RenderPipeline::invalid(self.clone(), desc.label.to_string()),
                Some(e),
            ),
        };
        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use crate::device::trace::IntoTrace;
            trace.add(trace::Action::CreateGeneralRenderPipeline {
                id: render_pipeline.to_trace(),
                desc: desc.to_trace(),
            });
        }
        api_log!(
            "Device::create_render_pipeline -> {:?}",
            Arc::as_ptr(&render_pipeline)
        );
        (render_pipeline, error)
    }

    pub fn create_render_pipeline_inner(
        self: &Arc<Self>,
        desc: pipeline::ResolvedGeneralRenderPipelineDescriptor,
    ) -> Result<Arc<pipeline::RenderPipeline>, pipeline::CreateRenderPipelineError> {
        use wgt::TextureFormatFeatureFlags as Tfff;

        self.check_is_valid()?;

        let mut minimum_binding_sizes = FastHashMap::default();

        let color_targets = desc
            .fragment
            .as_ref()
            .map_or(&[][..], |fragment| &fragment.targets);
        let depth_stencil_state = desc.depth_stencil.as_ref();

        check_color_attachment_count(color_targets.len(), self.limits.max_color_attachments)?;

        {
            let cts: ArrayVec<_, { hal::MAX_COLOR_ATTACHMENTS }> =
                color_targets.iter().filter_map(|x| x.as_ref()).collect();
            if !cts.is_empty() && {
                let first = &cts[0];
                cts[1..]
                    .iter()
                    .any(|ct| ct.write_mask != first.write_mask || ct.blend != first.blend)
            } {
                self.require_downlevel_flags(wgt::DownlevelFlags::INDEPENDENT_BLEND)?;
            }
        }

        let mut io = validation::StageIo::default();
        let mut validated_stages = wgt::ShaderStages::empty();

        let mut vertex_steps;
        let mut hal_vertex_buffer_layouts;
        let mut total_attributes;
        let mut dual_source_blending = false;
        let mut has_depth_attachment = false;
        if let pipeline::RenderPipelineVertexProcessor::Vertex(ref vertex) = desc.vertex {
            if vertex.buffers.len() > self.limits.max_vertex_buffers as usize {
                return Err(pipeline::CreateRenderPipelineError::TooManyVertexBuffers {
                    given: vertex.buffers.len() as u32,
                    limit: self.limits.max_vertex_buffers,
                });
            }

            vertex_steps = Vec::with_capacity(vertex.buffers.len());
            hal_vertex_buffer_layouts = Vec::with_capacity(vertex.buffers.len());
            total_attributes = 0;
            for (i, vb_state) in vertex.buffers.iter().enumerate() {
                let Some(vb_state) = vb_state else {
                    vertex_steps.push(None);
                    hal_vertex_buffer_layouts.push(None);
                    continue;
                };

                // https://gpuweb.github.io/gpuweb/#abstract-opdef-validating-gpuvertexbufferlayout

                if vb_state.array_stride > self.limits.max_vertex_buffer_array_stride as u64 {
                    return Err(pipeline::CreateRenderPipelineError::VertexStrideTooLarge {
                        index: i as u32,
                        given: vb_state.array_stride as u32,
                        limit: self.limits.max_vertex_buffer_array_stride,
                    });
                }
                if vb_state.array_stride % wgt::VERTEX_ALIGNMENT != 0 {
                    return Err(pipeline::CreateRenderPipelineError::UnalignedVertexStride {
                        index: i as u32,
                        stride: vb_state.array_stride,
                    });
                }

                let max_stride = if vb_state.array_stride == 0 {
                    self.limits.max_vertex_buffer_array_stride as u64
                } else {
                    vb_state.array_stride
                };
                let mut last_stride = 0;
                for attribute in vb_state.attributes.iter() {
                    let attribute_stride = attribute.offset + attribute.format.size();
                    if attribute_stride > max_stride {
                        return Err(
                            pipeline::CreateRenderPipelineError::VertexAttributeStrideTooLarge {
                                location: attribute.shader_location,
                                given: attribute_stride as u32,
                                limit: max_stride as u32,
                            },
                        );
                    }

                    let required_offset_alignment = attribute.format.size().min(4);
                    if attribute.offset % required_offset_alignment != 0 {
                        return Err(
                            pipeline::CreateRenderPipelineError::InvalidVertexAttributeOffset {
                                location: attribute.shader_location,
                                offset: attribute.offset,
                            },
                        );
                    }

                    if attribute.shader_location >= self.limits.max_vertex_attributes {
                        return Err(
                            pipeline::CreateRenderPipelineError::VertexAttributeLocationTooLarge {
                                given: attribute.shader_location,
                                limit: self.limits.max_vertex_attributes,
                            },
                        );
                    }

                    last_stride = last_stride.max(attribute_stride);
                }

                vertex_steps.push(Some(pipeline::VertexStep {
                    stride: vb_state.array_stride,
                    last_stride,
                    mode: vb_state.step_mode,
                }));
                hal_vertex_buffer_layouts.push(if vb_state.attributes.is_empty() {
                    None
                } else {
                    Some(hal::VertexBufferLayout {
                        array_stride: vb_state.array_stride,
                        step_mode: vb_state.step_mode,
                        attributes: vb_state.attributes.as_ref(),
                    })
                });

                for attribute in vb_state.attributes.iter() {
                    if attribute.offset >= 0x10000000 {
                        return Err(
                            pipeline::CreateRenderPipelineError::InvalidVertexAttributeOffset {
                                location: attribute.shader_location,
                                offset: attribute.offset,
                            },
                        );
                    }

                    if let wgt::VertexFormat::Float64
                    | wgt::VertexFormat::Float64x2
                    | wgt::VertexFormat::Float64x3
                    | wgt::VertexFormat::Float64x4 = attribute.format
                    {
                        self.require_features(wgt::Features::VERTEX_ATTRIBUTE_64BIT)?;
                    }

                    let previous = io.varyings.insert(
                        attribute.shader_location,
                        validation::InterfaceVar::vertex_attribute(attribute.format),
                    );

                    if previous.is_some() {
                        return Err(pipeline::CreateRenderPipelineError::ShaderLocationClash(
                            attribute.shader_location,
                        ));
                    }
                }
                total_attributes += vb_state.attributes.len();
            }

            if total_attributes > self.limits.max_vertex_attributes as usize {
                return Err(
                    pipeline::CreateRenderPipelineError::TooManyVertexAttributes {
                        given: total_attributes as u32,
                        limit: self.limits.max_vertex_attributes,
                    },
                );
            }
        } else {
            vertex_steps = Vec::new();
            hal_vertex_buffer_layouts = Vec::new();
        };

        if desc.primitive.strip_index_format.is_some() && !desc.primitive.topology.is_strip() {
            return Err(
                pipeline::CreateRenderPipelineError::StripIndexFormatForNonStripTopology {
                    strip_index_format: desc.primitive.strip_index_format,
                    topology: desc.primitive.topology,
                },
            );
        }

        if desc.primitive.unclipped_depth {
            self.require_features(wgt::Features::DEPTH_CLIP_CONTROL)?;
        }

        if desc.primitive.polygon_mode == wgt::PolygonMode::Line {
            self.require_features(wgt::Features::POLYGON_MODE_LINE)?;
        }
        if desc.primitive.polygon_mode == wgt::PolygonMode::Point {
            self.require_features(wgt::Features::POLYGON_MODE_POINT)?;
        }

        if desc.primitive.conservative {
            self.require_features(wgt::Features::CONSERVATIVE_RASTERIZATION)?;
        }

        if desc.primitive.conservative && desc.primitive.polygon_mode != wgt::PolygonMode::Fill {
            return Err(
                pipeline::CreateRenderPipelineError::ConservativeRasterizationNonFillPolygonMode,
            );
        }

        let mut target_specified = false;
        let mut required_color_outputs = 0u64;
        const _: () = {
            assert!(hal::MAX_COLOR_ATTACHMENTS <= 64);
        };

        for (i, cs) in color_targets.iter().enumerate() {
            if let Some(cs) = cs.as_ref() {
                target_specified = true;
                let error = 'error: {
                    // This is expected to be the operative check for illegal write mask
                    // values (larger than 15), because WebGPU requires that it be validated
                    // on the device timeline.
                    if cs.write_mask.contains_unknown_bits() {
                        break 'error Some(ColorStateError::InvalidWriteMask(cs.write_mask));
                    } else if cs.write_mask != ColorWrites::NONE {
                        required_color_outputs |= 1 << i;
                    }

                    let format_features = self.describe_format_features(cs.format)?;
                    if !format_features
                        .allowed_usages
                        .contains(wgt::TextureUsages::RENDER_ATTACHMENT)
                    {
                        break 'error Some(ColorStateError::FormatNotRenderable(cs.format));
                    }
                    if cs.blend.is_some() && !format_features.flags.contains(Tfff::BLENDABLE) {
                        break 'error Some(ColorStateError::FormatNotBlendable(cs.format));
                    }
                    if !hal::FormatAspects::from(cs.format).contains(hal::FormatAspects::COLOR) {
                        break 'error Some(ColorStateError::FormatNotColor(cs.format));
                    }

                    if desc.multisample.count > 1
                        && !format_features
                            .flags
                            .sample_count_supported(desc.multisample.count)
                    {
                        break 'error Some(ColorStateError::InvalidSampleCount(
                            desc.multisample.count,
                            cs.format,
                            cs.format
                                .guaranteed_format_features(self.features)
                                .flags
                                .supported_sample_counts(),
                            self.adapter
                                .get_texture_format_features(cs.format)
                                .flags
                                .supported_sample_counts(),
                        ));
                    }

                    if let Some(blend_mode) = cs.blend {
                        for component in [&blend_mode.color, &blend_mode.alpha] {
                            for factor in [component.src_factor, component.dst_factor] {
                                if factor.uses_second_blend_source() {
                                    self.require_features(wgt::Features::DUAL_SOURCE_BLENDING)?;
                                    if i == 0 {
                                        dual_source_blending = true;
                                    } else {
                                        break 'error Some(
                                            ColorStateError::BlendFactorOnUnsupportedTarget {
                                                factor,
                                                target: i as u32,
                                            },
                                        );
                                    }
                                }

                                if [wgt::BlendOperation::Min, wgt::BlendOperation::Max]
                                    .contains(&component.operation)
                                    && factor != wgt::BlendFactor::One
                                {
                                    break 'error Some(ColorStateError::InvalidMinMaxBlendFactor {
                                        factor,
                                        target: i as u32,
                                    });
                                }
                            }
                        }
                    }

                    break 'error None;
                };
                if let Some(e) = error {
                    return Err(pipeline::CreateRenderPipelineError::ColorState(i as u8, e));
                }
            }
        }

        if dual_source_blending && color_targets.len() > 1 {
            return Err(
                pipeline::CreateRenderPipelineError::DualSourceBlendingWithMultipleColorTargets {
                    count: color_targets.len(),
                },
            );
        }

        validation::validate_color_attachment_bytes_per_sample(
            color_targets.iter().flatten().map(|cs| cs.format),
            self.limits.max_color_attachment_bytes_per_sample,
        )
        .map_err(pipeline::CreateRenderPipelineError::ColorAttachment)?;

        if let Some(ds) = depth_stencil_state {
            // See <https://gpuweb.github.io/gpuweb/#abstract-opdef-validating-gpudepthstencilstate>.
            target_specified = true;
            let error = 'error: {
                if !ds.format.is_depth_stencil_format() {
                    // This error case is not redundant with the aspect check below when
                    // neither depth nor stencil is enabled at all.
                    break 'error Some(pipeline::DepthStencilStateError::FormatNotDepthOrStencil(
                        ds.format,
                    ));
                }

                let format_features = self.describe_format_features(ds.format)?;
                if !format_features
                    .allowed_usages
                    .contains(wgt::TextureUsages::RENDER_ATTACHMENT)
                {
                    break 'error Some(pipeline::DepthStencilStateError::FormatNotRenderable(
                        ds.format,
                    ));
                }

                let aspect = hal::FormatAspects::from(ds.format);
                if aspect.contains(hal::FormatAspects::DEPTH) {
                    has_depth_attachment = true;
                } else if ds.is_depth_enabled() {
                    break 'error Some(pipeline::DepthStencilStateError::FormatNotDepth(ds.format));
                }
                if has_depth_attachment {
                    let Some(depth_write_enabled) = ds.depth_write_enabled else {
                        break 'error Some(
                            pipeline::DepthStencilStateError::MissingDepthWriteEnabled(ds.format),
                        );
                    };

                    let depth_compare_required = depth_write_enabled
                        || ds.stencil.front.depth_fail_op != wgt::StencilOperation::Keep
                        || ds.stencil.back.depth_fail_op != wgt::StencilOperation::Keep;
                    if depth_compare_required && ds.depth_compare.is_none() {
                        break 'error Some(pipeline::DepthStencilStateError::MissingDepthCompare(
                            ds.format,
                        ));
                    }
                }

                if ds.stencil.is_enabled() && !aspect.contains(hal::FormatAspects::STENCIL) {
                    break 'error Some(pipeline::DepthStencilStateError::FormatNotStencil(
                        ds.format,
                    ));
                }
                if desc.multisample.count > 1
                    && !format_features
                        .flags
                        .sample_count_supported(desc.multisample.count)
                {
                    break 'error Some(pipeline::DepthStencilStateError::InvalidSampleCount(
                        desc.multisample.count,
                        ds.format,
                        ds.format
                            .guaranteed_format_features(self.features)
                            .flags
                            .supported_sample_counts(),
                        self.adapter
                            .get_texture_format_features(ds.format)
                            .flags
                            .supported_sample_counts(),
                    ));
                }

                break 'error None;
            };
            if let Some(e) = error {
                return Err(pipeline::CreateRenderPipelineError::DepthStencilState(e));
            }

            if ds.bias.clamp != 0.0 {
                self.require_downlevel_flags(wgt::DownlevelFlags::DEPTH_BIAS_CLAMP)?;
            }

            if (ds.bias.is_enabled() || ds.bias.clamp != 0.0)
                && !desc.primitive.topology.is_triangles()
            {
                return Err(pipeline::CreateRenderPipelineError::DepthStencilState(
                    pipeline::DepthStencilStateError::DepthBiasWithIncompatibleTopology(
                        desc.primitive.topology,
                    ),
                ));
            }
        }

        if !target_specified {
            return Err(pipeline::CreateRenderPipelineError::NoTargetSpecified);
        }

        let is_auto_layout = desc.layout.is_none();

        // Get the pipeline layout from the desc if it is provided.
        let pipeline_layout = match desc.layout {
            Some(pipeline_layout) => {
                pipeline_layout.same_device(self)?;
                pipeline_layout.check_valid()?;
                Some(pipeline_layout)
            }
            None => None,
        };

        let mut binding_layout_source = match pipeline_layout {
            Some(pipeline_layout) => validation::BindingLayoutSource::Provided(pipeline_layout),
            None => validation::BindingLayoutSource::new_derived(&self.limits),
        };

        let samples = {
            let sc = desc.multisample.count;
            if sc == 0 || sc > 32 || !sc.is_power_of_two() {
                return Err(pipeline::CreateRenderPipelineError::InvalidSampleCount(sc));
            }
            sc
        };

        let mut vertex_stage = None;
        let mut task_stage = None;
        let mut mesh_stage = None;
        let mut _vertex_entry_point_name = String::new();
        let mut _task_entry_point_name = String::new();
        let mut _mesh_entry_point_name = String::new();
        let mut passthrough_stages = wgt::ShaderStages::empty();
        match desc.vertex {
            pipeline::RenderPipelineVertexProcessor::Vertex(ref vertex) => {
                vertex_stage = {
                    let stage_desc = &vertex.stage;
                    let stage = validation::ShaderStageForValidation::Vertex {
                        topology: desc.primitive.topology,
                        compare_function: desc.depth_stencil.as_ref().and_then(|d| d.depth_compare),
                    };
                    let stage_bit = stage.to_wgt_bit();
                    let stage_err = |error| pipeline::CreateRenderPipelineError::Stage {
                        stage: stage_bit,
                        error,
                    };

                    let vertex_shader_module = &stage_desc.module;
                    let vertex_shader_module_state = vertex_shader_module
                        .state()
                        .map_err(Into::into)
                        .map_err(stage_err)?;
                    vertex_shader_module.same_device(self)?;

                    if vertex_shader_module_state.interface.interface().is_none() {
                        passthrough_stages |= stage_bit;
                    }

                    _vertex_entry_point_name = vertex_shader_module
                        .finalize_entry_point_name(
                            stage.to_naga(),
                            stage_desc.entry_point.as_ref().map(|ep| ep.as_ref()),
                        )
                        .map_err(stage_err)?;

                    if let Some(interface) = vertex_shader_module_state.interface.interface() {
                        io = interface
                            .check_stage(
                                &mut binding_layout_source,
                                &mut minimum_binding_sizes,
                                &_vertex_entry_point_name,
                                stage,
                                io,
                                Some(desc.primitive.topology),
                            )
                            .map_err(stage_err)?;
                        validated_stages |= stage_bit;
                    }
                    Some(hal::ProgrammableStage {
                        module: vertex_shader_module_state.raw.as_ref(),
                        entry_point: &_vertex_entry_point_name,
                        constants: &stage_desc.constants,
                        zero_initialize_workgroup_memory: stage_desc
                            .zero_initialize_workgroup_memory,
                    })
                };
            }
            pipeline::RenderPipelineVertexProcessor::Mesh(ref task, ref mesh) => {
                self.require_features(wgt::Features::EXPERIMENTAL_MESH_SHADER)?;

                task_stage = if let Some(task) = task {
                    let stage_desc = &task.stage;
                    let stage = validation::ShaderStageForValidation::Task;
                    let stage_bit = stage.to_wgt_bit();
                    let stage_err = |error| pipeline::CreateRenderPipelineError::Stage {
                        stage: stage_bit,
                        error,
                    };

                    let task_shader_module = &stage_desc.module;
                    let task_shader_module_state = task_shader_module
                        .state()
                        .map_err(Into::into)
                        .map_err(stage_err)?;
                    task_shader_module.same_device(self)?;

                    if task_shader_module_state.interface.interface().is_none() {
                        passthrough_stages |= stage_bit;
                    }

                    _task_entry_point_name = task_shader_module
                        .finalize_entry_point_name(
                            stage.to_naga(),
                            stage_desc.entry_point.as_ref().map(|ep| ep.as_ref()),
                        )
                        .map_err(stage_err)?;

                    if let Some(interface) = task_shader_module_state.interface.interface() {
                        io = interface
                            .check_stage(
                                &mut binding_layout_source,
                                &mut minimum_binding_sizes,
                                &_task_entry_point_name,
                                stage,
                                io,
                                Some(desc.primitive.topology),
                            )
                            .map_err(stage_err)?;
                        validated_stages |= stage_bit;
                    }
                    Some(hal::ProgrammableStage {
                        module: task_shader_module_state.raw.as_ref(),
                        entry_point: &_task_entry_point_name,
                        constants: &stage_desc.constants,
                        zero_initialize_workgroup_memory: stage_desc
                            .zero_initialize_workgroup_memory,
                    })
                } else {
                    None
                };
                mesh_stage = {
                    let stage_desc = &mesh.stage;
                    let stage = validation::ShaderStageForValidation::Mesh;
                    let stage_bit = stage.to_wgt_bit();
                    let stage_err = |error| pipeline::CreateRenderPipelineError::Stage {
                        stage: stage_bit,
                        error,
                    };

                    let mesh_shader_module = &stage_desc.module;
                    let mesh_shader_module_state = mesh_shader_module
                        .state()
                        .map_err(Into::into)
                        .map_err(stage_err)?;
                    mesh_shader_module.same_device(self)?;

                    if mesh_shader_module_state.interface.interface().is_none() {
                        passthrough_stages |= stage_bit;
                    }

                    _mesh_entry_point_name = mesh_shader_module
                        .finalize_entry_point_name(
                            stage.to_naga(),
                            stage_desc.entry_point.as_ref().map(|ep| ep.as_ref()),
                        )
                        .map_err(stage_err)?;

                    if let Some(interface) = mesh_shader_module_state.interface.interface() {
                        io = interface
                            .check_stage(
                                &mut binding_layout_source,
                                &mut minimum_binding_sizes,
                                &_mesh_entry_point_name,
                                stage,
                                io,
                                Some(desc.primitive.topology),
                            )
                            .map_err(stage_err)?;
                        validated_stages |= stage_bit;
                    }
                    Some(hal::ProgrammableStage {
                        module: mesh_shader_module_state.raw.as_ref(),
                        entry_point: &_mesh_entry_point_name,
                        constants: &stage_desc.constants,
                        zero_initialize_workgroup_memory: stage_desc
                            .zero_initialize_workgroup_memory,
                    })
                };
            }
        }

        let fragment_entry_point_name;
        let fragment_stage = match desc.fragment {
            Some(ref fragment_state) => {
                let stage = validation::ShaderStageForValidation::Fragment {
                    dual_source_blending,
                    has_depth_attachment,
                };
                let stage_bit = stage.to_wgt_bit();
                let stage_err = |error| pipeline::CreateRenderPipelineError::Stage {
                    stage: stage_bit,
                    error,
                };

                let shader_module = &fragment_state.stage.module;
                let shader_module_state = shader_module
                    .state()
                    .map_err(Into::into)
                    .map_err(stage_err)?;
                shader_module.same_device(self)?;

                if shader_module_state.interface.interface().is_none() {
                    passthrough_stages |= stage_bit;
                }

                fragment_entry_point_name = shader_module
                    .finalize_entry_point_name(
                        stage.to_naga(),
                        fragment_state
                            .stage
                            .entry_point
                            .as_ref()
                            .map(|ep| ep.as_ref()),
                    )
                    .map_err(stage_err)?;

                if let Some(interface) = shader_module_state.interface.interface() {
                    io = interface
                        .check_stage(
                            &mut binding_layout_source,
                            &mut minimum_binding_sizes,
                            &fragment_entry_point_name,
                            stage,
                            io,
                            Some(desc.primitive.topology),
                        )
                        .map_err(stage_err)?;
                    validated_stages |= stage_bit;
                }

                Some(hal::ProgrammableStage {
                    module: shader_module_state.raw.as_ref(),
                    entry_point: &fragment_entry_point_name,
                    constants: &fragment_state.stage.constants,
                    zero_initialize_workgroup_memory: fragment_state
                        .stage
                        .zero_initialize_workgroup_memory,
                })
            }
            None => None,
        };

        if !passthrough_stages.is_empty() && is_auto_layout {
            return Err(pipeline::CreateRenderPipelineError::Implicit(
                pipeline::ImplicitLayoutError::Passthrough(passthrough_stages),
            ));
        }

        let mut active_color_outputs = 0u64;
        if validated_stages.contains(wgt::ShaderStages::FRAGMENT) {
            for (i, output) in io.varyings.iter() {
                active_color_outputs |= 1 << i;
                match color_targets.get(*i as usize) {
                    Some(Some(state)) => {
                        validation::check_color_attachment_compatibility(state, output.ty)
                            .map_err(|err| {
                                pipeline::CreateRenderPipelineError::ColorState(*i as u8, err)
                            })?;
                    }
                    _ => {
                        log::debug!(
                            "The fragment stage {:?} output @location({}) values are ignored",
                            fragment_stage
                                .as_ref()
                                .map_or("", |stage| stage.entry_point),
                            i
                        );
                    }
                }
            }

            let missing_color_outputs = required_color_outputs & !active_color_outputs;
            if missing_color_outputs != 0 {
                return Err(pipeline::CreateRenderPipelineError::ColorState(
                    missing_color_outputs.trailing_zeros() as u8,
                    ColorStateError::OutputNotPresent,
                ));
            }
        }

        let last_stage = match desc.fragment {
            Some(_) => wgt::ShaderStages::FRAGMENT,
            None => wgt::ShaderStages::VERTEX,
        };
        if is_auto_layout && !validated_stages.contains(last_stage) {
            return Err(pipeline::ImplicitLayoutError::ReflectionError(last_stage).into());
        }

        let pipeline_layout = match binding_layout_source {
            validation::BindingLayoutSource::Provided(pipeline_layout) => pipeline_layout,
            validation::BindingLayoutSource::Derived(entries) => {
                self.create_derived_pipeline_layout(entries, io.immediates.size())?
            }
        };

        let naga::valid::ImmediateUsage::Valid {
            slots: immediate_slots_required,
            size: _,
        } = io.immediates
        else {
            unreachable!("Immediates exceeding maxImmediateSize should have been rejected");
        };

        if let pipeline::RenderPipelineVertexProcessor::Vertex(ref vertex) = desc.vertex {
            let bind_groups_plus_vertex_buffers =
                u32::try_from(pipeline_layout.bind_group_layouts.len() + vertex.buffers.len())
                    .unwrap();
            if bind_groups_plus_vertex_buffers > self.limits.max_bind_groups_plus_vertex_buffers {
                return Err(
                    pipeline::CreateRenderPipelineError::TooManyBindGroupsPlusVertexBuffers {
                        given: bind_groups_plus_vertex_buffers,
                        limit: self.limits.max_bind_groups_plus_vertex_buffers,
                    },
                );
            }

            let given = pipeline_layout
                .buffers_and_acceleration_structures_in_vertex_stage
                .saturating_add(vertex.buffers.len() as u32);
            if !self
                .instance_flags
                .contains(wgt::InstanceFlags::STRICT_WEBGPU_COMPLIANCE)
            {
                let limit = self
                    .limits
                    .max_buffers_and_acceleration_structures_per_shader_stage;
                if given > limit {
                    return Err(
                    pipeline::CreateRenderPipelineError::TooManyBuffersAndAccelerationStructuresInVertexStage {
                        given,
                        limit,
                    },
                );
                }
            }
        }

        // Multiview is only supported if the feature is enabled
        if let Some(mv_mask) = desc.multiview_mask {
            self.require_features(wgt::Features::MULTIVIEW)?;
            if !(mv_mask.get() + 1).is_power_of_two() {
                self.require_features(wgt::Features::SELECTIVE_MULTIVIEW)?;
            }
        }

        if !self
            .downlevel
            .flags
            .contains(wgt::DownlevelFlags::BUFFER_BINDINGS_NOT_16_BYTE_ALIGNED)
        {
            for (binding, size) in minimum_binding_sizes.iter() {
                if size.get() % 16 != 0 {
                    return Err(pipeline::CreateRenderPipelineError::UnalignedShader {
                        binding: binding.binding,
                        group: binding.group,
                        size: size.get(),
                    });
                }
            }
        }

        let late_sized_buffer_groups =
            Device::make_late_sized_buffer_groups(&minimum_binding_sizes, &pipeline_layout);

        let cache = match desc.cache {
            Some(cache) => {
                cache.check_is_valid()?;
                cache.same_device(self)?;
                Some(cache)
            }
            None => None,
        };

        let is_mesh = mesh_stage.is_some();
        let has_task_shader = task_stage.is_some();
        let raw = {
            let pipeline_desc = hal::RenderPipelineDescriptor {
                label: desc.label.to_hal(self.instance_flags),
                layout: pipeline_layout.raw()?,
                vertex_processor: match vertex_stage {
                    Some(vertex_stage) => hal::VertexProcessor::Standard {
                        vertex_buffers: &hal_vertex_buffer_layouts,
                        vertex_stage,
                    },
                    None => hal::VertexProcessor::Mesh {
                        task_stage,
                        mesh_stage: mesh_stage.unwrap(),
                    },
                },
                primitive: desc.primitive,
                depth_stencil: desc.depth_stencil.clone(),
                multisample: desc.multisample,
                fragment_stage,
                color_targets,
                multiview_mask: desc.multiview_mask,
                cache: cache.as_ref().map(|it| it.raw()).transpose()?,
            };
            unsafe { self.raw().create_render_pipeline(&pipeline_desc) }.map_err(
                |err| match err {
                    hal::PipelineError::Device(error) => {
                        pipeline::CreateRenderPipelineError::Device(self.handle_hal_error(error))
                    }
                    hal::PipelineError::Linkage(stage, msg) => {
                        pipeline::CreateRenderPipelineError::Internal { stage, error: msg }
                    }
                    hal::PipelineError::EntryPoint(stage) => {
                        pipeline::CreateRenderPipelineError::Internal {
                            stage: hal::auxil::map_naga_stage(stage),
                            error: ENTRYPOINT_FAILURE_ERROR.to_string(),
                        }
                    }
                    hal::PipelineError::PipelineConstants(stage, error) => {
                        pipeline::CreateRenderPipelineError::PipelineConstants { stage, error }
                    }
                },
            )?
        };

        let pass_context = RenderPassContext {
            attachments: AttachmentData {
                colors: color_targets
                    .iter()
                    .map(|state| state.as_ref().map(|s| s.format))
                    .collect(),
                resolves: ArrayVec::new(),
                depth_stencil: depth_stencil_state.as_ref().map(|state| state.format),
            },
            sample_count: samples,
            multiview_mask: desc.multiview_mask,
        };

        let mut flags = pipeline::PipelineFlags::empty();
        for state in color_targets.iter().filter_map(|s| s.as_ref()) {
            if let Some(ref bs) = state.blend {
                if bs.color.uses_constant() | bs.alpha.uses_constant() {
                    flags |= pipeline::PipelineFlags::BLEND_CONSTANT;
                }
            }
        }
        if let Some(ds) = depth_stencil_state.as_ref() {
            if ds.stencil.is_enabled() && ds.stencil.needs_ref_value() {
                flags |= pipeline::PipelineFlags::STENCIL_REFERENCE;
            }
            if !ds.is_depth_read_only() {
                flags |= pipeline::PipelineFlags::WRITES_DEPTH;
            }
            if !ds.is_stencil_read_only(desc.primitive.cull_mode) {
                flags |= pipeline::PipelineFlags::WRITES_STENCIL;
            }
        }
        let shader_modules = {
            let mut shader_modules = ArrayVec::new();
            match desc.vertex {
                pipeline::RenderPipelineVertexProcessor::Vertex(vertex) => {
                    shader_modules.push(vertex.stage.module)
                }
                pipeline::RenderPipelineVertexProcessor::Mesh(task, mesh) => {
                    if let Some(task) = task {
                        shader_modules.push(task.stage.module);
                    }
                    shader_modules.push(mesh.stage.module);
                }
            }
            shader_modules.extend(desc.fragment.map(|f| f.stage.module));
            shader_modules
        };

        let pipeline = pipeline::RenderPipeline {
            state: ResourceState::Valid(pipeline::RenderPipelineState {
                raw: ManuallyDrop::new(raw),
                layout: pipeline_layout.clone(),
            }),
            device: self.clone(),
            pass_context,
            _shader_modules: shader_modules,
            flags,
            topology: desc.primitive.topology,
            strip_index_format: desc.primitive.strip_index_format,
            vertex_steps,
            late_sized_buffer_groups,
            immediate_slots_required,
            label: desc.label.to_string(),
            tracking_data: TrackingData::new(self.tracker_indices.render_pipelines.clone()),
            is_mesh,
            has_task_shader,
        };

        let pipeline = Arc::new(pipeline);

        if is_auto_layout {
            for bgl in pipeline_layout.bind_group_layouts.iter() {
                let Some(bgl) = bgl else {
                    continue;
                };

                // `bind_group_layouts` might contain duplicate entries, so we need to ignore the
                // result.
                let _ = bgl.exclusive_pipeline.set((&pipeline).into());
            }
        }

        Ok(pipeline)
    }

    /// # Safety
    /// The `data` field on `desc` must have previously been returned from
    /// [`pipeline::PipelineCache::get_data`]
    pub unsafe fn create_pipeline_cache(
        self: &Arc<Self>,
        desc: &pipeline::PipelineCacheDescriptor,
    ) -> (
        Arc<pipeline::PipelineCache>,
        Option<pipeline::CreatePipelineCacheError>,
    ) {
        profiling::scope!("Device::create_pipeline_cache");
        let (cache, error) = match unsafe { self.create_pipeline_cache_inner(desc) } {
            Ok(cache) => (cache, None),
            Err(e) => (
                pipeline::PipelineCache::invalid(self.clone(), desc),
                Some(e),
            ),
        };
        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use trace::IntoTrace;
            trace.add(trace::Action::CreatePipelineCache {
                id: cache.to_trace(),
                desc: desc.clone(),
            });
        }
        api_log!("Device::create_pipeline_cache -> {:?}", Arc::as_ptr(&cache));
        (cache, error)
    }

    /// # Safety
    /// The `data` field on `desc` must have previously been returned from
    /// [`pipeline::PipelineCache::get_data`]
    pub(crate) unsafe fn create_pipeline_cache_inner(
        self: &Arc<Self>,
        desc: &pipeline::PipelineCacheDescriptor,
    ) -> Result<Arc<pipeline::PipelineCache>, pipeline::CreatePipelineCacheError> {
        use crate::pipeline_cache;

        self.check_is_valid()?;

        self.require_features(wgt::Features::PIPELINE_CACHE)?;
        let data = if let Some((data, validation_key)) = desc
            .data
            .as_ref()
            .zip(self.raw().pipeline_cache_validation_key())
        {
            let data = pipeline_cache::validate_pipeline_cache(
                data,
                &self.adapter.raw.info,
                validation_key,
            );
            match data {
                Ok(data) => Some(data),
                Err(e) if e.was_avoidable() || !desc.fallback => return Err(e.into()),
                // If the error was unavoidable and we are asked to fallback, do so
                Err(_) => None,
            }
        } else {
            None
        };
        let cache_desc = hal::PipelineCacheDescriptor {
            data,
            label: desc.label.to_hal(self.instance_flags),
        };
        let raw = match unsafe { self.raw().create_pipeline_cache(&cache_desc) } {
            Ok(raw) => raw,
            Err(e) => match e {
                hal::PipelineCacheError::Device(e) => return Err(self.handle_hal_error(e).into()),
            },
        };
        let cache = pipeline::PipelineCache {
            device: self.clone(),
            label: desc.label.to_string(),
            // This would be none in the error condition, which we don't implement yet
            raw: ResourceState::Valid(raw),
        };

        let cache = Arc::new(cache);

        Ok(cache)
    }

    fn get_texture_format_features(&self, format: TextureFormat) -> wgt::TextureFormatFeatures {
        // Variant of adapter.get_texture_format_features that takes device features into account
        use wgt::TextureFormatFeatureFlags as tfsc;
        let mut format_features = self.adapter.get_texture_format_features(format);
        if (format == TextureFormat::R32Float
            || format == TextureFormat::Rg32Float
            || format == TextureFormat::Rgba32Float)
            && !self.features.contains(wgt::Features::FLOAT32_FILTERABLE)
        {
            format_features.flags.set(tfsc::FILTERABLE, false);
        }
        format_features
    }

    pub(crate) fn describe_format_features(
        &self,
        format: TextureFormat,
    ) -> Result<wgt::TextureFormatFeatures, MissingFeatures> {
        self.require_features(format.required_features())?;

        let using_device_features = self
            .features
            .contains(wgt::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES);
        // If we're running downlevel, we need to manually ask the backend what
        // we can use as we can't trust WebGPU.
        let downlevel = !self
            .downlevel
            .flags
            .contains(wgt::DownlevelFlags::WEBGPU_TEXTURE_FORMAT_SUPPORT);

        if using_device_features || downlevel {
            Ok(self.get_texture_format_features(format))
        } else {
            Ok(format.guaranteed_format_features(self.features))
        }
    }

    #[cfg(feature = "replay")]
    pub(crate) fn wait_for_submit(
        &self,
        submission_index: crate::SubmissionIndex,
    ) -> Result<(), DeviceError> {
        let last_done_index = unsafe { self.raw().get_fence_value(self.fence.as_ref()) }
            .map_err(|e| self.handle_hal_error(e))?;
        if last_done_index < submission_index {
            unsafe { self.raw().wait(self.fence.as_ref(), submission_index, None) }
                .map_err(|e| self.handle_hal_error(e))?;
            if let Some(queue) = self.get_queue() {
                let closures = queue.lock_life().triage_submissions(submission_index);
                assert!(
                    closures.is_empty(),
                    "wait_for_submit is not expected to work with closures"
                );
            }
        }
        Ok(())
    }

    pub fn create_query_set(
        self: &Arc<Self>,
        desc: &resource::QuerySetDescriptor,
    ) -> (Arc<QuerySet>, Option<resource::CreateQuerySetError>) {
        profiling::scope!("Device::create_query_set");
        let (query_set, error) = match self.create_query_set_inner(desc) {
            Ok(query_set) => (query_set, None),
            Err(e) => (QuerySet::invalid(Arc::clone(self), desc), Some(e)),
        };
        #[cfg(feature = "trace")]
        if let Some(ref mut trace) = *self.trace.lock() {
            use trace::IntoTrace;
            trace.add(trace::Action::CreateQuerySet {
                id: query_set.to_trace(),
                desc: desc.clone(),
            });
        }
        api_log!("Device::create_query_set -> {:?}", Arc::as_ptr(&query_set));
        (query_set, error)
    }

    pub(crate) fn create_query_set_inner(
        self: &Arc<Self>,
        desc: &resource::QuerySetDescriptor,
    ) -> Result<Arc<QuerySet>, resource::CreateQuerySetError> {
        use resource::CreateQuerySetError as Error;

        self.check_is_valid()?;

        match desc.ty {
            wgt::QueryType::Occlusion => {}
            wgt::QueryType::Timestamp => {
                self.require_features(wgt::Features::TIMESTAMP_QUERY)?;
            }
            wgt::QueryType::PipelineStatistics(..) => {
                self.require_features(wgt::Features::PIPELINE_STATISTICS_QUERY)?;
            }
        }

        if desc.count == 0 {
            return Err(Error::ZeroCount);
        }

        if desc.count > wgt::QUERY_SET_MAX_QUERIES {
            return Err(Error::TooManyQueries {
                count: desc.count,
                maximum: wgt::QUERY_SET_MAX_QUERIES,
            });
        }

        let hal_desc = desc.map_label(|label| label.to_hal(self.instance_flags));

        let raw = unsafe { self.raw().create_query_set(&hal_desc) }
            .map_err(|e| self.handle_hal_error_with_nonfatal_oom(e))?;

        let query_set = QuerySet {
            state: ResourceState::Valid(QuerySetState {
                raw: Snatchable::new(raw),
            }),
            device: self.clone(),
            label: desc.label.to_string(),
            tracking_data: TrackingData::new(self.tracker_indices.query_sets.clone()),
            desc: desc.map_label(|_| ()),
            initialized_slots: Mutex::new(
                rank::QUERY_SET_INITIALIZED_SLOTS,
                bit_vec::BitVec::from_elem(desc.count as usize, false),
            ),
        };

        let query_set = Arc::new(query_set);

        Ok(query_set)
    }

    fn lose(&self, message: &str) {
        // Follow the steps at https://gpuweb.github.io/gpuweb/#lose-the-device.

        // Mark the device explicitly as invalid. This is checked in various
        // places to prevent new work from being submitted.
        self.valid.store(false, Ordering::Release);

        // 1) Resolve the GPUDevice device.lost promise.
        if let Some(device_lost_closure) = self.device_lost_closure.lock().take() {
            device_lost_closure(DeviceLostReason::Unknown, message.to_string());
        }

        // 2) Complete any outstanding mapAsync() steps.
        // 3) Complete any outstanding onSubmittedWorkDone() steps.

        // These parts are passively accomplished by setting valid to false,
        // since that will prevent any new work from being added to the queues.
        // Future calls to poll_devices will continue to check the work queues
        // until they are cleared, and then drop the device.
    }

    fn release_gpu_resources(&self) {
        // This is called when the device is lost, which makes every associated
        // resource invalid and unusable. This is an opportunity to release all of
        // the underlying gpu resources, even though the objects remain visible to
        // the user agent. We purge this memory naturally when resources have been
        // moved into the appropriate buckets, so this function just needs to
        // initiate movement into those buckets, and it can do that by calling
        // "destroy" on all the resources we know about.

        // During these iterations, we discard all errors. We don't care!
        let trackers = self.trackers.lock();
        for buffer in trackers.buffers.used_resources() {
            if let Some(buffer) = Weak::upgrade(buffer) {
                buffer.destroy();
            }
        }
        for texture in trackers.textures.used_resources() {
            if let Some(texture) = Weak::upgrade(texture) {
                texture.destroy();
            }
        }
    }

    pub(crate) fn new_usage_scope(&self) -> UsageScope<'_> {
        UsageScope::new_pooled(
            &self.usage_scopes,
            &self.tracker_indices,
            self.ordered_buffer_usages,
            self.ordered_texture_usages,
        )
    }

    /// `device_lost_closure` might never be called.
    pub fn set_device_lost_closure(&self, device_lost_closure: DeviceLostClosure) {
        self.device_lost_closure.lock().replace(device_lost_closure);
    }

    pub fn destroy(self: &Arc<Self>) {
        api_log!("Device::destroy {:?}", Arc::as_ptr(self));

        // Follow the steps at
        // https://gpuweb.github.io/gpuweb/#dom-gpudevice-destroy.
        // It's legal to call destroy multiple times, but if the device
        // is already invalid, there's nothing more to do. There's also
        // no need to return an error.
        if !self.is_valid() {
            return;
        }

        // The last part of destroy is to lose the device. The spec says
        // delay that until all "currently-enqueued operations on any
        // queue on this device are completed." This is accomplished by
        // setting valid to false, and then relying upon maintain to
        // check for empty queues and a DeviceLostClosure. At that time,
        // the DeviceLostClosure will be called with "destroyed" as the
        // reason.
        self.valid.store(false, Ordering::Release);
    }

    pub fn get_internal_counters(&self) -> wgt::InternalCounters {
        wgt::InternalCounters {
            hal: self.get_hal_counters(),
            core: wgt::CoreCounters {},
        }
    }

    pub fn get_hal_counters(&self) -> wgt::HalCounters {
        self.raw().get_internal_counters()
    }

    pub fn generate_allocator_report(&self) -> Option<wgt::AllocatorReport> {
        self.raw().generate_allocator_report()
    }
}

crate::impl_resource_type!(Device);
crate::impl_labeled!(Device);
crate::impl_storage_item!(Device);