gfx-hal Texture操作源码简析 2019.1.21

文档列表见：Rust 移动端跨平台复杂图形渲染项目开发系列总结（目录）

Texture的使用已在以OpenGL/ES视角介绍gfx-hal(Vulkan) Texture接口使用介绍，本文档分析Vulkan接口如何映射到Metal，内容安排完全参考前一个文档，建议结合阅读。

源码路径入口: src/backend/metal/src。
git commit id: c6fbead。

Metal 仿 Vulkan API 的纹理创建流程分析

通过create_image()创建n::Image(即MTLTextureDescriptor)

create_image()根据参数创建并配置MTLTextureDescriptor。usage: image::Usage参数决定了应该申请什么内存类型的纹理。内存类型影响性能。按照Metal开发流程，之后由此描述符向MTLDevice申请MTLTexture对象，MTLTextureDescriptor对象可以复用。

// 位置：gfx/src/backend/metal/src/device.rs
// 调用示例：
// (width, height, 1 /* image::Layer */, 1 /* image::NumSamples */),
// 1 /* image::Level */,
// Rgba8Srgb::SELF,
// image::Tiling::Optimal,
// image::Usage::TRANSFER_DST | image::Usage::SAMPLED,
// image::ViewCapabilities::empty(),
unsafe fn create_image(
    &self,
    kind: image::Kind,
    mip_levels: image::Level,
    format: format::Format,
    tiling: image::Tiling,
    usage: image::Usage,
    view_caps: image::ViewCapabilities,
) -> Result<n::Image, image::CreationError> {
    debug!("create_image {:?} with {} mips of {:?} {:?} and usage {:?}",
        kind, mip_levels, format, tiling, usage);

    let is_cube = view_caps.contains(image::ViewCapabilities::KIND_CUBE);
    let mtl_format = self.shared.private_caps
        .map_format(format)
        .ok_or(image::CreationError::Format(format))?;

    // 创建MTLTextureDescriptor
    let descriptor = metal::TextureDescriptor::new();

    // 2D纹理(不是纹理数组)对应的值为
    // mtl_type = MTLTextureType::D2, num_layers = None
    let (mtl_type, num_layers) = match kind {
        image::Kind::D1(_, 1) => {
            assert!(!is_cube);
            (MTLTextureType::D1, None)
        }
        image::Kind::D1(_, layers) => {
            assert!(!is_cube);
            (MTLTextureType::D1Array, Some(layers))
        }
        image::Kind::D2(_, _, layers, 1) => {
            if is_cube && layers > 6 {
                assert_eq!(layers % 6, 0);
                (MTLTextureType::CubeArray, Some(layers / 6))
            } else if is_cube {
                assert_eq!(layers, 6);
                (MTLTextureType::Cube, None)
            } else if layers > 1 {
                (MTLTextureType::D2Array, Some(layers))
            } else {
                // 2D纹理(不是纹理数组)执行此分支
                (MTLTextureType::D2, None)
            }
        }
        image::Kind::D2(_, _, 1, samples) if !is_cube => {
            descriptor.set_sample_count(samples as u64);
            (MTLTextureType::D2Multisample, None)
        }
        image::Kind::D2(..) => {
            error!("Multi-sampled array textures or cubes are not supported: {:?}", kind);
            return Err(image::CreationError::Kind)
        }
        image::Kind::D3(..) => {
            assert!(!is_cube && !view_caps.contains(image::ViewCapabilities::KIND_2D_ARRAY));
            (MTLTextureType::D3, None)
        }
    };

    descriptor.set_texture_type(mtl_type);
    // 处理纹理数组
    if let Some(count) = num_layers {
        descriptor.set_array_length(count as u64);
    }
    let extent = kind.extent();
    descriptor.set_width(extent.width as u64);
    descriptor.set_height(extent.height as u64);
    descriptor.set_depth(extent.depth as u64);
    descriptor.set_mipmap_level_count(mip_levels as u64);
    descriptor.set_pixel_format(mtl_format);
    // Metal文档：
    // The default value is `MTLTextureUsageShaderRead`. You should 
    // always set specific texture usage values; do not rely on the 
    // `MTLTextureUsageUnknown` value for the best performance.
    // Set this value to `MTLTextureUsageRenderTarget` if you 
    // intend to use the resulting `MTLTexture` object as a render target. 
    // This may significantly improve your app’s performance (with certain hardware).
    //
    // iOS 9才支持此API，gfx没适配iOS 8
    descriptor.set_usage(conv::map_texture_usage(usage, tiling));

    let base = format.base_format();
    let format_desc = base.0.desc();
    let mip_sizes = (0 .. mip_levels)
        .map(|level| {
            // pitches为某一level的二维图像大小，以字节表示。
            // 比如，640x480 RGBA在mipmap level 1 的pitches = 1228800
            let pitches = n::Image::pitches_impl(extent.at_level(level), format_desc);
            num_layers.unwrap_or(1) as buffer::Offset * pitches[3]
        })
        .collect();

    let host_usage = image::Usage::TRANSFER_SRC | image::Usage::TRANSFER_DST;
    // 通过Tiling::Optimal参数创建的2D纹理只为GPU读取优化，
    // Tiling::Linear参数创建的2D纹理为CPU/GPU读取优化，
    // 最终得到host_visible = false
    let host_visible = mtl_type == MTLTextureType::D2 &&
        mip_levels == 1 && num_layers.is_none() &&
        format_desc.aspects.contains(format::Aspects::COLOR) &&
        tiling == image::Tiling::Linear &&
        host_usage.contains(usage);

    Ok(n::Image {
        like: n::ImageLike::Unbound {
            descriptor,
            mip_sizes,
            host_visible,
        },
        kind,
        format_desc,
        shader_channel: base.1.into(),
        mtl_format,
        mtl_type,
    })
}
复制代码

