Metal Performans Primitifleri

Tüm OS: Metal 4 ve sonrası Metal Performance Primitives desteği sağlar.

Metal Performance Primitives, Apple silicon üzerinde verimlilik ve performans için tasarlanmış optimize edilmiş ilkel işlemlerden oluşan bir kütüphanedir.

Başlık dosyası:

<MetalPerformancePrimitives/MetalPerformancePrimitives.h>

mpp isim alanı içinde bu fonksiyonları tanımlar.

tensor_ops isim alanı (mpp içinde) matris çarpımı ve konvolüsyon gibi tensor işlemlerini içerir. Bu fonksiyonlar tensor ve cooperative_tensors üzerinde çalışır ve Apple silicon GPU'ları için ayarlanmıştır.

Desteklenen GPU aileleri developer.apple.com üzerindeki Metal Feature Set Tables içinde listelenmiştir.

---

7.1 Çalıştırma Kapsamları

Tüm OS: Metal 4 ve sonrası çalıştırma kapsamlarını destekler.

TensorOps gibi işlemler şu şekillerde çalıştırılabilir:

• Tek bir thread üzerinde
• Bir SIMD‑group içindeki thread'ler arasında işbirliğiyle
• Birden fazla SIMD‑group arasında

Çalıştırma kapsamları işbirliği düzeyini belirtir.

Çalıştırma Kapsamı Türleri

• execution_thread — İşbirliği kapsamı tek bir thread'dir.
• simdgroups_per_threadgroup — İşbirliği kapsamı N SIMD‑group'tur.

TensorOps N = 1 veya simdgroups_per_threadgroup değerini destekler. N = 1 olduğunda execution_simdgroup kullanabilirsiniz.

---

7.2 Tensor İşlemleri (TensorOps)

Tüm OS: Metal 4 ve sonrası tensor işlemlerini destekler.

TensorOps, tensor ve cooperative_tensors üzerinde çalışan GPU hızlandırmalı fonksiyonlardır.

Bu yapılar çalıştırma kapsamı gibi özelliklerle örneklenen sınıf şablonlarıdır.

Örnek şablon:

template <
    matmul2d_descriptor Desc,
    typename Scope,
    class... Args>
matmul2d

Davranış

Bir TensorOp run metodu çağrıldığında, belirtilen kapsam içindeki tüm thread'ler bu metodu çağırmalıdır, aksi halde sonuç tanımsızdır.

Örnek:

• Kapsam execution_simdgroup ise, SIMD‑group içindeki her thread run çağrısını yapmalıdır.
• Farklı SIMD‑group'lar birbirlerinden farklı yürütme akışına sahip olabilir.

TensorOps dahili olarak çalıştırma kapsamı seviyesinde bir barrier kullanabilir. Örneğin kapsam tüm bir threadgroup ise, TensorOp uygulaması içinde bir barrier kullanılması durumunda kodunuzun doğru çalışması gerekir.

TensorOps sonucu ElementType değeri device veya threadgroup adres alanında olan bir tensor içine yazıyorsa, sonuçları okumadan önce uygun thread kapsamına sahip bir barrier (bkz. bölüm 6.9.1) eklemeli ve uygun bellek bayraklarını ayarlamalısınız. Belleği thread adres alanında olan tensorlar veya cooperative_tensors için barrier kullanmanız gerekmez.

Örneğin TensorOp run metodu sonucu ElementType değeri threadgroup belleğinde olan bir tensor içine yazıyor ve kapsam execution_simdgroups<2> ise:

threadgroup_barrier(mem_flags::mem_threadgroup)

çağrısını tensor sonucunu okumadan önce yapın.

Başka bir örnek olarak TensorOp run metodu sonucu ElementType değeri device belleğinde olan bir tensor içine yazıyor ve kapsam execution_simdgroup ise:

simdgroup_barrier(mem_flags::mem_device)

çağrısını tensor sonucunu okumadan önce yapın.

---

Tablo 7.2 TensorOps

Genelleştirilmiş bir matris çarpımı gerçekleştirmek için bir nesne tanımlar:

C = A * B + C

A ve B host‑bound, origin‑shifted veya shader tarafından ayrılmış tensorlar olabilir.

C host‑bound, origin‑shifted, shader tarafından ayrılmış tensorlar veya cooperative_tensor olabilir.

TensorOp Şablon Sınıfları

template <convolution2d_descriptor Desc, typename Scope, typename... ConvArgs> convolution2d

