Veri Türleri

Bu bölüm Metal veri türlerini açıklar; bunlar vektörler, matrisler, atomik türler, tamponlar (buffer), dokular (texture), örnekleyiciler (sampler), diziler, kullanıcı tanımlı yapılar, tür hizalaması ve tür dönüşümünü içerir.

---

2.1 Skaler Veri Türleri

Metal aşağıda listelenen skaler türleri destekler. Aşağıdaki C türlerini desteklemez:

• double
• long long
• unsigned long long
• long double

Tablo 2.1 Metal Skaler Veri Türleri

bool
true veya false değerini alan koşullu bir veri türüdür.
true değeri tamsayı sabiti 1 olarak genişletilir ve false değeri 0 olarak genişletilir.

char / int8_t
İşaretli iki'nin tümleyeni (two’s complement) 8‑bit tamsayı.

unsigned char / uchar / uint8_t
İşaretsiz 8‑bit tamsayı.

short / int16_t
İşaretli iki'nin tümleyeni 16‑bit tamsayı.

unsigned short / ushort / uint16_t
İşaretsiz 16‑bit tamsayı.

int / int32_t
İşaretli iki'nin tümleyeni 32‑bit tamsayı.

unsigned int / uint / uint32_t
İşaretsiz 32‑bit tamsayı.

long / int64_t
İşaretli iki'nin tümleyeni 64‑bit tamsayı.

unsigned long / uint64_t
İşaretsiz 64‑bit tamsayı.

half
IEEE 754 binary16 standardına uyan 16‑bit kayan nokta.

bfloat (Metal 3.1 ve sonrası)
Makine öğrenimi için kullanılan 16‑bit brain floating point türüdür. float türünün 8‑bit mantis kullanan kısaltılmış bir sürümüdür.

float
IEEE 754 tek duyarlıklı standardına uyan 32‑bit kayan nokta.

size_t
sizeof operatörünün sonucunu temsil eden işaretsiz tamsayı türü (64‑bit).

ptrdiff_t
İki işaretçinin çıkarılması sonucunda oluşan işaretli tamsayı türü (64‑bit).

void
Herhangi bir değeri temsil etmeyen boş bir tür.

Sabit Değer (Literal) Sonekleri

Metal aşağıdaki literal soneklerini destekler:

• Tek duyarlıklı kayan nokta için f veya F (örn. 0.5f)
• Half duyarlık için h veya H (örn. 0.5h)
• Brain floating point için bf veya BF (örn. 0.5bf)
• İşaretsiz tamsayılar için u veya U
• İşaretli long tamsayılar için l veya L

---

Vektör Veri Türleri

Metal, sistem vektör matematik kütüphanesindeki vektör türlerinin bir alt kümesini destekler.

Vektör türü adları şu deseni kullanır; burada n = 2, 3 veya 4:

• booln
• charn
• shortn
• intn
• longn
• ucharn
• ushortn
• uintn
• ulongn
• halfn
• bfloatn (Metal 3.1 ve sonrası)
• floatn

Metal ayrıca aşağıdaki şablon biçimini destekler:

vec<T,n>

Burada:

• T bir skaler türdür
• n 2, 3 veya 4 değerlerinden biridir

---

Vektör Bileşenlerine Erişim

Vektör bileşenlerine dizi indeksleme veya bileşen seçiciler kullanılarak erişilebilir.

Örnek

int4 test = int4(0, 1, 2, 3);

• test.x → 0
• test.y → 1
• test.z → 2
• test.w → 3

Renk takma adları da desteklenir:

• test.r → 0
• test.g → 1
• test.b → 2
• test.a → 3

Swizzling

float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);

float4 swiz = pos.wzyx; // (4.0f, 3.0f, 2.0f, 1.0f)
float4 dup  = pos.xxyy; // (1.0f, 1.0f, 2.0f, 2.0f)

Birden fazla bileşene atama:

c.xy   = float2(3.0f, 4.0f);
c.xyz  = float3(3.0f, 4.0f, 5.0f);
pos.wx = float2(7.0f, 8.0f);
pos.xyz = float3(3.0f, 5.0f, 9.0f);

Eğer iki iş parçacığı aynı vektörün farklı bileşenlerine eşzamanlı olarak yazarsa sonuç tanımsızdır.

---

Geçersiz Vektör Bileşeni Erişimi

Aşağıdaki işlemler derleme zamanında hataya yol açar:

Var Olmayan Bileşenlere Erişim

float2 pos;
pos.z = 1.0f; // geçersiz

Çok Fazla Bileşene Erişim

2 bileşenli vektörler yalnızca .xy veya .rg erişimine izin verir.
3 bileşenli vektörler yalnızca .xyz veya .rgb erişimine izin verir.

Adlandırma Kurallarını Karıştırma

Aynı erişimde .rgba ve .xyzw kullanımını karıştırmak yasaktır.

Örnek:

float3 pos;
pos.z = 1.0f; // geçerli
pos.w = 1.0f; // geçersiz

Sol tarafta aynı bileşene iki kez erişmek belirsizdir ve hataya yol açar:

// 'x' iki kez kullanıldığı için geçersizdir.
pos.xx = float2(3.0f, 4.0f);

// float2 ile float4 uyuşmazlığı nedeniyle geçersizdir.
pos.xy = float4(1.0f, 2.0f, 3.0f, 4.0f);

pos.xg = float2(3.0f, 4.0f);

float3 coord = pos.ryz;

Swizzle içeren bir vektöre işaretçi veya referans almak hataya yol açar:

float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);

my_func(&pos.xy); // Bu, swizzle için geçersiz bir işaretçidir.