Image 相关结构定义如下：

#[derive(Debug)]
pub struct Image {
    pub(crate) like: ImageLike,
    pub(crate) kind: image::Kind,
    pub(crate) format_desc: FormatDesc,
    pub(crate) shader_channel: Channel,
    pub(crate) mtl_format: metal::MTLPixelFormat,
    pub(crate) mtl_type: metal::MTLTextureType,
}

#[derive(Debug)]
pub enum ImageLike {
    /// This image has not yet been bound to memory.
    Unbound {
        descriptor: metal::TextureDescriptor,
        mip_sizes: Vec<buffer::Offset>,
        host_visible: bool,
    },
    /// This is a linearly tiled HOST-visible image, which is represented by a buffer.
    Buffer(Buffer),
    /// This is a regular image represented by a texture.
    Texture(metal::Texture),
}

/// Specifies the kind of an image to be allocated.
#[derive(Clone, Copy, Debug, Eq, Hash, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serde", derive(Serialize, Deserialize))]
pub enum Kind {
    /// A single one-dimensional row of texels.
    D1(Size, Layer),
    /// Two-dimensional image.
    D2(Size, Size, Layer, NumSamples),
    /// Volumetric image.
    D3(Size, Size, Size),
}
复制代码

get_image_requirements()配置n::Image(即MTLTextureDescriptor)的存储参数

get_image_requirements(&n::Image)根据前面创建的n::Image即MTLTextureDescriptor属性及当前GPU能力创建合适的Requirements。

unsafe fn get_image_requirements(&self, image: &n::Image) -> memory::Requirements {
    let (descriptor, mip_sizes, host_visible) = match image.like {
        n::ImageLike::Unbound { ref descriptor, ref mip_sizes, host_visible } =>
            (descriptor, mip_sizes, host_visible),
        n::ImageLike::Texture(..) |
        n::ImageLike::Buffer(..) => panic!("Expected Image::Unbound"),
    };

    // fixme 有待确认 MacBook Pro (Retina, 15-inch, Mid 2015) 是否真不支持Resource Heaps
    if self.shared.private_caps.resource_heaps {
        // We don't know what memory type the user will try to allocate the image with, 
        // so we test them all get the most stringent ones. 
        // Note we don't check Shared because heaps can't use it
        let mut max_size = 0;
        let mut max_alignment = 0;
        let types = if host_visible {
            MemoryTypes::all()
        } else {
            MemoryTypes::PRIVATE
        };
        for (i, _) in self.memory_types.iter().enumerate() {
            if !types.contains(MemoryTypes::from_bits(1 << i).unwrap()) {
                continue
            }
            let (storage, cache_mode) = MemoryTypes::describe(i);
            // Metal文档：
            // iOS 9才支持，gfx没做适配
            // In iOS and tvOS, the default value is `MTLStorageModeShared`. 
            // In macOS, the default value is `MTLStorageModeManaged`.
            descriptor.set_storage_mode(storage);
            // API版本同上
            // The CPU cache mode used for the CPU mapping of the texture.
            // The default value is MTLCPUCacheModeDefaultCache.
            descriptor.set_cpu_cache_mode(cache_mode);

            let requirements = self.shared.device
                .lock()
                .heap_texture_size_and_align(descriptor);
            max_size = cmp::max(max_size, requirements.size);
            max_alignment = cmp::max(max_alignment, requirements.align);
        }
        memory::Requirements {
            size: max_size,
            alignment: max_alignment,
            type_mask: types.bits(),
        }
    } else if host_visible {
        assert_eq!(mip_sizes.len(), 1);
        let mask = self.shared.private_caps.buffer_alignment - 1;
        memory::Requirements {
            size: (mip_sizes[0] + mask) & !mask,
            alignment: self.shared.private_caps.buffer_alignment,
            type_mask: MemoryTypes::all().bits(),
        }
    } else {
        // 对于2D纹理，`create_image()`执行后host_visible = false，则macOS执行此分支
        memory::Requirements {
            size: mip_sizes.iter().sum(),
            alignment: 4,
            type_mask: MemoryTypes::PRIVATE.bits(),
        }
    }
}
复制代码