Sinir ağlarında gerçekleşen bir 2D konvolüsyon gerçekleştiren bir nesne tanımlar. 2D ifadesi genişlik × yükseklik olan iki uzamsal boyutu temsil eder. Bu işlem tarafından tüketilen tensor 4 boyutludur.

Şu anda desteklenen tek Scope değeri:

execution_simdgroups<N>

burada N, simdgroups_per_threadgroup değeridir.

Daha fazla ayrıntı için 7.2.2 bölümüne bakın.

Daha fazla ayrıntı için 7.2.1 bölümüne bakın.

---

7.2.1 Matris Çarpımı

matmul2d şablon sınıfı iki tensorun genelleştirilmiş matris çarpımını gerçekleştirir:

C = A * B

veya sonucu bir tensor içine biriktirilen matris çarpımı:

C = A * B + C

Bu işlem M × K boyutlu A tensorunu K × N boyutlu B tensoru ile çarpar ve sonucu M × N boyutlu C tensoruna biriktirir.

A ve B host‑bound, origin‑shifted veya shader tarafından ayrılmış tensorlar olabilir. C ise host‑bound, origin‑shifted, shader tarafından ayrılmış tensorlar veya cooperative_tensor olabilir.

Tablo 7.3 desteklenen veri tipi kombinasyonlarını gösterir.

Tablo 7.4 OS 26.1 ve sonrası için desteklenen ek veri tiplerini gösterir.

---

Tablo 7.4 OS 26.1 ve Sonrasında Desteklenen Ek MatMul2D Veri Tipleri

Tensor Tipi Kombinasyonları

• bfloat / bfloat / bfloat
• bfloat / bfloat / float
• bfloat / float / float
• bfloat / char / bfloat
• bfloat / char / float
• float / bfloat / float
• char / bfloat / bfloat
• char / bfloat / float
• bfloat / half / bfloat
• bfloat / half / half
• bfloat / half / float
• half / bfloat / bfloat
• half / bfloat / half
• half / bfloat / float

---

matmul2d Tanımlayıcısı Oluşturma

matmul2d oluşturmak için önce aşağıdaki yapıcıyı kullanarak bir tanımlayıcı oluşturun:

matmul2d_descriptor(
    int M,
    int N,
    int K = dynamic_length_v<int>,
    bool transpose_left = false,
    bool transpose_right = false,
    bool relaxed_precision = false,
    mode matmul_mode = mode::multiply
);

---

Tablo 7.5 MatMul2D Tanımlayıcı Parametreleri

M, N, K
Tensor boyutları:

• M × K tensor A
• K × N tensor B
• M × N tensor C

relaxed_precision
İşlemin float veri tipi için gevşetilmiş hassasiyet kullanıp kullanamayacağını belirtir. Gevşetilmiş hassasiyet, çarpma işleminden önce mantissa'nın kısaltılmasına izin verir. Varsayılan değer false'tur.

matmul_mode
multiply veya multiply_accumulate işlemlerinden hangisinin yapılacağını belirtir. Varsayılan değer multiply'dır.

---

Tablo 7.6 MatMul2D Üye Fonksiyonları

Run

template <
    typename LeftOperandType,
    typename RightOperandType,
    typename DestinationOperandType>
void run(
    thread LeftOperandType &left,
    thread RightOperandType &right,
    thread DestinationOperandType &destination);

Bir matris çarpımı gerçekleştirir:

C = A * B

• C hedef tensor
• A sol tensor
• B sağ tensor

---

Hedef Cooperative Tensor Alma

template <
    typename LeftOperandType,
    typename RightOperandType,
    typename ElementType,
    typename CoordType = int>
cooperative_tensor<...>
get_destination_cooperative_tensor() thread const;

Matris çarpımı sonucunu saklayabilecek bir cooperative_tensor döndürür.

---

Satır Azaltma Hedefi

template <
    typename LeftOperandType,
    typename RightOperandType,
    typename ElementType,
    typename CoordType = int>
cooperative_tensor<...>
get_row_reduction_destination_cooperative_tensor() thread const;

Matris çarpımı sonucunun satır azaltma sonucunu saklayabilecek bir cooperative_tensor döndürür.

---

Sütun Azaltma Hedefi

template <
    typename LeftOperandType,
    typename RightOperandType,
    typename ElementType,
    typename CoordType = int>
cooperative_tensor<...>
get_column_reduction_destination_cooperative_tensor() thread const;

