Adres Alanları

• device (bkz. bölüm 4.1)
• constant (bkz. bölüm 4.2)
• thread (bkz. bölüm 4.3)
• threadgroup (bkz. bölüm 4.4)
• threadgroup_imageblock (bkz. bölüm 4.5)
• ray_data (bkz. bölüm 4.6)
• object_data (bkz. bölüm 4.7)

Örnek:

threadgroup int *l_data,
constant float *c_data)
{ … }

Metal bellek modeli, MSL programlarında bellek nesnelerinin davranışını ve yapısını tanımlar. Bir adres uzayı özniteliği, tampon bellek nesnelerinin hangi bellek bölgesinden ayrıldığını belirtir. Bu öznitelikler, erişim kısıtlamalarını da belirtebilen ayrık adres uzaylarını tanımlar.

Tüm işletim sistemleri: Metal 1 ve sonrası device, threadgroup, constant ve thread özniteliklerini destekler. Metal 2.3 ve sonrası ray_data özniteliklerini destekler. Metal 3 ve sonrası object_data özniteliklerini destekler.

• iOS: Metal 2 ve sonrası threadgroup_imageblock özniteliğini destekler.
• macOS: Metal 2.3 ve sonrası threadgroup_imageblock özniteliğini destekler.

Bir grafik veya kernel fonksiyonuna verilen ve bir türe işaretçi veya referans olan tüm argümanlar bir adres uzayı özniteliği ile bildirilmelidir.

• Grafik fonksiyonları için işaretçi veya referans argümanları device veya constant adres uzayında bildirilmelidir.
• Kernel fonksiyonları için işaretçi veya referans argümanları device, threadgroup, threadgroup_imageblock veya constant adres uzayında bildirilmelidir.

Birden fazla adres uzayı özniteliğini tanıtan örnek:

void foo(device int *g_data,

Program kapsamındaki bir değişkenin adres uzayı constant olmalıdır.

İşaretçi veya referans olan herhangi bir değişken, bu bölümde ele alınan adres uzayı özniteliklerinden biriyle bildirilmelidir. Bir işaretçi veya referans türü bildiriminde adres uzayı özniteliği eksikse derleme hatası oluşur.

4.1 Device Adres Uzayı

device adres uzayı, cihaz bellek havuzundan ayrılan ve hem okunabilir hem de yazılabilir olan tampon bellek nesnelerini ifade eder.

Bir tampon bellek nesnesi, bir skaler, vektör veya kullanıcı tanımlı yapı türüne işaretçi ya da referans olarak bildirilebilir. Bir uygulamada Metal API çağrıları tampon nesnesi için belleği ayırır ve bu da tampon belleğinin gerçek boyutunu belirler.

Örnekler:

// Dört bileşenli bir float vektör dizisi.
device float4 *color;

struct Foo {
    float a[3];
    int b[2];
};

// Foo elemanlarından oluşan bir dizi.
device Foo *my_info;

4.2 Constant Adres Uzayı

Doku nesneleri her zaman device adres uzayından ayrıldığı için, doku türleri için device adres özniteliğini kullanmanız gerekmez. Bir doku nesnesinin elemanlarına doğrudan erişemezsiniz; bu nedenle bir doku nesnesinden okumak ve ona yazmak için yerleşik fonksiyonları kullanın (bkz. bölüm 6.12).

constant adres uzayı, cihaz bellek havuzundan ayrılan ve salt okunur olan tampon bellek nesnelerini ifade eder.

Program kapsamındaki değişkenleri constant adres uzayında bildirmeli ve bildirim sırasında başlatmalısınız. Başlatıcı ifade bir çekirdek sabit ifade olmalıdır (bkz. C++17 belirtiminin bölüm 5.20). Derleyici bir çekirdek sabit ifadeyi derleme zamanında değerlendirebilir.

Program kapsamındaki değişkenler programla aynı yaşam süresine sahiptir ve değerleri hesaplama veya grafik fonksiyonlarına yapılan çağrılar arasında korunur.

constant adres uzayına işaretçi veya referanslar fonksiyonlara argüman olarak kullanılabilir.

constant adres uzayında bildirilen değişkenlere yazmak derleme zamanı hatasıdır. Böyle bir değişkeni başlatmadan bildirmek de derleme zamanı hatasıdır.

Kernel, vertex ve fragment fonksiyonlarına geçirilen constant adres uzayındaki tamponlar GPU’ya bağlı olarak asgari hizalama gereksinimlerine sahiptir. Daha fazla bilgi için Metal Feature Set Tables içindeki Minimum constant buffer offset alignment bölümüne bakın.

4.3 Thread Adres Uzayı

thread adres uzayı iş parçacığı başına belleği ifade eder. Bu adres uzayında ayrılan değişkenler diğer iş parçacıkları tarafından görülemez. Bir grafik veya kernel fonksiyonu içinde bildirilen değişkenler thread adres uzayında ayrılır.

Örnek:

// İş parçacığı başına adres uzayında ayrılmış bir float
float x;

// İş parçacığı başına adres uzayındaki x değişkenine işaretçi
thread float *p = &x;

Bir GPU hesaplama birimi, bir threadgroup içinde birden fazla iş parçacığını eşzamanlı olarak çalıştırabilir ve bir GPU, her hesaplama birimi için ayrı bir threadgroup çalıştırabilir.

4.4 Threadgroup Adres Uzayı

Bir threadgroup içindeki iş parçacıkları, genellikle device belleğinden daha hızlı olan threadgroup belleğinde veri paylaşarak birlikte çalışabilir.

threadgroup adres uzayını şu amaçlarla kullanın:

• Bir kernel, mesh veya object fonksiyonunda bir threadgroup değişkeni ayırmak.
• Bir threadgroup adresine işaretçi olan bir kernel, fragment veya object fonksiyonu parametresini tanımlamak.

Bir kernel, mesh veya object fonksiyonundaki threadgroup değişkenleri yalnızca kerneli çalıştıran threadgroup’un yaşam süresi boyunca var olur.

Örnek kernel:

kernel void
my_kernel(threadgroup float *sharedParameter [[threadgroup(0)]],

[[kernel]] void
my_kernel(...)
{
    // threadgroup adres uzayında 10 float'tan oluşan bir dizi ayır
    threadgroup float sharedFloatArray[10];
}

4.4.1 SIMD-Grupları ve Quad-Grupları

• macOS: Metal 2 ve sonrası SIMD-group fonksiyonlarını destekler. Metal 2.1 ve sonrası quad-group fonksiyonlarını destekler.
• iOS: Metal 2.2 ve sonrası bazı SIMD-group fonksiyonlarını destekler. Metal 2 ve sonrası quad-group fonksiyonlarını destekler.

Bir threadgroup içinde iş parçacıkları, eşzamanlı yürütülen iş parçacıkları toplulukları olan SIMD-groups olarak bölünebilir. SIMD-groups’a eşleme kernel yürütmesi sırasında değişmeden kalır.

Bir threadgroup içindeki tüm SIMD-groups aynı boyutta olmalıdır; ancak en büyük indekse sahip olan grup, threadgroup boyutu SIMD-group boyutuna tam bölünmüyorsa daha küçük olabilir.

Bir quad-group, iş parçacığı yürütme genişliği 4 olan bir SIMD-group’tur.

Daha fazla bilgi için:

• İş parçacığı senkronizasyon fonksiyonları: bölüm 6.9.1
• SIMD-group fonksiyonları: bölüm 6.9.2
• Quad-group fonksiyonları: bölüm 6.9.3

4.5 Threadgroup Imageblock Adres Uzayı

threadgroup_imageblock adres uzayı, yalnızca bir imageblock<T, L> nesnesi kullanılarak erişilebilen ve threadgroup belleğinde ayrılan nesneleri ifade eder.

threadgroup_imageblock adres uzayında ayrılmış kullanıcı tanımlı bir türe işaretçi, bir tile shading fonksiyonuna argüman olabilir.

• Her tile için tam olarak bir threadgroup vardır.
• Her threadgroup, kendi tile’ı ile ilişkilendirilmiş threadgroup belleğine ve imageblock’a erişebilir.

Bir kernel fonksiyonunda bu adres uzayında ayrılan değişkenler:

• Kerneli çalıştıran her threadgroup için ayrılır
• Threadgroup içindeki tüm iş parçacıkları tarafından paylaşılır
• Yalnızca o threadgroup’un yaşam süresi boyunca var olur

Her iş parçacığı imageblock’lara erişmek için açık 2B koordinatlar kullanır. İş parçacıkları ile imageblock arasında herhangi bir uzamsal ilişki varsaymayın. Threadgroup boyutları tile boyutundan daha küçük olabilir.

4.6 Ray Data Adres Uzayı

Tüm işletim sistemleri: Metal 2.3 ve sonrası ray_data adres uzayını destekler.

ray_data adres uzayı, yalnızca ışın izleme için kullanılan bir intersection function içinde erişilebilen belleği ifade eder.

Intersection fonksiyonları, ray_data adres uzayında bulunan [[payload]] özniteliğini kullanarak özel bir yükten okuma ve yazma yapabilir. Bir shader intersect() fonksiyonunu bir payload ile çağırdığında:

1. Sistem payload’u ray_data adres uzayına kopyalar.
2. Intersection fonksiyonu çalıştırılır.
3. Fonksiyon döndüğünde payload tekrar dışarı kopyalanır.

4.7 Object Data Adres Uzayı

Tüm işletim sistemleri: Metal 3 ve sonrası object_data adres uzayını destekler.

Object fonksiyonları, bir mesh fonksiyonuna payload iletmek için object_data adres uzayını kullanır. Programlama modeli threadgroup içinde iş birliğine dayanır ve threadgroup adres uzayına benzer.

Threadgroup içindeki iş parçacıkları payload’u ve mesh_grid_properties::set_threadgroups_per_grid değerini hesaplar. object_data içindeki payload açıkça bağlanmaz veya başlatılmaz ve uygulama bunun yaşam süresini yönetir.

4.8 Bellek Tutarlılığı

Metal’deki bellek işlemleri bir tutarlılık kapsamı tanımlar.

Bir store işlemi için kapsam, uygun şekilde senkronize edildiğinde sonucu gözlemleyebilecek iş parçacıkları kümesini tanımlar. Bir load işlemi için ise yüklemenin, uygun şekilde senkronize edildiğinde hangi iş parçacıklarının store işlemlerini gözlemleyebileceğini tanımlar.

Metal aşağıdaki tutarlılık kapsamlarını sağlar:

• Thread coherence — bellek yazmaları yalnızca aynı iş parçacığı tarafından görülebilir.
• Threadgroup coherence — bellek yazmaları aynı threadgroup içindeki iş parçacıkları tarafından görülebilir.
• Device coherence — bellek yazmaları cihaz üzerindeki tüm threadgroup’lar boyunca tüm iş parçacıkları tarafından görülebilir.

thread adres uzayındaki bellek thread tutarlılığına, threadgroup adres uzayındaki bellek ise threadgroup tutarlılığına sahiptir. Varsayılan olarak device adres uzayındaki bellek threadgroup tutarlılığına sahiptir.

Metal 3.2 ve sonrası, tamponlar için coherent(device) niteleyicilerini ve dokular için memory_coherence_device özelliğini destekler; bu, nesnenin device tutarlılığına sahip olduğunu gösterir. Yani bellek işlemleri uygun şekilde senkronize edildiğinde cihaz üzerindeki iş parçacıkları arasında görünür olur.

[[kernel]] void example(
    coherent device float *dptr1,
    coherent(device) device float4 *dptr2,
    texture2d<float, access::read, memory_coherence_device> tex,
    texture2d<float, access::read,
    memory_coherence::memory_coherence_device> tex2)
{ … }