Requirements包含大小、对齐、内存类型信息。

/// Memory requirements for a certain resource (buffer/image).
#[derive(Clone, Copy, Debug)]
pub struct Requirements {
    /// Size in the memory.
    pub size: u64,
    /// Memory alignment.
    pub alignment: u64,
    /// Supported memory types.
    pub type_mask: u64,
}
复制代码

条件分配n::Image(即MTLTextureDescriptor)的后台Buffer

let memory_types = adapter.physical_device.memory_properties().memory_types;
// 2D纹理对应的device_type = 0
let device_type = get_memory_type_id_from_properties(
    memory_types,
    image_req,
    memory::Properties::DEVICE_LOCAL,
);

fn get_memory_type_id_from_properties(
    memory_types: &[hal::MemoryType],
    requirements: memory::Requirements,
    memory_properties: memory::Properties,
) -> hal::MemoryTypeId {
    memory_types
        .iter()
        .enumerate()
        .position(|(id, mem_type)| {
            // type_mask is a bit field where each bit represents a memory type. 
            // If the bit is set to 1 it means we can use that type for our buffer. 
            // So this code finds the first memory type that has a `1` (or, is allowed), 
            // and is visible to the CPU.
            requirements.type_mask & (1 << id) != 0
                && mem_type.properties.contains(memory_properties)
        })
        .unwrap()
        .into()
}
复制代码

allocate_memory() 根据 MTLStorageMode 类型分配 MTLBuffer 或 MTLHeap（gfx目前用&& false直接关闭了 MTLHeapDescriptor 分支。我问过他们，原因是他们还没找到 MTLHeap 合适的使用场合。），如果 MTLStorageMode 为MTLStorageMode::Private则在后面的bind_image_memory()分配 MTLTexture 对象。

unsafe fn allocate_memory(&self, memory_type: hal::MemoryTypeId, size: u64) 
    -> Result<n::Memory, AllocationError> {
    // 由于memory_type在前面执行为0，则2D纹理在macOS上执行结果为
    // storage = Private, cache = DefaultCache
    let (storage, cache) = MemoryTypes::describe(memory_type.0);
    let device = self.shared.device.lock();
    debug!("allocate_memory type {:?} of size {}", memory_type, size);

    // Heaps cannot be used for CPU coherent resources
    //TEMP: MacOS supports Private only, iOS and tvOS can do private/shared
    let heap = if self.shared.private_caps.resource_heaps 
                && storage != MTLStorageMode::Shared 
                && false {
        // iOS 10开始支持MTLHeapDescriptor，gfx没做适配
        let descriptor = metal::HeapDescriptor::new();
        descriptor.set_storage_mode(storage);
        descriptor.set_cpu_cache_mode(cache);
        descriptor.set_size(size);
        let heap_raw = device.new_heap(&descriptor);
        n::MemoryHeap::Native(heap_raw)
    } else if storage == MTLStorageMode::Private {
        // 由于前一分支有`&& false`且`storage = Private`，Image挂载的Memory执行此分支。
        // 后续方法调用才分配实际存储空间
        n::MemoryHeap::Private
    } else {
        // Buffer挂载的Memory执行此分支，分配实际存储空间
        let options = conv::resource_options_from_storage_and_cache(storage, cache);
        let cpu_buffer = device.new_buffer(size, options);
        debug!("\tbacked by cpu buffer {:?}", cpu_buffer.as_ptr());
        n::MemoryHeap::Public(memory_type, cpu_buffer)
    };

    Ok(n::Memory::new(heap, size))
}
复制代码

metal-backend的MemoryHeap与Memory定义如下所示：