Bir vektör türü üzerinde kullanılan sizeof operatörü vektörün boyutunu döndürür. Bu genellikle bileşen sayısının her bileşenin boyutuyla çarpılmasıdır; ancak 3 bileşenli vektörlerde boyut, 4 bileşenli vektör ile aynıdır (Tablo 2.3'e bakınız). Örneğin sizeof(float4) değeri 16 döndürür ve sizeof(half4) değeri 8 döndürür.

Skaler veya vektör kümelerinden vektör oluşturmak için yapıcıları kullanabilirsiniz. Parametre imzası, vektörün nasıl oluşturulacağını ve başlatılacağını belirler. Örneğin, vektör yalnızca tek bir skaler parametre ile başlatılırsa, oluşturulan vektörün tüm bileşenleri bu skaler değere ayarlanır.

Bir vektörü birden fazla skalerden, bir veya daha fazla vektörden veya skaler ve vektör karışımından oluşturursanız, Metal vektörün bileşenlerini verilen argümanların bileşenlerinden sırayla tüketir. Metal argümanları soldan sağa doğru tüketir. Metal, bir sonraki argümana geçmeden önce bir argümanın tüm bileşenlerini sırayla tüketir.

float3(float3 x);

Bu, float2 için yapıcıların bir listesidir:

float2(float x);
float2(float x, float y);
float2(float2 x);

Aşağıdaki örnekler yapıcıların kullanımını gösterir:

float x = 1.0f, y = 2.0f, z = 3.0f, w = 4.0f;

float4 a = float4(0.0f);
float4 b = float4(x, y, z, w);

float2 c = float2(5.0f, 6.0f);

float2 a = float2(x, y);
float2 b = float2(z, w);

float4 x = float4(a.xy, b.xy);

Bir vektör yapıcısını eksik başlatmak derleme zamanında hataya neden olur.

2.2.3 Packed Vektör Türleri

Desteklenen packed vektör türleri:

• packed_charn
• packed_shortn
• packed_intn
• packed_ucharn
• packed_ushortn
• packed_uintn
• packed_halfn
• packed_bfloatn (Metal 3.1 ve sonrası)
• packed_floatn
• packed_longn (Metal 2.3 ve sonrası)

Bölüm 2.2'de açıklanan vektör veri türlerini vektörün boyutuna göre hizalamanız gerekir. Ayrıca vektör verisinin sıkı biçimde paketlenmesini isteyebilirsiniz; örneğin konum, normal, teğet vektörleri ve doku koordinatları içeren bir vertex yapısının sıkı şekilde paketlenip bir vertex fonksiyonuna buffer olarak aktarılması gibi.

Desteklenen packed vektör türü adları n değeri 2, 3 veya 4 olacak şekilde tanımlanır ve sırasıyla 2, 3 veya 4 bileşenli vektör türünü temsil eder. (packed_booln vektör türü adları ayrılmıştır.)

Metal ayrıca şunu destekler:

packed_vec<T, n>

Burada T geçerli bir skaler türdür ve n 2, 3 veya 4 olup 2, 3 veya 4 bileşenli packed vektörü temsil eder.

Tablo 2.4 packed vektör veri türlerinin boyutunu ve hizalamasını listeler.

Tablo 2.4. Packed Vektör Veri Türlerinin Boyutu ve Hizalaması

Packed vektör veri türleri genellikle veri depolama biçimi olarak kullanılır. Metal, packed vektör veri türleri için atama, aritmetik, mantıksal, ilişkisel ve kopya yapıcı operatörlerini destekler. Metal ayrıca packed vektör veri türü ile hizalanmış vektör veri türü arasında yükleme ve saklama işlemlerini de destekler.

Örnekler

device float4 *buffer;
device packed_float4 *packed_buffer;

int i;

packed_float4 f(buffer[i]);

packed_buffer[i] = buffer[i];

// packed_float4 türünden float4 türüne dönüştürme.
buffer[i] = float4(packed_buffer[i]);

Packed vektör veri türünün bileşenlerine erişmek için dizi indeksi kullanabilirsiniz. Metal 2.1 ve sonrasında packed vektör veri türünün bileşenlerine erişmek için .xyzw veya .rgba seçim söz dizimini kullanabilirsiniz. Packed vektörler için swizzle semantiği ve kısıtlamaları vektör türleriyle aynıdır.

Örnek:

packed_float4 f;

f[0] = 1.0f; // Geçerli
f.x = 1.0f;  // Geçerli, Metal 2.1 ve sonrası.

2.3 Matris Veri Türleri

Metal, sistem matematik kütüphanesi tarafından sağlanan matris veri türlerinin bir alt kümesini destekler.

Desteklenen matris türü adları n ve m değerlerinin sütun ve satır sayısını temsil ettiği şekilde tanımlanır. Hem n hem m 2, 3 veya 4 olmalıdır.

floatnxm türündeki bir matris n adet floatm vektöründen oluşur. Benzer şekilde halfnxm türündeki bir matris n adet halfm vektöründen oluşur.

Metal ayrıca şunu destekler:

matrix<T, c, r>

Burada T geçerli bir kayan nokta türüdür ve c ile r değerleri 2, 3 veya 4 olabilir.

Tablo 2.5. Matris Veri Türlerinin Boyutu ve Hizalaması

2.3.1 Matris Bileşenlerine Erişim

Bir matrisin bileşenlerine erişmek için dizi alt indis söz dizimini kullanabilirsiniz. Tek bir alt indis uygulamak matrisi sütun vektörleri dizisi olarak ele alır. İki alt indis kullanımı önce bir sütunu sonra bir satırı seçer. En üst sütun sütun 0'dır.

float4x4 m;

// Bu, 2. sütunu tamamen 2.0 yapar.
m[1] = float4(2.0f);

// Bu, 1. sütunun 1. elemanını 1.0 yapar.
m[0][0] = 1.0f;

// Bu, 3. sütunun 4. elemanını 3.0 yapar.
m[2][3] = 3.0f;

floatnxm ve halfnxm matrislerine n adet floatm veya n adet halfm girişinden oluşan bir dizi olarak erişin.

Sabit olmayan bir ifade ile matris sınırları dışında bir bileşene erişmek tanımsız davranışa yol açar. Sabit bir ifade ile matris sınırları dışında bir bileşene erişmek derleme zamanı hatası üretir.

2.3.2 Matris Yapıcıları

Skalerler, vektörler veya matrislerden oluşan bir kümeden matris oluşturmak için yapıcıları kullanın. Parametre imzası matrisin nasıl oluşturulacağını ve başlatılacağını belirler.

Bir matrisi yalnızca tek bir skaler parametre ile başlatırsanız, elde edilen matris köşegen boyunca bu skaler değeri içerir ve kalan tüm bileşenler 0.0 olarak başlatılır.

Aynı satır ve sütun sayısına sahip başka bir matristen de bir matris oluşturabilirsiniz.

Metal matris bileşenlerini sütun öncelikli (column‑major) sırayla oluşturur ve tüketir. Yapıcı, matrisin her bileşenini başlatmak için tam olarak yeterli sayıda değer sağlamalıdır. Fazla argüman verilmesi hataya neden olur ve eksik başlatma derleme zamanı hatasına yol açar.

T türünde n sütun ve m satıra sahip bir matrisi, m bileşenli T türünde n vektörden oluşturabilirsiniz.

Geçerli yapıcılar:

float2x2(float2, float2);
float3x3(float3, float3, float3);
float3x2(float2, float2, float2);

Metal 2 ve sonrasında matrisler n × m adet skalerden de oluşturulabilir:

float2x2(float, float, float, float);
float3x2(float, float, float, float, float, float);

Aşağıdaki yapıcı desteklenmez, çünkü vektör ve skalerleri karıştırır:

// Desteklenmez
float2x3(float2 a, float b, float2 c, float d);

2.4 SIMD-group Matris Veri Türleri

SIMD‑group matris türleri Metal 2.3 ve sonrasında desteklenir.

Desteklenen türler şunları içerir:

• simdgroup_half8x8
• simdgroup_bfloat8x8 (Metal 3.1 ve sonrası)

Metal şu şablonu tanımlar:

simdgroup_matrix<T, Cols, Rows>

Bu tanım <metal_simdgroup_matrix> içinde yer alır.

SIMD‑group matrisler üzerindeki işlemler bir SIMD grubundaki iş parçacıkları tarafından işbirliği içinde yürütülür. Bu nedenle tüm işlemler SIMD grubu içinde uniform control flow altında yürütülmelidir; aksi halde davranış tanımsızdır.

Matris elemanlarının SIMD grubundaki iş parçacıklarına nasıl eşlendiği belirtilmemiştir. SIMD‑group matrislerde desteklenen fonksiyonlar için bölüm 6.7'ye bakınız.

Bir türün veya nesnenin hizalama gereksinimini belirtmek için alignas hizalama belirtecini kullanabilirsiniz. Metal derleyicisi veri öğelerini türlerinin gereksinimlerine göre otomatik olarak hizalar.

Bir grafik veya kernel fonksiyonuna, bir veri türüne işaretçi olarak bildirilen argümanlar için Metal derleyicisi başvurulan nesnenin doğru şekilde hizalanmış olduğunu varsayar.

Atomik türlere ait nesneler veri yarışlarından arındırılmıştır. Bir iş parçacığı atomik bir nesneye yazarken başka bir iş parçacığı ondan okuma yaparsa davranış iyi tanımlıdır.

2.6 Atomik Veri Türleri

Metal şunu destekler:

atomic<T>

Burada T şu türlerden biri olabilir:

• int
• uint
• bool (Metal 2.4 ve sonrası)
• ulong (Metal 2.4 ve sonrası)
• float (Metal 3 ve sonrası)

Metal atomik fonksiyonları yalnızca Metal atomik veri türleri üzerinde çalışabilir. Bu fonksiyonlar C++17 atomik ve senkronizasyon fonksiyonlarının bir alt kümesidir.

Metal şu takma adları sağlar:

• atomic_int — atomic<int> için takma ad (Metal 1 ve sonrası)
• atomic_uint — atomic<uint> için takma ad (Metal 1 ve sonrası)
• atomic_bool — atomic<bool> için takma ad (Metal 2.4 ve sonrası)
• atomic_ulong — atomic<ulong> için takma ad (Metal 2.4 ve sonrası)
• atomic_float — atomic<float> için takma ad (Metal 3 ve sonrası)

2.7 Piksel Veri Türleri

Desteklenen sürümler:

• iOS: Metal 2 ve sonrası
• macOS: Metal 2.3 ve sonrası

Metal piksel veri türü, piksel formatını ve buna karşılık gelen ALU türünü tanımlayan şablon tabanlı bir türdür. <metal_pixel> başlık dosyası Metal piksel veri türlerini tanımlar.

Piksel veri türleri ile bunlara karşılık gelen ALU türleri arasında yalnızca atama ve eşitlik/eşit değil karşılaştırmalarına izin verilir.

Tablo 2.6. Metal Piksel Veri Türleri

Örnek:

kernel void my_kernel(device rgba8unorm<half4> *p [[buffer(0)]],
                      uint gid [[thread_position_in_grid]])
{
    rgba8unorm<half4> x = p[index];
    half4 val = p[gid];

    p[gid] = val;
    p[index] = x;
}

Başka bir örnek:

struct Foo {
    rgba8unorm<half4> a;
};

kernel void my_kernel(device Foo *p [[buffer(0)]],
                      uint gid [[thread_position_in_grid]])
{
    half4 a = p[gid].a;
    p[gid].a = a;
}

2.8 Tamponlar

MSL, bir arabelleği device, constant veya threadgroup adres alanında bulunan yerleşik ya da kullanıcı tanımlı bir veri türüne işaret eden bir gösterici olarak uygular.

Sıradan Metal arabellekleri şunları içerebilir:

• float ve int gibi temel türler
• Vektör ve matris türleri
• Arabellek türlerinden oluşan diziler
• Arabellek türlerinden oluşan birleşimler

Örnek:

vertex ColorInOut
phong_vertex(const device packed_float3* vertices [[buffer(0)]],
             const device packed_float3* normals [[buffer(1)]],
             constant AAPL::uniforms_t& uniforms [[buffer(2)]],
             unsigned int vid [[vertex_id]])
{
    ...
}

2.9 Dokular

Doku veri türü, bir dokunun tek bir mipmap seviyesinin tamamına ya da bir kısmına karşılık gelen bir, iki veya üç boyutlu doku verisine yönelik bir tanıtıcıdır.

Metal 3.2 ve sonrasında, dokular isteğe bağlı bir bellek tutarlılığı parametresini destekler.

Örnek şablon tanımları:

// Metal 3.2 ve sonrası
texture1d<T, access a = access::sample,
          memory_coherence c = memory_coherence_threadgroup>

texture1d_array<T, access a = access::sample,
                memory_coherence c = memory_coherence_threadgroup>

texture2d<T, access a = access::sample,
          memory_coherence c = memory_coherence_threadgroup>

texture2d_array<T, access a = access::sample,
                memory_coherence c = memory_coherence_threadgroup>

texture3d<T, access a = access::sample,
          memory_coherence c = memory_coherence_threadgroup>

texturecube<T, access a = access::sample,
            memory_coherence c = memory_coherence_threadgroup>

texturecube_array<T, access a = access::sample,
                  memory_coherence c = memory_coherence_threadgroup>

texture2d_ms<T, access a = access::read,
             memory_coherence c = memory_coherence_threadgroup>

Tüm işletim sistemi sürümleri Metal 3.2 ile başlayarak Apple silicon üzerinde memory_coherence desteği sunar.

enum class access { sample, read, write, read_write };

enum memory_coherence {
    memory_coherence_threadgroup,
    memory_coherence_device
};

Metal 3.2 öncesinde:

texture1d<T, access a = access::sample>

Aşağıdaki şablonlar belirli doku veri türlerini tanımlar:

texture2d_ms_array<T, access a = access::read,
memory_coherence c = memory_coherence_threadgroup>

Bir derinlik biçimi ile sample_compare kullanmak için aşağıdaki doku türlerinden birini bildirmeniz gerekir:

depth2d<T, access a = access::sample,
memory_coherence c = memory_coherence_threadgroup>

depth2d_array<T, access a = access::sample,
memory_coherence c = memory_coherence_threadgroup>

depthcube<T, access a = access::sample,
memory_coherence c = memory_coherence_threadgroup>

depthcube_array<T, access a = access::sample,
memory_coherence c = memory_coherence_threadgroup>

• macOS: texture2d_ms_array ve depth2d_ms_array türlerini Metal 2 ve sonrasında destekler. Diğer tüm türler Metal 1 ve sonrasında desteklenir.
• iOS: texture2d_ms_array ve depth2d_ms_array dışında tüm türleri Metal 1 ve sonrasında destekler.

T, bir dokudan okuma yapıldığında döndürülen bileşenlerden birinin renk türünü veya dokuya yazılırken belirtilen bileşenlerden birinin renk türünü belirtir.

Doku türleri için (derinlik doku türleri hariç) T şu türlerden biri olabilir:

• half
• float
• short
• ushort
• int
• uint

Derinlik doku türleri için T mutlaka float olmalıdır.

Eğer T int veya short ise, doku ile ilişkili veri işaretli tamsayı biçimi kullanmalıdır. Eğer T uint veya ushort ise, doku ile ilişkili veri işaretsiz tamsayı biçimi kullanmalıdır.

Eğer T half ise, doku ile ilişkili veri ya normalize edilmiş (işaretli veya işaretsiz tamsayı) ya da yarı hassasiyetli biçimde olmalıdır. Eğer T float ise, doku ile ilişkili veri ya normalize edilmiş (işaretli veya işaretsiz tamsayı), yarı hassasiyetli ya da tek hassasiyetli biçimde olmalıdır.

Erişim Öznitelikleri

• sample — Bir grafik veya çekirdek fonksiyonu doku nesnesinden örnekleme yapabilir. sample, bir örnekleyici ile veya örnekleyici olmadan dokudan okuma yeteneğini ifade eder.

• read — Bir örnekleyici olmadan, bir grafik veya çekirdek fonksiyonu yalnızca doku nesnesini okuyabilir.

• write — Bir grafik veya çekirdek fonksiyonu doku nesnesine yazabilir.

• read_write — Bir grafik veya çekirdek fonksiyonu doku nesnesini okuyabilir ve yazabilir.

• Tüm OS: Metal 1.2 ve sonrası read_write erişimini destekler. Metal 1 ve sonrası diğer erişim niteleyicilerini destekler.

Çoklu örneklemeli dokular yalnızca read özniteliğini destekler. Derinlik dokuları yalnızca sample ve read özniteliklerini destekler. Seyrek dokular write veya read_write özniteliklerini desteklemez.

Bu erişim öznitelikleri bir dokuya erişim desteğini tanımlar.

Erişim Niteleyicileri Kullanımına Örnek

void foo (texture2d<float> imgA [[texture(0)]],
          texture2d<float, access::read> imgB [[texture(1)]],
          texture2d<float, access::write> imgC [[texture(2)]])
{
    texture2d<float> x = imgA;                // OK
    texture2d<float, access::read> y = imgB;  // OK
    texture2d<float, access::write> z;        // Bu geçersizdir.
}

constant texture2d<float, access::sample,
memory_coherence_device> gtex [[texture(2)]];

constant texture2d<int, access::write,
memory_coherence::memory_coherence_device>
gtex2 [[texture(8)]];

(texture özniteliğinin açıklaması için bölüm 5.2.1'e bakın.)

Bir doku türünü, bir fonksiyon içinde bildirilen herhangi bir değişken için değişken türü olarak kullanabilirsiniz. Bir fonksiyon içinde bildirilen doku türündeki değişkenlerin erişim özniteliği access::read veya access::sample olmalıdır. Bir fonksiyon içinde access::read veya access::sample niteleyicileri kullanılmadan doku türünde değişken bildirimi yapılması derleme hatasına yol açar.

Örnekler

void foo (texture2d<float> imgA [[texture(0)]],
          texture2d<float, access::read> imgB [[texture(1)]])

Metal 3.2 ve sonrasında, doku işlemlerinin cihaz genelinde tutarlı olup olmadığını belirtebilirsiniz; bu, uygun şekilde senkronize edildiğinde doku işlemlerinin farklı threadgroup'lardaki diğer iş parçacıkları tarafından görülebileceği anlamına gelir.

---

Texture Buffers

Bir texture buffer, büyük bir 1D piksel veri dizisine erişebilen ve bu veri üzerinde piksel biçimleri arasında dinamik tür dönüşümünü optimize edilmiş performansla gerçekleştirebilen bir doku türüdür.

Texture buffer'lar tür dönüşümünü diğer tekniklere göre daha verimli biçimde gerçekleştirir; daha fazla sayıda elemana erişim sağlar ve sınır dışı okuma erişimlerini yönetir.

Benzer tür dönüşümü texture buffer kullanmadan şu yöntemlerle de sağlanabilir:

• Piksel verisini (herhangi bir başka dizi gibi) bir doku nesnesinden okuyup piksel dönüşümünü istenen biçime dönüştürmek.
• Bir doku nesnesini bir arabellek nesnesinin verisi etrafında sarmalayarak, paylaşılan arabellek verisine doku üzerinden erişmek. Bu sarma tekniği piksel dönüşümü sağlar ancak ek bir işlem adımı gerektirir ve doku boyutunu sınırlar.

Aşağıdaki şablon, diğer dokular gibi kullanabileceğiniz opak texture_buffer türünü tanımlar:

texture_buffer<T, access a = access::read>

T, bir texture buffer'dan okuma yapıldığında döndürülen bileşen türünü veya yazma sırasında belirtilen türü ifade eder. Bir texture buffer için T şu türlerden biri olabilir:

• half
• float
• short
• ushort
• int
• uint

Alfa kanalı olmayan bir biçim için (R, RG veya RGB gibi), sınır dışı bir okuma (0, 0, 0, 1) döndürür. Alfa içeren bir biçim için (RGBA gibi), sınır dışı bir okuma (0, 0, 0, 0) döndürür. Bazı aygıtlarda sınır dışı okuma performans cezasına yol açabilir.

Bir texture buffer, maksimum genişliği 16384 olan genel bir 1D dokudan daha fazla doku verisini destekleyebilir. Ancak bir texture buffer'dan örnekleme yapılamaz.

Bir texture buffer ayrıca veriyi dönüştürür ve kaynak biçim ne olursa olsun talep edilen doku biçiminde sunar. Örneğin verisi RGBA8Unorm olarak depolanmış bir texture buffer kullanabilir ve daha sonra bir shader içinde bunu float4 gibi dönüştürülmüş bir tür olarak okuyabilirsiniz.

Bunun sonucunda tek bir pipeline state object, yeniden derleme yapılmadan farklı piksel biçimlerinde depolanan verilere erişebilir.

Bir texture buffer, diğer doku türleri gibi, bir shader fonksiyonu içinde yerel değişken türü olarak bildirilebilir.

• Texture buffer dizileri için 2.12.1 bölümüne bakın.
• Texture buffer hakkında daha fazla bilgi için 6.12.16 bölümüne bakın.

access şu değerlerden biri olabilir:

• read
• write
• read_write

Metal, sınır dışı bir yazma işlemini yok sayar.

---

2.10 Samplers

sampler türü, bir dokunun nasıl örnekleneceğini tanımlar. Metal API, bir sampler nesnesi oluşturmanıza ve bunu bir grafik veya çekirdek fonksiyonuna argüman olarak iletmenize olanak tanır.

Bir sampler nesnesini doğrudan program kaynak kodunda da tanımlayabilirsiniz. Bu durumda yalnızca sampler durumunun bir alt kümesi belirtilebilir:

• adresleme modu
• filtre modu
• normalize edilmiş koordinatlar
• karşılaştırma fonksiyonu

Tablo 2.7, desteklenen sampler durum numaralandırmalarını ve bunlara karşılık gelen değerleri (ve varsayılanları) listeler.

Sampler Durumu Numaralandırma Değerleri

coord

Geçerli değerler:

• normalized (varsayılan)
• pixel

Doku koordinatlarının normalize edilmiş değerler olup olmadığını belirtir.

address

Geçerli değerler:

• repeat
• mirrored_repeat
• clamp_to_edge (varsayılan)
• clamp_to_zero
• clamp_to_border

Doku koordinatları için adresleme modunu ayarlar.

s_address, t_address, r_address

address ile aynı geçerli değerlere sahiptir. Tek tek doku koordinatları için adresleme modunu ayarlar.

border_color

Desteklenen platformlar:

• macOS: Metal 1.2
• iOS: Metal 2.3

Geçerli değerler:

• transparent_black (varsayılan)
• opaque_black
• opaque_white

clamp_to_border adresleme modu ile kullanılan sınır rengini belirtir.

filter

Geçerli değerler:

• nearest (varsayılan)
• linear

Büyütme ve küçültme filtreleme modlarını ayarlar.

mag_filter

Geçerli değerler:

• nearest (varsayılan)
• linear

min_filter

Geçerli değerler:

• nearest (varsayılan)
• linear

mip_filter

Geçerli değerler:

• none (varsayılan)
• nearest
• linear

none ise, doku tek bir mip seviyesi varmış gibi örneklenir. Tüm örnekler seviye 0'dan okunur.

compare_func

Geçerli değerler:

• never (varsayılan)
• less
• less_equal
• greater
• greater_equal
• equal
• not_equal
• always

sample_compare ve gather_compare doku fonksiyonları tarafından kullanılır.

reduction

Tüm işletim sistemlerinde Metal 2.3 ile desteklenir.

Geçerli değerler:

• weighted_average
• minimum
• maximum

Filtrelenmiş piksel değerinin nasıl hesaplandığını kontrol eder.

bias

Tüm işletim sistemlerinde Metal 4.0 ile desteklenir.

Bir float değeri kabul eder ve örnekleme yapılmadan önce seviye‑ayrıntı (LOD) yanlılığı uygular.

Adresleme Notları

clamp_to_border kullanıldığında, bir dokunun dışından örnekleme yapılması yalnızca sınır rengini kullanır ve dokunun kenarındaki renkleri kullanmaz.

clamp_to_zero, transparent_black (0.0, 0.0, 0.0) sınır rengi ve dokudan alınan alfa bileşeni değeri ile clamp_to_border ile eşdeğerdir.

Eğer coord değeri pixel olarak ayarlanmışsa:

• min_filter ve mag_filter aynı olmalıdır
• mip_filter none olmalıdır
• adresleme modları clamp_to_zero, clamp_to_border veya clamp_to_edge olmalıdır

Sampler Bildirim Örnekleri

Metal Shading Language (MSL) kaynak kodunda başlatılan sampler'ları bildirmek için constexpr kullanmanız gerekir.

constexpr sampler a(coord::normalized);
constexpr sampler b(address::repeat);

constexpr sampler s(address::clamp_to_zero,
                    filter::linear,
                    compare_func::less);

constexpr sampler s(address::clamp_to_zero,
                    filter::linear,
                    compare_func::less,
                    max_anisotropy(10),
                    lod_clamp(0.0f, MAXFLOAT));

kernel void
my_kernel(device float4 *p [[buffer(0)]],
          texture2d<float> img [[texture(0)]],
          sampler smp [[sampler(3)]])
{
    ...
}

---

Imageblocks

Bir imageblock, threadgroup belleğinde ayrılan (genişlik, yükseklik ve örnek sayısı ile temsil edilen) iki boyutlu bir veri yapısıdır. 2D görüntü verisini işlemek için verimli bir mekanizma sağlar.

Yapının her bir elemanı şunlardan biri olabilir:

• skaler veya vektör tamsayı türleri
• skaler veya vektör kayan nokta türleri
• piksel veri türleri (bkz. Tablo 2.6)
• bu türlerin dizileri
• bu türlerden oluşturulmuş yapılar

Imageblock'un veri düzeni opaktır. Elemanlara (x, y) koordinatları ve isteğe bağlı olarak bir örnek indeksi kullanılarak erişilir.

• 5.6 bölümü [[imageblock_data(type)]] gibi imageblock özniteliklerini açıklar.
• 6.13 bölümü imageblock'lar için yerleşik fonksiyonları listeler.

Imageblock'lar yalnızca fragment ve kernel fonksiyonlarıyla kullanılır.

• Fragment fonksiyonları yalnızca fragmanın imageblock verisine erişebilir.
• Kernel fonksiyonları threadgroup içindeki tüm iş parçacıklarının imageblock'a erişmesine izin verir.

Bir imageblock slice, imageblock içinde belirli bir elemanın tüm piksel konumları veya iş parçacıkları için değerlerini temsil eden bir bölgeyi ifade eder.

Imageblock slice'ın depolama türü hedef doku biçimiyle uyumlu olmalıdır.

Imageblock Slice Uyumluluğu

Tamsayı Türleri

• int, short → R32Sint, R16Sint, R8Sint
• int2, short2 → RG32Sint, RG16Sint, RG8Sint
• int4, short4 → RGBA32Sint, RGBA16Sint, RGBA8Sint

İşaretsiz Tamsayı Türleri

• uint, ushort → R32Uint, R16Uint, R8Uint
• uint2, ushort2 → RG32Uint, RG16Uint, RG8Uint
• uint4, ushort4 → RGBA32Uint, RGBA16Uint, RGBA8Uint

Kayan Nokta Türleri

• float, half → R32Float, R16Float, R8Unorm, R8Snorm, R16Unorm, R16Snorm
• float2, half2 → RG32Float, RG16Float, RG8Unorm, RG8Snorm, RG16Unorm, RG16Snorm
• float4, half4 → RGBA32Float, RGBA16Float, RGBA8Unorm, RGBA8Snorm, RGBA16Unorm, RGBA16Snorm, RGB10A2Unorm, RG11B10Float, RGB9E5Float

Özel Piksel Türleri

• r8unorm<T> → A8Unorm, R8Unorm
• r8snorm<T> → R8Snorm
• r16unorm<T> → R16Unorm
• r16snorm<T> → R16Snorm
• rg8unorm<T> → RG8Unorm
• rg8snorm<T> → RG8Snorm
• rgba8snorm<T> → RGBA8Snorm, BGRA8Unorm
• rgba16unorm<T> → RGBA16Unorm
• rgba16snorm<T> → RGBA16Snorm
• rgb10a2<T> → RGB10A2Unorm
• rg11b10f<T> → RG11B10Float
• rgb9e5<T> → RGB9E5Float

---

2.12 Bileşik Türler

Bir yapı üyesi gösterici türünde değilse, adres alanı özniteliği ile belirtilmemelidir. Bir bileşik türdeki tüm üyeler aynı adres alanına ait olmalıdır.

Platform Desteği

• iOS:

  • Metal 1.2 ve sonrası doku dizilerini destekler.
  • Metal 2 ve sonrası sampler dizilerini destekler.
  • Metal 2.1 ve sonrası texture buffer dizilerini destekler.

• macOS:

  • Metal 2 ve sonrası doku ve sampler dizilerini destekler.
  • Metal 2.1 ve sonrası texture buffer dizilerini destekler.

Doku Dizileri

Bir doku dizisini şu şekilde bildirin:

typename

typename, access::read veya access::sample ile bildirilmiş bir doku türü olmalıdır.

• macOS Metal 2 ve sonrası yazılabilir dokuların (access::write) dizilerini destekler.
• iOS Metal 2.2 ve sonrası, Apple GPU Family 5 ve sonrası ile birlikte bunu destekler.

Texture Buffer Dizileri

Bir texture buffer dizisi access::read niteleyicisi kullanılarak oluşturulur.

Sampler Dizileri

Örnekleyici dizileri benzer şekilde bildirilebilir ve grafik, kernel veya kullanıcı fonksiyonlarına argüman olarak geçirilebilir.

Şunları yapabilirsiniz:

• Doku veya örnekleyici dizilerini fonksiyon argümanları olarak geçirmek.
• Doku veya örnekleyici dizilerini yerel değişkenler olarak bildirmek.
• Program kapsamı içinde örnekleyici dizileri bildirmek.

Bir argument buffer içinde kullanılmadığı sürece (bkz. bölüm 2.13), bir yapı içinde array<T, N>—yani doku, doku arabelleği veya örnekleyici dizisi—bildiremezsiniz.

array_ref<T>

MSL ayrıca array_ref<T> türünü de destekler. Bu, T türünde size() adet öğeden oluşan değiştirilemez bir diziyi temsil eder.

Gereksinimler:

• T bir sampler türü veya desteklenen bir doku türü olmalıdır (doku arabellekleri dahil).
• Dizinin depolama alanı array_ref<T> nesnesine ait değildir.

metal::array gibi bitişik yineleyicilere sahip türlerden örtük dönüşümler sağlanır.

array_ref<T> için yaygın bir kullanım, fonksiyonlara doku dizilerini geçirmek ve böylece birden fazla dizi türünü kabul edebilmelerini sağlamaktır.

array_ref<T> türü grafik ve kernel fonksiyonlarına argüman olarak geçirilemez. Ancak array_ref<T> türü kullanıcı fonksiyonlarına argüman olarak geçirilebilir. array_ref<T> türü fonksiyonların içinde yerel değişken olarak bildirilemez.

Bölüm 2.12.1.1 ile 2.12.1.3 arasında listelenen üye fonksiyonlar; doku dizileri, örnekleyici dizileri ve array_ref<T> türleri için kullanılabilir.

---

2.12.1.1 Operatörü ile Dizi Elemanına Erişim

Doku, doku arabelleği veya örnekleyici dizilerinin elemanlarına [] operatörü kullanılarak erişilebilir:

reference operator[] (size_t pos);

Doku, doku arabelleği veya örnekleyici dizilerinin elemanlarına ya da şablon türü array_ref<T>'ye aşağıdaki [] operatörü varyantı kullanılarak erişilebilir:

constexpr const_reference operator[] (size_t pos) const;

---

2.12.1.2 Dizi Kapasitesi

size(), doku, doku arabelleği veya örnekleyici dizisindeki eleman sayısını döndürür:

constexpr size_t size();
constexpr size_t size() const;

Örnek:

kernel void
my_kernel(const array<texture2d<float>, 10> src [[texture(0)]],
          texture2d<float, access::write> dst [[texture(10)]],
          ...)
{
    for (int i = 0; i < src.size(); i++) {
        if (is_null_texture(src[i]))
            break;
        process_image(src[i], dst);
    }
}

---

2.12.1.3 Şablon Dizileri İçin Yapıcılar

constexpr array_ref();
constexpr array_ref(const array_ref &);
array_ref & operator=(const array_ref &);
constexpr array_ref(const T * array, size_t length);

template<size_t N>
constexpr array_ref(const T (&a)[N]);

template<typename T>
constexpr array_ref<T> make_array_ref(const T * array, size_t length);

template<typename T, size_t N>
constexpr array_ref<T> make_array_ref(const T (&a)[N]);

Dizi oluşturma örnekleri:

float4 foo(array_ref<texture2d<float>> src)
{
    float4 clr(0.0f);
    for (int i = 0; i < src.size(); i++) {
        clr += process_texture(src[i]);
    }
    return clr;
}

kernel void
my_kernel_A(const array<texture2d<float>, 10> src [[texture(0)]],
            texture2d<float, access::write> dst [[texture(10)]],
            ...)
{
    float4 clr = foo(src);
    ...
}

kernel void
my_kernel_B(const array<texture2d<float>, 20> src [[texture(0)]],
            texture2d<float, access::write> dst [[texture(10)]],
            ...)
{
    float4 clr = foo(src);
    ...
}

Program kapsamı içinde bildirilen bir örnekleyici dizisi örneği:

constexpr array<sampler, 2> samplers = {
    sampler(address::clamp_to_zero),
    sampler(coord::pixel)
};

---

2.12.2 Buffer, Texture ve Sampler Yapıları

Bir grafik, kernel, visible veya kullanıcı fonksiyonuna verilen argüman; yalnızca buffer, texture veya sampler üyeleri içeren bir yapı veya iç içe geçmiş bir yapı olabilir. Böyle bir yapıyı değer olarak geçirmeniz gerekir. Grafik veya kernel fonksiyonuna argüman türü olarak geçirilen bu tür bir yapının her üyesi, konumunu belirtmek için bir özniteliğe sahip olabilir (bölüm 5.2.1'de açıklandığı gibi).

Argüman olarak geçirilen bir yapı örneği:

struct Foo {
    texture2d<float> a [[texture(0)]];
    depth2d<float> b [[texture(1)]];
};

[[kernel]] void
my_kernel(Foo f)
{
    ...
}

Aşağıdaki örnekte gösterildiği gibi yapıları iç içe de kullanabilirsiniz:

struct Foo {
    texture2d<float> a [[texture(0)]];
    depth2d<float> b [[texture(1)]];
};

struct Bar {
    Foo f;
    sampler s [[sampler(0)]];
};

Geçersiz örnek:

struct MyResources {
    texture2d<float> a [[texture(0)]];
    depth2d<float> b [[texture(1)]];
    int c;
};

[[kernel]] void
my_kernel(MyResources r) // Bu geçersiz bir kullanımdır.
{
    ...
}

Sabit bir buffer kullanan örnek:

struct Foo {
    texture2d<float, access::write> a;
    depth2d<float> b;
    sampler c;
    texture2d<float> d;
    device float4* e;
    texture2d<float> f;
    texture_buffer<float> g;
    int h;
};

kernel void
my_kernel(constant Foo &f [[buffer(0)]])
{
    ...
}

---

2.13 Argument Buffers

Tüm OS: Metal 2 ve sonrası argument buffer desteği sağlar.

Argument buffer'lar temel buffer türlerini; işaretçiler (buffer'lar), texture'lar, texture buffer'lar ve sampler'ları içerecek şekilde genişletir. Ancak argument buffer'lar union içeremez.

Texture ve sampler dizileri mevcut array<T, N> şablon türü kullanılarak bildirilebilir. Diğer tüm geçerli buffer türlerinin dizileri ise C tarzı dizi sözdizimi kullanılarak bildirilebilir.

Argument buffer üyelerine genel bir [[id(n)]] özniteliği atanabilir; burada n, Metal API üzerinden buffer elemanını tanımlamak için kullanılabilen 32 bit işaretsiz bir tam sayıdır.

Argument buffer'lar; buffer, texture, sampler içermeleri veya [[id]] özniteliğine sahip herhangi bir öğe barındırmaları durumunda normal buffer'lardan ayırt edilebilir.

Aynı indeks bir argument buffer'ın birden fazla üyesine atanamaz. Elle atanan indekslerin ardışık olması gerekmez ancak monotonik olarak artmaları gerekir.

Örnek

struct Foo {
    Foo foo1; // foo1'e idx 0 atanır, t1 ve t2'ye idx 1 ve 3 atanır
    Foo foo2; // foo2'ye idx 4 atanır, t1 ve t2'ye idx 5 ve 7 atanır
};

Otomatik ID Atama Kuralları

1. Yapı üyeleri: Metal, yapı üyelerine sırayla ID atar ve önceki yapı üyesinin maksimum ID değerine 1 ekler.

2. Dizi elemanları: Metal, dizi elemanlarına sırayla ID atar ve önceki dizi elemanının maksimum ID değerine 1 ekler.

3. İç içe yapılar veya diziler: Eğer bir yapı üyesi veya dizi elemanı E kendisi de bir yapı veya dizi ise, Metal E'ye atanan ID'den başlayarak üyelerine özyineli olarak indeks atar.

4. En üst düzey argument buffer argümanları: Metal, önceki kurallara göre 0'dan başlayarak ID atar.

Örnekler

struct MaterialTexture {
    texture2d<float> tex;        // Atanan indeks 0
    float4 uvScaleOffset;        // Atanan indeks 1
};

struct Material {
    float4 diffuse;                          // Atanan indeks 0
    array<texture2d<float>, 3> texs1;        // Atanan indeksler 1–3
    MaterialTexture materials[3];            // Atanan indeksler 4–9
    int constants [[id(20)]] [4];            // Atanan indeksler 20–23
};

struct Material {
    MaterialTexture diffuse;                 // Atanan indeksler 0,1
    MaterialTexture normal [[id(4)]];        // Atanan indeksler 4,5
    MaterialTexture specular;                // Atanan indeksler 6,7
};

---

2.13.1 Argument Buffer'lar İçin Tier 2 Donanım Desteği

Tier 2 donanım ile argument buffer'lar, Tier 1 donanımda bulunmayan ek yeteneklere sahiptir.

Argument buffer'lara işaretçi indeksleme yoluyla erişebilirsiniz. Aşağıda gösterilen sözdizimi, ardışık ve bağımsız olarak kodlanmış argument buffer'lardan oluşan bir diziyi ifade eder:

constant Resources &resources = resArray[3];

struct TStruct {
    texture2d<float> tex;
};

kernel void
kern(constant TStruct *textures [[buffer(0)]]);

Kaynakların Kopyalanması

GPU tarafından yönlendirilen pipeline'ları ve dolaylı draw çağrılarını ve dispatch işlemlerini desteklemek için, bir fonksiyon içinde yapılar ve diziler arasında kaynakları kopyalayabilirsiniz:

dst.a = src.d;
...

Sampler'lar thread adres alanından device adres alanına kopyalanamaz. Bu nedenle sampler'lar bir argument buffer'a yalnızca başka bir argument buffer'dan doğrudan kopyalanabilir.

struct Resources {
    sampler sam;
};

constexpr sampler sam2;

dst->sam = src->sam; // Geçerli: device -> device

dst->sam = sam1;     // Geçersiz: thread -> device

dst->sam = sam2;     // Geçersiz: thread -> device

Argument buffer'lar başka argument buffer'lara işaretçiler içerebilir.

struct Textures {
    texture2d<float> diffuse;
    texture2d<float> specular;
};

struct Material {
    Textures textures;
};

fragment float4
fragFunc(device Material &material);

---

2.14 Uniform Türü

Tüm OS: Metal 2 ve sonrası uniform türlerini destekler.

Aşağıdaki fonksiyon örneğinde i değişkeni, texInput tarafından verilen doku dizisini indekslemek için kullanılır. i değişkeni nonuniform'dur; yani draw veya dispatch çağrısı sırasında grafik veya kernel fonksiyonunu çalıştıran thread'ler için farklı değerler alabilir.

kernel void
my_kernel(array<texture2d<float>, 10> texInput,
          array<texture2d<float>, 10> texOutput,
          sampler s,
          ...,
          uint2 gid [[thread_position_in_grid]])
{
    int i = ...;
    float4 color = texInput[i].sample(s, float2(gid));
    ...
    texOutput[i].write(color, float2(gid));
}

Eğer i değişkeni fonksiyonu çalıştıran tüm thread'ler için aynı değere sahipse, donanım optimizasyonlar uygulayabilir.

Bir değişkenin tüm thread'lerde uniform olduğunu belirtmek için MSL uniform adlı bir şablon sınıfı ekler (metal_uniform başlık dosyasında bulunur). Bu şablon yalnızca aritmetik veya vektör türleriyle örneklenebilir.

uint2 gid [[thread_position_in_grid]]
{
    uniform<int> i = ...;
    float4 color = texInput[i].sample(s, float2(gid));
    ...
    texOutput[i].write(color, float2(gid));
}

Bir değişken uniform tür olarak bildirilmiş ancak thread'ler arasında aynı değere sahip değilse, davranış tanımsızdır.

Uniform değişkenler örtük olarak nonuniform türlere dönüşebilir. Nonuniform bir değişkeni uniform bir değişkene atamak derleme zamanı hatasıyla sonuçlanır.

int j = ...;

if (i < j) { // Nonuniform sonuç
    ...
    i++; // Derleme zamanı hatasına neden olur
}

bool p = ...; // Nonuniform koşul
uniform<int> a = ..., b = ...;
uniform<int> c = p ? a : b; // Derleme zamanı hatası

Uniform bir türe işaretçi bildirmek yasaldır, ancak uniform bir işaretçi bildirmek geçerli değildir.

---

2.14.3 Uniform Kontrol Akışı

Bir kontrol akışı koşullu testi uniform bir niceliğe dayanıyorsa, tüm program örnekleri bu koşullu testte aynı yolu izler. Uniform niceliklere dayanan kod genellikle nonuniform niceliklere dayanan koda göre daha verimlidir.

---

2.15 Visible Function Table

Tüm OS: Metal 2.3 ve sonrası visible function table desteği sağlar.

<metal_visible_function_table> başlık dosyasında tanımlanan visible_function_table türü, device belleğinde depolanan görünür fonksiyonlara ait fonksiyon işaretçilerinden oluşan bir tabloyu temsil eder.

• Metal 2.3 içinde compute (kernel) fonksiyonlarında kullanılabilir.
• Metal 2.4 ve sonrası sürümlerde fragment, vertex ve tile fonksiyonlarında da kullanılabilir.

Bu opak bir türdür ve GPU tablonun içeriğini değiştiremez.

visible_function_table türünü bir argument buffer içinde kullanabilir veya bir buffer bağlama noktası kullanarak doğrudan nitelikli bir fonksiyona geçirebilirsiniz.

Bildirim

visible_function_table<typename T>

Burada T, tabloda saklanan fonksiyonun imzasıdır.

Şununla uyumlu örnek: [[visible]] int func(float f)

visible_function_table<int(float)> functions;

Üye Fonksiyonlar

using fnptr = T (*)(...) [[visible]];
fnptr operator[](uint index) const;
uint size() const;
bool empty() const;

Bir tablonun null olup olmadığını belirlemek için:

bool is_null_visible_function_table(visible_function_table<T>);

Kullanım Örneği

using TFuncSig = void(float, int);

kernel void F(uint tid [[thread_position_in_grid]],
              device float* buf [[buffer(0)]],
              visible_function_table<TFuncSig> table [[buffer(1)]])
{
    uint tsize = table.size();
    table[tid % tsize](buf[tid], tid);
}

[[function_groups]] Özniteliği

İsteğe bağlı [[function_groups]] özniteliği, dolaylı çağrı yoluyla çağrılabilecek olası fonksiyon gruplarını belirtir. Bu, derleyici için bir optimizasyon ipucu görevi görür.

Bu öznitelik şu yerlere uygulanabilir:

• Başlatıcı ifadeleri olan değişken bildirimleri
• İfade deyimleri
• Return deyimleri

Örnek:

float h(visible_function_table<float(float)> table,
        float (*fnptr[3])(float))
{
    // Dolaylı çağrı "group1" ile sınırlandırılmıştır
    [[function_groups("group1")]] float x = table[0](1.0f);
}

---

2.16 Function Groups Özniteliği

Tüm OS: Metal 2.3 ve sonrası [[function_groups]] desteği sağlar.

// `fnptr[0]` için dolaylı çağrı herhangi bir fonksiyonu çağırabilir.
x += fnptr[0](2.0f);

// `fnptr[1]` için dolaylı çağrı
// "group2" + "group3" ile sınırlandırılmıştır.
[[function_groups("group2", "group3")]]
return x + fnptr[1](3.0f);

2.17 Ray-Tracing Türleri

Tüm OS: Metal 2.3 ve sonrası ray-tracing türlerini destekler.

<metal_raytracing> başlık dosyası bu türleri metal::raytracing ad alanında tanımlar. Metal 2.3 ve sonrası sürümlerde bu türler yalnızca compute fonksiyonlarında (kernel fonksiyonlarında) desteklenir; aşağıda belirtilen durumlar hariç. Metal 2.4 ve sonrası sürümlerde vertex, fragment ve tile fonksiyonlarında da desteklenir. Metal 3.1 ve sonrası sürümlerde ray tracing eğrileri ve çok seviyeli instancing desteği sağlar.

2.17.1 Ray-Tracing Intersection Etiketleri

Tüm OS: Metal 2.3 ve sonrası ray-tracing intersection etiketlerini destekler.

Bu etiketler ray tracing sırasında aşağıdakileri tanımlarken kullanılır:

• intersection fonksiyonları ([[intersection]], bölüm 5.1.6)
• intersection fonksiyon tabloları (intersection_function_table, bölüm 2.17.3)
• intersection sonuçları (intersection_result, bölüm 2.17.4)
• intersector türleri ve ilişkili fonksiyonlar (intersector, bölüm 6.18.2)
• intersection sorguları (intersection_query, bölüm 6.18.5)

<metal_raytracing> başlık dosyası metal::raytracing ad alanında intersection_tags tanımlar. Bunlar Tablo 2.9'da listelenmiştir ve ayrıca aşağıdakiler tanımlanırken de kullanılır:

• hızlandırma yapısı türleri (acceleration_structure, bölümler 2.17.7 ve 6.18.1)

Bu etiketler ray-tracing sürecini yapılandırır ve farklı türlerin ve tabloların davranışını ve semantiğini kontrol eder. Şunları tanımlar:

• kesiştirilen hızlandırma yapısının türü
• intersection fonksiyonları için mevcut yerleşik parametreler
• intersection fonksiyon tablosundaki intersection fonksiyonunun türü
• intersector türleri veya intersection_query nesneleri üzerinde kullanılabilir yöntemler
• intersection_result türünde döndürülen veriler

intersection_tags, aşağıdaki ilgili kullanımlar arasında hem etiket türü hem de sıra bakımından eşleşmelidir:

• intersection_function_table
• intersection_result
• intersector
• intersection_query

Aksi takdirde derleyici bir hata üretir.

Kesiştirilen hızlandırma yapısı türü, eğer hızlandırma yapısını kesiştirmek için kullanılan diğer ray‑tracing türlerinde instancing, primitive_motion ve instance_motion etiketleri bulunuyorsa bunların sıralamasıyla eşleşmelidir.

intersection_function_table içinde kesişim fonksiyonları çağrılırken, aynı sıralı etiket kümesini kullandıklarından emin olun. Aksi takdirde sonuç tanımsızdır.

Tablo 2.9 — Intersection Tags

instancing

Kesişim fonksiyonlarının, bölüm 5.2.3.7’de açıklandığı gibi yerleşik instance_id ve/veya user_instance_id değerlerini okuyabildiğini ve hızlandırma yapısının bir instance hızlandırma yapısı olduğunu belirtir.

intersector<intersection_tags...>::intersect() fonksiyonu ve intersection_query<intersection_tags...>, hızlandırma yapısının bir instance_acceleration_structure olduğunu varsayar ve instance_id değerini döndürür.

triangle_data

Üçgen kesişim fonksiyonlarının, bölüm 5.2.3.7’de açıklandığı gibi barycentric_coord veya front_facing özniteliklerine sahip giriş parametrelerini okuyabildiğini belirtir.

Bu etiket bir hızlandırma yapısı tanımlanırken kullanılamaz.

intersector<intersection_tags...>::intersect() fonksiyonu ve intersection_query<intersection_tags...> aşağıdaki değerleri döndürür:

• triangle_barycentric_coord
• triangle_front_facing

world_space_data

Bu etiketle bildirilen kesişim fonksiyonlarının aşağıdakileri sorgulayabildiğini belirtir:

• world_space_origin
• world_space_direction
• object_to_world_transform
• world_to_object_transform

(bölüm 5.2.3.7’de açıklandığı gibi).

Bu etiket bir hızlandırma yapısı veya intersection_query tanımlanırken kullanılamaz. intersector ve intersection_function_table içinde dünya uzayı verileri için destek sağlar.

primitive_motion

Tüm OS: Metal 2.4 ve sonrası.

Aşağıdakilerde primitive düzeyinde hareket desteğini etkinleştirir:

• intersector
• intersection_function_table
• hızlandırma yapıları

instance_motion

Tüm OS: Metal 2.4 ve sonrası.

Aşağıdakilerde instance düzeyinde hareket desteğini etkinleştirir:

• intersector
• intersection_function_table
• hızlandırma yapıları

extended_limits

Tüm OS: Metal 2.4 ve sonrası.

Kesişim fonksiyonlarına geçirilen hızlandırma yapılarının aşağıdakiler için genişletilmiş sınırlarla oluşturulduğunu belirtir:

• primitive sayısı
• geometri sayısı
• instance sayısı

Ayrıca görünürlük maskeleri için kullanılan bit sayısını da artırır.

Bu etiket bir hızlandırma yapısı tanımlanırken kullanılamaz.

curve_data

Tüm OS: Metal 3.1 ve sonrası.

Eğri kesişim noktasının curve_parameter değerini aşağıdaki biçimlerde erişilebilir yapar:

• intersection_query yöntemlerinden dönen intersection_result nesnesinin bir alanı olarak
• kesişim fonksiyonlarına bir giriş parametresi olarak

max_levels<Count>

Tüm OS: Metal 3.1 ve sonrası.

Aşağıdakilerde çok seviyeli instancing desteğini etkinleştirir:

• intersector
• intersection_query
• intersection_function_table

Bu etiket hızlandırma yapılarında kullanılamaz.

Count, geçilebilecek maksimum hızlandırma yapısı seviyesi sayısını belirten bir şablon parametresidir.

• intersection_query için: [2, 16] aralığında olmalıdır
• intersector için: [2, 32] aralığında olmalıdır

intersection_function_table için değer, intersection_query veya intersector ile kullanımına uygun olmalıdır.

intersection_function_buffer

Kesişim fonksiyonunun bir intersection function buffer içinde kullanılabilir olduğunu bildirir.

user_data

Tüm OS: Metal 4 ve sonrası.

user_data_buffer ile işaretlenen bir parametre olarak kullanıcı verisi işaretçisini erişilebilir hale getirir. Bu, bir intersection function buffer içinde kullanılmak üzere kesişim fonksiyonuna kaynakların veya diğer verilerin aktarılmasını sağlar.

Intersection Tag'lerinin Geçerli Kombinasyonları

Metal 2.3 ve sonrası:

• etiket yok
• triangle_data
• instancing
• instancing, triangle_data
• instancing, world_space_data

Ek geçerli kombinasyonlar:

• instancing, triangle_data, world_space_data
• triangle_data, primitive_motion
• instancing, triangle_data, primitive_motion
• instancing, world_space_data, primitive_motion
• instancing, triangle_data, instance_motion
• instancing, world_space_data, instance_motion
• instancing, primitive_motion, instance_motion

Metal 2.4 ve sonrası şunları ekler:

• primitive_motion
• instancing, primitive_motion
• instancing, triangle_data, world_space_data, primitive_motion
• instance_motion
• instancing, instance_motion
• instancing, triangle_data, world_space_data, instance_motion

Metal 3.1 ve sonrası:

• curve_data yukarıdaki herhangi bir kombinasyona eklenebilir.
• max_levels<Count> instancing içeren kombinasyonlara eklenebilir.

Metal 4 ve sonrası:

• intersection_function_buffer yukarıdaki tüm kombinasyonlara eklenebilir.
• user_data yalnızca intersection_function_buffer ile birlikte kullanılabilir.

2.17.2 Ray Type

ray yapısı, kesişim için gerekli bir ışının özelliklerini saklar.

struct ray
{
    ray(float3 origin = 0.0f, float3 direction = 0.0f,
        float min_distance = 0.0f, float max_distance = INFINITY);

    float3 origin;
    float3 direction;
    float min_distance;
    float max_distance;
};

Işının origin ve direction alanları dünya uzayındadır. Bir ray nesnesi özel bir kesişim veya üçgen kesişim fonksiyonuna geçirildiğinde min_distance ve max_distance alanları mevcut arama aralığını yansıtır.

Aday kesişimler bulunduğunda, max_distance daraltılmış arama aralığına uyacak şekilde azalır. Kesişim fonksiyonlarının içinde origin ve direction nesne uzayındadır.

Ray yönünün normalize edilmesi gerekmez, ancak sıfır olmamalıdır.

Geçersiz Ray Örnekleri

Bir ray geçersiz olabilir. Örnekler:

• origin veya direction INF ya da NaN içerir
• min_distance == NaN veya max_distance == NaN
• min_distance == INF (not: max_distance pozitif INF olabilir)
• length(ray.direction) == 0.0
• min_distance > max_distance
• min_distance < 0.0 veya max_distance < 0.0

2.17.3 Intersection Function Table

intersection_function_table<intersection_tags...> yapı türü, bölüm 5.1.6’da tanımlandığı gibi shader’a geçirilen özel kesişim fonksiyonlarından oluşan bir tabloyu tanımlar.

Kesişim etiketleri Tablo 2.9’da tanımlanmıştır. intersection_function_table türü üzerindeki etiketler ile kesişim fonksiyonlarının etiketleri eşleşmelidir.

Sorgu Fonksiyonları

bool is_null_intersection_function_table(
    intersection_function_table<intersection_tags...>)

Üye Fonksiyonlar

Girdi sayısını döndürür:

uint size() const

Tablonun boş olup olmadığını kontrol eder:

bool empty() const

Belirli bir indeksteki buffer’ı döndürür (T, device veya constant adres alanında bir işaretçi ya da referans olmalıdır):

template<typename T>
T get_buffer(uint index) const

Belirtilen indekste bir visible_function_table döndürür:

template<typename T>
visible_function_table<T> get_visible_function_table(uint index) const

Metal 3.1 Ek Fonksiyonlar

Bir device veya constant buffer ayarlar:

void set_buffer(const device void *buf, uint index)
void set_buffer(constant void *buf, uint index)

Bir visible function table ayarlar:

template<typename T>
void set_visible_function_table(visible_function_table<T> vft, uint index)

Intersection Result

Bir kesişimin sonucu intersection_result<intersection_tags...> yapısında döndürülür.

class intersection_type {
    none,
    triangle,
    bounding_box,
    curve // Metal 3.1 ve sonrasında kullanılabilir.
};

template <typename...intersection_tags>
struct intersection_result
{
    intersection_type type;
    float distance;
    uint primitive_id;
    uint geometry_id;

    const device void *primitive_data; // Metal 3 ve sonrasında kullanılabilir.

    // intersection_tags içinde max_levels<Count> olmadan instancing varsa kullanılabilir.
    uint instance_id;
    uint user_instance_id; // Metal 2.4 ve sonrasında kullanılabilir.

    // Metal 3.1 ve sonrasında instancing ve max_levels<Count> kullanıldığında.
    uint instance_count;
    uint instance_id[Count - 1];
    uint user_instance_id[Count - 1];

    // Yalnızca triangle_data dahil edildiğinde kullanılabilir.
    float2 triangle_barycentric_coord;
    bool triangle_front_facing;

    // world_space_data ve instancing dahil edildiğinde kullanılabilir.
    float4x3 world_to_object_transform;
    float4x3 object_to_world_transform;

    // curve_data dahil edildiğinde kullanılabilir.
    float curve_parameter;
};

Bir ray geçersizse intersection_type::none döndürülür.

Döndürülen mesafe dünya uzayında ölçülür.

v0, v1 ve v2 köşe öznitelikleri için, belirtilen üçgen barycentric noktasındaki öznitelik değeri:

v1 * triangle_barycentric_coord.x +
v2 * triangle_barycentric_coord.y +
v0 * (1.0f - (triangle_barycentric_coord.x + triangle_barycentric_coord.y))

intersection_result_ref

Tüm OS: Metal 3.2 ve sonrası Apple silicon için intersection_result_ref<intersection_tags...> desteği sağlar.

Bazı durumlarda intersection_result kopyalamaktan kaçınılabilir. Bunun yerine intersection_result_ref<intersection_tags...> kullanılabilir; bu türün yaşam süresi intersector::intersect fonksiyonuna geçirilen lambda fonksiyonunun süresiyle sınırlıdır.

template <typename...intersection_tags>
struct intersection_result_ref {
public:
    intersection_type get_type() const;
    float get_distance() const;
    uint get_primitive_id() const;
    uint get_geometry_id() const;
    const device void *get_primitive_data() const;

    float3 get_ray_origin() const;
    float3 get_ray_direction() const;
    float get_ray_min_distance() const;

    // max_levels<Count> olmadan
    uint get_instance_id() const;
    uint get_user_instance_id() const;

    // max_levels<Count> ile
    uint get_instance_count() const;
    uint get_instance_id(uint depth) const;
    uint get_user_instance_id(uint depth) const;

    float2 get_triangle_barycentric_coord() const;
    bool is_triangle_front_facing() const;

    float get_curve_parameter() const;

    float4x3 get_object_to_world_transform() const;
    float4x3 get_world_to_object_transform() const;
};

Intersector

intersector<intersection_tags...> yapısı, hızlandırma yapısı üzerinde dolaşımı kontrol eden ve ray kesişimleri gerçekleştiren (örneğin intersect() aracılığıyla) bir nesneyi tanımlar.

intersection_tags (Tablo 2.9) kullanılarak intersector’un hangi tür hızlandırma yapılarıyla çalışacağı ve hangi fonksiyonların kullanılabilir olacağı özelleştirilebilir.

intersector türündeki etiketler, ilgili kesişim fonksiyonunun (bölüm 5.1.6) etiketleriyle eşleşmelidir; aksi takdirde davranış tanımsızdır.

// Varsayılan bir intersector oluştur.
intersector<> primitiveIntersector;

// Üçgen ve dünya uzayı verisini destekleyen
// özelleştirilmiş bir intersector oluştur.
intersector<triangle_data, instancing, world_space_data>
instanceInter;

intersector<intersection_tags...> türü aşağıdaki kolaylık sağlayan sonuç türünü sunar:

intersector<intersection_tags...>::result

Hızlandırma yapısı türleri şunları içerir:

• primitive_acceleration_structure
• instance_acceleration_structure

Bunlar, buffer bağlama noktaları kullanılarak veya argument buffer aracılığıyla doğrudan bağlanabilen opak nesnelerdir.

Acceleration Structures

struct AccelerationStructs {
    primitive_acceleration_structure prim_accel;
    instance_acceleration_structure inst_accel;
    array<primitive_acceleration_structure, 2> prim_accel_array;
    array<instance_acceleration_structure, 2> inst_accel_array;
};

[[kernel]]
void intersectInstancesKernel(
    primitive_acceleration_structure prim_accel [[buffer(0)]],
    instance_acceleration_structure inst_accel [[buffer(1)]],
    device AccelerationStructs *accels [[buffer(3)]])
{ … }

Bu türlerin varsayılan olarak başlatılmış değişkenlerini oluşturmak mümkündür ve varsayılan değer hızlandırma yapıları için null değerdir.

Metal 2.4 ve sonrasında hızlandırma yapısı, acceleration_structure<intersection_tags…> adlı şablonlanmış bir türle değiştirilmiştir. intersection_tags şablon parametresi boş olabilir veya Tablo 2.9’da (Intersection tags) tanımlanan instancing, primitive_motion veya instance_motion kombinasyonlarından oluşabilir.

Örneğin aşağıdaki kod primitive motion desteği olan bir instance hızlandırma yapısını tanımlar:

acceleration_structure<instancing, primitive_motion> accel_struct;

Primitive Acceleration Structure Etiket Kombinasyonları

• etiket yok
• primitive_motion

Instance Acceleration Structure Etiket Kombinasyonları

• instancing

Geriye dönük uyumluluğu korumak için:

• primitive_acceleration_structure, acceleration_structure<> için bir takma addır
• instance_acceleration_structure, acceleration_structure<instancing> için bir takma addır

Daha önce olduğu gibi bunlar buffer bağlama noktaları kullanılarak veya argument buffer aracılığıyla doğrudan bağlanabilen opak nesnelerdir.

struct AccelerationMotionStructs {
    acceleration_structure<primitive_motion> prim_motion_accel;
    acceleration_structure<instancing, instance_motion> inst_motion_accel;
    array<acceleration_structure<>, 2> prim_accel_array;
    array<acceleration_structure<instancing>, 2> inst_accel_array;
};

[[kernel]]
void intersectMotionKernel(
    acceleration_structure<primitive_motion> prim [[buffer(15)]],
    acceleration_structure<instancing, primitive_motion, instance_motion> inst [[buffer(16)]],
    device AccelerationMotionStructs *accels [[buffer(17)]])
{ … }

Bu hızlandırma yapıları Metal API üzerinden compute veya grafik fonksiyonlarına bağlanırken, hızlandırma yapısının türü shader’da tanımlanan türle eşleşmelidir.

Instance hızlandırma yapıları için, primitive_motion desteği olmayan instance hızlandırma yapıları, primitive_motion bekleyen bir shader’a bağlanabilir.

Örneğin acceleration_structure<instancing> bekleyen bir shader’a geçirilebilen bir instance hızlandırma yapısını içeren bir Metal buffer, acceleration_structure<instancing, primitive_motion> bekleyen bir shader’a da verilebilir.

Bu özellik, primitive motion olup olmamasına bakılmaksızın tek bir shader fonksiyonunun hızlandırma yapısını işleyebilmesini sağlar; ancak bunun karşılığında ray tracing çalışma zamanı primitive motion için kontrol yapar. Bu maliyetten kaçınmak için iki ayrı fonksiyon yazılabilir—biri primitive_motion kullanan bir hızlandırma yapısıyla, diğeri ise kullanmadan.

Intersection Query

intersection_query<intersection_tags...> türü, kullanıcıların ray tracing sürecini ve özel kesişim kodunun ne zaman çağrılacağını tamamen kontrol etmesini sağlayan bir nesne tanımlar.

Kesişim sorgusu nesnesi, sorguyu bir hızlandırma yapısı içinde ilerletmek ve sorgu dolaşım bilgilerini sorgulamak için işlevler sağlar. Hızlandırma yapısının türünü özelleştirmek ve hangi işlevlerin kullanılabilir olduğunu belirlemek için intersection_query<intersection_tags...> türünü oluştururken intersection_tags (Tablo 2.9'da tanımlanmıştır) kullanın (bkz. Bölüm 6.18.5).

Desteklenen Etiket Kombinasyonları

• etiket yok
• triangle_data
• instancing, triangle_data
• instancing, max_levels<Count>
• instancing, triangle_data, max_levels<Count>
• instancing

Metal 3.1 ek kombinasyonları destekler. Metal 3.1 ve sonrasında, yukarıda listelenen tüm kombinasyonlara curve_data eklenebilir.

Kısıtlamalar

• Atama yapılamaz.
• Bir yapı veya union üyeleri için kullanılamaz.

Bu kısıtlamalar, kesişim sorgusu nesnesinin kopyalanmasını engeller.

Interpolant Türü

<metal_interpolate> içinde tanımlanan interpolant türü interpolant<T, P>, çekme modeli enterpolasyonu için bir fragment shader girişini kapsayan şablonlaştırılmış bir türdür (Bölüm 6.11).

• T, girişin veri türünü temsil eder.
• P, perspektif doğruluğunu temsil eder.

Desteklenen değerler:

• T: skaler ve vektör kayan nokta türleri
• P: interpolation::perspective, interpolation::no_perspective

Kullanım:

• [[stage_in]] ile bir fragment shader giriş argümanı olarak.

Böyle bir bildirim, aynı isimde veya [[user(name)]] özniteliğine sahip T türünde karşılık gelen bir vertex shader çıkış argümanı ile eşleşmelidir. Bildirim bir örnekleme-ve-enterpolasyon özniteliğine sahip olamaz (Bölüm 5.4).

Bir interpolant<T, P> değişkeni otomatik olarak T türüne dönüştürülemez. Bunun yerine, enterpolasyon yöntemleri çağrılarak bir değer alınır (bkz. Bölüm 6.11).

Eğer P interpolation::perspective ise, enterpolasyon perspektif açısından doğru şekilde yapılır.

2.19 Vertex Başına Değerler

Tüm OS: Metal 4 ve sonrasında vertex başına değerler desteklenir.

<metal_vertex_value> içinde tanımlanan vertex değer türü vertex_value<T>, fragment shader içinde vertex başına değerlere (raster öncesi vertex başına üçgen öznitelikleri) erişim sağlar.

Fragment shader giriş argümanı olarak vertex_value<T> türünde bir değişken bildirebilirsiniz; burada T türü vertex çıkışındaki karşılık gelen türle eşleşmelidir.

enum class vertex_index { first, second, third };

T get(vertex_index i);

Yukarıdaki fonksiyonu çağırarak i indeksindeki vertex başına (enterpolasyon uygulanmamış) değeri döndürün.

Örnek

struct vertex_in {
    float3 position [[attribute(0)]];
    float4 color [[attribute(1)]];
};

struct vertex_out {
    float4 position [[position]];
    float4 color;
};

[[vertex]]
vertex_out vert(vertex_in vert_in [[stage_in]]) { … }

struct fragment_in {
    float4 position [[position]];
    float3 barycentric_coords [[barycentric_coord, center_no_perspective]];
    vertex_value<float4> color;
};

struct fragment_out {
    float4 color;
};

[[fragment]]
fragment_out frag(fragment_in frag_in [[stage_in]]) {
    fragment_out frag_out;

    auto bc = frag_in.barycentric_coords;
    auto c1 = frag_in.color.get(vertex_index::first);
    auto c2 = frag_in.color.get(vertex_index::second);
    auto c3 = frag_in.color.get(vertex_index::third);

    frag_out.color = c1 * bc.x + c2 * bc.y + c3 * bc.z;
    return frag_out;
}

2.20 Mesh Shader Türleri

Tüm OS: Metal 3 ve sonrasında mesh shader türleri desteklenir. Metal bu türleri mesh pipeline içinde geometriyi işlemek için kullanır ve bunları <metal_mesh> başlık dosyasında tanımlar.

Bir nesne fonksiyonu (Bölüm 5.1.7), object aşamasından belirli bir threadgroup için gönderilen mesh ızgarasının boyutunu belirtmek amacıyla mesh_grid_properties türünü kullanabilir.

Object grid içindeki belirli bir threadgroup için set_threadgroups_per_grid üye fonksiyonu hiç çağrılmazsa, o object grid threadgroup için hiçbir mesh ızgarası gönderilmez.

set_threadgroups_per_grid çağrıları, her thread tarafından gerçekleştirilen threadgroup belleğine yazma işlemi gibi davranır.

Mesh Türü

Bir mesh fonksiyonu (Bölüm 5.1.8) aşağıdaki türde bir argüman kullanabilir:

mesh<V, P, NV, NP, t>

Parametreler:

• V — vertex türü
• P — primitive türü
• NV — maksimum vertex sayısı
• NP — maksimum primitive sayısı
• t — mesh topolojisi

Topoloji numaralandırması:

enum topology {
    point,
    line,
    triangle
};

Vertex Türü Gereksinimleri

Geçerli bir V vertex türü, vertex fonksiyonunun dönüş türü ile aynı kuralları izler (Bölüm 5.2.3.3) ve şu kısıtlamalara sahiptir:

• Skaler veya vektör tamsayı ya da kayan nokta alanları içerebilir.
• Mesh vertex özniteliklerini destekler.

Örnek:

struct VertexOut {
    float4 position [[position]];
};

struct PrimitiveOut {
    float color [[flat]];
};

using custom_mesh_t = metal::mesh<VertexOut, PrimitiveOut, 64, 64, metal::topology::triangle>;

Mesh Statik Üyeleri

void set_vertex(uint I, V v);
void set_primitive(uint I, P p);
void set_primitive_count(uint c);
void set_index(uint I, uchar v);
void set_indices(uint I, uchar2 v);
void set_indices(uint I, uchar4 v);

Ek üyeler:

• uint max_primitive — maksimum primitive sayısı (NP)
• uint indices_per_primitive — t topolojisine göre primitive başına indeks sayısı
• uint max_indices — maksimum indeks sayısı (max_primitive * indices_per_primitive)

Kullanım Kuralları

• set_vertex(I, v), [0, max_vertices) aralığında I indeksindeki vertex'i ayarlar.
• set_primitive(I, p), [0, max_primitive) aralığında I indeksindeki primitive'i ayarlar.
• set_primitive_count(c), c değeri [0, max_primitive] aralığında olacak şekilde primitive sayısını ayarlar.
• set_index(I, v), [0, max_indices) aralığında I indeksindeki bir indeksi ayarlar.

set_indices aşırı yüklemeleri, indeks arabelleği konumu geçerli ve uygun şekilde hizalanmış olduğunda kullanılabilir.

2.21 Tensor Türleri

Tüm OS: Metal 4 ve sonrasında tensor türleri desteklenir.

Tensorler, makine öğreniminde yaygın olarak kullanılan çok boyutlu veri yapılarıdır.

Bir tensor şunlara sahiptir:

• Bir veri türü (tüm elemanlar aynı türü paylaşır).
• Boyut sayısını temsil eden bir rank.
• Boyut genişliklerini ve stride değerlerini temsil eden bir yerleşim (layout).

Tensor Varyantları

Metal iki tensor türü tanımlar:

• tensor<…> — shaderlara argümanlar, global bağlamalar, argüman arabellekleri aracılığıyla iletilir veya shader içinde ayrılır.
• cooperative_tensor<…> — depolama thread belleğindedir ve katılan threadler arasında bölüştürülür.

Extents Türü

<metal_tensor> içinde tanımlanır:

extents<IndexType, size_t… Extents>

Parametreler:

• IndexType — boyut büyüklükleri ve indeksleme için kullanılan tamsayı türü.
• Extents — derleme zamanında belirlenen boyutlar veya dynamic_extent.

Takma ad:

dextents<class IndexType, size_t Rank>

Bu, tüm boyutları dinamik olan extents yapısını temsil eder.

Üye Fonksiyonlar

static constexpr rank_type rank();
static constexpr rank_type rank_dynamic();
static constexpr size_t static_extent(rank_type r);
constexpr index_type extent(rank_type r);

Açıklamalar:

• rank() — boyut sayısı.
• rank_dynamic() — dinamik boyut sayısı.
• static_extent(r) — r boyutunun derleme zamanındaki genişliği.
• extent(r) — r boyutunun çalışma zamanındaki genişliği.

<metal_tensor> başlık dosyası tensor<ElementType, Extents, DescriptorType, class… Tag> türünü tanımlar. Bu türü tensorleri shaderlara argümanlar, global bağlamalar veya argüman arabellekleri aracılığıyla iletmek için kullanın. Ayrıca shader içinde tensor oluşturmak için de bu türü kullanabilirsiniz. Tablo 2.16, şablonu örneklerken belirtebileceğiniz şablon parametrelerini açıklar.

Tablo 2.16 Tensor şablon parametreleri

ElementType

Değer türü şu türlerden biri olabilir: half, bfloat, float, char, uchar, short, ushort, int, uint, long veya ulong.

Adres alanı constant, device, threadgroup veya thread olabilir (bkz. bölüm 4).

Değer const, volatile veya coherent(device) olabilir (bkz. bölüm 4.8).

Extents

Extents, tensorun boyutlarını extents<…> kullanarak tanımlar (bkz. bölüm 2.21.1). IndexType genişlik türü short, ushort, int, uint, long veya ulong olabilir.

DescriptorType

DescriptorType, descriptor'un nerede bulunduğunu açıklar. Şunlardan biri olabilir:

• tensor_handle: tensor, tensor descriptor'una bir handle içerir
• tensor_inline: tensor, tensor descriptor'unu doğrudan tutar

Varsayılan değer tensor_handle'dır.

Tags

Tags, ek derleme zamanı özelliklerini içerir. Desteklenen tek etiket tensor_offset'tir ve yalnızca DescriptorType tensor_handle ise kullanılabilir. Bu etiketle işaretlenmiş bir tensor, tensorun başlangıç noktasını kaydıran bir dizi offset içerir (bkz. bölüm 2.21.2.2).

Tablo 2.17, tensor<ElementType, Extents, DescriptorType, Tags…> tarafından tanımlanan üye türlerini açıklar.

Tablo 2.17 Tensor üye türü tanımı

element_type

Tensor türünü özelleştirirken kullandığınız tam nitelikli eleman türü.

value_type

element_type'ın niteliksiz eşdeğeri.

index_type

Extents, stride ve indeksler için kullanılan tür.

size_type

index_type'ın işaretsiz eşdeğeri.

rank_type

Tensorun rank değeri için kullanılan tür.

extents_type

Tensor türünü özelleştirirken kullandığınız extents türü (bölüm 2.21.1).

Yapıcılar

Tüm tensorler aşağıdaki yapıcıları destekler:

tensor() thread;

// Kopya yapıcılar
tensor(const thread tensor &) thread;
tensor(const device tensor &) thread;
tensor(const device coherent(device) tensor &) thread;

// Dönüşüm yapıcısı extent <-> dextent.
tensor(const thread tensor<element_type, OtherExtents, tensor_handle, Tags...> &other) thread;
tensor(const device tensor<element_type, OtherExtents, tensor_handle, Tags...> &other) thread;
tensor(const device device(coherent) tensor<element_type, OtherExtents, tensor_handle, Tags...> &other) thread;
tensor(const constant tensor<element_type, OtherExtents, tensor_handle, Tags...> &other) thread;

// Dönüşüm yapıcısı tensor_handle, tensor_offset <- tensor_handle.
tensor(const thread tensor<element_type, OtherExtents, tensor_handle> &other) thread;
tensor(const device tensor<element_type, OtherExtents, tensor_handle> &other) thread;
tensor(const device device(coherent) tensor<element_type, OtherExtents, tensor_handle> &other) thread;
tensor(const constant tensor<element_type, OtherExtents, tensor_handle> &other) thread;

Üye Fonksiyonlar

Tensorun rank değerini (boyut sayısını) almak için aşağıdaki üye fonksiyonu çağırın:

static constexpr size_t get_rank();

Tensorun r boyutu boyunca statik extent değerini (boyut büyüklüğünü) almak için aşağıdaki üye fonksiyonu çağırın:

static constexpr size_t get_static_extent(rank_type r);

Örneğin, tensorun extent değeri extents<int, 32, 64> ise, get_static_extent(0) 32 döndürür ve get_static_extent(1) 64 döndürür.

r boyutu boyunca extent değerinin statik olup olmadığını belirlemek için aşağıdaki üye fonksiyonu çağırın:

static constexpr bool has_static_extent(rank_type r);

Tensorun tüm extent değerlerinin statik olup olmadığını belirlemek için aşağıdaki üye fonksiyonu çağırın:

static constexpr bool has_all_static_extent();

Tensorun r boyutu boyunca stride değerini almak için aşağıdaki üye fonksiyonu çağırın:

index_type get_stride(rank_type r);

Eleman Erişimi

template<class… OtherIndexTypes>
reference operator[](OtherIndexTypes… index);

template<class OtherIndexTypes>
reference operator[](thread const array<OtherIndexTypes, get_rank()> &index);

template<class… OtherIndexTypes>
value_type get(OtherIndexTypes… index);

template<class OtherIndexTypes>
value_type get(thread const array<OtherIndexTypes, get_rank()> &index);

template<class… OtherIndexTypes>
void set(value_type v, OtherIndexTypes… index);

template<class OtherIndexTypes>
void set(value_type v, thread const array<OtherIndexTypes, get_rank()> &index);

Dilim (Slice)

template<size_t… SliceExtents, class… OtherIndexTypes>
tensor<ElementType, SliceExtents, DescriptorType, SliceTags…>
slice(OtherIndexTypes… index);

Bir tensorun çok boyutlu bir indeksindeki elemanın referansını almak için [] operatörünü çağırın. İndeks tensor sınırlarının dışındaysa, eleman erişimi tanımsız davranışa yol açar.

Bir indeks konumundaki tensor elemanını yüklemek için get fonksiyonunu çağırın. Bu fonksiyon yayın (broadcast) semantiğini destekler; çok boyutlu indeksin i konumundaki değeri sıfırdan büyük ve get_extent(i) değerinden büyükse, etkin indeks 0 olarak kabul edilir. Etkin indeks tensor sınırlarının dışındaysa, yükleme işlemi varsayılan değeri döndürür.

Bir indeks konumundaki tensor elemanına v değerini yazmak için set fonksiyonunu çağırın. İndeks tensor sınırlarının dışındaysa, GPU yazma işlemini yok sayar.

Başlangıç noktası index kadar kaydırılmış ve extents değerleri SliceExtents olan bir tensor dilimi elde etmek için slice üye fonksiyonunu çağırın. Döndürülen dilim tensoru, özgün tensor ile aynı DescriptorType değerine sahiptir ve ya başlangıcı kaydırılmış bir tensor (bkz. bölüm 2.21.2.2) ya da shader içinde ayrılmış bir tensor (bkz. bölüm 2.21.2.3) olur. Eğer OtherExtents dynamic_extent ise, slice fonksiyonu index konumundan başlayarak kalan elemanları döndürür. Bu durum tensorun giriş tensorunun sınırlarının dışına çıkmasına neden olursa, sonuç tanımsız davranıştır.

2.21.2.1 Host'a Bağlı Tensorler

Host'a bağlı tensorler, host tarafında ayrılan ve yapılandırılan tensorlerdir. Host'a bağlı bir tensor bildirmek için DescriptorType şablon parametresi olarak tensor_handle belirtin. ElementType, device veya constant adres alanlarından biri ile nitelendirilebilir.

[[kernel]]
void gemm(
    tensor<device half, dextents<int, 2>, tensor_handle> ta [[buffer(0)]],
    tensor<constant float, dextents<int, 2>> tb [[buffer(1)]])
{ … }

2.21.2.2 Başlangıcı Kaydırılmış Tensorler

Yukarıdaki örnek, ta değişkenini değer türü half olan ve device belleğinde ayrılmış bir tensor olarak tanımlar. tb değişkenini ise değer türü float olan ve constant bellekte ayrılmış bir tensor olarak tanımlar. Varsayılan DescriptorType değeri tensor_handle olduğu için bu durumda ayrıca belirtilmesine gerek olmadığını unutmayın.

Başlangıcı kaydırılmış tensorler, tensor_offset etiketi ile işaretlenmiş host'a bağlı tensorlerdir. Bu tensorlerin başlangıç noktası, temel tensor'a göre bir dizi offset (eleman sayısı cinsinden) kadar kaydırılmıştır.

Örneğin, tensor türünün slice üye fonksiyonunu kullanarak başlangıcı kaydırılmış bir tensor elde edebilirsiniz. Döndürülen tensor, temel tensorun belleğini takma ad (alias) olarak kullanır. slice fonksiyonunun ilk çağrısı, tensor içinde kalan eleman sayısı özgün tensor ve kaydırılmış başlangıca bağlı olduğu için dinamik extents değerine sahip bir tensor döndürür. İkinci çağrı, başlangıcı temel tensorun (32, 32, 32) konumunda olan 16x16x16 boyutunda bir tensor döndürür. Son çağrı ise başlangıcı temel tensorun (16, 16, 32) konumunda olan 16x16x16 boyutunda bir tensor döndürür.

[[kernel]]
void offsetTensor(tensor<device float, extents<int, 64, 128, 256>> tbase) {

    // Başlangıcı kaydırılmış tensor
    tensor<device float, dextents<int,3>, tensor_handle, tensor_offset> t3 =
        tbase.slice(8, 16, 32);

// Orijini kaydırılmış 16x16x16 tensör tensor<device float, extents<int, 16, 16, 16>, tensor_handle, tensor_offset> t4 = tbase.slice<16, 16, 16>(32, 32, 32);

// Orijini kaydırılmış tensör
auto t5 = tbase.slice<16, 16, 16>(16, 16, 32);

}

## 2.21.2.3 Shader-Tarafından Ayrılan Tensörler

Shader tarafından ayrılan (inline) tensörler, doğrudan bir shader içinde ayrılan tensörlerdir. Shader tarafından ayrılan bir tensörü tanımlamak için `DescriptorType` şablon parametresi olarak `tensor_inline` belirtin. `ElementType` değerini `device`, `constant`, `threadgroup` veya `thread` adres alanlarından biriyle nitelendirebilirsiniz. Shader tarafından ayrılan tensör türlerini bir birleşik tür (aggregate type) içinde tanımlayamazsınız (bkz. bölüm 2.12).

Shader tarafından ayrılan tensörler aşağıdaki ek kurucuları destekler:

// İşaretçi, extents ve strides ile ham kurucu tensor(data_handle_type ptr, thread const OtherExtentsType &extents, thread const array<OtherStrideType, get_rank()> &strides) thread;

// İşaretçi ve extents ile ham kurucu (strides için örtük sıkı paketlenmiş yerleşim kullanılır) template <class OtherExtentsType> tensor(data_handle_type ptr, thread const OtherExtentsType &extents) thread;

Örnek:

[[kernel]] void func1(threadgroup half *buf) {

tensor<threadgroup half, dextents<int, 3>, tensor_inline>
    t1(buf, dextents<int, 3>(16, 32, 64));

auto t2 = tensor(buf, dextents<int, 3>(16, 32, 64));

…

}

## 2.21.3 İşbirlikçi Tensör Türü

`<metal_cooperative_tensor>` başlık dosyası `cooperative_tensor<ElementType, Extents, Layout>` türünü tanımlar. `cooperative_tensor`, elemanları thread belleğinde katılan thread’ler arasında bölüştürülmüş olan bir tensörü temsil eder. Her thread yalnızca kendi bölümündeki elemanlara erişebilir. Bu thread’ler aynı threadgroup’a aittir ve ardışık SIMD-gruplarına yayılmış olabilir. Bir `cooperative_tensor` türünü birleşik bir tür içinde tanımlayamazsınız (bkz. bölüm 2.12).

## Tablo 2.18 İşbirlikçi tensör şablon parametreleri

### ElementType
`ElementType`, tensördeki temel türün tipidir. İşbirlikçi tensörler için adres alanı `thread`’dir.

### Extents
`Extents`, tensörün boyutlarını `extents<…>` kullanarak tanımlar (bkz. bölüm 2.21.1).

### Layout
`Layout`, tensörün çok boyutlu koordinat uzayının her thread için önceden bölümlenmiş depolamaya nasıl eşleneceğini belirtir.

Genellikle `cooperative_tensor` doğrudan oluşturulmaz; çünkü `Layout` aygıta özgüdür. Bunun yerine Metal Performance Primitives gibi kütüphaneler kullanılır (bkz. bölüm 7). Bu kütüphaneler, matris çarpımı ve evrişim gibi tensörler üzerinde çalışan işleçleri içeren optimize edilmiş ilkel işlemlerden oluşur. Tensör işlemleri kullanılarak elde edilirler ve bu işlemler ara sonuçları depolamak için bu yapıları kullanır. Tensör işlemi, kullanım şekline ve GPU’ya bağlı olarak bir `cooperative_tensor` için verimli ve yüksek performanslı bir `Layout` belirler.

`Layout`, `cooperative_tensor` yapılandırmasını tanımlayan aşağıdaki arayüzü sağlayan opak bir nesnedir. Yerleşim, `cooperative_tensor` tarafından çeşitli işlevlerini gerçekleştirmek için kullanılır. Genellikle bu işlevleri doğrudan çağırmanız gerekmez.

### Layout Arayüzü

Her thread’in `cooperative_tensor` için ayırması gereken depolama miktarını döndürür:

static size_t thread_storage_size();

Her thread’in ayırması gereken depolamanın hizalamasını döndürür:

static constexpr size_t thread_storage_align();

`idx` konumundaki elemanın geçerli olup olmadığını belirler:

static bool is_valid_element(const thread void *, uint16 idx);

`idx` konumundaki elemana işaretçiyi alır. Eğer `idx` geçersizse sonuç tanımsızdır.

static thread void * get_element_pointer(const thread void *, uint16_t idx);

Elemana ait işaretçi verilerek indeks değerini döndürür:

static uint16_t get_element_index(const thread void *storage, const thread void *element);

`idx` konumundaki çok boyutlu indeksi döndürür:

template <class OtherIndexType, size_t Rank> static array<OtherIndexType, Rank> get_multidimensional_index(const thread void *, uint16_t idx);

Bu thread’e ait elemanları thread başına depolamaya yükler:

template <class T, class E, class D, class... Tags> static void load(thread void *storage, const thread tensor<T, E, D, Tags...> &);

Bu thread’e ait elemanları thread başına depolamadan hedef tensöre kaydeder:

template <class T, class E, class D, class... Tags> static void store(const thread void *storage, const thread tensor<T, E, D, Tags...> &);

İteratörler uyumlu olduğunda indeksleri eşler:

template <class FromIterator, class ToIterator> static uint16_t map_index(const thread void *from_storage, uint16_t from_idx, const thread void *to_storage);

## 2.21.3.2 İşbirlikçi Tensör

### Tablo 2.19 İşbirlikçi tensör türü tanımı

**value_type**  
`element_type` için niteliksiz eşdeğer.

**extents_type**  
`cooperative_tensor` türünü özelleştirmek için kullandığınız extents türü (bölüm 2.21.1).

**index_type**  
Extents için kullandığınız indeks türü.

**size_type**  
`index_type` türünün işaretsiz eşdeğeri.

**thread_index_type**  
Thread başına depolamayı indekslemek için kullanılan indeks türü.

**thread_size_type**  
`thread_index_type` türünün işaretsiz eşdeğeri.

**data_handle_type**  
`element_type` için işaretçi.

**reference**  
`element_type` için referans.

**const_reference**  
`reference` türünün const eşdeğeri.

**iterator**  
`element_type` için rastgele erişimli iteratör.

**const_iterator**  
`iterator` türünün const eşdeğeri.

**element_type**  
`cooperative_tensor` türünü özelleştirirken kullandığınız tam nitelikli eleman türü.

**rank_type**  
`cooperative_tensor` derecesi (rank) için kullanılan tür (`extents` aracılığıyla).

### Cooperative Load

template<class T, class E, class D, class…> void load(const thread tensor<T, E, D, …> &t) thread;

Bu üye işlev, bir `t` tensöründeki tüm elemanları işbirlikçi biçimde cooperative tensöre yükler. Fonksiyon, bir tensörün uyumlu bir cooperative tensöre genişletildiği broadcast semantiğini destekler. İki tensörün broadcast için uyumlu olması için aynı dereceye sahip olmaları ve boyutlar üzerinde yineleme yapılırken boyutların eşit olması veya yüklenen tensörün ilgili boyutta `1` boyutuna sahip olması gerekir.

Örneğin, `64x1` boyutundaki bir tensörü `64x2` boyutundaki bir cooperative tensöre yükleyebilirsiniz.

### Cooperative Store

template<class T, class E, class D, class…> void store(thread tensor<T, E, D, …> &t) thread const;

Bu üye işlev, bir cooperative tensördeki tüm elemanları işbirlikçi biçimde `t` tensörüne kaydeder. Fonksiyon, `load` için açıklanan aynı broadcast semantiğini destekler. Örneğin, `64x1` boyutundaki bir cooperative tensörü `64x2` boyutundaki bir tensöre kaydedebilirsiniz.

Aşağıdaki üye işlevi çağırarak bu thread’e özel olan maksimum eleman sayısını döndürebilirsiniz. Bu değer cooperative tensöre katılan tüm thread’lerde aynıdır.

Bir cooperative tensörde `idx` konumundaki elemana referans almak için `[]` operatörünü çağırın. Eğer `idx` cooperative tensör sınırlarının dışındaysa elemana erişim tanımsız davranışa yol açar.

reference operator[](thread_index_type idx); const_reference operator[](thread_index_type idx) const;

Aşağıdaki üye işlevi çağırarak bu thread’in sahip olduğu bellekten `it`, `idx` veya `ptr` konumundaki değeri alabilirsiniz.

Cooperative tensörün derecesini (rank) almak için aşağıdaki üye işlevi çağırın:

static constexpr rank_type get_rank(); thread_size_type get_capacity() thread const;

Aşağıdaki üye işlevleri çağırarak bu thread’in sahip olduğu bellekten `it`, `idx` veya `ptr` konumuna değer atayabilirsiniz:

void set(iterator it, value_type v) thread;

void set(thread_index_type idx, value_type v) thread;

void set(thread element_type *ptr, value_type v) thread;

Aşağıdaki üye işlevi çağırarak `it`, `idx` veya `ptr` konumundaki elemana karşılık gelen mantıksal çok boyutlu indeksi alabilirsiniz:

array<index_type, get_rank()> get_multidimensional_index(const_iterator it) thread const;

array<index_type, get_rank()> get_multidimensional_index(thread_index_type idx) thread const;

array<index_type, get_rank()> get_multidimensional_index( const thread element_type *ptr) thread const;

iterator begin() thread; const_iterator begin() thread const;

Aşağıdaki üye işlevi çağırarak `it`, `idx` veya `ptr` tarafından işaret edilen elemanın geçerli olup olmadığını belirleyebilirsiniz. Döndürülen değer false ise eleman geçersizdir ve erişim tanımsız davranışa yol açar.

bool is_valid_element(const_iterator it) const;

bool is_valid_element(thread_index_type idx) const;

bool is_valid_element(const thread element_type *ptr) const;

Aşağıdaki üye işlevler, başlangıca işaret eden ve indeks `0` ile aynı elemana karşılık gelen bir iteratör döndürür.

Aşağıdaki üye işlevler sona işaret eden bir iteratör döndürür.

Aşağıdaki üye işlevler `idx` veya `ptr` ile ilişkili elemana karşılık gelen bir iteratör döndürür:

iterator get_iterator(thread_index_type idx) thread; const_iterator get_iterator(thread_index_type idx) thread const; iterator get_iterator(const thread element_type *ptr) thread; const_iterator get_iterator( const thread element_type *ptr) thread const;

Aşağıdaki işlevler, başka bir `cooperative_tensor` içindeki `it` tarafından işaret edilen elemana karşılık gelen bu `cooperative_tensor` içindeki elemana işaret eder. Bu işlevler, iki `cooperative_tensor` yerleşimi uyumlu olduğunda erişilebilir olabilir.

template<class OtherIterator> iterator map_iterator(const thread OtherIterator &it);

template<class OtherIterator> const_iterator map_iterator( const thread OtherIterator &it) const;

## 2.22 Tür Dönüşümleri ve Veriyi Yeniden Yorumlama

`static_cast` operatörü, bir skaler veya vektör türünden başka bir skaler veya vektör türüne varsayılan yuvarlama modu kullanılarak ve doyum uygulanmadan dönüşüm yapar. Kayan nokta türüne dönüştürürken eşitlik durumlarında en yakın çift sayıya yuvarlanır; tamsayıya dönüştürürken sıfıra doğru yuvarlanır.

Kaynak türü skaler veya vektör Boolean ise `false` değeri sıfıra, `true` değeri ise bire dönüştürülür.

Metal, herhangi bir skaler veya vektör veri türünün (işaretçi olmayan) aynı boyuttaki başka bir skaler veya vektör veri türü olarak yeniden yorumlanabilmesine izin veren `as_type<type-id>` operatörünü ekler. İşlenendeki bitler değiştirilmeden doğrudan yeni tür olarak döndürülür. Fonksiyon argümanları için olağan tür yükseltmeleri uygulanmaz.

Örneğin:

as_type<float>(0x3f800000)

`0x3f800000` bit deseninin IEEE-754 tek duyarlıklı değer olarak yorumlandığında karşılığı olan `1.0f` değerini döndürür.

`as_type<type-id>` operatörünü kullanarak veriyi farklı bayt sayısına sahip bir türe yeniden yorumlamak hataya yol açar.

### Geçerli ve Geçersiz Tür Dönüşümü Örnekleri

// Geçerli. İçerir:
// (int4)(0x3f800000, 0x40000000, 0x40400000, 0x40800000)
float4 f = float4(1.0f, 2.0f, 3.0f, 4.0f);
int4 i = as_type<int4>(f);

int i;

// Geçerli.
short2 j = as_type<short2>(i);

half4 f;

// Hata. Sonuç ve operand farklı boyutlara sahip
float4 g = as_type<float4>(f);

float4 f;

// Geçerli. g.xyz, f.xyz ile aynı değerlere sahiptir.
float3 g = as_type<float3>(f);

Yerleşik skaler türler arasında ( void hariç ) örtük dönüşümler desteklenir. Örtük dönüşüm gerçekleştiğinde bu işlem yalnızca ifadenin değerinin yeniden yorumlanması değildir; değer yeni türde eşdeğer bir değere dönüştürülür. Örneğin, 5 tamsayı değeri 5.0 kayan nokta değerine dönüştürülür.

bfloat, yalnızca daha yüksek kayan nokta derecesine sahip bir türe örtük dönüşüme izin veren genişletilmiş bir kayan nokta türüdür. bfloat örtük olarak float türüne dönüştürülebilirken half türüne dönüştürülemez; ayrıca float veya half türleri örtük olarak bfloat türüne dönüştürülemez.

Tüm vektör türlerinin dönüşüm derecesi skaler türlerden daha yüksek kabul edilir. Bir vektör türünden başka bir vektör veya skaler türe örtük dönüşüme izin verilmez ve derleme hatası oluşur.

Örneğin, dört bileşenli bir tamsayı vektörünü dört bileşenli bir kayan nokta vektörüne dönüştürme girişimi başarısız olur.

Skaler türden vektör türüne örtük dönüşümler desteklenir. Skaler değer vektörün her elemanında çoğaltılır. Skaler değer ayrıca vektörün eleman türüne göre olağan aritmetik dönüşümlere de tabi olabilir.

Skalerden matrise ve vektörden matrise örtük dönüşümler desteklenmez ve derleme hatası oluşur. Bir matris türünden başka bir matris, vektör veya skaler türe örtük dönüşümlere izin verilmez ve derleme hatası oluşur.

İşaretçi türleri için örtük dönüşümler C++17 Spesifikasyonu'nda açıklanan kuralları izler.