cubecl_cpp/metal/
binding.rs

1use crate::{
2    Dialect,
3    metal::AddressSpace,
4    shared::{Builtin, Component, KernelArg},
5};
6
7pub fn format_global_binding_arg<D: Dialect>(
8    binding: &KernelArg<D>,
9    attr_idx: &mut usize,
10    f: &mut core::fmt::Formatter<'_>,
11) -> core::fmt::Result {
12    let comma = if *attr_idx > 0 { "," } else { "" };
13    let address_space = AddressSpace::from(binding);
14    let ty = binding.value.item();
15    let name = binding.value;
16    let attribute = address_space.attribute();
17
18    write!(f, "{comma}\n    {ty} {name}",)?;
19    // attribute
20    attribute.indexed_fmt(*attr_idx, f)?;
21    *attr_idx += 1;
22    Ok(())
23}
24
25pub fn format_metal_builtin_binding_arg<D: Dialect>(
26    f: &mut core::fmt::Formatter<'_>,
27    builtin: &Builtin<D>,
28    comma: bool,
29) -> core::fmt::Result {
30    let ty = builtin.item();
31    let attribute = builtin.attribute();
32    let comma = if comma { "," } else { "" };
33    write!(f, "{comma}\n    {ty} {builtin} {attribute}",)?;
34    Ok(())
35}
cubecl_cpp/metal/binding.rs

cubecl_cpp/metal/
binding.rs