#[derive(Debug)]
pub struct Memory {
    pub(crate) heap: MemoryHeap,
    pub(crate) size: u64,
}

#[derive(Debug)]
pub(crate) enum MemoryHeap {
    Private,
    Public(MemoryTypeId, metal::Buffer),
    Native(metal::Heap),
}
复制代码

分配MTLTexture对象

bind_image_memory(memory: &n::Memory, offset: u64, image: &mut n::Image)根据 MemoryHeap 决定从 MTLDevice 分配 MTLTexture对象或基于 MTLHeap 创建 MTLTexture 别名，即共享 MTLHeap 存储空间。

unsafe fn bind_image_memory(
    &self, memory: &n::Memory, offset: u64, image: &mut n::Image
) -> Result<(), BindError> {
    let like = {
        let (descriptor, mip_sizes) = match image.like {
            n::ImageLike::Unbound { ref descriptor, ref mip_sizes, .. } =>
                (descriptor, mip_sizes),
            n::ImageLike::Texture(..) |
            n::ImageLike::Buffer(..) => panic!("Expected Image::Unbound"),
        };

        match memory.heap {
            // 目前没执行此分支
            n::MemoryHeap::Native(ref heap) => {
                let resource_options = conv::resource_options_from_storage_and_cache(
                    heap.storage_mode(),
                    heap.cpu_cache_mode());
                descriptor.set_resource_options(resource_options);
                n::ImageLike::Texture(
                    heap.new_texture(descriptor)
                        .unwrap_or_else(|| {
                            // heap.new_texture执行失败的兜底方案
                            // 目前gfx逻辑不执行此 match 分支
                            // TODO: disable hazard tracking?
                            self.shared.device
                                .lock()
                                .new_texture(&descriptor)
                        })
                )
            },
            n::MemoryHeap::Public(_memory_type, ref cpu_buffer) => {
                assert_eq!(mip_sizes.len(), 1);
                n::ImageLike::Buffer(
                    n::Buffer::Bound {
                        raw: cpu_buffer.clone(),
                        range: offset .. offset + mip_sizes[0] as u64,
                        options: MTLResourceOptions::StorageModeShared,
                    }
                )
            }
            // 由前面n::MemoryHeap::Private可知，将执行下面支持，分配出MTLTexture对象。
            n::MemoryHeap::Private => {
                descriptor.set_storage_mode(MTLStorageMode::Private);
                n::ImageLike::Texture(
                    self.shared.device
                        .lock()
                        .new_texture(descriptor)
                )
            }
        }
    };

    Ok(image.like = like)
}
复制代码

上传CPU数据到纹理

创建Staging Buffer

此部分参考Buffer的相关操作。

创建Fence

创建用于数据拷贝的Submmit

创建带类型的Command Pool

let mut staging_pool = device.borrow().device.create_command_pool_typed(
    &device.borrow().queues,
    pool::CommandPoolCreateFlags::empty(),
    16,
);
复制代码

/// Create a new command pool for a given queue family.
///
/// *Note*: the family has to be associated by one as the `Gpu::queue_groups`.
fn create_command_pool(&self, family: QueueFamilyId, create_flags: CommandPoolCreateFlags) 
-> B::CommandPool;

/// Create a strongly typed command pool wrapper.
fn create_command_pool_typed<C>(
    &self,
    group: &QueueGroup<B, C>,
    flags: CommandPoolCreateFlags,
    max_buffers: usize,
) -> CommandPool<B, C> {
    let raw = self.create_command_pool(group.family(), flags);
    let mut pool = unsafe { CommandPool::new(raw) };
    pool.reserve(max_buffers);
    pool
}
复制代码

fn create_command_pool(
    &self, _family: QueueFamilyId, _flags: CommandPoolCreateFlags
) -> command::CommandPool {
    command::CommandPool::new(&self.shared, self.online_recording.clone())
}
复制代码

pub struct CommandPool {
    shared: Arc<Shared>,
    allocated: Vec<CommandBufferInnerPtr>,
    pool_shared: PoolSharedPtr,
}

impl CommandPool {
    pub(crate) fn new(
        shared: &Arc<Shared>,
        online_recording: OnlineRecording,
    ) -> Self {
        let pool_shared = PoolShared {
            #[cfg(feature = "dispatch")]
            dispatch_queue: match online_recording {
                OnlineRecording::Immediate |
                OnlineRecording::Deferred => None,
                OnlineRecording::Remote(priority) => Some(dispatch::Queue::global(priority.clone())),
            },
            online_recording,
        };
        CommandPool {
            shared: Arc::clone(shared),
            allocated: Vec::new(),
            pool_shared: Arc::new(RefCell::new(pool_shared)),
        }
    }
}
复制代码