Matris çarpımı sonucunun sütun azaltma sonucunu saklayabilecek bir cooperative_tensor döndürür.

---

matmul2d Örnekleme

matmul2d şablonunu örneklemek için tanımlayıcıyı ve çalıştırma kapsamını geçirin.

Matris çarpımını yürütmek için run metodunu şu parametrelerle çağırın:

• sol tensor (A)
• sağ tensor (B)
• hedef tensor (C)

template <
    typename LeftOperandType,
    typename RightOperandType,
    typename DestinationOperandType>
void run(
    thread LeftOperandType &left,
    thread RightOperandType &right,
    thread DestinationOperandType &destination);

---

Örnek: Matris Çarpımı

#include <metal_tensor>
#include <MetalPerformancePrimitives/MetalPerformancePrimitives.h>

using namespace metal;
using namespace mpp;

[[ kernel ]] void matrixMultiply(
    tensor<device half, dextents<int, 2>> a [[ buffer(0) ]],
    tensor<device half, dextents<int, 2>> b [[ buffer(1) ]],
    tensor<device half, dextents<int, 2>> c [[ buffer(2) ]],
    uint2 tgid [[thread_position_in_grid]])
{
    // 4 SIMD-group'tan oluşan bir threadgroup için matmul işlemi oluştur.
    constexpr auto matmulDescriptor =
        tensor_ops::matmul2d_descriptor(64, 32, 0);

    tensor_ops::matmul2d<matmulDescriptor, execution_simdgroups<4>> matmulOp;

    // Bu threadgroup'un üzerinde çalışacağı uygun dilimi oluştur.
    auto mA = a.slice(0, tgid.y * 64);
    auto mB = b.slice(tgid.x * 32, 0);
    auto mC = c.slice(tgid.x * 32, tgid.y * 64);

    // C'nin sıfır ile başlatıldığı varsayılarak işlemi yürüt.
    matmulOp.run(mA, mB, mC);
}

---

Hedef Olarak cooperative_tensor Kullanma

Bir matmul2d TensorOp'un hedefi olarak cooperative_tensor kullanmak için aşağıdaki üye fonksiyonu kullanın:

template <
    typename LeftOperandType,
    typename RightOperandType,
    typename ElementType,
    typename CoordType = int>
cooperative_tensor<...>
get_destination_cooperative_tensor() thread const;

Bu fonksiyon, depolaması matmul2d kapsamındaki thread'ler arasında bölünmüş olan bir cooperative_tensor döndürür.

Tensor A, B ve C için desteklenen element tipleri için Tablo 7.3 ve Tablo 7.4'e bakın.

---

Örnek: cooperative_tensor ile matmul2d

#include <metal_tensor>
#include <MetalPerformancePrimitives/MetalPerformancePrimitives.h>

using namespace metal;
using namespace mpp;

[[ kernel ]] void gemmBias(
    tensor<device float, dextents<int, 2>> a [[ buffer(0) ]],
    tensor<device float, dextents<int, 2>> b [[ buffer(1) ]],
    tensor<device float, dextents<int, 2>> c [[ buffer(2) ]],
    device float* bufBias [[buffer(3)]],
    uint2 tgid [[thread_position_in_grid]])
{
    // Bias tensorunu buffer'dan oluştur.
    array<int,1> stride = {1};

    tensor<device float, dextents<int, 1>, tensor_inline>
        tBias(bufBias, dextents<int,1>(64), stride);

    // 4 SIMD-group'tan oluşan bir threadgroup için matmul işlemi oluştur.
    constexpr auto matmulDescriptor =
        tensor_ops::matmul2d_descriptor(
            64, 32, 0, false, false, false,
            tensor_ops::matmul2d_descriptor::mode::multiply_accumulate);

    tensor_ops::matmul2d<matmulDescriptor, execution_simdgroups<4>> matmulOp;

    // Cooperative tensor oluştur.
    auto cTc = matmulOp.get_destination_cooperative_tensor<
        decltype(a), decltype(b), float>();

    // Bias değerini yükle, matris çarpımını çalıştır ve sonucu sakla.
    cTc.load(tBias);
    matmulOp.run(a, b, cTc);
    cTc.store(c);
}

---

cooperative_tensor Üzerinde Azaltma İşlemleri

cooperative_tensor üzerinde satır veya sütun toplamı, maksimum veya minimum azaltmaları yapılabilir ve sonuç hedef bir 1D cooperative_tensor içine yazılabilir; bunun için matmul2d kapsamı şu olmalıdır:

execution_simdgroup

---

Tablo 7.7 Cooperative Tensorlar İçin Azaltma İlgili Fonksiyonlar

Satır Azaltma

template <
    class ElementType,
    class SrcExtents,
    class DstExtents,
    class SrcLayout,
    class DstLayout>
inline void reduce_rows(
    thread metal::cooperative_tensor<
        ElementType, SrcExtents, SrcLayout> &source,
    thread metal::cooperative_tensor<
        ElementType, DstExtents, DstLayout> &destination,
    reduction_operation op = reduction_operation::sum,
    ElementType identity =
        reduction_operation_identity<ElementType>::sum_identity);

Her satırın azaltma sonucunu döndürür ve sonucu hedef cooperative_tensor içine saklar. Varsayılan işlem sum'dır.

---

Sütun Azaltma

template <
    class ElementType,
    class SrcExtents,
    class DstExtents,
    class SrcLayout,
    class DstLayout>
inline void reduce_columns(
    thread metal::cooperative_tensor<
        ElementType, SrcExtents, SrcLayout> &source,
    thread metal::cooperative_tensor<
        ElementType, DstExtents, DstLayout> &destination,
    reduction_operation op = reduction_operation::sum,
    ElementType identity =
        reduction_operation_identity<ElementType>::sum_identity);

Her sütunun azaltma sonucunu döndürür ve sonucu hedef cooperative_tensor içine saklar. Varsayılan işlem sum'dır.

---

Azaltma İşlemi Numaralandırması

enum class reduction_operation {
    sum, // Satır/sütundaki elemanların toplamı
    max, // Satır/sütundaki maksimum değer
    min  // Satır/sütundaki minimum değer
};

---

Azaltma Kimlik Değerleri

template <typename ElementType>
struct reduction_operation_identity
{
    static const constant ElementType sum_identity;
    static const constant ElementType max_identity;
    static const constant ElementType min_identity;
};

---

Örnek: Satır Azaltma

[[ kernel ]] void rowReduction(
    tensor<device float, dextents<int, 2>> aT [[ buffer(0) ]],
    tensor<device float, dextents<int, 2>> bT [[ buffer(1) ]],
    tensor<device float, dextents<int, 2>> cT [[ buffer(2) ]],
    tensor<device float, dextents<int, 1>> dR [[ buffer(3) ]],
    uint2 tgid [[thread_position_in_grid]])
{
    constexpr auto matmulDescriptor =
        tensor_ops::matmul2d_descriptor(64, 32, 0);

    tensor_ops::matmul2d<matmulDescriptor, execution_simdgroup> matmulOp;

    auto cTdest = matmulOp.get_destination_cooperative_tensor<
        decltype(aT), decltype(bT), float>();

    // Matris çarpımını çalıştır
    matmulOp.run(aT, bT, cTdest);

    // Satır azaltma
    auto cTred =
        matmulOp.get_row_reduction_destination_cooperative_tensor<
            decltype(aT), decltype(bT), float>();

    reduce_rows(
        cTdest,
        cTred,
        tensor_ops::reduction_operation::sum,
        0.0f);

    cTred.store(dR);
}

---

Iterator Uyumluluğunu Kontrol Etme

İki cooperative tensorun iterator paylaşabilip paylaşamayacağını kontrol etmek için:

template <
    class SrcElementType,
    class DstElementType,
    class SrcExtents,
    class DstExtents,
    class SrcLayout,
    class DstLayout>
inline bool is_iterator_compatible(
    const thread metal::cooperative_tensor<
        SrcElementType,
        SrcExtents,
        SrcLayout> &source,
    const thread metal::cooperative_tensor<
        DstElementType,
        DstExtents,
        DstLayout> &destination);

Azaltma sonucunun map_iterator kullanan başka bir tensor ile kullanılıp kullanılamayacağını belirler. Uygunsa true döndürür.

---

Örnek: is_iterator_compatible ve map_iterator