创建Command Buffer

/// Get a primary command buffer for recording.
///
/// You can only record to one command buffer per pool at the same time.
/// If more command buffers are requested than allocated, new buffers will be reserved.
/// The command buffer will be returned in 'recording' state.
pub fn acquire_command_buffer<S: Shot>(
    &mut self, allow_pending_resubmit: bool
) -> CommandBuffer<B, C, S> {
    self.reserve(1);

    let buffer = &mut self.buffers[self.next_buffer];
    let mut flags = S::FLAGS;
    if allow_pending_resubmit {
        flags |= CommandBufferFlags::SIMULTANEOUS_USE;
    }
    buffer.begin(flags, CommandBufferInheritanceInfo::default());
    self.next_buffer += 1;
    unsafe {
        CommandBuffer::new(buffer)
    }
}
复制代码

创建Barrier

复制代码

向Command Buffer提交Barrier

复制代码

向Command Buffer提交Copy Buffer to Image命令

pub trait Backend: 'static + Sized + Eq + Clone + Hash + fmt::Debug + Any + Send + Sync {
    type CommandBuffer: command::RawCommandBuffer<Self>;
}
复制代码

/// A strongly-typed command buffer that will only implement methods 
/// that are valid for the operations it supports.
pub struct CommandBuffer<B: Backend, C, S = OneShot, L = Primary, R = <B as Backend>::CommandBuffer>
{
    pub(crate) raw: R,
    pub(crate) _marker: PhantomData<(B, C, S, L)>,
}
复制代码

trait RawCommandBuffer 定义了一系列具体后端需要实现的接口。

/// Copies regions from the source buffer to the destination image.
unsafe fn copy_buffer_to_image<T>(
    &mut self,
    src: &B::Buffer,
    dst: &B::Image,
    dst_layout: Layout,
    regions: T,
) where
    T: IntoIterator,
    T::Item: Borrow<BufferImageCopy>;
复制代码

• hal/command/compute.rs
• hal/command/graphics.rs
• hal/command/transfer.rs

分别给 hal/command/mod.rs 中定义的 CommandBuffer 添加相应用途的API。传输数据用 copy_*** 系列API，故调用了 hal/command/transfer.rs 的 copy_buffer_to_image() ，它内部转发给具体图形后端的同名方法。

/// Identical to the `RawCommandBuffer` method of the same name.
pub unsafe fn copy_buffer_to_image<T>(
    &mut self,
    src: &B::Buffer,
    dst: &B::Image,
    dst_layout: image::Layout,
    regions: T,
) where
    T: IntoIterator,
    T::Item: Borrow<BufferImageCopy>,
{
    self.raw.copy_buffer_to_image(src, dst, dst_layout, regions)
}
复制代码

metal-backend的 copy_buffer_to_image() 实现如下：

fn copy_buffer_to_image<T>(
    &mut self,
    src: &native::Buffer,
    dst: &native::Image,
    _dst_layout: Layout,
    regions: T,
) where
    T: IntoIterator,
    T::Item: Borrow<com::BufferImageCopy>,
{
    match dst.like {
        native::ImageLike::Texture(ref dst_raw) => {
            // 由前面可知，从StagingBuffer通过BlitCommandBuffer拷贝数据到Texture执行此分支
            let commands = regions.into_iter().filter_map(|region| {
                let r = region.borrow();
                if r.image_extent.is_empty() {
                    None
                } else {
                    Some(soft::BlitCommand::CopyBufferToImage {
                        src: AsNative::from(src.raw.as_ref()),
                        dst: AsNative::from(dst_raw.as_ref()),
                        dst_desc: dst.format_desc,
                        region: com::BufferImageCopy {
                            buffer_offset: r.buffer_offset + src.range.start,
                            .. r.clone()
                        },
                    })
                }
            });
            self.inner
                .borrow_mut()
                .sink()
                .blit_commands(commands);
        }
        native::ImageLike::Buffer(ref dst_buffer) => {
            self.copy_buffer(src, dst_buffer, regions.into_iter().map(|region| {
                let r = region.borrow();
                com::BufferCopy {
                    src: r.buffer_offset,
                    dst: dst.byte_offset(r.image_offset),
                    size: dst.byte_extent(r.image_extent),
                }
            }))
        }
    }
}
复制代码

由上可知，Buffer到Buffer的拷贝执行 copy_buffer() 。

unsafe fn copy_buffer<T>(
    &mut self,
    src: &native::Buffer,
    dst: &native::Buffer,
    regions: T,
) where
    T: IntoIterator,
    T::Item: Borrow<com::BufferCopy>,
{
    let pso = &*self.shared.service_pipes.copy_buffer;
    let wg_size = MTLSize {
        width: pso.thread_execution_width(),
        height: 1,
        depth: 1,
    };

    let (src_raw, src_range) = src.as_bound();
    let (dst_raw, dst_range) = dst.as_bound();

    let mut compute_datas = Vec::new();
    let mut inner = self.inner.borrow_mut();
    let mut blit_commands = Vec::new();
    let mut compute_commands = vec![ //TODO: get rid of heap
        soft::ComputeCommand::BindPipeline(pso),
    ];

    for region in regions {
        let r = region.borrow();
        if r.size % WORD_SIZE as u64 == 0 &&
            r.src % WORD_SIZE as u64 == 0 &&
            r.dst % WORD_SIZE as u64 == 0
        {
            blit_commands.push(soft::BlitCommand::CopyBuffer {
                src: AsNative::from(src_raw),
                dst: AsNative::from(dst_raw),
                region: com::BufferCopy {
                    src: r.src + src_range.start,
                    dst: r.dst + dst_range.start,
                    size: r.size,
                },
            });
        } else {
            // not natively supported, going through a compute shader
            assert_eq!(0, r.size >> 32);
            let src_aligned = r.src & !(WORD_SIZE as u64 - 1);
            let dst_aligned = r.dst & !(WORD_SIZE as u64 - 1);
            let offsets = (r.src - src_aligned) | ((r.dst - dst_aligned) << 16);
            let size_and_offsets = [r.size as u32, offsets as u32];
            compute_datas.push(Box::new(size_and_offsets));

            let wg_count = MTLSize {
                width: (r.size + wg_size.width - 1) / wg_size.width,
                height: 1,
                depth: 1,
            };

            compute_commands.push(soft::ComputeCommand::BindBuffer {
                index: 0,
                buffer: AsNative::from(dst_raw),
                offset: dst_aligned + dst_range.start,
            });
            compute_commands.push(soft::ComputeCommand::BindBuffer {
                index: 1,
                buffer: AsNative::from(src_raw),
                offset: src_aligned + src_range.start,
            });
            compute_commands.push(soft::ComputeCommand::BindBufferData {
                index: 2,
                words: unsafe {
                    // Rust doesn't see that compute_datas will not lose this item
                    // and the boxed contents can't be moved otherwise.
                    mem::transmute(&compute_datas.last().unwrap()[..])
                },
            });
            compute_commands.push(soft::ComputeCommand::Dispatch {
                wg_size,
                wg_count,
            });
        }
    }

    let sink = inner.sink();
    if !blit_commands.is_empty() {
        sink.blit_commands(blit_commands.into_iter());
    }
    if compute_commands.len() > 1 { // first is bind PSO
        sink.quick_compute("copy_buffer", compute_commands.into_iter());
    }
}
复制代码

结束Command Buffer编码

复制代码

下载纹理数据到CPU

假设用Vulkan接口模拟glReadPixels()，大概分成以下几步：

Sampler

采样器表达了Shader读取纹理(Image)的方式，比如U、V、W坐标超出范围时是采取重复还是截断策略。

创建Sampler

使用Vulkan接口实现

VkSamplerCreateInfo samplerInfo;
samplerInfo.sType = VK_STRUCTURE_TYPE_SAMPLER_CREATE_INFO;
samplerInfo.pNext = nullptr;
samplerInfo.flags = 0;

samplerInfo.mipmapMode = mVkMipmapMode;
samplerInfo.minFilter = mVkMinFilter;
samplerInfo.magFilter = mVkMagFilter;
samplerInfo.addressModeU = mVkAddressModeU;
samplerInfo.addressModeV = mVkAddressModeV;
samplerInfo.addressModeW = mVkAddressModeW;

samplerInfo.minLod = mMinLod;
samplerInfo.maxLod = mMaxLod;
samplerInfo.mipLodBias = mMipLodBias;
samplerInfo.anisotropyEnable = mAnisotropyEnabled;
samplerInfo.maxAnisotropy = mMaxAnisotropy;
samplerInfo.compareEnable = mCompareEnabled;
samplerInfo.compareOp = mVkCompareOp;
samplerInfo.borderColor = mVkBorderColor;
samplerInfo.unnormalizedCoordinates = mUnnormalizedCoordinates;