[[ kernel ]] void gemm_map(
    tensor<device float, dextents<int, 2>> aT [[ buffer(0) ]],
    tensor<device float, dextents<int, 2>> bT [[ buffer(1) ]],
    tensor<device float, dextents<int, 2>> dT [[ buffer(2) ]])
{
    constexpr auto matmulDescriptor =
        tensor_ops::matmul2d_descriptor(64, 32, 0);

    tensor_ops::matmul2d<matmulDescriptor, execution_simdgroup> matmulOp;

    auto cTdest = matmulOp.get_destination_cooperative_tensor<
        decltype(aT), decltype(bT), float>();

    matmulOp.run(aT, bT, cTdest);

auto cTred =
        matmulOp.get_row_reduction_destination_cooperative_tensor<
            decltype(aT), decltype(bT), float>();

    auto identity = metal::numeric_limits<float>::lowest();

    reduce_rows(
        cTdest,
        cTred,
        tensor_ops::reduction_operation::min,
        identity);

    if (tensor_ops::is_iterator_compatible(cTdest, cTred)) {
        for (auto it = cTdest.begin(); it != cTdest.end(); it++) {
            auto cTred_it = cTred.map_iterator(it);
            *it += *cTred_it;
        }
    } else {
        // Başka bir işlem yap
    }

    cTdest.store(dT);
}

## 7.2.2 Konvolüsyon

`convolution2d` şablon sınıfı, 2B konvolüsyon gerçekleştirir; burada 2B, genişlik × yükseklik olmak üzere iki uzamsal boyutu ifade eder. İşlem, Tablo 7.8'de açıklandığı gibi bir tensör veya `cooperative_tensor` üretmek için bir aktivasyon ve bir ağırlık tensörü alır.

enum class convolution2d_activation_layout { nhwc, };

enum class convolution2d_weights_layout { hwio, };

convolution2d_descriptor( int4 destination_dimensions, int4 source_dimensions, int2 kernel_dimensions, convolution2d_activation_layout activation_layout = convolution2d_activation_layout::nhwc, convolution2d_weights_layout weight_layout = convolution2d_weights_layout::hwio, int2 strides = int2(1, 1), int2 dilations = int2(1, 1), int groups = 1, bool relaxed_precision = false, mode convolution2d_mode = mode::multiply);

### Parametreler

- **destination_dimensions** — Çıkış tensörünün boyutunu belirtir.
- **source_dimensions** — Giriş tensörünün boyutunu belirtir.
- **kernel_dimensions** — Konvolüsyon penceresinin boyutunu belirtir.
- **strides** — Konvolüsyonun adım uzunluğunu belirtir.
- **dilations** — Çekirdek elemanları arasındaki aralığı belirtir.

### Tablo 7.8. Convolution2d Parametreleri

template < convolution2d_descriptor Desc, typename Scope, typename... ConvArgs> convolution2d;

- **groups** — Girdinin kanal ekseni boyunca kaç gruba bölündüğünü belirtir.
- **relaxed_precision** — İşlemin float veri türü için gevşetilmiş hassasiyet kullanıp kullanamayacağını belirtir. Gevşetilmiş hassasiyet, işlemin çarpma işleminden önce mantissa'yı kısaltmasına izin verir.
- **convolution2d_mode** — `multiply` veya `multiply_accumulate` işlemlerinden hangisinin gerçekleştirileceğini belirtir.

`convolution2d` şablonunu başlatmak için descriptor ve scope geçirilir. Şu anda desteklenen tek scope `execution_simdgroups<N>`'dir; burada `N`, `simdgroups_per_threadgroup` değeridir.

### Konvolüsyonun Çalıştırılması

Konvolüsyonu çalıştırmak için `convolution2d` sınıfının `run` yöntemi çağrılır:

template <typename ActivationTensorType, typename WeightsTensorType, typename DestinationTensorType, typename... RunArgs> void run(thread ActivationTensorType &activation, thread WeightsTensorType &weights, thread DestinationTensorType &destination) const;

### Tablo 7.9. Convolution Run Parametreleri

**activation**

**NHWC** düzenine sahip aktivasyon tensörü:

- N = batch (en yavaş ilerleyen boyut)
- H = yükseklik
- W = genişlik
- C = giriş kanalları (en hızlı ilerleyen boyut)

**weights**

**HWIO** düzenine sahip ağırlık tensörü:

- H = çekirdek yüksekliği
- W = çekirdek genişliği
- I = giriş kanalları
- O = çıkış kanalları (en hızlı ilerleyen boyut)

**destination**

`tensor` veya `cooperative_tensor` olabilen hedef tensör. Eğer bir tensör ise biçimi **NHWO** düzenindedir:

- N = batch (en yavaş ilerleyen boyut)
- H = yükseklik
- W = genişlik
- O = çıkış kanalları (en hızlı ilerleyen boyut)

---