VkResult err = vkCreateSampler(mVkDevice, &samplerInfo, nullptr, &mVkSampler);
// deal with (err != VK_ERROR_OUT_OF_HOST_MEMORY && err != VK_ERROR_OUT_OF_DEVICE_MEMORY && err != VK_ERROR_TOO_MANY_OBJECTS);
复制代码

使用gfx-hal接口实现

let sampler = match device.create_sampler(image::SamplerInfo::new(image::Filter::Linear, image::WrapMode::Clamp)) {
    Ok(sampler) => Some(sampler),
    Err(msg) => {
        error!("{}", msg);
        None
    };
}
复制代码

使用gfx-hal接口源码

/// Specifies how to sample from an image.
// TODO: document the details of sampling.
#[derive(Clone, Debug, Eq, Hash, PartialEq)]
#[cfg_attr(feature = "serde", derive(Serialize, Deserialize))]
pub struct SamplerInfo {
    /// Minification filter method to use.
    pub min_filter: Filter,
    /// Magnification filter method to use.
    pub mag_filter: Filter,
    /// Mip filter method to use.
    pub mip_filter: Filter,
    /// Wrapping mode for each of the U, V, and W axis (S, T, and R in OpenGL
    /// speak).
    pub wrap_mode: (WrapMode, WrapMode, WrapMode),
    /// This bias is added to every computed mipmap level (N + lod_bias). For
    /// example, if it would select mipmap level 2 and lod_bias is 1, it will
    /// use mipmap level 3.
    pub lod_bias: Lod,
    /// This range is used to clamp LOD level used for sampling.
    pub lod_range: Range<Lod>,
    /// Comparison mode, used primary for a shadow map.
    pub comparison: Option<Comparison>,
    /// Border color is used when one of the wrap modes is set to border.
    pub border: PackedColor,
    /// Anisotropic filtering.
    pub anisotropic: Anisotropic,
}

impl SamplerInfo {
    /// Create a new sampler description with a given filter method for all filtering operations
    /// and a wrapping mode, using no LOD modifications.
    pub fn new(filter: Filter, wrap: WrapMode) -> Self {
        SamplerInfo {
            min_filter: filter,
            mag_filter: filter,
            mip_filter: filter,
            wrap_mode: (wrap, wrap, wrap),
            lod_bias: Lod(0),
            lod_range: Lod(-8000)..Lod(8000),
            comparison: None,
            border: PackedColor(0),
            anisotropic: Anisotropic::Off,
        }
    }
}
复制代码

gfx-hal Vulkan模块源码

fn create_sampler(&self, sampler_info: image::SamplerInfo) -> Result<n::Sampler, d::AllocationError> {
    use hal::pso::Comparison;

    let (anisotropy_enable, max_anisotropy) = match sampler_info.anisotropic {
        image::Anisotropic::Off => (vk::VK_FALSE, 1.0),
        image::Anisotropic::On(aniso) => {
            if self.raw.1.contains(Features::SAMPLER_ANISOTROPY) {
                (vk::VK_TRUE, aniso as f32)
            } else {
                warn!("Anisotropy({}) was requested on a device with disabled feature", aniso);
                (vk::VK_FALSE, 1.0)
            }
        },
    };
    let info = vk::SamplerCreateInfo {
        s_type: vk::StructureType::SamplerCreateInfo,
        p_next: ptr::null(),
        flags: vk::SamplerCreateFlags::empty(),
        mag_filter: conv::map_filter(sampler_info.mag_filter),
        min_filter: conv::map_filter(sampler_info.min_filter),
        mipmap_mode: conv::map_mip_filter(sampler_info.mip_filter),
        address_mode_u: conv::map_wrap(sampler_info.wrap_mode.0),
        address_mode_v: conv::map_wrap(sampler_info.wrap_mode.1),
        address_mode_w: conv::map_wrap(sampler_info.wrap_mode.2),
        mip_lod_bias: sampler_info.lod_bias.into(),
        anisotropy_enable,
        max_anisotropy,
        compare_enable: if sampler_info.comparison.is_some() { vk::VK_TRUE } else { vk::VK_FALSE },
        compare_op: conv::map_comparison(sampler_info.comparison.unwrap_or(Comparison::Never)),
        min_lod: sampler_info.lod_range.start.into(),
        max_lod: sampler_info.lod_range.end.into(),
        border_color: match conv::map_border_color(sampler_info.border) {
            Some(bc) => bc,
            None => {
                error!("Unsupported border color {:x}", sampler_info.border.0);
                vk::BorderColor::FloatTransparentBlack
            }
        },
        unnormalized_coordinates: vk::VK_FALSE,
    };

    let result = unsafe {
        self.raw.0.create_sampler(&info, None)
    };

    match result {
        Ok(sampler) => Ok(n::Sampler(sampler)),
        Err(vk::Result::ErrorTooManyObjects) => Err(d::AllocationError::TooManyObjects),
        Err(vk::Result::ErrorOutOfHostMemory) => Err(d::OutOfMemory::OutOfHostMemory.into()),
        Err(vk::Result::ErrorOutOfDeviceMemory) => Err(d::OutOfMemory::OutOfDeviceMemory.into()),
        _ => unreachable!(),
    }
}
复制代码

gfx-hal Metal模块源码

To create a sampler, first create a MTLSamplerDescriptor object and configure the descriptor’s properties. Then call the makeSamplerState(descriptor:) method on the MTLDevice object that will use this sampler. Once the sampler is created, the descriptor can be disposed of or reconfigured to create other sampler objects.

developer.apple.com/documentati…

fn create_sampler(&self, info: image::SamplerInfo) -> Result<n::Sampler, AllocationError> {
    // todo 缺少supportArgumentBuffers API iOS 11
    let descriptor = metal::SamplerDescriptor::new();

    descriptor.set_min_filter(conv::map_filter(info.min_filter));
    descriptor.set_mag_filter(conv::map_filter(info.mag_filter));
    descriptor.set_mip_filter(match info.mip_filter {
        // Note: this shouldn't be required, but Metal appears to be confused when mipmaps
        // are provided even with trivial LOD bias.
        image::Filter::Nearest if info.lod_range.end < image::Lod::from(0.5) =>
            MTLSamplerMipFilter::NotMipmapped,
        image::Filter::Nearest => MTLSamplerMipFilter::Nearest,
        image::Filter::Linear => MTLSamplerMipFilter::Linear,
    });

    if let image::Anisotropic::On(aniso) = info.anisotropic {
        descriptor.set_max_anisotropy(aniso as _);
    }

    let (s, t, r) = info.wrap_mode;
    descriptor.set_address_mode_s(conv::map_wrap_mode(s));
    descriptor.set_address_mode_t(conv::map_wrap_mode(t));
    descriptor.set_address_mode_r(conv::map_wrap_mode(r));

    unsafe { descriptor.set_lod_bias(info.lod_bias.into()) };
    descriptor.set_lod_min_clamp(info.lod_range.start.into());
    descriptor.set_lod_max_clamp(info.lod_range.end.into());
    
    let caps = &self.private_caps;
    // TODO: Clarify minimum macOS version with Apple (43707452)
    if (caps.os_is_mac && caps.has_version_at_least(10, 13)) ||
        (!caps.os_is_mac && caps.has_version_at_least(9, 0)) {
        descriptor.set_lod_average(true); // optimization
    }

    if let Some(fun) = info.comparison {
        // iOS 9.0 API
        // The sampler comparison function used when sampling texels from a depth texture.
        descriptor.set_compare_function(conv::map_compare_function(fun));
    }
    if [r, s, t].iter().any(|&am| am == image::WrapMode::Border) {
        descriptor.set_border_color(match info.border.0 {
            0x00000000 => MTLSamplerBorderColor::TransparentBlack,
            0x000000FF => MTLSamplerBorderColor::OpaqueBlack,
            0xFFFFFFFF => MTLSamplerBorderColor::OpaqueWhite,
            other => {
                error!("Border color 0x{:X} is not supported", other);
                MTLSamplerBorderColor::TransparentBlack
            }
        });
    }

    Ok(n::Sampler(
        self.shared.device
        .lock()
        .new_sampler(&descriptor)
    ))
}
复制代码

销毁Sampler

使用Vulkan接口实现

if (mVkSampler != VK_NULL_HANDLE) {
    vkDestroySampler(mVkDevice, mVkSampler, nullptr);
    mVkSampler = VK_NULL_HANDLE;
}
复制代码

使用gfx-hal接口实现

device.destroy_sampler(sampler);
复制代码

gfx-hal Vulkan模块源码

fn destroy_sampler(&self, sampler: n::Sampler) {
    unsafe { self.raw.0.destroy_sampler(sampler.0, None); }
}
复制代码

gfx-hal Metal模块源码

fn destroy_sampler(&self, _sampler: n::Sampler) {
}
复制代码