METAL SHADING LANGUAGE · BÖLÜM 6
Metal Standart Kütüphanesi
ASIL METİN Apple — Metal Shading Language Specification, Sürüm 4
6.1 Ad Alanı ve Başlık Dosyaları
Bu bölüm Metal Standart Kütüphanesi (MSLib) içindeki fonksiyonları açıklar.
Metal, tüm MSLib fonksiyonlarını ve numaralandırmalarını metal ad alanı içinde tanımlar. MSLib fonksiyonlarında açıklanan başlık dosyalarına ek olarak, <metal_stdlib> başlık dosyası da mevcuttur ve MSLib tarafından desteklenen tüm fonksiyonlara erişebilir.
Yerleşik Ortak Fonksiyonlar
clamp
fmin(fmax(x, minval), maxval) değerini döndürür.
minval > maxval ise sonuçlar tanımsızdır.
mix
x ve y değerlerinin doğrusal harmanını döndürür ve şu şekilde uygulanır:
x + (y - x) * a
veya:
(1 - a) * x + a * y
a değeri 0.0 ile 1.0 aralığında olmalıdır. Eğer a değeri 0.0 ile 1.0 aralığında değilse, döndürülen değerler tanımsızdır.
6.2 Ortak Fonksiyonlar
<metal_common> başlık dosyası Tablo 6.1'deki fonksiyonları tanımlar. T, skaler veya vektör türünde half ya da float kayan nokta tiplerinden biridir.
Tablo 6.1. Metal standart kütüphanesindeki ortak fonksiyonlar
T clamp(T x, T minval, T maxval)fmin(fmax(x, minval), maxval) değerini döndürür. minval > maxval ise sonuçlar tanımsızdır.
T mix(T x, T y, T a)x ve y değerlerinin doğrusal harmanını döndürür.
T sign(T x)x > 0 ise 1.0, x = -0.0 ise -0.0, x = +0.0 ise +0.0, veya x < 0 ise -1.0 döndürür. x değeri NaN ise 0.0 döndürür.
T smoothstep(T edge0, T edge1, T x)x <= edge0 ise 0.0, x >= edge1 ise 1.0 döndürür ve edge0 < x < edge1 olduğunda 0 ile 1 arasında düzgün bir Hermite enterpolasyonu gerçekleştirir. Bu, düzgün geçişli bir eşik fonksiyonu gerektiğinde kullanışlıdır.
Şuna eşdeğerdir:
t = clamp((x - edge0) / (edge1 - edge0), 0, 1);
return t * t * (3 - 2 * t);
edge0 >= edge1 ise veya x, edge0 ya da edge1 değerlerinden biri NaN ise sonuçlar tanımsızdır.
T step(T edge, T x)x < edge ise 0.0, aksi halde 1.0 döndürür.
Tek duyarlıklı kayan nokta için Metal, şu ortak fonksiyonların hem hassas hem de hızlı varyantlarını destekler: clamp ve saturate. Hızlı ve hassas fonksiyon varyantları arasındaki fark NaN değerlerini ele alma biçimleridir. Hızlı varyantta NaN davranışı tanımsızdır, hassas varyantlar ise NaN işlemesi için IEEE 754 kurallarını izler.
-ffast-math derleyici seçeneği (bkz. bölüm 1.6.3) Metal kaynağı derlenirken uygun varyantı seçer. Ayrıca metal::precise ve metal::fast iç içe ad alanları, bu ortak fonksiyonların hızlı veya hassas varyantını açıkça seçmek için bir yöntem sağlar.
6.3 Tamsayı Fonksiyonları
<metal_integer> başlık dosyası Tablo 6.2'deki tamsayı fonksiyonlarını tanımlar. T, skaler veya vektör tamsayı tiplerinden biridir. Tu, karşılık gelen işaretsiz skaler veya vektör tamsayı tipidir. T32, skaler veya vektör 32‑bit int veya uint tiplerinden biridir.
Tablo 6.2. Metal standart kütüphanesindeki tamsayı fonksiyonları
T extract_bits(T x, uint offset, uint bits)
Tüm işletim sistemleri: Metal 1.2 ve sonrası
x değerinden [offset, offset + bits - 1] bitlerini çıkarır ve sonucu en düşük anlamlı bitlerde döndürür.
- İşaretsiz veri tipleri için sonucun en yüksek anlamlı bitleri sıfır yapılır.
- İşaretli veri tipleri için en yüksek anlamlı bitler
offset + bits - 1bitinin değeriyle ayarlanır.
bits sıfır ise sonuç sıfırdır. offset ile bits toplamı, işlenenin saklanması için kullanılan bit sayısından büyükse sonuç tanımsızdır.
T insert_bits(T base, T insert, uint offset, uint bits)
Tüm işletim sistemleri: Metal 1.2 ve sonrası
insert değerinin en düşük anlamlı bitlerinin base içine yerleştirilmesini döndürür.
Sonuçta [offset, offset + bits - 1] bitleri insert değerinin [0, bits - 1] bitlerinden alınır ve diğer tüm bitler doğrudan base değerinin karşılık gelen bitlerinden alınır.
bits sıfır ise sonuç base olur. offset ile bits toplamı, işlenenin saklanması için kullanılan bit sayısından büyükse sonuç tanımsızdır.
T clamp(T x, T minval, T maxval)min(max(x, minval), maxval) değerini döndürür. minval > maxval ise sonuçlar tanımsızdır.
T clz(T x)x içindeki en yüksek anlamlı bit konumundan başlayarak öndeki 0 bitlerinin sayısını döndürür. Eğer x değeri 0 ise, x tipinin bit cinsinden boyutunu (veya x bir vektörse bileşen tipinin boyutunu) döndürür.
T ctz(T x)x içindeki sondaki 0 bitlerinin sayısını döndürür. Eğer x değeri 0 ise, x tipinin bit cinsinden boyutunu veya x bir vektörse bileşen tipinin boyutunu döndürür.
T hadd(T x, T y)(x + y) >> 1 değerini döndürür. Ara toplam modulo taşma yapmaz.
T max(T x, T y)x < y ise y, aksi halde x döndürülür.
T max3(T x, T y, T z)
Tüm işletim sistemleri: Metal 2.1 ve sonrasımax(x, max(y, z)) değerini döndürür.
T median3(T x, T y, T z)
Tüm işletim sistemleri: Metal 2.1 ve sonrasıx, y ve z değerlerinin ortadaki değerini döndürür.
T min(T x, T y)y < x ise y, aksi halde x döndürür.
T min3(T x, T y, T z)
Tüm işletim sistemleri: Metal 2.1 ve sonrasımin(x, min(y, z)) değerini döndürür.
T32 mul24(T32 x, T32 y)
Tüm işletim sistemleri: Metal 2.1 ve sonrası
İki adet 24‑bit tamsayı değeri x ve y ile çarpma işlemi yapar ve 32‑bit tamsayı sonucu döndürür. x ve y 32‑bit tamsayıdır ancak yalnızca en düşük 24 bit çarpma işlemine katılır.
T mulhi(T x, T y)x * y işlemini hesaplar ve çarpımın yüksek yarısını döndürür.
T popcount(T x)x içindeki sıfır olmayan bitlerin sayısını döndürür.
T reverse_bits(T x)
Tüm işletim sistemleri: Metal 2.1 ve sonrası
x değerinin bitlerinin ters çevrilmiş halini döndürür. Sonucun n numaralı biti, x değerinin (bits − 1) − n bitinden alınır; burada bits, x değerini temsil etmek için kullanılan toplam bit sayısıdır.
T rhadd(T x, T y)(x + y + 1) >> 1 değerini döndürür. Ara toplam modulo taşma yapmaz.
T rotate(T v, T i)v içindeki her eleman için bitler, i içindeki karşılık gelen elemanın belirttiği bit sayısı kadar sola kaydırılır. Soldan taşan bitler sağ taraftan yeniden eklenir.
T subsat(T x, T y)x − y değerini döndürür ve sonucu doygunluğa getirir.
mul24 fonksiyonu yalnızca x ve y değerleri [-2^23, 2^23 - 1] aralığında işaretli tamsayılar olduğunda veya [0, 2^24 - 1] aralığında işaretsiz tamsayılar olduğunda açıklanan şekilde çalışır. x ve y bu aralıkta değilse çarpma sonucu uygulamaya bağlıdır.
<metal_relational> başlık dosyası Tablo 6.3'teki ilişkisel fonksiyonları tanımlar.
T,bfloattipleri dahil olmak üzere skaler veya vektör kayan nokta tiplerinden biridir.Ti, skaler veya vektör tamsayı ya da Boolean tiplerinden biridir.Tb, yalnızca skaler veya vektör Boolean tiplerini ifade eder.
Tablo 6.3. Metal standart kütüphanesindeki ilişkisel fonksiyonlar
Tb isfinite(T x)
Sonlu değer testi yapar.
Tb isinf(T x)
Sonsuzluk değerini (pozitif veya negatif) test eder.
Tb isnan(T x)
Bir NaN değeri olup olmadığını test eder.
Tb isnormal(T x)
Normal bir değer olup olmadığını test eder.
bool all(Tb x)
Yalnızca x içindeki tüm bileşenler true ise true döndürür.
bool any(Tb x)x içindeki herhangi bir bileşen true ise true döndürür.
Tb isordered(T x, T y)
Argümanların sıralı olup olmadığını test eder. (x == x) && (y == y) döndürür.
Tb isunordered(T x, T y)x veya y değeri NaN ise true, aksi halde false döndürür.
Tb not(Tb x)x değerinin bileşen bazında mantıksal tümleyenini döndürür.
T select(T a, T b, Tb c)
Ti select(Ti a, Ti b, Tb c)
Vektör tipleri için:
result[i] = c[i] ? b[i] : a[i]
Skaler tipler için:
result = c ? b : a
Tb signbit(T x)x içindeki kayan nokta değerinin işaret biti ayarlanmışsa true, aksi halde false döndürür.
6.5 Matematik Fonksiyonları
<metal_math> başlık dosyası Tablo 6.4'teki matematik fonksiyonlarını tanımlar.
T, skaler veya vektör half ya da float kayan nokta tiplerinden biridir.Ti, yalnızca skaler veya vektör tamsayı tiplerini ifade eder.
Tablo 6.4. Metal standart kütüphanesindeki matematik fonksiyonları
T asinh(T x) — x değerinin ters hiperbolik sinüsünü hesaplar.
T atan(T y_over_x) — x değerinin ark tanjantını hesaplar.
T atan2(T y, T x) — y / x değerinin ark tanjantını hesaplar.
T atanh(T x) — x değerinin hiperbolik ark tanjantını hesaplar.
T ceil(T x) — x değerini pozitif sonsuza yuvarlayarak tam sayıya dönüştürür.
T copysign(T x, T y) — x değerini y değerinin işaretine göre işaret değiştirerek döndürür.
T cos(T x) — x değerinin kosinüsünü hesaplar.
T cosh(T x) — hiperbolik kosinüs hesaplar.
T cospi(T x) — cos(πx) hesaplar.
T divide(T x, T y) — x / y hesaplar.
T exp(T x) — tabanı e olan üstel fonksiyon.
T exp2(T x) — tabanı 2 olan üstel fonksiyon.
T exp10(T x) — tabanı 10 olan üstel fonksiyon.
T fdim(T x, T y) — x > y ise x − y, aksi halde +0 döndürür.
T floor(T x) — x değerini negatif sonsuza doğru yuvarlar.
T fma(T a, T b, T c)a ve b değerlerinin sonsuz hassasiyetteki çarpımı ile c toplamının doğru yuvarlanmış kayan nokta gösterimini döndürür. Ara yuvarlama yapılmaz. Kenar durumları IEEE 754‑2008 standardını izler.
T fabs(T x) / T abs(T x) — mutlak değeri hesaplar.
T fmax3(T x, T y, T z) / T max3(T x, T y, T z)
Tüm işletim sistemleri: Metal 2.1 ve sonrasıfmax(x, fmax(y, z)) döndürür.
T fmedian3(T x, T y, T z)
Tüm işletim sistemleri: Metal 1 ve sonrası
T median3(T x, T y, T z)
Tüm işletim sistemleri: Metal 2.1 ve sonrasıx, y ve z değerlerinin ortadaki değerini döndürür.
T fmin(T x, T y) / T min(T x, T y)
Daha küçük değeri döndürür.
T fmin3(T x, T y, T z) / T min3(T x, T y, T z)
Tüm işletim sistemleri: Metal 2.1 ve sonrasıfmin(x, fmin(y, z)) döndürür.
T fmod(T x, T y) — x − y * trunc(x / y) döndürür.
T fract(T x) — x değerinin [0, 1) aralığındaki kesirli kısmını döndürür.
T frexp(T x, Ti &exponent)x değerinden mantis ve üs çıkarır. x değerinin her bileşeni mantissa * 2^exponent eşitliğini sağlar.
Ti ilogb(T x) — üssü tamsayı olarak döndürür.
T ldexp(T x, Ti k) — x değerini 2^k ile çarpar.
T log(T x) — doğal logaritma.
T log2(T x) — tabanı 2 olan logaritma.
T log10(T x) — tabanı 10 olan logaritma.
T modf(T x, T &intval)x değerini tam ve kesirli parçalara ayırır. Kesirli kısmı döndürür ve tam kısmı intval içinde saklar.
T pow(T x, T y) — x^y hesaplar.
T powr(T x, T y) — x ≥ 0 koşulu altında x^y hesaplar.
T rint(T x) — en yakın tam sayıya yuvarlar (eşitlik durumunda çift sayıya yuvarlama).
T round(T x) — en yakın tam sayıya yuvarlar; yarım durumlar sıfırdan uzaklaştırılır.
T rsqrt(T x) — ters karekök.
T sin(T x) — sinüs hesaplar.
T sincos(T x, T &cosval)x değerinin sinüsünü döndürür ve cosval değişkenine kosinüsü yazar.
T sinh(T x) — hiperbolik sinüs.
T sinpi(T x) — sin(πx) hesaplar.
T sqrt(T x) — karekök.
T tan(T x) — tanjant.
T tanh(T x) — hiperbolik tanjant.
T tanpi(T x) — tan(πx) hesaplar.
T trunc(T x) — sıfıra doğru yuvarlar.
Eğer hem x hem de y kayan nokta sıfır değerlerini temsil ediyorsa sonuç y değerinin işaretine sahiptir. x veya y değerlerinden biri NaN ise sonuç NaN olur.
fmedian3 için, tüm değerler NaN ise NaN döndürülür. Aksi halde NaN eksik veri olarak değerlendirilir ve kümeden çıkarılır. İki değer NaN ise NaN olmayan değer döndürülür. Bir değer NaN ise fonksiyon iki NaN olmayan değerden herhangi birini döndürebilir.
Tek duyarlıklı kayan nokta için Metal, çoğu matematik fonksiyonu için iki varyant destekler: precise ve fast. -ffast-math derleyici seçeneği derleme sırasında uygun varyantı seçer. metal::precise ve metal::fast ad alanları da açıkça kullanılabilir.
Örnek:
float a = sin(x); // -ffast-math ayarına bağlı olarak fast veya precise
float b = fast::sin(x); // Fast sürüm
float c = precise::cos(x); // Precise sürüm
Tüm işletim sistemleri: Metal 1.2 ve sonrası Tablo 6.5 ve Tablo 6.6 içindeki sabitleri destekler.
Tablo 6.5. Tek duyarlıklı kayan nokta matematik fonksiyonları için sabitler
MAXFLOAT — Sonsuz olmayan en büyük tek duyarlıklı kayan nokta sayısı.
HUGE_VALF — +infinity değerine eşit pozitif float sabit ifadesi.
INFINITY — Pozitif sonsuzluğu temsil eden float sabiti.
NAN — Sessiz bir NaN değerini temsil eden float sabiti.
M_E_F — e değeri.
M_LOG2E_F — log2(e) değeri.
M_LOG10E_F — log10(e) değeri.
M_LN2_F — ln(2) değeri.
M_LN10_F — ln(10) değeri.
M_PI_F — π değeri.
M_PI_2_F — π / 2 değeri.
M_PI_4_F — π / 4 değeri.
M_1_PI_F — 1 / π değeri.
M_2_PI_F — 2 / π değeri.
M_2_SQRTPI_F — 2 / √π değeri.
Tablo 6.6. Half duyarlıklı kayan nokta matematik fonksiyonları için sabitler
MAXHALF — Sonsuz olmayan en büyük half duyarlıklı kayan nokta sayısı.
HUGE_VALH — +infinity değerine eşit pozitif half sabit ifadesi.
M_E_H — e değeri.
M_LOG2E_H — log2(e) değeri.
M_LOG10E_H — log10(e) değeri.
M_LN2_H — ln(2) değeri.
M_LN10_H — ln(10) değeri.
M_PI_H — π değeri.
M_PI_2_H — π / 2 değeri.
M_PI_4_H — π / 4 değeri.
M_1_PI_H — 1 / π değeri.
M_2_PI_H — 2 / π değeri.
M_SQRT2_H — √2 değeri.
M_SQRT1_2_H — 1 / √2 değeri.
M_2_SQRTPI_H — 2 / √π değeri.
Tablo 6.7. bfloat matematik fonksiyonları için sabitler
M_E_BF — e değeri.
M_LOG2E_BF — log2(e) değeri.
M_LOG10E_BF — log10(e) değeri.
M_LN2_BF — ln(2) değeri.
M_LN10_BF — ln(10) değeri.
M_PI_BF — π değeri.
M_PI_2_BF — π / 2 değeri.
M_PI_4_BF — π / 4 değeri.
M_1_PI_BF — 1 / π değeri.
M_2_PI_BF — 2 / π değeri.
M_2_SQRTPI_BF — 2 / √π değeri.
M_SQRT2_BF — √2 değeri.
M_SQRT1_2_BF — 1 / √2 değeri.
6.6 Matris Fonksiyonları
float determinant(floatnxn)
half determinant(halfnxn)
HUGE.VALBF
Pozitif bir half sabit ifadesidir. HUGE.VALBF değeri +infinity olarak değerlendirilir.
<metal_matrix> başlık dosyası Tablo 6.8'deki fonksiyonları tanımlar. T float veya half tipidir.
Tablo 6.8. Metal Standart Kütüphanesinde Matris Fonksiyonları
Yerleşik Matris Fonksiyonları
Açıklama
- Matrisin determinantını hesaplar. Matrisin kare matris olması gerekir.
- Bir matrisi transpoze eder.
Örnek:
float4x4 mA;
float det = determinant(mA);
6.7 SIMD-Group Matris Fonksiyonları
<metal_simdgroup_matrix> başlık dosyası SIMD-group matris fonksiyonlarını tanımlar.
6.7.1 Matris Elemanlarını Oluşturma, Yükleme ve Saklama
Metal Shading Library, bir SIMD-group matrisini bir değerle başlatmak, threadgroup veya device belleğinden veri yüklemek ve veriyi threadgroup veya device belleğine yazmak için aşağıdaki fonksiyonları destekler.
Tablo 6.9. SIMD-Group Matris Yükleme ve Saklama Fonksiyonları
Fonksiyonlar ve Açıklamalar
simdgroup_matrix<T,Cols,Rows> make_filled_simdgroup_matrix(T value)
- Verilen değer ile köşegen bir matris oluşturur.
- Verilen değerle doldurulmuş bir SIMD-group matrisini başlatır.
simdgroup_matrix<T,Cols,Rows>(T dval)
void simdgroup_load( thread simdgroup_matrix<T,Cols,Rows>& d, const threadgroup T *src, ulong elements_per_row = Cols, ulong2 matrix_origin = 0, bool transpose_matrix = false)} }```
- Verileri threadgroup belleğinden bir SIMD-group matrisine yükler.
elements_per_rowparametresi, kaynak bellek düzenindeki eleman sayısını belirtir.
void simdgroup_load( thread simdgroup_matrix<T,Cols,Rows>& d, const device T *src, ulong elements_per_row = Cols, ulong2 matrix_origin = 0, bool transpose_matrix = false)
- Verileri device belleğinden bir SIMD-group matrisine yükler.
elements_per_rowparametresi, kaynak bellek düzenindeki eleman sayısını belirtir.
void simdgroup_store( thread simdgroup_matrix<T,Cols,Rows> a, threadgroup T *dst, ulong elements_per_row = Cols, ulong2 matrix_origin = 0, bool transpose_matrix = false)
- Verileri bir SIMD-group matrisinden threadgroup belleğine kaydeder.
elements_per_rowparametresi, hedef bellek düzenindeki eleman sayısını belirtir.
void simdgroup_store( thread simdgroup_matrix<T,Cols,Rows> a, device T *dst, ulong elements_per_row = Cols, ulong2 matrix_origin = 0, bool transpose_matrix = false)
- Verileri bir SIMD-group matrisinden device belleğine kaydeder.
elements_per_rowparametresi, hedef bellek düzenindeki eleman sayısını belirtir.
İşlemler
void simdgroup_multiply_accumulate( thread simdgroup_matrix<T,Cols,Rows>& d, thread simdgroup_matrix<T,K,Rows>& a, thread simdgroup_matrix<T,Cols,K>& b, thread simdgroup_matrix<T,Cols,Rows>& c)
void simdgroup_multiply( thread simdgroup_matrix<T,Cols,Rows>& d, thread simdgroup_matrix<T,K,Rows>& a, thread simdgroup_matrix<T,Cols,K>& b)
Örnek
kernel void float_matmad(device float *pMatA, device float *pMatB,
device float *pMatC, device float *pMatR)
{
simdgroup_float8x8 sgMatA;
simdgroup_float8x8 sgMatB;
simdgroup_float8x8 sgMatC;
simdgroup_float8x8 sgMatR;
simdgroup_load(sgMatA, pMatA);
simdgroup_load(sgMatB, pMatB);
simdgroup_load(sgMatC, pMatC);
simdgroup_multiply_accumulate(sgMatR, sgMatA, sgMatB, sgMatC);
simdgroup_store(sgMatR, pMatR);
}
Tablo 6.11’deki fonksiyonlar <metal_geometric> başlık dosyasında tanımlıdır. T bir vektör kayan nokta türüdür (floatn veya halfn). Ts ise karşılık gelen skaler türü ifade eder. Eğer T floatn ise, skaler tür Ts float olur. Eğer T halfn ise, Ts half olur.
6.8 Geometrik Fonksiyonlar
Tablo 6.11. Metal Standart Kütüphanesindeki Geometrik Fonksiyonlar
T cross(T x, T y)
x ve y’nin vektörel çarpımını döndürür. T’nin 3 bileşenli bir vektör türü olması gerekir.
Ts distance(T x, T y)
x ve y arasındaki mesafeyi döndürür; bu değer length(x - y) ifadesine eşittir.
Ts distance_squared(T x, T y)
x ve y arasındaki mesafenin karesini döndürür.
Ts dot(T x, T y)
x ve y’nin skaler çarpımını döndürür; bu işlem şu şekildedir:
x[0] * y[0] + x[1] * y[1] + ...
Ts length(T x)
x vektörünün uzunluğunu döndürür; bu değer şu şekildedir:
sqrt(x[0]^2 + x[1]^2 + ...)
Ts length_squared(T x)
x vektörünün uzunluğunun karesini döndürür; bu değer şu şekildedir:
(x[0]^2 + x[1]^2 + ...)
T normalize(T x)
x ile aynı yönde fakat uzunluğu 1 olan bir vektör döndürür.
T reflect(T I, T N)
Gelen vektör I ve yüzey yönelimi N için normalize edilmiş N (NN) hesaplanır ve yansıma yönü döndürülür:
I – 2 * dot(NN, I) * NN
T refract(T I, T N, Ts eta)
Gelen vektör I, yüzey normali N ve kırılma indislerinin oranı eta için kırılma vektörünü döndürür.
İstenen sonuçları elde etmek için gelen vektör I ve yüzey normali N giriş parametrelerinin önceden normalize edilmiş olması gerekir.
T faceforward(T N, T I, T Nref)
Eğer dot(Nref, I) < 0.0 ise N, aksi durumda –N döndürülür.
Tek duyarlıklı kayan nokta için Metal ayrıca aşağıdaki geometrik fonksiyonların hem hassas hem de hızlı bir varyantını destekler: distance, length ve normalize. Metal kaynak kodu derlenirken uygun varyantı seçmek için ffast-math derleyici seçeneğini kullanın (bkz. bölüm 1.6.3).
Buna ek olarak, metal::precise ve metal::fast iç içe ad alanları bu geometrik fonksiyonların hızlı veya hassas varyantını açıkça seçmek için bir yöntem sağlar.
6.9 Senkronizasyon ve SIMD-Group Fonksiyonları
Senkronizasyon ve SIMD-group fonksiyonlarını şu yerlerde kullanabilirsiniz:
[[kernel]]fonksiyonları[[fragment]]fonksiyonları- kernel veya fragment fonksiyonlarının çağırdığı
[[visible]]fonksiyonları
6.9.1 Threadgroup ve SIMD-Group Senkronizasyon Fonksiyonları
<metal_compute> başlık dosyası, Tablo 6.12’de listelenen threadgroup ve SIMD-group senkronizasyon fonksiyonlarını tanımlar.
Tablo 6.12. Metal Standart Kütüphanesindeki Senkronizasyon Hesaplama Fonksiyonları
void simdgroup_barrier(mem_flags flags)
- macOS: Metal 2 ve sonrası
- iOS: Metal 1.2 ve sonrası
Kernel, fragment, mesh veya object çalıştıran bir threadgroup içindeki tüm thread’lerin, herhangi bir thread bariyerin ötesinde yürütmeye devam edebilmeden önce bu fonksiyonu çalıştırması gerekir.
Kernel, fragment, mesh veya object çalıştıran bir SIMD-group içindeki tüm thread’lerin de herhangi bir thread simdgroup_barrier ötesinde yürütmeye devam edebilmeden önce bu fonksiyonu çalıştırması gerekir.
Bir bariyer fonksiyonu (threadgroup_barrier veya simdgroup_barrier) hem yürütme hem de bellek bariyeri görevi görür. Kernel’i çalıştıran bir threadgroup (veya SIMD-group) içindeki tüm thread’lerin bariyer fonksiyonuna ulaşması gerekir.
Apple silicon üzerinde, yürütmesi sona ermiş bir thread artık bariyere katılmaz ve kalan thread’leri engellemez.
Eğer threadgroup_barrier (veya simdgroup_barrier) bir koşul ifadesi içinde yer alıyorsa ve herhangi bir thread koşul bloğuna girip bariyer fonksiyonunu çalıştırıyorsa, threadgroup (veya SIMD-group) içindeki tüm thread’lerin de koşul bloğuna girip bariyer fonksiyonunu çalıştırması gerekir.
Eğer threadgroup_barrier (veya simdgroup_barrier) bir döngü içinde yer alıyorsa, döngünün her yinelemesinde threadgroup (veya SIMD-group) içindeki herhangi bir thread bariyeri çalıştırıyorsa, diğer tüm thread’lerin de yürütmenin devam edebilmesi için bariyer fonksiyonunu çalıştırması gerekir.
threadgroup_barrier (veya simdgroup_barrier) fonksiyonu ayrıca threadgroup veya device belleğine yapılan bellek işlemlerinin doğru sıralamasını sağlamak için bir bellek çiti (okuma ve yazma için) kuyruğa alabilir.
Tablo 6.13. Bariyer Fonksiyonları için Bellek Bayrağı Enum Değerleri
Bellek Bayrakları (mem_flags)
mem_none
Bu bayrak, threadgroup_barrier veya simdgroup_barrier fonksiyonunun yalnızca yürütme bariyeri olarak davranmasını sağlar ve bir bellek çiti uygulamaz.
mem_device
GPU’nun threadgroup veya SIMD-group içindeki thread’ler için device belleğine yapılan bellek işlemlerini doğru sırada gerçekleştirmesini sağlar.
mem_threadgroup
GPU’nun threadgroup veya SIMD-group içindeki thread’ler için threadgroup belleğine yapılan bellek işlemlerini doğru sırada gerçekleştirmesini sağlar.
mem_texture
- macOS: Metal 1.2 ve sonrası
- iOS: Metal 2 ve sonrası
read_write erişim niteleyicisine sahip dokular için GPU’nun threadgroup veya SIMD-group içindeki thread’ler tarafından yapılan texture belleği işlemlerini doğru sırada gerçekleştirmesini sağlar.
mem_threadgroup_imageblock
GPU’nun threadgroup veya SIMD-group içindeki thread’ler için threadgroup imageblock belleğine yapılan bellek işlemlerini doğru sırada gerçekleştirmesini sağlar.
mem_object_data
GPU’nun threadgroup veya SIMD-group içindeki thread’ler için object_data belleğine yapılan bellek işlemlerini doğru sırada gerçekleştirmesini sağlar.
<metal_simdgroup> başlık dosyası kernel ve fragment fonksiyonları için SIMD-group fonksiyonlarını tanımlar.
- macOS, Metal 2 ve sonrasında SIMD-group fonksiyonlarını destekler.
- iOS, çoğu SIMD-group fonksiyonunu Metal 2.2 ve sonrasında destekler.
SIMD-group fonksiyonları, threadgroup belleği kullanmadan veya bariyer gibi senkronizasyon işlemlerine ihtiyaç duymadan SIMD-group içindeki thread’lerin veri paylaşmasına olanak tanır.
Aktif thread, yürütülmekte olan bir thread’dir. Pasif thread, yürütülmeyen bir thread’dir. Örneğin bir thread, akış kontrolü nedeniyle veya bir görevin grubu dolduracak kadar işi olmaması nedeniyle aktif olmayabilir. Bir thread yalnızca SIMD-group içindeki başka bir aktif thread’den veri okumalıdır.
Yardımcı thread’ler de aktif veya pasif olabilir. Örneğin bir yardımcı thread yürütmesini tamamladığında pasif bir yardımcı thread haline gelir. Bir thread’in yardımcı thread olup olmadığını incelemek için simd_is_helper_thread() kullanılır.
Tablo 6.14’te yer tutucu T, aşağıdakiler hariç herhangi bir tamsayı veya kayan nokta türünde skaler veya vektörü temsil eder:
- bool
- long
- ulong
- void
- size_t
- ptrdiff_t
Bit düzeyindeki işlemler için Ti bir tamsayı skaler veya vektörü olmalıdır.
Tablo 6.14. Metal Standart Kütüphanesindeki SIMD-Group Fonksiyonları
simd_active_threads_mask()
Aktif thread’leri temsil eden bir simd_vote maskesi döndürür. Bu fonksiyon simd_ballot(true) ile eşdeğerdir ve aktif thread’leri temsil eden bitleri 1, pasif thread’leri temsil eden bitleri 0 olarak ayarlar.
bool simd_all(bool expr)
Tüm aktif thread’ler expr ifadesini true olarak değerlendirirse true döndürür.
bool simd_any(bool expr)
En az bir aktif thread expr ifadesini true olarak değerlendirirse true döndürür.
T simd_broadcast(T data, ushort broadcast_lane_id)
SIMD lane kimliği broadcast_lane_id olan thread’den veriyi tümüne yayınlar.
T simd_broadcast_first(T data)
SIMD-group içindeki ilk aktif thread’den (en küçük indekse sahip aktif thread) veriyi tüm aktif thread’lere yayınlar.
bool simd_is_first()
Geçerli thread mevcut SIMD-group içindeki ilk aktif thread ise true, aksi halde false döndürür.
T simd_shuffle(T data, ushort simd_lane_id)
SIMD lane kimliği simd_lane_id olan thread’den veriyi döndürür.
T simd_shuffle_and_fill_down(T data, T filling_data, ushort delta)
T simd_shuffle_and_fill_down(T data, T filling_data, ushort delta, ushort modulo)
Çağıran thread’in SIMD lane kimliği ile delta toplamına eşit lane kimliğine sahip thread’den veri veya filling_data döndürür. Eğer bu toplam SIMD-group boyutunu aşarsa, filling_data’nın alt delta lane’lerindeki değerler data’nın üst delta lane’lerine kopyalanır.
T simd_shuffle_and_fill_up(T data, T filling_data, ushort delta)
T simd_shuffle_and_fill_up(T data, T filling_data, ushort delta, ushort modulo)
Çağıran thread’in SIMD lane kimliği ile delta farkına eşit lane kimliğine sahip thread’den veri veya filling_data döndürür. Eğer fark negatifse, filling_data’nın üst delta lane’lerindeki değerler data’nın alt delta lane’lerine kopyalanır.
T simd_shuffle_down(T data, ushort delta)
Çağıran thread’in SIMD lane kimliği ile delta toplamına eşit lane kimliğine sahip thread’den veriyi döndürür. Bu fonksiyon değerleri SIMD-group etrafında dolaştırmaz.
Yerleşik SIMD-group Fonksiyonları
Shuffle Fonksiyonları
T simd_shuffle_rotate_up(T data, ushort delta)
Veriyi, SIMD lane kimliği çağıran thread’in SIMD lane kimliğinden delta çıkarılmasıyla elde edilen thread’den döndürür. delta değeri bir SIMD-group içindeki tüm thread’ler için aynı olmalıdır. Bu fonksiyon değerleri SIMD-group içinde dolaştırır.
T simd_shuffle_up(T data, ushort delta)
Veriyi, SIMD lane kimliği çağıran thread’in SIMD lane kimliğinden delta çıkarılmasıyla elde edilen thread’den döndürür. delta değeri bir SIMD-group içindeki tüm thread’ler için aynı olmalıdır. Bu fonksiyon değerleri SIMD-group içinde dolaştırmadığı için data’nın alt delta lane’lerini değiştirmez.
Ti simd_shuffle_xor(Ti value, ushort mask)
Veriyi, SIMD lane kimliği çağıran thread’in SIMD lane kimliği ile mask değerinin bit düzeyinde XOR (^) işlemine eşit olan thread’den döndürür. mask değeri bir SIMD-group içindeki tüm thread’ler için aynı olmalıdır.
SIMD-group Reduction Fonksiyonları
Tablo 6.15. Metal Standart Kütüphanesindeki SIMD-Group Reduction Fonksiyonları
Ti simd_and(Ti data)
SIMD-group içindeki tüm aktif thread’ler boyunca data değerlerinin bit düzeyinde AND (&) sonucunu döndürür ve sonucu SIMD-group içindeki tüm aktif thread’lere yayınlar.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
bool simd_is_helper_thread()
Geçerli thread bir yardımcı thread ise true, aksi halde false döndürür.
Bu fonksiyonu bir fragment fonksiyonundan veya fragment fonksiyonunuzun çağırdığı başka bir fonksiyondan çağırmanız gerekir; aksi halde derleme zamanında hata oluşabilir.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
T simd_min(T data)
SIMD-group içindeki tüm aktif thread’ler arasından en düşük değere sahip veriyi döndürür ve bu değeri SIMD-group içindeki tüm aktif thread’lere yayınlar.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
Ti simd_or(Ti data)
SIMD-group içindeki tüm aktif thread’ler boyunca data değerlerinin bit düzeyinde OR (|) sonucunu döndürür ve sonucu SIMD-group içindeki tüm aktif thread’lere yayınlar.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
T simd_prefix_exclusive_product(T data)
Belirli bir thread için, SIMD-group içindeki kendisinden daha düşük indekse sahip tüm aktif thread’lerin data giriş değerlerinin çarpımını döndürür. Gruptaki ilk thread T(1) döndürür.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
T simd_prefix_exclusive_sum(T data)
Belirli bir thread için, SIMD-group içindeki kendisinden daha düşük indekse sahip tüm aktif thread’lerin data giriş değerlerinin toplamını döndürür. Gruptaki ilk thread T(0) döndürür.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
T simd_prefix_inclusive_product(T data)
Belirli bir thread için, SIMD-group içindeki kendisinden daha düşük veya aynı indekse sahip tüm aktif thread’lerin data giriş değerlerinin çarpımını döndürür.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
T simd_prefix_inclusive_sum(T data)
Belirli bir thread için, SIMD-group içindeki kendisinden daha düşük veya aynı indekse sahip tüm aktif thread’lerin data giriş değerlerinin toplamını döndürür.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
T simd_product(T data)
SIMD-group içindeki tüm aktif thread’ler boyunca data giriş değerlerinin çarpımını döndürür ve sonucu SIMD-group içindeki tüm aktif thread’lere yayınlar.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
T simd_sum(T data)
SIMD-group içindeki tüm aktif thread’ler boyunca data giriş değerlerinin toplamını döndürür ve sonucu SIMD-group içindeki tüm aktif thread’lere yayınlar.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
Ti simd_xor(Ti data)
SIMD-group içindeki tüm aktif thread’ler boyunca data değerlerinin bit düzeyinde XOR (^) sonucunu döndürür ve sonucu SIMD-group içindeki tüm aktif thread’lere yayınlar.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.3 ve sonrası
6.9.2.1 Örnekler
Başlangıç SIMD-group Durumu
SIMD Lane Kimliği:
0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
Veri:
a b c d e f g h i j k l m n o p
simd_shuffle_up()
simd_shuffle_up() fonksiyonu her SIMD-group’u delta kadar thread yukarı kaydırır.
delta = 2 için örnek:
Hesaplanan SIMD Lane Kimliği:
-2 -1 0 1 2 3 4 5 6 7 8 9 10 11 12 13
Geçerli:
0 0 1 1 1 1 1 1 1 1 1 1 1 1 1 1
Veri:
a b a b c d e f g h i j k l m n
simd_shuffle_up() fonksiyonu, değerleri SIMD-group içinde dolaştırmayan bir işlemdir ve alt delta lane’lerini etkilemez.
simd_shuffle_down()
Benzer şekilde, simd_shuffle_down() fonksiyonu her SIMD grubunu delta iş parçacığı kadar aşağı kaydırır.
delta = 2 için örnek:
Hesaplanan SIMD Lane ID:
2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17
Geçerli:
1 1 1 1 1 1 1 1 1 1 1 1 1 1 0 0
Veri:
c d e f g h i j k l m n o p o p
simd_shuffle_down() fonksiyonu sarma (wrap) yapmayan bir işlemdir ve üstteki delta lane’lerini etkilemez.
Shuffle ve Fill İşlemleri
Başlangıç durumu örneği:
SIMD Lane ID:
0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
Veri:
a b c d e f g h s t u v w x y z
Doldurma:
fa fb fc fd fe ff fg fh fs ft fu fv fw fx fy fz
simd_shuffle_and_fill_up()
Her SIMD grubunu delta iş parçacığı kadar yukarı kaydırır (simd_shuffle_up() ile benzer) ve SIMD lane kimliklerini sararak üst doldurma lane’lerindeki değerleri alt veri lane’lerine atar.
delta = 2 için örnek:
Hesaplanan SIMD Lane ID:
-2 -1 0 1 2 3 4 5 6 7 8 9 10 11 12 13
Veri:
fy fz a b c d e f g h s t u v w x
modulo ile
modulo parametresi SIMD grubunu modulo boyutunda vektörlere böler ve her vektörü delta iş parçacığı kadar kaydırır.
modulo = 8 ve delta = 2 için örnek:
Hesaplanan SIMD Lane ID:
-2 -1 0 1 2 3 4 5 -2 -1 0 1 2 3 4 5
Veri:
fg fh a b c d e f fy fz s t u v w x
simd_shuffle_and_fill_down()
Her SIMD grubunu delta iş parçacığı kadar aşağı kaydırır (simd_shuffle_down() ile benzer) ve alt doldurma lane’lerindeki değerleri üst veri lane’lerine atar.
delta = 2 için örnek:
Hesaplanan SIMD Lane ID:
2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17
Veri:
c d e f g h s t u v w x y z fa fb
modulo ile
modulo = 8 ve delta = 2 için örnek:
Hesaplanan SIMD Lane ID:
2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17
Veri:
c d e f g h fa fb u v w x y z fs ft
Reduction Örneği
kernel void
reduce(const device int *input [[buffer(0)]],
device atomic_int *output [[buffer(1)]],
threadgroup int *ldata [[threadgroup(0)]],
uint gid [[thread_position_in_grid]],
uint lid [[thread_position_in_threadgroup]],
uint lsize [[threads_per_threadgroup]],
uint simd_size [[threads_per_simdgroup]],
uint simd_lane_id [[thread_index_in_simdgroup]],
uint simd_group_id [[simdgroup_index_in_threadgroup]])
{
// İlk seviyede reduction işlemini gerçekleştir.
// Aygıt belleğinden oku, threadgroup belleğine yaz.
int val = input[gid] + input[gid + lsize];
for (uint s = lsize / simd_size; s > simd_size; s /= simd_size)
{
// SIMD başına kısmi reduction işlemi gerçekleştir.
for (uint offset = simd_size / 2; offset > 0; offset /= 2)
val += simd_shuffle_down(val, offset);
// SIMD başına kısmi reduction değerini threadgroup belleğine yaz.
if (simd_lane_id == 0)
ldata[simd_group_id] = val;
// Tüm kısmi reduction işlemlerinin tamamlanmasını bekle.
threadgroup_barrier(mem_flags::mem_threadgroup);
val = (lid < s) ? ldata[lid] : 0;
}
// Son SIMD başına kısmi reduction işlemini gerçekleştir.
for (uint offset = simd_size / 2; offset > 0; offset /= 2)
val += simd_shuffle_down(val, offset);
// Reduction sonucunu atomik olarak güncelle.
if (lid == 0)
atomic_fetch_add_explicit(output, val, memory_order_relaxed);
}
simd_vote Sarmalayıcı Türü
simd_active_threads_mask ve simd_ballot fonksiyonları, açıkça temel türü olan vote_t türüne dönüştürülebilen simd_vote sarmalayıcı türünü kullanır.
class simd_vote {
public:
explicit constexpr simd_vote(vote_t v = 0);
explicit constexpr operator vote_t() const;
bool all() const;
bool any() const;
private:
uint64_t v;
};
simd_all(expr), tüm aktif iş parçacıklarıexprifadesinitrueolarak değerlendiriyorsatruedöndürür.simd_ballot(expr).all(), tüm iş parçacıkları aktifse veexprifadesinitrueolarak değerlendiriyorsatruedöndürür.
simd_vote::all() hangi iş parçacıklarının aktif olduğunu kontrol etmez.
SIMD grubunda 64’ten daha az iş parçacığı bulunan donanımlarda simd_vote değerinin üst bitleri tanımsızdır. Bunların sıfır olduğunu varsaymayın.
6.9.3 Quad-Group Fonksiyonları
macOS: Metal 2.1 ve sonrası quad-group fonksiyonlarını destekler.
Bir quad-group fonksiyonu, yürütme genişliği 4 olan bir SIMD-group fonksiyonudur. Aktif ve aktif olmayan iş parçacıklarının davranışı SIMD gruplarıyla aynı kuralları izler.
Yardımcı iş parçacıkları yalnızca fragment shader içinde quad-group’lar için gradyanları hesaplamak amacıyla çalıştırılır ve ardından pasif hale gelir.
Kernel ve fragment fonksiyonları aşağıdaki tablolarda listelenen quad-group fonksiyonlarını çağırabilir. İş parçacıkları yalnızca quad-group içindeki başka bir aktif iş parçacığından veri okuyabilir.
Yer tutucu T, aşağıdakiler hariç herhangi bir tamsayı veya kayan noktalı skaler ya da vektör türünü temsil eder:
boolvoidsize_tptrdiff_t
Bit düzeyi işlemler için T bir tamsayı skaler veya vektörü olmalıdır.
Tablo 6.16. Quad-group Ballot Fonksiyonu
quad_vote quad_ballot(bool expr)
Her bitin quad-group içindeki aktif iş parçacıklarında expr Boolean ifadesinin true olarak değerlendirildiği konumu gösterdiği bir quad_vote bit maskesi döndürür. Aktif olmayan iş parçacıklarına karşılık gelen bitler 0 olarak ayarlanır.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.2 ve sonrası
Tablo 6.17. Quad-group Permute Fonksiyonları
T quad_broadcast(T data, ushort broadcast_lane_id)
data değerini quad lane kimliği broadcast_lane_id olan iş parçacığından yayınlar. Bu değer quad-group içindeki tüm iş parçacıkları için aynı olmalıdır.
macOS: Metal 2 ve sonrası
iOS: Metal 2 ve sonrası
T quad_broadcast_first(T data)
Quad-group içindeki ilk aktif iş parçacığından (en küçük indeksli aktif iş parçacığı) alınan veriyi tüm aktif iş parçacıklarına yayınlar.
macOS: Metal 2.1 ve sonrası
iOS: Metal 2.2 ve sonrası
T quad_shuffle(T data, ushort quad_lane_id)
Quad lane kimliği quad_lane_id olan iş parçacığından veriyi döndürür. Değer iş parçacıkları arasında farklı olabilir ancak geçerli bir lane kimliği olmalıdır.
macOS: Metal 2 ve sonrası
iOS: Metal 2 ve sonrası
T quad_shuffle_and_fill_down(T data, T filling_data, ushort delta)
Quad lane kimliği çağıran iş parçacığının quad lane kimliği artı delta olan iş parçacığından data veya filling_data değerini döndürür. Eğer toplam quad-group boyutunu aşarsa, filling_data içindeki alt delta lane’lerindeki değerler data içindeki üst lane’lere kopyalanır.
T quad_shuffle_and_fill_down(T data, T filling_data, ushort delta, ushort modulo)
modulo parametresi quad-group’u genişliği modulo olan vektörlere böler (2 veya 4 olmalıdır).
T quad_shuffle_and_fill_up(T data, T filling_data, ushort delta)
Quad lane kimliği çağıran iş parçacığının quad lane kimliği eksi delta olan iş parçacığından data veya filling_data değerini döndürür. Sonuç negatifse, filling_data içindeki üst delta lane’lerindeki değerler data içindeki alt lane’lere kopyalanır.
T quad_shuffle_and_fill_up(T data, T filling_data, ushort delta, ushort modulo)
quad_shuffle_and_fill_up ile benzerdir ancak modulo (2 veya 4) ile tanımlanan vektörler üzerinde çalışır.
Tüm işletim sistemleri: Metal 2.4 ve sonrası.
Yerleşik Quad-Group Fonksiyonları
Shuffle Fonksiyonları
T quad_shuffle_down(T data, ushort delta)
- macOS: Metal 2 ve sonrası
- iOS: Metal 2 ve sonrası
Quad lane kimliği, çağıran iş parçacığının quad lane kimliği ile delta değerinin toplamı olan iş parçacığından veriyi döndürür.
delta değeri quad-group içindeki tüm iş parçacıkları için aynı olmalıdır.
Bu fonksiyon quad-group içinde değerleri sarmadığı için data içindeki üst delta lane’lerini değiştirmez.
T quad_shuffle_rotate_down(T data, ushort delta)
- macOS: Metal 2.1 ve sonrası
- iOS: Metal 2.2 ve sonrası
Quad lane kimliği, çağıran iş parçacığının quad lane kimliği ile delta değerinin toplamı olan iş parçacığından veriyi döndürür.
delta değeri quad-group içindeki tüm iş parçacıkları için aynı olmalıdır.
Bu fonksiyon değerleri quad-group içinde sarar.
T quad_shuffle_rotate_up(T data, ushort delta)
- macOS: Metal 2.1 ve sonrası
- iOS: Metal 2.2 ve sonrası
Quad lane kimliği, çağıran iş parçacığının quad lane kimliği eksi delta farkı olan iş parçacığından veriyi döndürür.
delta değeri quad-group içindeki tüm iş parçacıkları için aynı olmalıdır.
Bu fonksiyon değerleri quad-group içinde sarar.
T quad_shuffle_up(T data, ushort delta)
Quad lane kimliği, çağıran iş parçacığının quad lane kimliği eksi delta farkı olan iş parçacığından veriyi döndürür.
delta değeri quad-group içindeki tüm iş parçacıkları için aynı olmalıdır.
Bu fonksiyon değerleri quad-group içinde sarmadığı için data içindeki alt delta lane’lerini değiştirmez.
T quad_shuffle_xor(T value, ushort mask)
Quad lane kimliği, çağıran iş parçacığının quad lane kimliği ile mask değerinin bit düzeyi XOR (^) işlemi sonucu olan iş parçacığından veriyi döndürür. mask değeri quad-group içindeki tüm iş parçacıkları için aynı olmalıdır.
Quad-Group Reduction Fonksiyonları (Metal Standard Library)
quad_vote quad_active_threads_mask()
- macOS: Metal 2 ve sonrası
- iOS: Metal 2 ve sonrası
Aktif iş parçacıklarını temsil eden bir quad_vote maskesi döndürür.
Fonksiyon quad_ballot(true) ile eşdeğerdir ve aktif iş parçacıklarını temsil eden bitleri 1, aktif olmayanları 0 olarak ayarlar.
bool quad_all(bool expr)
Tüm aktif iş parçacıkları expr ifadesini true olarak değerlendiriyorsa true döndürür.
bool quad_any(bool expr)
En az bir aktif iş parçacığı expr ifadesini true olarak değerlendiriyorsa true döndürür.
bool quad_is_first()
Mevcut iş parçacığı, geçerli quad-group içindeki ilk aktif iş parçacığıysa (en küçük indeksli aktif iş parçacığı) true, aksi halde false döndürür.
bool quad_is_helper_thread()
Mevcut iş parçacığı bir yardımcı iş parçacığıysa true, aksi halde false döndürür.
Bu fonksiyonu bir fragment fonksiyonundan veya fragment fonksiyonunun çağırdığı başka bir fonksiyondan çağırmanız gerekir; aksi halde derleme zamanında hata oluşabilir.
T quad_and(T data)
Quad-group içindeki tüm aktif iş parçacıklarında data değerlerinin bit düzeyi AND (&) sonucunu döndürür ve sonucu quad-group içindeki tüm aktif iş parçacıklarına yayınlar.
T quad_max(T data)
Quad-group içindeki tüm aktif iş parçacıkları arasındaki en yüksek değere sahip data değerini döndürür ve bu değeri quad-group içindeki tüm aktif iş parçacıklarına yayınlar.
T quad_min(T data)
Quad-group içindeki tüm aktif iş parçacıkları arasındaki en düşük değere sahip data değerini döndürür ve bu değeri quad-group içindeki tüm aktif iş parçacıklarına yayınlar.
T quad_or(T data)
- macOS: Metal 2.1 ve sonrası
- iOS: Metal 2.3 ve sonrası
Quad-group içindeki tüm aktif iş parçacıklarında data değerlerinin bit düzeyi OR (|) sonucunu döndürür ve sonucu quad-group içindeki tüm aktif iş parçacıklarına yayınlar.
T quad_xor(T data)
- macOS: Metal 2.1 ve sonrası
- iOS: Metal 2.3 ve sonrası
Quad-group içindeki tüm aktif iş parçacıklarında data değerlerinin bit düzeyi XOR (^) sonucunu döndürür ve sonucu quad-group içindeki tüm aktif iş parçacıklarına yayınlar.
Prefix İşlemleri
T quad_prefix_exclusive_product(T data)
Belirli bir iş parçacığı için, quad-group içindeki kendisinden daha düşük indeksli tüm aktif iş parçacıklarının data giriş değerlerinin çarpımını döndürür. Grup içindeki ilk iş parçacığı için T(1) döndürür.
T quad_prefix_inclusive_product(T data)
Belirli bir iş parçacığı için, quad-group içindeki kendisinden daha düşük veya aynı indeksli tüm aktif iş parçacıklarının data giriş değerlerinin çarpımını döndürür.
T quad_prefix_exclusive_sum(T data)
- macOS: Metal 2.1 ve sonrası
- iOS: Metal 2.3 ve sonrası
Belirli bir iş parçacığı için, quad-group içindeki kendisinden daha düşük indeksli tüm aktif iş parçacıklarının data giriş değerlerinin toplamını döndürür. Grup içindeki ilk iş parçacığı için T(0) döndürür.
T quad_prefix_inclusive_sum(T data)
- macOS: Metal 2.1 ve sonrası
- iOS: Metal 2.3 ve sonrası
Belirli bir iş parçacığı için, quad-group içindeki kendisinden daha düşük veya aynı indeksli tüm aktif iş parçacıklarının data giriş değerlerinin toplamını döndürür.
Grup İşlemleri
T quad_product(T data)
- macOS: Metal 2.1 ve sonrası
- iOS: Metal 2.3 ve sonrası
Quad-group içindeki tüm aktif iş parçacıkları boyunca data giriş değerlerinin çarpımını döndürür ve sonucu quad-group içindeki tüm aktif iş parçacıklarına yayınlar.
T quad_sum(T data)
- macOS: Metal 2.1 ve sonrası
- iOS: Metal 2.3 ve sonrası
Quad-group içindeki tüm aktif iş parçacıkları boyunca data giriş değerlerinin toplamını döndürür ve sonucu quad-group içindeki tüm aktif iş parçacıklarına yayınlar.
Quad Lane Kimlikleri
Bir kernel fonksiyonunda quad’lar SIMD-group boyunca bölünür. Bir fragment fonksiyonunda lane kimliği, 2 × 2 boyutlu bir quad içindeki fragment konumunu temsil eder:
- Lane ID 0 sol üst piksel
- Lane ID 1 sağ üst piksel
- Lane ID 2 sol alt piksel
- Lane ID 3 sağ alt piksel
Shuffle Fonksiyonu Örnekleri
Shuffle fonksiyonlarını göstermek için, aşağıdaki quad-group başlangıç durumuyla başlayın:
data
a
b
c
d
quad_shuffle_up()
quad_shuffle_up() fonksiyonu her quad-group’u delta iş parçacığı kadar yukarı kaydırır.
delta = 2 için örnek:
- Quad lane kimliklerini iki aşağı kaydırır
- Alt iki lane’i geçersiz olarak işaretler
quad_shuffle_up() fonksiyonu sarma yapmayan bir işlemdir ve alt delta lane’lerini etkilemez.
quad_shuffle_down()
quad_shuffle_down() fonksiyonu her quad-group’u delta iş parçacığı kadar aşağı kaydırır.
delta = 2 için örnek:
- Quad lane kimliklerini iki yukarı kaydırır
- Üst iki lane’i geçersiz olarak işaretler
quad_shuffle_down() fonksiyonu sarma yapmayan bir işlemdir ve üst delta lane’lerini etkilemez.
Shuffle-and-Fill Fonksiyonları
Başlangıç quad-group durumu:
data a b c d
filling fa fb fc fd
quad_shuffle_and_fill_up()
Her quad-group’u delta iş parçacığı kadar yukarı kaydırır (quad_shuffle_up() ile benzer) ve quad lane kimliklerini sararak üst doldurma lane’lerindeki değerleri alt veri lane’lerine atar.
delta = 2 için örnek:
fillingiçindeki üst iki lane’idataiçindeki alt iki lane’e atar
modulo parametresi ile
Quad-group, her biri modulo boyutunda vektörlere bölünür ve her vektör delta iş parçacığı kadar kaydırılır.
modulo = 2 ve delta = 1 için örnek:
- Quad lane kimliklerini bir aşağı kaydırır
fillingiçindeki her vektörün üst lane’inidataiçindeki her vektörün alt lane’ine atar
quad_shuffle_and_fill_down()
Her quad-group’u delta iş parçacığı kadar aşağı kaydırır (quad_shuffle_down() ile benzer) ve quad lane kimliklerini sararak alt doldurma lane’lerindeki değerleri üst veri lane’lerine atar.
delta = 2 için örnek:
- Quad lane kimliklerini iki yukarı kaydırır
fillingiçindeki alt iki lane’idataiçindeki üst iki lane’e atar
Ortaya çıkan veri örneği:
data
c
d
fa
fb
modulo parametresi ile
Quad-group, her biri modulo boyutunda vektörlere bölünür ve her vektör delta iş parçacığı kadar kaydırılır.
modulo = 2 ve delta = 1 için örnek:
- Quad lane kimliklerini bir yukarı kaydırır
fillingiçindeki her vektörün alt lane’inidataiçindeki her vektörün üst lane’ine atar
Ortaya çıkan veri örneği:
data
b
fa
d
fc
quad_vote Türü
quad_ballot fonksiyonu, temel türüne açıkça dönüştürülebilen quad_vote sarmalayıcı türünü kullanır.
class quad_vote {
public:
typedef XXX vote_t;
explicit constexpr quad_vote(vote_t v = 0);
explicit constexpr operator vote_t() const;
// Quad-gruptaki iş parçacıklarına karşılık gelen
// tüm bitler (en alttaki dört bit) ayarlanmışsa true döndürür.
bool all() const;
// Quad-gruptaki bir iş parçacığına karşılık gelen
// herhangi bir bit ayarlanmışsa true döndürür.
bool any() const;
};
quad_vote yapıcısı üst bitleri (yani en alttaki dört bit dışındakileri) maskeler. Bu nedenle Metal üst bitleri temizler ve vote_t türüne dönüştürdüğünüzde en alttaki dört bit değişmez.
Grafik Fonksiyonları
Bu bölümde ve alt bölümlerinde yer alan grafik fonksiyonları şu başlık dosyasında tanımlanır:
<metal_graphics>
Bu grafik fonksiyonlarını yalnızca bir fragment fonksiyonundan çağırabilirsiniz.
Tablo 6.19, Tablo 6.20 ve Tablo 6.21’de listelenen fonksiyonlar yalnızca bir fragment fonksiyonunun içinde veya bir fragment fonksiyonundan çağrılan bir fonksiyonun içinde çağrılabilir. Aksi durumda davranış tanımsızdır ve derleme zamanında hataya yol açabilir.
Türevlerin (açık veya örtük) değerlendirilmesine yardımcı olmak için fragment fonksiyonu yardımcı iş parçacıkları oluşturulabilir. Yardımcı iş parçacıkları diğer fragment iş parçacıklarıyla aynı kodu çalıştırır, ancak render hedeflerini veya fragment fonksiyonunun erişebildiği diğer belleği değiştiren yan etkilere sahip değildir.
Özellikle:
- Yardımcı iş parçacıklarına karşılık gelen fragmentler, fragment yürütmesi tamamlandığında render hedeflerinde herhangi bir güncelleme yapılmadan atılır.
- Yardımcı iş parçacıkları tarafından tamponlara ve dokulara yapılan store ve atomik işlemler alttaki belleğe herhangi bir etki yapmaz.
Türev Fragment Fonksiyonları
Metal, türevleri hesaplamak için aşağıdaki fonksiyonları içerir. T aşağıdakilerden biridir:
float, float2, float3, float4, half, half2, half3 veya half4.
Türevler, uniform olmayan kontrol akışı içinde tanımsızdır.
Not: Metal 2.2 ve önceki sürümlerde discard_fragment, kontrol akışını uniform olmayan hale getirebilirdi. Metal 2.3 ve sonraki sürümlerde discard_fragment, kontrol akışının uniform olmayan olarak kabul edilip edilmediğini etkilemez.
Tablo 6.19 — Türev Fragment Fonksiyonları
T dfdx(T p)
Belirtilen değerin ekran uzayı x koordinatına göre yüksek hassasiyetli kısmi türevini döndürür.
T dfdy(T p)
Belirtilen değerin ekran uzayı y koordinatına göre yüksek hassasiyetli kısmi türevini döndürür.
T fwidth(T p)
p için yerel fark alma kullanarak x ve y yönlerindeki mutlak türevlerin toplamını döndürür:
fabs(dfdx(p)) + fabs(dfdy(p))
Fragment Fonksiyonları — Örnekler
Metal aşağıda listelenen örnek başına (per-sample) fonksiyonları içerir.
Tablo 6.20 — Örnek Fragment Fonksiyonları
uint get_num_samples()
Çoklu örneklemeli renk bağlayıcısı için örnek sayısını döndürür.
float2 get_sample_position(uint index)
Belirli bir index örnek indeksi için normalize edilmiş örnek ofsetini (x, y) döndürür. x ve y değerleri [0.0 … 1.0] aralığındadır.
Eğer MTLRenderPassDescriptor nesnesinin setSamplePositions:count: yöntemi kullanılarak örnek konumlarını özelleştirdiyseniz, get_sample_position(index) belirtilen indeks için programlanmış konumu döndürür.
Fragment Akış Kontrolü
Tablo 6.21 — Fragment Akış Kontrol Fonksiyonu
void discard_fragment(void)
Geçerli fragment’i sonlandırılmış olarak işaretler ve bu fragment’in fragment fonksiyonundan çıkan çıktısını atar.
discard_fragment çağrılmadan önce bir fragment iş parçacığından bir tampon veya dokuya yapılan yazmalar atılmaz.
Bir fragment iş parçacığıyla ilişkili birden fazla fragment iş parçacığı veya yardımcı iş parçacığı türevleri hesaplamak için birlikte yürütülür.
- Metal 2.2 ve önceki sürümlerde, bu iş parçacıklarından herhangi biri (ancak hepsi değil)
discard_fragmentçalıştırırsa türev hesaplamaları tanımsızdır. - Metal 2.3 ve sonraki sürümlerde,
discard_fragmentfragment’i sonlandırılmış olarak işaretler, paralel yürütmeyi sürdürür ve türevlerin tanımlı olup olmadığını etkilemez.
Yürütme devam etse bile fragment şu öğeleri atar:
- Fragment çıktısı
discard_fragmentçağrısından sonra tamponlara veya dokulara yapılan tüm yazmalar
Pull-Model Interpolation
Metal 2.3 ve sonraki sürümleri destekleyen tüm işletim sistemleri pull-model interpolasyonu destekler.
interpolant<T,P> interpolant türü ve ilişkili yöntemler şu dosyada tanımlanır:
<metal_interpolate>
Bir fragment fonksiyonunda, interpolant<T,P> değerlerini yöntemlerini çağırarak açıkça interpolate edersiniz. Interpolant farklı modlarda birden fazla kez örneklenebilir ve diğer fonksiyonlara aktarılabilir.
Perspektif doğruluğu, tür içindeki P değerine göre tüm interpolasyonlar boyunca sabittir.
Tablo 6.22 — Pull-Model Interpolant Yöntemleri
T interpolate_at_center()
Shader girdisini bir pikselin merkezinde örnekler ve giriş T türünde olup şu niteliklerden biriyle tanımlanmış olsaydı elde edilecek aynı değeri döndürür:
[[center_perspective]]veya[[center_no_perspective]]
T interpolate_at_centroid()
Shader girdisini pikselin kapsanan alanı içinde örnekler ve giriş T türünde olup şu niteliklerden biriyle tanımlanmış olsaydı elde edilecek aynı değeri döndürür:
[[centroid_perspective]]veya[[centroid_no_perspective]]
T interpolate_at_offset(float2 offset)
Shader girdisini bir pikselin sol üst köşesinden itibaren belirtilen pencere koordinatı ofsetinde örnekler. İzin verilen ofset bileşenleri 1/16 piksel ızgarası boyunca [0.0, 1.0) aralığındadır.
T interpolate_at_sample(uint sample)
Shader girdisini belirtilen örnek indeksinin konumunda örnekler ve giriş T türünde olup şu niteliklerden biriyle tanımlanmış olsaydı elde edilecek aynı değeri döndürür:
[[sample_perspective]]veya[[sample_no_perspective]]
Verilen indeksle bir örnek mevcut değilse interpolasyon konumu tanımsızdır.
⚠️ Uyarı: Orijinal metnin bu kısmı (28. parça) API içerik filtresi veya bağlantı hatası nedeniyle işlenememiştir.
Tv atomic_exchange(ushort coord, ushort array, Tv desired) const
Tv atomic_fetch_op(uint coord, uint array, Tv operand)
Tv atomic_fetch_op(ushort coord, ushort array, Tv operand) const
void atomic_min(uint coord, uint array, ulong4 operand)
void atomic_min(ushort coord, ushort array, ulong4 operand)
void atomic_max(uint coord, uint array, ulong4 operand)
void atomic_max(ushort coord, ushort array, ulong4 operand)
6.12.3 2D Doku
Tv sample(sampler s, float2 coord, int2 offset = int2(0)) const
Tv sample(sampler s, float2 coord, lod_options options, int2 offset = int2(0)) const
Tv sample(sampler s, float2 coord, bias bias_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
Tv sample(sampler s, float2 coord, gradient2d grad_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
Bu üye fonksiyonlar 2D bir dokudan sampler kullanmadan okuma gerçekleştirir:
Tv read(ushort2 coord, ushort lod = 0) const
Tüm işletim sistemleri: Metal 1.2 ve sonrası.
Bu üye fonksiyonlar 2D bir dokuya yazma işlemi yapar. macOS’te Metal 2.2 ve önceki sürümlerde lod sıfır sabiti olmalıdır. Metal 2.3 ve sonraki sürümlerde Apple silicon için bu kısıtlama kaldırılmıştır.
void write(Tv color, ushort2 coord, ushort lod = 0)
Tüm işletim sistemleri: Metal 1.2 ve sonrası.
Bu üye fonksiyon, 2D bir dokudan örnekleme yapılırken bilinear interpolasyon için dört örnek toplar:
Tv gather(sampler s, float2 coord, int2 offset = int2(0), component c = component::x) const
Seyrek 2D Doku Örnekleme
Bu üye fonksiyonlar iOS’ta Metal 2.2 ve sonraki sürümlerde, macOS’te ise Metal 2.3 ve sonraki sürümlerde seyrek bir 2D dokudan örnekleme yapar:
Tv read(uint2 coord, uint lod = 0) const
void write(Tv color, uint2 coord, uint lod = 0)
enum class component { x, y, z, w };
Doku Sorguları
uint get_width(uint lod = 0) const
uint get_height(uint lod = 0) const
Seyrek Örnekleme
sparse_color<Tv> sparse_sample(sampler s, float2 coord, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, bias options, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, level options, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, bias bias_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, gradient2d grad_options, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, gradient2d grad_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
Bu üye fonksiyonlar iOS’ta Metal 2.2 ve sonraki sürümlerde, macOS’te ise Metal 2.3 ve sonraki sürümlerde seyrek bir 2D dokudan sampler kullanılmadan okuma yapar:
sparse_color<Tv> sparse_read(ushort2 coord, ushort lod = 0) const
sparse_color<Tv> sparse_read(uint2 coord, uint lod = 0) const
Bu üye fonksiyon iOS’ta Metal 2.2 ve sonraki sürümlerde, macOS’te ise Metal 2.3 ve sonraki sürümlerde seyrek bir 2D dokudan bilinear interpolasyon için dört örnek toplar:
sparse_color<Tv> sparse_gather(sampler s, float2 coord, int2 offset = int2(0), component c = component::x) const
Bu üye fonksiyonlar iOS’ta Metal 2.3 ve sonraki sürümlerde, macOS’te ise Metal 2.2 ve sonraki sürümlerde bir doku getirme işlemini simüle eder ve LOD (mip seviyesi) hesaplama sonucunu döndürür.
Atomik İşlemler (Metal 3.1 ve sonrası)
2D bir dokudan atomik yükleme.
2D bir dokuya atomik yazma.
Atomik karşılaştır ve değiştir:
bool atomic_compare_exchange_weak(uint2 coord, thread Tv *expected, Tv desired) const
bool atomic_compare_exchange_weak(ushort2 coord, thread Tv *expected, Tv desired) const
Atomik değişim:
Tv atomic_exchange(uint2 coord, Tv desired) const
Tv atomic_exchange(ushort2 coord, Tv desired) const
Atomik getir ve değiştir (op: int ve uint renk türleri için add, and, max, min, or, sub veya xor):
Tv atomic_fetch_op(uint2 coord, Tv operand)
Tv atomic_fetch_op(ushort2 coord, Tv operand) const
Atomik yazma:
void atomic_store(Tv color, uint2 coord) const
void atomic_store(Tv color, ushort2 coord) const
6.12.3.1 2D Doku Örnekleme Örneği
texture2d<float> tex;
sampler s;
float2 coord;
int2 offset;
float lod;
// İsteğe bağlı argüman yok.
float4 clr = tex.sample(s, coord);
// Bir mipmap seviyesi kullanarak örnekleme.
clr = tex.sample(s, coord, level(lod));
// Bir ofset kullanarak örnekleme.
clr = tex.sample(s, coord, offset);
// Bir mipmap seviyesi ve bir ofset kullanarak örnekleme.
clr = tex.sample(s, coord, level(lod), offset);
Ek aşırı yüklemeler:
Tv sample(sampler s, float2 coord, uint array, int2 offset = int2(0)) const
Tv sample(sampler s, float2 coord, uint array, gradient2d grad_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
Aşağıdaki kod, 2D doku örnekleme fonksiyonunun argümanlarına bağlı olarak çeşitli kullanım biçimlerini gösterir.
6.12.4 2D Doku Dizisi
Bu bölümdeki fonksiyonlar için aşağıdaki veri türleri ve karşılık gelen yapıcı fonksiyonlar örnekleme seçeneklerini (lod_options) belirtmek için kullanılabilir:
bias(float value)
level(float lod)
gradient2d(float2 dPdx, float2 dPdy)
min_lod_clamp(float lod) — Tüm işletim sistemleri: Metal 2.2 ve sonrası.
Örnekleme
Tv sample(sampler s, float2 coord, uint array, lod_options options, int2 offset = int2(0)) const
Tv sample(sampler s, float2 coord, uint array, bias bias_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
Sampler Olmadan Okuma
Tv read(ushort2 coord, ushort array, ushort lod = 0) const
Tüm işletim sistemleri: Metal 1.2 ve sonrası.
Yazma
void write(Tv color, ushort2 coord, ushort array, ushort lod = 0)
Tüm işletim sistemleri: Metal 1.2 ve sonrası.
Tv read(uint2 coord, uint array, uint lod = 0) const
Bu üye fonksiyonlar bir 2D doku dizisine yazma işlemi yapar. macOS’te Metal 2.2 ve önceki sürümlerde lod sıfır sabiti olmalıdır. Metal 2.3 ve sonraki sürümlerde Apple silicon için bu kısıtlama kaldırılmıştır.
void write(Tv color, uint2 coord, uint array, uint lod = 0)
Gather
Tv gather(sampler s, float2 coord, uint array, int2 offset = int2(0), component c = component::x) const
Bu üye fonksiyon, bir 2D doku dizisinden örnekleme yapılırken bilinear interpolasyon için dört örnek toplar.
Seyrek 2D Doku Dizisi
sparse_color<Tv> sparse_sample(sampler s, float2 coord, uint array, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, uint array, bias options, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, uint array, level options, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, uint array, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, uint array, gradient2d options, int2 offset = int2(0)) const
sparse_color<Tv> sparse_sample(sampler s, float2 coord, uint array, gradient2d grad_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
Sampler olmadan seyrek okuma:
sparse_color<Tv> sparse_read(ushort2 coord, ushort array, ushort lod = 0) const
sparse_color<Tv> sparse_read(uint2 coord, uint array, uint lod = 0) const
Seyrek gather:
sparse_color<Tv> sparse_gather(sampler s, float2 coord, uint array, int2 offset = int2(0), component c = component::x) const
Atomik İşlemler (Metal 3.1 ve sonrası)
void atomic_store(Tv color, ushort2 coord, ushort array) const
bool atomic_compare_exchange_weak(uint2 coord, uint array, thread Tv *expected, Tv desired) const
bool atomic_compare_exchange_weak(ushort2 coord, ushort array, thread Tv *expected, Tv desired) const
Tv atomic_exchange(uint2 coord, uint array, Tv desired) const
Tv atomic_exchange(ushort2 coord, ushort array, Tv desired) const
Tv atomic_fetch_op(uint2 coord, uint array, Tv operand)
Tv atomic_fetch_op(ushort2 coord, ushort array, Tv operand) const
void atomic_min(uint2 coord, uint array, ulong4 operand)
void atomic_min(ushort2 coord, ushort array, ulong4 operand)
void atomic_max(uint2 coord, uint array, ulong4 operand)
void atomic_max(ushort2 coord, ushort array, ulong4 operand)
6.12.5 3D Doku
Örnekleme seçenekleri için yapıcılar:
min_lod_clamp(float lod) — Tüm işletim sistemleri: Metal 2.2 ve sonrası.
gradient3d(float3 dPdx, float3 dPdy)
Örnekleme
Tv sample(sampler s, float3 coord, int3 offset = int3(0)) const
Tv sample(sampler s, float3 coord, lod_options options, int3 offset = int3(0)) const
Tv sample(sampler s, float3 coord, bias bias_options, min_lod_clamp min_lod_clamp_options, int3 offset = int3(0)) const
Tv sample(sampler s, float3 coord, gradient3d grad_options, min_lod_clamp min_lod_clamp_options, int3 offset = int3(0)) const
Sampler Olmadan Okuma
Tv read(ushort3 coord, ushort lod = 0) const
Tüm işletim sistemleri: Metal 1.2 ve sonrası.
Yazma
void write(Tv color, ushort3 coord, ushort lod = 0)
Tüm işletim sistemleri: Metal 1.2 ve sonrası.
Ek fonksiyonlar:
Tv read(uint3 coord, uint lod = 0) const
void write(Tv color, uint3 coord, uint lod = 0)
Doku Sorguları
uint get_width(uint lod = 0) const
uint get_height(uint lod = 0) const
uint get_depth(uint lod = 0) const
Seyrek 3D Doku
sparse_color<Tv> sparse_sample(sampler s, float3 coord, bias options, int3 offset = int3(0)) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, level options, int3 offset = int3(0)) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, min_lod_clamp min_lod_clamp_options, int3 offset = int3(0)) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, bias bias_options, min_lod_clamp min_lod_clamp_options, int3 offset = int3(0)) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, gradient3d grad_options, int3 offset = int3(0)) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, gradient3d grad_options, min_lod_clamp min_lod_clamp_options, int3 offset = int3(0)) const
Sampler kullanılmadan yapılan seyrek okumalar:
sparse_color<Tv> sparse_read(uint3 coord, uint lod = 0) const
sparse_color<Tv> sparse_read(ushort3 coord, ushort lod = 0) const
Atomik İşlemler (Metal 3.1 ve sonrası)
bool atomic_compare_exchange_weak(uint3 coord, thread Tv *expected, Tv desired) const
bool atomic_compare_exchange_weak(ushort3 coord, thread Tv *expected, Tv desired) const
void atomic_store(Tv color, ushort3 coord) const
Tv atomic_exchange(uint3 coord, Tv desired) const
Tv atomic_exchange(ushort3 coord, Tv desired) const
Tv atomic_fetch_op(uint3 coord, Tv operand)
Tv atomic_fetch_op(ushort3 coord, Tv operand) const
void atomic_min(uint3 coord, ulong4 operand)
void atomic_min(ushort3 coord, ulong4 operand)
void atomic_max(uint3 coord, ulong4 operand)
void atomic_max(ushort3 coord, ulong4 operand)
Küp Doku Örnekleme
Bu üye fonksiyonlar bir küp dokudan örnekleme yapar:
Tv sample(sampler s, float3 coord) const
Tv sample(sampler s, float3 coord, lod_options options) const
Tv sample(sampler s, float3 coord, bias bias_options, min_lod_clamp min_lod_clamp_options) const
Tv sample(sampler s, float3 coord, gradientcube grad_options, min_lod_clamp min_lod_clamp_options) const
Küp Yüz Numaraları
| Yüz numarası | Küp yüzü |
|---|---|
| 0 | Pozitif X |
| 1 | Negatif X |
| 2 | Pozitif Y |
| 3 | Negatif Y |
| 4 | Pozitif Z |
| 5 | Negatif Z |
Tablo 6.22. Küp yüz numarası.
Bu üye fonksiyon, bir küp dokudan örnekleme yapılırken bilineer enterpolasyon için dört örnek toplar.
Bu üye fonksiyonlar bir küp dokuya yazma işlemi gerçekleştirir. macOS'ta, Metal 2.2 ve önceki sürümlerde lod sıfır sabiti olmalıdır. Metal 2.3 ve sonraki sürümler Apple silicon için bu kısıtlamayı kaldırır.
void write(Tv color, ushort2 coord, ushort face, ushort lod = 0) // Tüm OS: Metal 1.2 ve sonrası.
Bu üye fonksiyonlar iOS'ta Metal 2.2 ve sonrasında, macOS'ta Metal 2.3 ve sonrasında seyrek bir küp dokudan örnekleme yapar:
sparse_color<Tv> sparse_sample(sampler s, float3 coord, bias options) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, level options) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, min_lod_clamp min_lod_clamp_options) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, bias bias_options, min_lod_clamp min_lod_clamp_options) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, gradientcube grad_options) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, gradientcube grad_options, min_lod_clamp min_lod_clamp_options) const
Bu üye fonksiyonlar iOS'ta Metal 2.2 ve sonrasında, macOS'ta Metal 2.3 ve sonrasında seyrek bir küp dokudan sampler kullanılmadan okuma işlemi gerçekleştirir:
sparse_color<Tv> sparse_read(ushort2 coord, ushort face, ushort lod = 0) const
Bu üye fonksiyon, iOS'ta Metal 2.2 ve sonrasında, macOS'ta Metal 2.3 ve sonrasında seyrek bir küp dokudan bilineer enterpolasyon için dört örnek toplar:
sparse_color<Tv> sparse_gather(sampler s, float3 coord, component c = component::x) const
Küp Doku Sorguları
uint get_width(uint lod = 0) const
uint get_height(uint lod = 0) const
uint get_num_mip_levels() const
LOD Hesaplama
Bu üye fonksiyonlar bir doku getirme işlemini simüle eder ve LOD (mip seviyesi) hesaplama sonucunu iOS'ta Metal 2.3 ve sonrasında, macOS'ta Metal 2.2 ve sonrasında döndürür:
float calculate_clamped_lod(sampler s, float3 coord);
float calculate_unclamped_lod(sampler s, float3 coord);
Küp Doku Üzerinde Atomik İşlemler (Metal 4 ve sonrası)
Atomik karşılaştır ve değiştir:
bool atomic_compare_exchange_weak(uint2 coord, uint face, thread Tv *expected, Tv desired) const
bool atomic_compare_exchange_weak(ushort2 coord, ushort face, thread Tv *expected, Tv desired) const
Atomik yükleme:
Tv atomic_load(uint2 coord, uint face) const
Tv atomic_load(ushort2 coord, ushort face) const
Atomik saklama:
void atomic_store(Tv color, uint2 coord, uint face) const
void atomic_store(Tv color, ushort2 coord, ushort face) const
Atomik değişim:
Tv atomic_exchange(uint2 coord, uint face, Tv desired) const
Tv atomic_exchange(ushort2 coord, ushort face, Tv desired) const
Atomik min/maks:
void atomic_min(uint2 coord, uint face, ulong4 operand)
void atomic_min(ushort2 coord, ushort face, ulong4 operand)
void atomic_max(uint2 coord, uint face, ulong4 operand)
void atomic_max(ushort2 coord, ushort face, ulong4 operand)
6.12.7 Küp Doku Dizisi
Bu bölümdeki fonksiyonlar için aşağıdaki veri türleri ve karşılık gelen kurucu fonksiyonlar örnekleme seçeneklerini (lod_options) belirtmek için kullanılabilir:
bias(float value)
level(float lod)
gradientcube(float3 dPdx, float3 dPdy)
min_lod_clamp(float lod) // Tüm OS: Metal 2.2 ve sonrası.
Bir Küp Doku Dizisinden Örnekleme
Tv sample(sampler s, float3 coord, uint array, lod_options options) const
Tv sample(sampler s, float3 coord, uint array, gradientcube grad_options, min_lod_clamp min_lod_clamp_options) const
Tv sample(sampler s, float3 coord, uint array) const
Tv sample(sampler s, float3 coord, uint array, bias bias_options, min_lod_clamp min_lod_clamp_options) const
Örnek Toplama
Tv gather(sampler s, float3 coord, uint array, component c = component::x) const
Okuma
Tv read(ushort2 coord, ushort face, ushort array, ushort lod = 0) const // Tüm OS: Metal 1.2 ve sonrası.
Yazma
Bu üye fonksiyonlar bir küp doku dizisine yazma işlemi gerçekleştirir. macOS'ta Metal 2.2 ve önceki sürümlerde lod sıfır sabiti olmalıdır. Metal 2.3 ve sonraki sürümler Apple silicon için bu kısıtlamayı kaldırır.
void write(Tv color, uint2 coord, uint face, uint array, uint lod = 0)
void write(Tv color, ushort2 coord, ushort face, ushort array, ushort lod = 0) // Tüm OS: Metal 1.2 ve sonrası.
Seyrek Küp Doku Dizisi Örnekleme
iOS'ta Metal 2.2 ve sonrasında, macOS'ta Metal 2.3 ve sonrasında desteklenir.
sparse_color<Tv> sparse_sample(sampler s, float3 coord, uint array) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, uint array, bias options) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, uint array, level options) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, uint array, min_lod_clamp min_lod_clamp_options) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, uint array, bias bias_options, min_lod_clamp min_lod_clamp_options) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, uint array, gradientcube options) const
sparse_color<Tv> sparse_sample(sampler s, float3 coord, uint array, gradientcube grad_options, min_lod_clamp min_lod_clamp_options) const
Sampler kullanılmadan okuma:
sparse_color<Tv> sparse_read(uint2 coord, uint face, uint array, uint lod = 0) const
Sorgular
uint get_width(uint lod = 0) const
uint get_height(uint lod = 0) const
uint get_array_size() const
uint get_num_mip_levels() const
Seyrek Gather
sparse_color<Tv> sparse_gather(sampler s, float3 coord, uint array, component c = component::x) const
Atomik İşlemler (Metal 4 ve sonrası)
Atomik karşılaştır ve değiştir:
bool atomic_compare_exchange_weak(uint2 coord, uint face, uint array, thread Tv *expected, Tv desired) const
bool atomic_compare_exchange_weak(ushort2 coord, ushort face, ushort array, thread Tv *expected, Tv desired) const
Atomik yükleme:
Tv atomic_load(uint2 coord, uint face, uint array) const
Tv atomic_load(ushort2 coord, ushort face, ushort array) const
Atomik saklama:
void atomic_store(Tv color, uint2 coord, uint face, uint array) const
void atomic_store(Tv color, ushort2 coord, ushort face, ushort array) const
Atomik değişim:
Tv atomic_exchange(uint2 coord, uint face, uint array, Tv desired) const
Tv atomic_exchange(ushort2 coord, ushort face, ushort array, Tv desired) const
Atomik getir ve değiştir (burada op, int ve uint renk türleri için add, and, max, min, or, sub veya xor olabilir):
Tv atomic_fetch_op(uint2 coord, uint face, uint array, Tv operand)
Tv atomic_fetch_op(ushort2 coord, ushort face, ushort array, Tv operand) const
Atomik min/maks:
void atomic_min(uint2 coord, uint face, uint array, ulong4 operand)
void atomic_max(uint2 coord, uint face, uint array, ulong4 operand)
void atomic_min(ushort2 coord, ushort face, ushort array, ulong4 operand)
void atomic_max(ushort2 coord, ushort face, ushort array, ulong4 operand)
Eğer özelleştirilmiş örnek konumları kullanıyorsanız (MTLRenderPassDescriptor sınıfının setSamplePositions:count: yöntemi ile ayarlanır), read(coord, sample) programlanan örnek konumundaki örneğe ait veriyi döndürür.
Seyrek 2D Çoklu Örneklemeli Doku Okumaları
Bu üye fonksiyonlar iOS'ta Metal 2.2 ve sonrasında, macOS'ta Metal 2.3 ve sonrasında seyrek bir 2D çoklu örneklemeli dokudan sampler kullanılmadan okuma işlemi gerçekleştirir:
sparse_color<Tv> sparse_read(ushort2 coord, ushort sample) const
sparse_color<Tv> sparse_read(uint2 coord, uint sample) const
6.12.9 2D Çoklu Örneklemeli Doku Dizisi
Aşağıdaki üye fonksiyonlar bir 2D çoklu örneklemeli doku dizisinden sampler kullanılmadan okuma gerçekleştirebilir.
Seyrek okumalar iOS'ta Metal 2.2 ve sonrasında, macOS'ta Metal 2.3 ve sonrasında desteklenir.
2D Derinlik Dokusu
Bu bölümdeki fonksiyonlar için aşağıdaki veri türleri ve karşılık gelen kurucu fonksiyonlar örnekleme seçeneklerini (lod_options) belirtmek için kullanılabilir.
bias(float value)
level(float lod)
gradient2d(float2 dPdx, float2 dPdy)
min_lod_clamp(float lod) // Tüm OS: Metal 2.2 ve sonrası.
Örnekleme
T sample(sampler s, float2 coord, lod_options options, int2 offset = int2(0)) const
T sample(sampler s, float2 coord, bias bias_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
T sample(sampler s, float2 coord, gradient2d grad_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
T sample(sampler s, float2 coord, int2 offset = int2(0)) const
Karşılaştırmalı Örnekleme
T sample_compare(sampler s, float2 coord, float compare_value, int2 offset = int2(0)) const
T sample_compare(sampler s, float2 coord, float compare_value, lod_options options, int2 offset = int2(0)) const
sample_compare, compare_value değerini piksel değeriyle karşılaştırır (karşılaştırma başarılıysa 1.0, başarısızsa 0.0). Piksel başına bu karşılaştırma sonuçları normal doku filtrelemesinde olduğu gibi birleştirilir ve 0.0 ile 1.0 arasında bir sonuç değeri döndürülür.
macOS'ta Metal 2.2 ve önceki sürümler lod_options değerleri olan level ve min_lod_clamp seçeneklerini desteklemez; lod sıfır sabiti olmalıdır. Metal 2.3 ve sonraki sürümler Apple silicon için bu kısıtlamayı kaldırır.
Okuma
T read(ushort2 coord, ushort lod = 0) const // Tüm OS: Metal 1.2 ve sonrası.
T read(uint2 coord, uint lod = 0) const
Gather İşlemleri
Tv gather(sampler s, float2 coord, int2 offset = int2(0)) const
Tv gather_compare(sampler s, float2 coord, float compare_value, int2 offset = int2(0)) const
T bir float türü olmalıdır.
Sorgular
uint get_width(uint lod = 0) const
uint get_height(uint lod = 0) const
uint get_num_mip_levels() const
Seyrek 2D Derinlik Doku Örnekleme
iOS'ta Metal 2.2 ve sonrasında, macOS'ta Metal 2.3 ve sonrasında desteklenir.
sparse_color<T> sparse_sample(sampler s, float2 coord, int2 offset = int2(0)) const
sparse_color<T> sparse_sample(sampler s, float2 coord, gradient2d grad_options, int2 offset = int2(0)) const
sparse_color<T> sparse_sample(sampler s, float2 coord, bias options, int2 offset = int2(0)) const
sparse_color<T> sparse_sample(sampler s, float2 coord, level options, int2 offset = int2(0)) const
sparse_color<T> sparse_sample(sampler s, float2 coord, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
sparse_color<T> sparse_sample(sampler s, float2 coord, bias bias_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
sparse_color<T> sparse_sample(sampler s, float2 coord, gradient2d grad_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
Seyrek Karşılaştırmalı Örnekleme
sparse_color<T> sparse_sample_compare(sampler s, float2 coord, float compare_value, int2 offset = int2(0)) const
sparse_color<T> sparse_sample_compare(sampler s, float2 coord, float compare_value, bias options, int2 offset = int2(0)) const
sparse_color<T> sparse_sample_compare(sampler s, float2 coord, float compare_value, level options, int2 offset = int2(0)) const
sparse_color<T> sparse_sample_compare(sampler s, float2 coord, float compare_value, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
sparse_color<T> sparse_sample_compare(sampler s, float2 coord, float compare_value, bias bias_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
sparse_color<T> sparse_sample_compare(sampler s, float2 coord, float compare_value, gradient2d grad_options, int2 offset = int2(0)) const
sparse_color<T> sparse_sample_compare(sampler s, float2 coord, float compare_value, gradient2d grad_options, min_lod_clamp min_lod_clamp_options, int2 offset = int2(0)) const
Sampler kullanılmadan seyrek okumalar:
sparse_color<T> sparse_read(ushort2 coord, ushort lod = 0) const
sparse_color<T> sparse_read(uint2 coord, uint lod = 0) const
Seyrek gather:
sparse_color<Tv> sparse_gather(sampler s, float2 coord, int2 offset = int2(0), component c = component::x) const
Karşılaştırmalı seyrek gather:
sparse_color<Tv> sparse_gather_compare(sampler s, float2 coord, float compare_value, int2 offset = int2(0)) const
6.12.11 2D Derinlik Doku Dizisi
Bu bölümdeki üye fonksiyonlar, lod_options parametrelerinin örnekleme seçenek alanlarını ayarlamak için aşağıdaki veri türlerini ve kurucu fonksiyonları kullanır:
bias(float value)
level(float lod)
gradient2d(float2 dPdx, float2 dPdy)
min_lod_clamp(float lod) // Tüm OS: Metal 2.2 ve sonrası.
⚠️ Uyarı: Orijinal metnin bu kısmı (31. parça) API içerik filtresi veya bağlantı hatası nedeniyle işlenememiştir.
Seyrek Küp Derinlik Doku Dizisi Örnekleme
sparse_color<T> sparse_sample(sampler s, float3 coord, uint array, gradientcube grad_options) const;
sparse_color<T> sparse_sample(
sampler s,
float3 coord,
uint array,
gradientcube grad_options,
min_lod_clamp min_lod_clamp_options
) const;
Bu üye fonksiyonlar, iOS’ta Metal 2.2 ve sonrası ile macOS’ta Metal 2.3 ve sonrası sürümlerde seyrek bir küp derinlik doku dizisinden örnekleme yapar ve tek bir bileşeni bir karşılaştırma değeriyle karşılaştırır.
sparse_color<T> sparse_sample_compare(sampler s, float3 coord, uint array, float compare_value) const;
sparse_color<T> sparse_sample_compare(
sampler s,
float3 coord,
uint array,
float compare_value,
bias bias_options
) const;
sparse_color<T> sparse_sample_compare(
sampler s,
float3 coord,
uint array,
float compare_value,
level options
) const;
sparse_color<T> sparse_sample_compare(
sampler s,
float3 coord,
uint array,
float compare_value,
min_lod_clamp min_lod_clamp_options
) const;
sparse_color<T> sparse_sample_compare(
sampler s,
float3 coord,
uint array,
float compare_value,
bias bias_options,
min_lod_clamp min_lod_clamp_options
) const;
sparse_color<T> sparse_sample_compare(
sampler s,
float3 coord,
uint array,
float compare_value,
gradient2d grad_options
) const;
sparse_color<T> sparse_sample_compare(
sampler s,
float3 coord,
uint array,
float compare_value,
gradient2d grad_options,
min_lod_clamp min_lod_clamp_options
) const;
Bu üye fonksiyonlar, iOS’ta Metal 2.2 ve sonrası ile macOS’ta Metal 2.3 ve sonrası sürümlerde seyrek bir küp derinlik doku dizisinden sampler kullanmadan okuma işlemi gerçekleştirir.
Sparse Gather
sparse_color<Tv> sparse_gather(sampler s, float3 coord, uint array) const;
Bu üye fonksiyon, iOS’ta Metal 2.2 ve sonrası ile macOS’ta Metal 2.3 ve sonrası sürümlerde seyrek bir küp derinlik doku dizisinden çift doğrusal enterpolasyon için dört örnek toplar.
Sparse Gather Compare
Bu üye fonksiyon, seyrek bir 2D derinlik dokusundan örnekleri toplar ve bunları bir karşılaştırma değeriyle karşılaştırır.
sparse_color<Tv> sparse_gather_compare(
sampler s,
float3 coord,
uint array,
float compare_value
) const;
iOS’ta Metal 2.2 ve sonrası ile macOS’ta Metal 2.3 ve sonrası sürümlerde desteklenir.
LOD Simülasyonu
Bu üye fonksiyonlar bir doku getirme işlemini simüle eder ve LOD (mip seviyesi) hesaplama sonucunu iOS’ta Metal 2.3 ve sonrası ile macOS’ta Metal 2.2 ve sonrası sürümlerde döndürür.
Texture Buffer İşlemleri
Aşağıdaki üye fonksiyonlar bir texture buffer içindeki bir elemandan okuma yapabilir ve ona yazabilir (bkz. bölüm 2.9.1).
Atomic Compare and Exchange (Metal 3.1+)
Tv atomic_compare_exchange(uint coord, Tv compare_value, Tv desired) const;
Tv atomic_compare_exchange(ushort coord, Tv compare_value, Tv desired) const;
Atomic Exchange (Metal 3.1+)
Tv atomic_exchange(uint coord, Tv desired) const;
Tv atomic_exchange(ushort coord, Tv desired) const;
Atomic Fetch and Modify (Metal 3.1+)
Burada op, int ve uint renk türleri için add, and, max, min, or, sub veya xor olabilir.
Tv atomic_fetch_op(uint coord, Tv operand) const;
Tv atomic_fetch_op(ushort coord, Tv operand) const;
Atomic Min / Max (Metal 3.1+)
void atomic_min(uint coord, ulong4 operand);
void atomic_min(ushort coord, ulong4 operand);
void atomic_max(uint coord, ulong4 operand);
void atomic_max(ushort coord, ulong4 operand);
Texture Buffer Okuma Örneği
kernel void myKernel(texture_buffer<float, access::read> myBuffer)
{
uint index = ...;
float4 value = myBuffer.read(index);
}
Texture Buffer Boyutunu Sorgulama
uint get_width() const;
Texture Fence Fonksiyonu
texture.fence() üye fonksiyonu, bir iş parçacığının dokuya yaptığı yazma işlemlerinin aynı iş parçacığı tarafından daha sonra gerçekleştirilen okuma işlemleri için görünür olmasını sağlar. access::read_write niteliğiyle bildirilen doku türleri (texture buffer dahil) fence fonksiyonunu destekler.
Örnek
kernel void
my_kernel(texture2d<float, access::read_write> texA,
...,
ushort2 gid [[thread_position_in_grid]])
{
float4 clr = ...;
texA.write(clr, gid);
// İş parçacığı tarafından yapılan yazma işlemlerinin
// aynı iş parçacığının daha sonraki okumaları için
// görünür olmasını sağlamak amacıyla fence kullanılır.
texA.fence();
clr_new = texA.read(gid);
}
macOS: Metal 2 ve sonrası, texture2d_ms_array ve depth2d_ms_array için null texture fonksiyonlarını destekler.
6.12.18 Null Texture Fonksiyonları
Tüm işletim sistemleri: Metal 1.2 ve sonrası null texture fonksiyonlarını destekler.
Bir dokunun null texture olup olmadığını belirlemek için aşağıdaki fonksiyonları kullanın. Doku null ise is_null_texture true, aksi durumda false döndürür.
bool is_null_texture(texture1d<T, access>);
bool is_null_texture(texture1d_array<T, access>);
bool is_null_texture(texture2d<T, access>);
bool is_null_texture(texture2d_array<T, access>);
bool is_null_texture(texture3d<T, access>);
bool is_null_texture(texturecube<T, access>);
bool is_null_texture(texturecube_array<T, access>);
bool is_null_texture(texture2d_ms<T, access>);
bool is_null_texture(depth2d_ms<T, access>);
// Yalnızca macOS (Metal 2 ve sonrası)
bool is_null_texture(depth2d_ms_array<T, access>);
Null bir doku ile herhangi bir doku üye fonksiyonunun çağrılmasının davranışı tanımsızdır.
6.13 Imageblock Fonksiyonları
macOS: Metal 2.3 ve sonrası, Apple silicon için imageblock desteği sağlar.
iOS: Metal 2 ve sonrası imageblock desteği sağlar.
Bu bölüm imageblock’lar için Metal üye fonksiyonlarını listeler. (imageblock veri türü hakkında daha fazla bilgi için 2.11 ve 5.6 bölümlerine bakın.)
Imageblock Bilgisini Sorgulama
ushort get_width() const;
ushort get_height() const;
ushort get_num_samples() const;
Benzersiz Renk Sayısını Sorgulama
ushort get_num_colors(ushort2 coord) const;
Renk Kapsama Maskesi
Renk kapsama maskesini döndürür (belirli bir rengin imageblock içindeki bir veya daha fazla örneği kapsayıp kapsamadığını gösterir). Her örnek, döndürülen değerdeki bit konumuyla tanımlanır.
Bir bit ayarlanmışsa, örnek belirtilen renk indeksini kullanır.
ushort get_color_coverage_mask(ushort2 coord, ushort color_index) const;
color_index değeri 0 ile get_num_colors() - 1 aralığındadır.
6.13.1 Örtük Yerleşime Sahip Imageblock Fonksiyonları
Belirli bir (x, y) koordinatı için imageblock’u piksel hızında okur veya yazar.
Örnek veya Renk Hızı Erişimi
T read(ushort2 coord, ushort index, imageblock_data_rate data_rate) const;
void write(
T data,
ushort2 coord,
ushort index,
imageblock_data_rate data_rate
);
Örnek
struct Foo {
float4 a [[color(0)]];
int4 b [[color(1)]];
};
kernel void
my_kernel(imageblock<Foo, imageblock_layout_implicit> img_blk,
ushort2 lid [[thread_position_in_threadgroup]] ...)
{
Foo f = img_blk.read(lid);
float4 r = f.a;
f.a = r;
img_blk.write(f, lid);
}
Imageblock Slice Erişimi
Bu üye fonksiyonlar imageblock içindeki belirli bir veri üyesi için bir slice bölgesi elde eder.
const imageblock_slice<E, imageblock_layout_implicit> slice(ushort color_index) const;
const imageblock_slice<E, imageblock_layout_implicit>
slice(ushort color_index, ushort2 size) const;
Kopyalanan bölgenin başlangıç noktası (0,0) olur. size belirtilmezse imageblock’un tüm genişliği ve yüksekliği kopyalanır.
6.13.2 Açık Yerleşime Sahip Imageblock Fonksiyonları
Bu fonksiyonları belirli bir (x, y) koordinatı için imageblock verisine referans almak amacıyla kullanın.
Piksel Hızı Erişimi
threadgroup_imageblock T* data(ushort2 coord);
const threadgroup_imageblock T* data(ushort2 coord) const;
Örnek veya Renk Hızı Erişimi
threadgroup_imageblock T* data(
ushort2 coord,
ushort index,
imageblock_data_rate data_rate
);
const threadgroup_imageblock T* data(
ushort2 coord,
ushort index,
imageblock_data_rate data_rate
) const;
coord imageblock boyutlarının dışındaysa veya index geçersizse data() fonksiyonunun davranışı tanımsızdır.
Örnek veya renk hızında piksel saklayan bir imageblock için data(coord) çağrısı aşağıdakiyle eşdeğerdir:
data(coord, 0, imageblock_data_rate::sample);
Örnek
struct Foo {
rgba8unorm<half4> a;
int b;
};
kernel void
my_kernel(imageblock<Foo> img_blk,
ushort2 lid [[thread_position_in_threadgroup]] ...)
{
threadgroup_imageblock Foo* f = img_blk.data(lid);
half4 r = f->a;
f->a = r;
}
Veri Hızı Numaralandırması
enum class imageblock_data_rate {
color,
sample
};
Açık Yerleşim İçin Slice Erişimi
const imageblock_slice<E, imageblock_layout_explicit>
slice(const threadgroup_imageblock E& data_member) const;
const imageblock_slice<E, imageblock_layout_explicit>
slice(const threadgroup_imageblock E& data_member, ushort2 size) const;
Kopyalanan bölgenin başlangıç noktası (0,0) olur. size belirtilmezse tüm imageblock kopyalanır.
6.13.3 Bir Imageblock Slice’ını Bir Dokudaki Bölgeye Yazma
Bir imageblock_slice ile ilişkili pikselleri coord konumundan başlayarak bir dokuya yazmak için doku türlerindeki aşağıdaki write(...) üye fonksiyonlarını kullanın.
Aşağıdaki durumlarda yazma işlemi sınır dışı kabul edilir:
- Erişilen koordinatlar dokunun dışındaysa.
- Ayrıntı seviyesi (LOD) argümanı sınır dışındaysa.
imageblock_slice’ın herhangi bir bölümü dokunun dışına erişiyorsa.
Sınır dışı yazma işlemlerinin davranışı tanımsızdır. Yazma işlemi, imageblock slice ile hedef doku arasında eşleşen MSAA modları gerektirir.
1D Texture
void write(imageblock_slice<E, imageblock_layout_explicit> slice, uint coord, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice, ushort coord, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice, uint coord, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice, ushort coord, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice, uint coord, uint array, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice, ushort coord, ushort array, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice, uint coord, uint array, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice, ushort coord, ushort array, ushort lod = 0);
2D Texture
void write(imageblock_slice<E, imageblock_layout_explicit> slice, uint2 coord, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice, ushort2 coord, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice, uint2 coord, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice, ushort2 coord, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice, uint2 coord, uint array, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice, ushort2 coord, ushort array, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice, uint2 coord, uint array, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice, ushort2 coord, ushort array, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice,
uint2 coord, uint face, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice,
ushort2 coord, ushort face, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice,
uint2 coord, uint face, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice,
ushort2 coord, ushort face, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice,
uint2 coord, uint face, uint array, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice,
ushort2 coord, ushort face, ushort array, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice,
uint2 coord, uint face, uint array, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice,
ushort2 coord, ushort face, ushort array, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice,
uint3 coord, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_explicit> slice,
ushort3 coord, ushort lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice,
uint3 coord, uint lod = 0);
void write(imageblock_slice<E, imageblock_layout_implicit> slice,
ushort3 coord, ushort lod = 0);
Bir küp doku dizisi için:
Bir 3D doku için:
Örnek
struct Foo {
half4 a;
int b;
float c;
};
kernel void my_kernel(texture2d<half> src [[ texture(0) ]],
texture2d<half, access::write> dst [[ texture(1) ]],
imageblock<Foo> img_blk,
ushort2 lid [[ thread_position_in_threadgroup ]],
ushort2 gid [[ thread_position_in_grid ]])
{
// İş parçacığı kimliğini kullanarak giriş görüntüsünden pikseli oku.
half4 clr = src.read(gid);
// Image slice elde et.
threadgroup_imageblock Foo* f = img_blk.data(lid);
// Threadgroup içindeki iş parçacığı kimliğini kullanarak imageblock içine piksel yaz.
f->a = clr;
// Tüm iş parçacıklarının imageblock’a yazmayı bitirdiğinden emin olmak için bariyer.
// Bu durumda her iş parçacığı imageblock içindeki kendi konumuna yazdığı için
// bariyer gerekli değildir.
threadgroup_barrier(mem_flags::mem_threadgroup_imageblock);
// Imageblock içindeki pikselleri işle ve slice içindeki elemanları güncelle.
process_pixels_in_imageblock(img_blk, gid, lid);
// Tüm iş parçacıklarının imageblock içindeki elemanlara yazmayı
// tamamladığından emin olmak için bariyer.
threadgroup_barrier(mem_flags::mem_threadgroup_imageblock);
// Imageblock içindeki belirli bir elemanı çıkış görüntüsüne yaz.
// Threadgroup içindeki yalnızca bir iş parçacığı imageblock yazma işlemini gerçekleştirir.
if (lid.x == 0 && lid.y == 0)
dst.write(img_blk.slice(f->a), gid);
}
6.14 Pack ve Unpack Fonksiyonları
Bu bölüm, <metal_pack> başlık dosyasında tanımlanan ve vektör kayan noktalı verileri paketlenmiş bir tamsayı değerine dönüştürmek ve bu değerden geri dönüştürmek için kullanılan Metal fonksiyonlarını listeler. 8‑, 10‑ veya 16‑bit işaretli ya da işaretsiz tamsayı değerlerinden normalize edilmiş tek hassasiyetli veya yarım hassasiyetli kayan noktalı değerlere dönüşümün ve tersinin nasıl yapılacağı hakkında ayrıntılar için bölüm 8.7’nin alt bölümlerine bakın.
Tablo 6.23, tek bir işaretsiz tamsayıdan birden fazla değeri çıkaran ve ardından bunları bir vektörde saklanan kayan noktalı değerlere dönüştüren fonksiyonları listeler.
Tablo 6.23. Unpack fonksiyonları
32 bit işaretsiz bir tamsayıyı dört adet 8 bit işaretli veya işaretsiz tamsayıya aç ve her 8 bit değeri normalize edilmiş tek hassasiyetli veya yarım hassasiyetli kayan nokta değerine dönüştürerek 4 bileşenli bir vektör üret.
Yerleşik fonksiyonlar:
float4 unpack_unorm4x8_to_float(uint x)
float4 unpack_snorm4x8_to_float(uint x)
half4 unpack_unorm4x8_to_half(uint x)
half4 unpack_snorm4x8_to_half(uint x)
Yukarıdaki ile aynıdır, ancak r, g ve b renk değerleri sRGB'den doğrusal RGB'ye dönüştürülür.
float4 unpack_unorm4x8_srgb_to_float(uint x)
half4 unpack_unorm4x8_srgb_to_half(uint x)
32 bit işaretsiz bir tamsayıyı iki adet 16 bit işaretli veya işaretsiz tamsayıya aç ve her değeri normalize edilmiş bir kayan nokta değerine dönüştürerek 2 bileşenli bir vektör üret.
float2 unpack_unorm2x16_to_float(uint x)
float2 unpack_snorm2x16_to_float(uint x)
half2 unpack_unorm2x16_to_half(uint x)
half2 unpack_snorm2x16_to_half(uint x)
Paketlenmiş renk biçimlerini normalize edilmiş kayan nokta vektörlerine dönüştür.
float4 unpack_unorm10a2_to_float(uint x)
float3 unpack_unorm565_to_float(ushort x)
half4 unpack_unorm10a2_to_half(uint x)
half3 unpack_unorm565_to_half(ushort x)
float4 unpack_snorm10a2_to_float(uint x)
half4 unpack_snorm10a2_to_half(uint x)
Tüm işletim sistemleri: Metal 4 ve sonrası.
16 bit işaretsiz normalize edilmiş veya işaretli normalize edilmiş bir değerden yarım hassasiyetli kayan nokta değerine dönüştürürken, unpack_unorm2x16_to_half ve unpack_snorm2x16_to_half fonksiyonları hassasiyet kaybına neden olabilir.
6.14.2 Kayan Nokta Vektörünü Tamsayılara Dönüştürme, Ardından Tamsayıları Paketleme
Tablo 6.24, bir kayan nokta vektörüyle başlayan, bileşenleri tamsayı değerlere dönüştüren ve ardından çoklu değerleri tek bir işaretsiz tamsayı içinde paketleyen fonksiyonları listeler.
Tablo 6.24. Paketleme fonksiyonları
Dört bileşenli normalize edilmiş tek hassasiyetli veya yarım hassasiyetli kayan nokta vektörünü dört adet 8 bit tamsayıya dönüştür ve bunları 32 bit işaretsiz bir tamsayı içine paketle.
uint pack_float_to_unorm4x8(float4 x)
uint pack_float_to_snorm4x8(float4 x)
uint pack_half_to_unorm4x8(half4 x)
uint pack_half_to_snorm4x8(half4 x)
Dört bileşenli normalize edilmiş bir vektörü dört adet 8 bit tamsayıya dönüştür ve bunları 32 bit işaretsiz bir tamsayı içine paketle. Renk değerleri doğrusal RGB'den sRGB'ye dönüştürülür.
uint pack_float_to_srgb_unorm4x8(float4 x)
uint pack_half_to_srgb_unorm4x8(half4 x)
İki bileşenli normalize edilmiş kayan nokta vektörünü iki adet 16 bit tamsayıya dönüştür ve bunları 32 bit işaretsiz bir tamsayı içine paketle.
uint pack_float_to_unorm2x16(float2 x)
uint pack_float_to_snorm2x16(float2 x)
uint pack_half_to_unorm2x16(half2 x)
uint pack_half_to_snorm2x16(half2 x)
Normalize edilmiş kayan nokta vektörlerini paketlenmiş renk biçimlerine dönüştür.
uint pack_float_to_unorm10a2(float4)
ushort pack_float_to_unorm565(float3)
uint pack_half_to_unorm10a2(half4)
ushort pack_half_to_unorm565(half3)
uint pack_float_to_snorm10a2(float4)
uint pack_half_to_snorm10a2(half4)
Tüm işletim sistemleri: Metal 4 ve sonrası.
6.15 Atomik Fonksiyonlar
Metal programlama dili, C++17 atomik işlemleri ve senkronizasyon işlemlerinin bir alt kümesini uygular. Metal atomik fonksiyonları, bölüm 2.6'da açıklandığı gibi Metal atomik verileri üzerinde çalışmalıdır.
Atomik işlemler, bir iş parçacığında yapılan atamaların başka bir iş parçacığı tarafından görünür hâle getirilmesinde özel bir rol oynar. Bir veya daha fazla bellek konumu üzerindeki bir senkronizasyon işlemi ya bir acquire işlemi, bir release işlemi ya da her ikisidir. İlişkili bir bellek konumu olmayan bir senkronizasyon işlemi bir fence olarak adlandırılır ve bir acquire fence, bir release fence veya her ikisi olabilir. Buna ek olarak, senkronizasyon işlemi olmayan relaxed atomik işlemler de vardır.
Atomik türler üzerinde yalnızca birkaç işlem türü vardır, ancak bu türlerin birçok örneği bulunur. Bu bölüm her genel türü tanımlar.
Atomik fonksiyonlar <metal_atomic> başlık dosyasında tanımlanır.
6.15.1 Bellek Sırası
memory_order numaralandırması, ayrıntılı normal (atomik olmayan) bellek senkronizasyon işlemlerini belirtir (bkz. C++17 belirtiminin bölüm 29.3'ü) ve işlem sıralaması sağlayabilir:
enum memory_order {
memory_order_relaxed,
memory_order_seq_cst
};
atomic_thread_fence dışındaki atomik işlemler için memory_order_relaxed tek numaralandırma değeridir. memory_order_relaxed ile senkronizasyon veya sıralama kısıtlaması yoktur; işlem yalnızca atomiklik gerektirir. Bu işlemler belleği sıralamaz, ancak atomiklik ve değişiklik sırası tutarlılığı sağlar.
Relaxed bellek sıralamasının tipik bir kullanım alanı sayaçları güncellemektir; örneğin referans sayaçları. Bunun nedeni yalnızca atomiklik gerektirmesi, ancak ne sıralama ne de senkronizasyon gerektirmemesidir.
Metal 3.2 ve sonrasında, bir iş parçacığındaki store işleminden önce gerçekleşen her şeyin load işlemini gerçekleştiren iş parçacığında görünür bir yan etki hâline geldiğini belirtmek ve tüm etiketlenmiş atomik işlemler için tek bir toplam değişiklik sırası kurmak amacıyla atomic_thread_fence üzerinde memory_order_seq_cst kullanabilirsiniz.
6.15.2 İş Parçacığı Kapsamı
Tüm işletim sistemleri: Metal 3.2 ve sonrası Apple silicon için thread_scope desteği sağlar.
thread_scope numaralandırması, memory_order tarafından sağlanan bellek sırası kısıtlaması için bir iş parçacığı kümesini belirtir:
enum thread_scope {
thread_scope_thread,
thread_scope_simdgroup,
thread_scope_threadgroup,
thread_scope_device
};
Gayriresmî olarak, bir senkronizasyon işlemi üzerindeki iş parçacığı kapsamı, bu işlemin senkronize olabileceği veya kendisiyle senkronize olabilecek iş parçacıkları kümesini tanımlar. Bunu atomic_thread_fence ile kullanırsınız.
6.15.3 Fence Fonksiyonları
Tüm işletim sistemleri: Metal 3.2 ve sonrası Apple silicon için atomic_thread_fence desteği sağlar.
atomic_thread_fence, ilişkili bir atomik fonksiyon olmadan, bellek sırası ve iş parçacığı kapsamına göre atomik olmayan ve relaxed atomik erişimlerin bellek senkronizasyon sıralamasını kurar:
void atomic_thread_fence(mem_flags flags,
memory_order order,
thread_scope scope = thread_scope_device);
Bir fence aşağıdaki adres alanı kapsamlarında çalışır:
mem_flagsiçindemem_threadgroupvarsathreadgroupmem_flagsiçindemem_threadgroup_imageblockvarsathreadgroup_imageblockmem_flagsiçindemem_object_datavarsaobject_datamem_flagsiçindemem_devicevarsadevicemem_flagsiçindemem_texturevarsatexture
order değerine bağlı olarak işlem:
order == memory_order_relaxedise hiçbir etkiye sahip değildirorder == memory_order_seq_cstise sıralı tutarlı bir acquire ve release fence olur
Bir atomic_thread_fence, aynı memory_order değerine sahip bir atomik store işleminden farklı senkronizasyon kısıtlamaları uygular. Bir atomik store‑release işlemi, önceki tüm yazmaların store‑release sonrasına taşınmasını engeller ve memory_order_seq_cst sıralamasına sahip bir atomic_thread_fence, o kapsam içindeki sonraki tüm store işlemlerinin ötesine önceki tüm yazmaların taşınmasını engeller.
Buna ek olarak, atomik nesnelere yapılan erişimler iş parçacıkları arası senkronizasyon kurabilir ve memory_order tarafından belirtildiği şekilde atomik olmayan bellek erişimlerini sıralayabilir.
Aşağıdaki alt bölümlerde açıklanan atomik fonksiyonlarda:
- A, atomik türlerden birini ifade eder.
- M, aritmetik işlemler için diğer argümanın türünü ifade eder. Atomik tamsayı türleri için M, C ile aynıdır.
- C, karşılık gelen atomik olmayan türü ifade eder.
Her atomik fonksiyonun yalnızca bazı türleri destekleyebileceğini unutmayın. Aşağıdaki bölümler Metal'in hangi A türünü desteklediğini belirtir.
Tüm işletim sistemleri: Aksi belirtilmedikçe, adları explicit ile biten fonksiyonlar (atomic_store_explicit veya atomic_load_explicit gibi) Metal 1 ve sonrasında desteklenir. Metal 3, yalnızca aygıt belleği için atomic_float desteği sağlar.
iOS: Metal 2 ve sonrası atomic_store, atomic_load, atomic_exchange, atomic_compare_exchange_weak ve atomic_fetch_* fonksiyonlarını destekler.
6.15.4.1 Atomik Store Fonksiyonları
Bu fonksiyonlar object tarafından işaret edilen değeri atomik olarak desired ile değiştirir. Bu fonksiyonlar atomic_int, atomic_uint, atomic_bool ve atomic_float atomik türlerini destekler. Atomik store, atomic_float için yalnızca aygıt belleğini destekler.
Tüm işletim sistemleri: memory_order_relaxed ile atomic_store_explicit fonksiyonu desteği.
void atomic_store_explicit(threadgroup A* object,
C desired,
memory_order order); // Tüm işletim sistemleri: Metal 2'den beri
void atomic_store_explicit(volatile threadgroup A* object,
C desired,
memory_order order); // Tüm işletim sistemleri: Metal 1'den beri
void atomic_store_explicit(device A* object,
C desired,
memory_order order); // Tüm işletim sistemleri: Metal 2'den beri
void atomic_store_explicit(volatile device A* object,
C desired,
memory_order order); // Tüm işletim sistemleri: Metal 1'den beri
6.15.4.2 Atomik Load Fonksiyonları
Bu fonksiyonlar object tarafından işaret edilen değeri atomik olarak elde eder. atomic_int, atomic_uint, atomic_bool ve atomic_float atomik türlerini destekler. Atomik load, atomic_float için yalnızca aygıt belleğini destekler.
Tüm işletim sistemleri: memory_order_relaxed ile atomic_load_explicit fonksiyonu desteği.
C atomic_load_explicit(const threadgroup A* object,
memory_order order); // Tüm işletim sistemleri: Metal 2'den beri
C atomic_load_explicit(const volatile threadgroup A* object,
memory_order order); // Tüm işletim sistemleri: Metal 1'den beri
C atomic_load_explicit(const device A* object,
memory_order order); // Tüm işletim sistemleri: Metal 2'den beri
C atomic_load_explicit(const volatile device A* object,
memory_order order); // Tüm işletim sistemleri: Metal 1'den beri
6.15.4.3 Atomik Exchange Fonksiyonları
Bu fonksiyonlar object tarafından işaret edilen değeri atomik olarak desired ile değiştirir ve daha önce object tarafından tutulan değeri döndürür. Bu fonksiyonlar atomic_int, atomic_uint, atomic_bool ve atomic_float atomik türlerini destekler.
Tüm işletim sistemleri: memory_order_relaxed ile atomic_exchange_explicit fonksiyonu desteği.
C atomic_exchange_explicit(threadgroup A* object,
C desired,
memory_order order); // Tüm işletim sistemleri: Metal 2'den beri
C atomic_exchange_explicit(volatile threadgroup A* object,
C desired,
memory_order order); // Tüm işletim sistemleri: Metal 1'den beri
C atomic_exchange_explicit(device A* object,
C desired,
memory_order order); // Tüm işletim sistemleri: Metal 2'den beri
C atomic_exchange_explicit(volatile device A* object,
C desired,
memory_order order); // Tüm işletim sistemleri: Metal 1'den beri
Kopyalama işlemi std::memcpy benzer şekilde gerçekleştirilir.
6.15.4.4 Atomik Karşılaştır ve Değiştir Fonksiyonları
Bir compare‑and‑exchange fonksiyonunun etkisi aşağıdaki bölümde açıklanmıştır.
⚠️ Uyarı: Orijinal metnin bu kısmı (34. parça) API içerik filtresi veya bağlantı hatası nedeniyle işlenememiştir.
Metal 4 aşağıdaki yapı ve numaralandırmaları tanımlar:
void set_fragment_buffer(device T *buffer, uint index);
template <typename T ...>
void set_fragment_buffer(constant T *buffer, uint index);
void draw_primitives(
primitive_type type,
uint vertex_start,
uint vertex_count,
uint instance_count,
uint base_instance
);
// index_buffer temel alınarak aşırı yüklenmiş draw_indexed_primitives.
void draw_indexed_primitives(
primitive_type type,
uint index_count,
device ushort *index_buffer,
uint instance_count,
uint base_vertex,
uint base_instance
);
void draw_indexed_primitives(
primitive_type type,
uint index_count,
device uint *index_buffer,
uint instance_count,
uint base_vertex,
uint base_instance
);
void draw_indexed_primitives(
primitive_type type,
uint index_count,
constant ushort *index_buffer,
uint instance_count,
uint base_vertex,
uint base_instance
);
void draw_indexed_primitives(
primitive_type type,
uint index_count,
constant uint *index_buffer,
uint instance_count,
uint base_vertex,
uint base_instance
);
// patch_index_buffer ve tessellation_factor_buffer temel alınarak aşırı yüklenmiş draw_patches.
void draw_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
uint instance_count,
uint base_instance,
const device MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
uint instance_count,
uint base_instance,
const device MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
uint instance_count,
uint base_instance,
constant MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
uint instance_count,
uint base_instance,
constant MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
uint instance_count,
uint base_instance,
const device MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
uint instance_count,
uint base_instance,
const device MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
uint instance_count,
uint base_instance,
constant MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
uint instance_count,
uint base_instance,
constant MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
// patch_index_buffer,
// control_point_index_buffer ve tessellation_factor_buffer'a dayalı aşırı yüklenmiş draw_indexed_patches.
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
const device void *control_point_index_buffer,
uint instance_count,
uint base_instance,
const device MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
const device void *control_point_index_buffer,
uint instance_count,
uint base_instance,
const device MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
const device void *control_point_index_buffer,
uint instance_count,
uint base_instance,
constant MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
const device void *control_point_index_buffer,
uint instance_count,
uint base_instance,
constant MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
constant void *control_point_index_buffer,
uint instance_count,
uint base_instance,
const device MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
constant void *control_point_index_buffer,
uint instance_count,
uint base_instance,
const device MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
constant void *control_point_index_buffer,
uint instance_count,
uint base_instance,
constant MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
const device uint *patch_index_buffer,
constant void *control_point_index_buffer,
uint instance_count,
uint base_instance,
constant MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
const device void *control_point_index_buffer,
uint instance_count,
uint base_instance,
const device MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
const device void *control_point_index_buffer,
uint instance_count,
uint base_instance,
const device MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
const device void *control_point_index_buffer,
uint instance_count,
uint base_instance,
constant MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
const device void *control_point_index_buffer,
uint instance_count,
uint base_instance,
constant MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
constant void *control_point_index_buffer,
uint instance_count,
uint base_instance,
const device MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
constant void *control_point_index_buffer,
uint instance_count,
uint base_instance,
const device MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
constant void *control_point_index_buffer,
uint instance_count,
uint base_instance,
constant MTLQuadTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
void draw_indexed_patches(
uint number_of_patch_control_points,
uint patch_start,
uint patch_count,
constant uint *patch_index_buffer,
constant void *control_point_index_buffer,
uint instance_count,
uint base_instance,
constant MTLTriangleTessellationFactorsHalf *tessellation_factor_buffer,
uint instance_stride = 0
);
// Komutun tamamını sıfırla. reset() çağrısından sonra, ek
// değişiklikler yapılmazsa bu komutun yürütülmesi
// herhangi bir işlem gerçekleştirmez.
void reset();
// `source` komutunun içeriğini bu komuta kopyala.
void copy_command(render_command source);
};
command_buffer erişilirken Metal erişimin sınırlar içinde olup olmadığını kontrol etmez. Eğer erişim tamponun kapasitesinin ötesine geçerse davranış tanımsızdır.
render_command içinde sunulan yöntemler MTLIndirectRenderCommand arayüzünü yansıtır ve MTLRenderCommandEncoder ile benzerdir. MTLRenderCommandEncoder ile dikkat çekici farklar şunlardır:
render_commandiçindekidraw*yöntemlerine yapılan çağrılar, komutun gerçekleştireceği eylemleri kodlar. Birden fazla çağrı yapılırsa yalnızca sonuncusu etkili olur.- Tessellation argümanları doğrudan
render_command::draw_patchesverender_command::draw_indexed_patchesiçinde geçirilir. Diğer çağrılar tessellation argümanlarını ayarlamaz.
6.16.2 Dolaylı Komut Tamponlarında Compute Komutlarını Kodlama
ICB’ler, daha sonra GPU üzerinde yürütülmek üzere bir Metal tamponu içine dispatch komutlarının kodlanmasına izin verir.
Bir gölgelendirme dili fonksiyonunda, compute_command yapısına indeksli erişim sağlayan bir Metal tampon nesnesi içine ICB komutlarını kodlamak için command_buffer türünü kullanın:
iOS: Metal 2.2 ve sonrası compute komutları için dolaylı komut tamponlarını destekler.
macOS: Metal 2.3 ve sonrası compute komutları için dolaylı komut tamponlarını destekler.
[[kernel]] void producer(device arguments &args,
ushort cmd_idx [[thread_position_in_grid]])
{
compute_command cmd(args.cmd_buffer, cmd_idx);
...
}
compute_command her türlü dispatch komut türünü kodlayabilir. compute_command için aşağıdaki genel arayüz <metal_command_buffer> başlık dosyasında tanımlıdır. compute_pipeline_state türü compute pipeline durumlarını temsil eder ve yalnızca argüman tamponları aracılığıyla shader’lara geçirilebilir. Bir argüman tamponu içinde pipeline durumu skaler olarak veya bir dizi içinde geçirilebilir.
struct compute_command {
public:
explicit compute_command(command_buffer icb, unsigned cmd_index);
void set_compute_pipeline_state(compute_pipeline_state pipeline);
template <typename T ...>
void set_kernel_buffer(device T *buffer, uint index);
template <typename T ...>
void set_kernel_buffer(constant T *buffer, uint index);
// Metal 3.1: kernel stride değerlerinin geçirilmesini destekler.
template <typename T ...>
void set_kernel_buffer(device T *buffer, size_t stride, uint index);
template <typename T ...>
void set_kernel_buffer(constant T *buffer, size_t stride, uint index);
void set_barrier();
void clear_barrier();
void concurrent_dispatch_threadgroups(
uint3 threadgroups_per_grid,
uint3 threads_per_threadgroup
);
void concurrent_dispatch_threads(
uint3 threads_per_grid,
uint3 threads_per_threadgroup
);
void set_threadgroup_memory_length(uint length, uint index);
void set_stage_in_region(uint3 origin, uint3 size);
// Komutun tamamını sıfırla. reset() çağrısından sonra, ek
// değişiklik yapılmazsa bu komutun yürütülmesi
// herhangi bir işlem gerçekleştirmez.
void reset();
// `source` komutunun içeriğini bu komuta kopyala.
void copy_command(compute_command source);
};
command_buffer erişilirken Metal erişimin sınırlar içinde olup olmadığını kontrol etmez. Eğer erişim tamponun kapasitesinin ötesine geçerse davranış tanımsızdır.
compute_command içinde sunulan yöntemler MTLIndirectComputeCommand arayüzünü yansıtır ve MTLComputeCommandEncoder ile benzerdir.
Bir ICB içinde dispatch işlemleri her zaman eşzamanlıdır. compute_command içindeki concurrent_dispatch* yöntemlerine yapılan çağrılar, komutun gerçekleştireceği eylemleri kodlar. Birden fazla çağrı yapılırsa yalnızca sonuncusu etkili olur.
Uygulama gerekli yerlere bariyerler yerleştirmekten sorumludur. Bir ICB içinde kodlanan bariyerler üst (parent) encoder’ı etkilemez.
CPU, command_buffer bir shader’a argüman olarak geçirilmeden önce bu tampon içindeki tekil komutları başlatmış olabilir. Eğer CPU bir komutu önceden başlatmamışsa, o komutu kullanmadan önce sıfırlamanız gerekir.
6.16.3 Bir Dolaylı Komut Tamponunun Komutlarını Kopyalama
Bir komut yapısını (render_command veya compute_command) operator= ile kopyalamak komutun içeriğini kopyalamaz; yalnızca hedef komutun kaynak komut ile aynı tamponu ve indeksi göstermesini sağlar. Komutun içeriğini kopyalamak için Bölüm 6.16.1 ve 6.16.2’de listelenen copy_command fonksiyonlarını çağırın.
Kopyalama yalnızca uyumlu komut tamponlarını gösteren komutlar arasında desteklenir. İki komut tamponu yalnızca ICB tanımlayıcıları (MTLIndirectCommandBufferDescriptor nesneleri) eşleşiyorsa uyumludur. Komutların kendileri de tamponlar içinde geçerli indekslere başvurmalıdır.
Aşağıdaki örnek, copy_command kullanarak bir render komutunun içeriğinin cmd0’dan cmd1’e nasıl kopyalanacağını gösterir:
struct arguments {
command_buffer cmd_buffer;
render_pipeline_state pipeline_state_0;
render_pipeline_state pipeline_state_1;
};
[[kernel]] void producer(device arguments &args) {
render_command cmd0(args.cmd_buffer, 0);
render_command cmd1(args.cmd_buffer, 1);
cmd0.set_render_pipeline_state(args.pipeline_state_0);
// 1 indeksindeki komutun 0 indeksindeki komutu göstermesini sağla.
cmd1 = cmd0;
}
// Tampondaki 0 indeksindeki komut için pipeline durumunu değiştir.
cmd1.set_render_pipeline_state(args.pipeline_state_0);
// Tampondaki 1 indeksindeki komut henüz değiştirilmedi.
cmd1 = render_command(args.cmd_buffer, 1);
// 0 indeksindeki komutun içeriğini 1 indeksindeki komuta kopyala.
cmd1.copy_command(cmd0);
Değişken rasterization oranı (VRR), ekran konumuna bağlı olarak fragment shader çağrılma oranını azaltarak yüksek çözünürlüklü işleme sırasında gölgelendirme maliyetini düşürebilir. VRR özellikle Augmented Reality (AR) / Virtual Reality (VR) uygulamalarında çevresel bilgilerin aşırı örneklenmesini önlemek için yararlıdır.
Bir gölgelendirme dili fonksiyonunda VRR’yi desteklemek için katman başına rasterization oranı verisinin eşlemesini tanımlamak üzere rasterization_rate_map_decoder yapısını kullanın. Her katman ekran uzayında minimum kalite değerleri içerir ve farklı fiziksel fragment uzayı boyutlarına sahip olabilir. AR/VR için bu kalite değerleri lens dönüşümü veya göz izleme bilgisine dayanır.
struct rasterization_rate_map_data;
struct rasterization_rate_map_decoder {
explicit rasterization_rate_map_decoder(
constant rasterization_rate_map_data &data) thread;
float2 map_screen_to_physical_coordinates(
float2 screen_coordinates,
uint layer_index = 0) const thread;
uint2 map_screen_to_physical_coordinates(
uint2 screen_coordinates,
uint layer_index = 0) const thread;
float2 map_physical_to_screen_coordinates(
float2 physical_coordinates,
uint layer_index = 0) const thread;
uint2 map_physical_to_screen_coordinates(
uint2 physical_coordinates,
uint layer_index = 0) const thread;
};
VRR haritası ekran uzayı ile fiziksel fragment uzayı arasındaki eşlemeyi tanımlar ve işleme sonuçlarının istenen ekran çözünürlüğüne geri dönüştürülmesini sağlar.
Shader içinde ekran uzayı ile fiziksel fragment uzayı arasında dönüşüm yapmak için uygulama, rasterization_rate_map_decoder yapısındaki dönüşüm fonksiyonlarından herhangi birini kullanmadan önce tamponu harita verisiyle doldurmak amacıyla MTLRasterizationRateMap sınıfının copyParameterDataToBuffer:offset: yöntemini çağırmalıdır. copyParameterDataToBuffer:offset: yöntemi tarafından dışa aktarılan veriye işaret eden bir işaretçi dışında başka bir şey geçirilmesi tanımsız davranışa yol açar.
6.17 Değişken Rasterization Oranı
iOS: Metal 2.2 ve sonrası değişken rasterization oranını ve rasterization oranı haritasını destekler.
macOS: Metal 2.3 ve sonrası değişken rasterization oranını ve rasterization oranı haritasını destekler.
Aşağıdaki örnek, uygulamanın rasterization_rate_map_data verisini shader bağlama noktasında rasterization_rate_map_decoder yapısının kurucusuna nasıl geçirmesi gerektiğini gösterir:
[[fragment]] float4 fragment_shader(/* diğer argümanlar */
constant rasterization_rate_map_data &data [[buffer(0)]]) {
float2 screen_coords = ...;
rasterization_rate_map_decoder map(data);
float2 physical_coords =
map.map_screen_to_physical_coordinates(screen_coords);
...
}
Alternatif olarak uygulama, derlenmiş verinin saklandığı ofseti hesaplayabilir ve geçerli bir rasterization_rate_map_data oluşturmak için açık bir cast veya işaretçi aritmetiği kullanabilir. rasterization_rate_map_data eksik bir tür olduğundan, üzerinde bazı işlemler doğası gereği yasaktır (örneğin işaretçi türü üzerinde işaretçi aritmetiği veya sizeof).
6.18 Ray‑Tracing Fonksiyonları
Tüm işletim sistemleri: Metal 2.3 ve sonrası ray‑tracing fonksiyonlarını destekler.
Metal, ışın izleme işlevlerini ve türlerini <metal_raytracing> içinde metal::raytracing ad alanında tanımlar. Metal 2.3 ve sonrası sürümler, aşağıda belirtilen durumlar dışında bu işlevleri yalnızca bir compute fonksiyonunda (kernel fonksiyonu) destekler. Metal 2.4 ve sonrası sürümler, bunlar için vertex, fragment ve tile fonksiyonlarında ek destek sunar.
Metal 2.3 ve sonrası sürümlerde, bir hızlandırma yapısının (bkz. bölüm 2.17.7) null olup olmadığını kontrol etmek için aşağıdaki fonksiyonlardan birini çağırabilirsiniz.
Metal 2.4 ve sonrası sürümlerde, bir hızlandırma yapısının null olup olmadığını kontrol etmek için aşağıdaki fonksiyonu çağırabilirsiniz:
bool is_null_acceleration_structure(
acceleration_structure<intersection_tags...>
);
Metal 3.1 ve sonrası sürümlerde, bir instance hızlandırma yapısı tarafından başvurulan hızlandırma yapısı üzerinde aşağıdaki fonksiyonları kullanarak yineleme yapabilirsiniz.
Bir instance hızlandırma yapısındaki instance sayısını sorgulamak için aşağıdaki fonksiyonu çağırın:
uint get_instance_count() const;
Bir instance hızlandırma yapısı içinde bulunan bir instance tarafından başvurulan hızlandırma yapısını elde etmek için aşağıdaki fonksiyonu çağırın. Döndürülen tür, bölüm 2.17.7'de tanımlanan şablonlaştırılmış türdür.
template <typename... intersection_tags>
acceleration_structure<intersection_tags...>
get_acceleration_structure(uint instance_id);
Olası hızlandırma yapısı türleri şunlardır:
acceleration_structure<instancing>acceleration_structure<instancing, instance_motion>acceleration_structure<instancing, primitive_motion>acceleration_structure<instancing, primitive_motion, instance_motion>
Bildirilmiş dönüş türü, bir instance hızlandırma yapısı içinde bulunan instance tarafından başvurulan hızlandırma yapısı türü ile eşleşmezse sonuçlar tanımsızdır.
Instance ve/veya primitive motion etiketlerini kullanmayan instance hızlandırma yapıları, bu etiketleri içeren bir hızlandırma yapısı türü olarak döndürülebilir. Örneğin, herhangi bir hareket (instance veya primitive) içermeyen bir instance hızlandırma yapısı, bu etiketleri içeren bir tür olarak döndürülebilir.
Bu yetenek, birden fazla hızlandırma yapısı türüyle çalışırken her etiket kümesi için ayrı bir intersector sağlama gerekliliğini ortadan kaldırmanıza olanak tanır; ancak bu durum, bu etiketlere ihtiyaç duymayan bir hızlandırma yapısının dolaşılmasının getirebileceği olası performans maliyetiyle birlikte gelir.
intersector<intersection_tags...> nesnesini oluşturduktan sonra (bkz. bölüm 2.17.6), intersection_tags değerine bağlı olarak aşağıdaki intersect fonksiyonlarından birini çağırabilirsiniz.
using result_type = intersection_result<intersection_tags...>;
Tüm intersect fonksiyonlarının ray ve accel_struct parametrelerine sahip olması gerekir. Diğer parametreler isteğe bağlıdır.
Parametreler
- ray — Işın özellikleri.
- accel_struct —
acceleration_structure<intersection_tags...>türünde hızlandırma yapısı. - mask — Metal API
MTLAccelerationStructureInstanceDescriptoriçinde tanımlanan instance maskesi ile AND işlemine tabi tutulacak kesişim maskesi. Maskeleri örtüşmeyen instance'lar atlanır. - time — Işınla ilişkili zaman. Bu parametre,
intersection_tagsiçindeprimitive_motionveyainstance_motionbulunduğunda mevcuttur. - func_table —
intersection_function_table<intersection_tags...>türünde kesişim fonksiyon tablosu. - payload — Referansla geçirilen kullanıcı payload nesnesi.
intersect()çağrıldığında payloadray_dataadres alanına kopyalanır ve kesişim fonksiyonuna iletilir. Sonuç, kesişim fonksiyonundan çıkışta (bölüm 5.1.6) kopyalanır ve payload nesnesi güncellenir.
Eğer intersection_tags içinde intersection_function_buffer bulunuyorsa, isteğe bağlı olarak intersection_function_buffer_arguments türünde bir nesne iletebilirsiniz (bkz. bölüm 6.18.8). ifba.intersection_function_buffer, çağrının SIMD grubu içinde uniform olmalıdır.
Eğer intersection_tags içinde user_data bulunuyorsa, kesişim fonksiyonu için kullanıcı verisini işaret eden bir buffer isteğe bağlı olarak iletilebilir. Bir buffer iletirseniz ifba da iletmeniz gerekir.
Intersect Fonksiyonları (Instancing Yok)
Yalnızca intersection_tags içinde instancing bulunmadığında kullanılabilir.
result_type intersect(
ray ray,
primitive_acceleration_structure accel_struct) const;
result_type intersect(
ray ray,
primitive_acceleration_structure accel_struct,
intersection_function_table<intersection_tags...> func_table) const;
template <typename T>
result_type intersect(
ray ray,
primitive_acceleration_structure accel_struct,
intersection_function_table<intersection_tags...> func_table,
thread T &payload) const;
Intersect Fonksiyonları (Instancing)
Yalnızca intersection_tags içinde instancing bulunduğunda kullanılabilir.
result_type intersect(
ray ray,
instance_acceleration_structure accel_struct,
uint mask = ~0U) const;
result_type intersect(
ray ray,
instance_acceleration_structure accel_struct,
intersection_function_table<intersection_tags...> func_table) const;
Etiketlerin hareket (primitive_motion veya instance_motion) içerip içermemesine ve intersection_function_buffer kullanılıp kullanılmamasına bağlı olarak ek varyantlar vardır.
Örnekler:
template <typename T>
result_type intersect(
ray ray,
instance_acceleration_structure accel_struct,
intersection_function_table<intersection_tags...> func_table,
thread T &payload) const;
result_type intersect(
ray ray,
instance_acceleration_structure accel_struct,
uint mask,
intersection_function_table<intersection_tags...> func_table) const;
template <typename T>
result_type intersect(
ray ray,
instance_acceleration_structure accel_struct,
uint mask,
intersection_function_table<intersection_tags...> func_table,
thread T &payload) const;
Zamana Dayalı Kesişim (Motion)
intersection_tags içinde primitive_motion veya instance_motion bulunduğunda kullanılabilir.
template <typename T, intersection_tags...>
result_type intersect(
ray ray,
acceleration_structure<intersection_tags...> accel_struct,
float time) const;
Ek aşırı yüklemeler maskeler, fonksiyon tabloları ve payload'lar için izin verir.
Callback Tabanlı Kesişim (Metal 3.2+)
Metal 3.2 ve sonrası sürümlerde, bir kopyalamadan kaçınmak ve kesişimin belleğine doğrudan erişmek için intersection_result_ref<intersection_tags...> ve bir callback içinde ray_data payload işaretçisi kullanılabilir.
template <typename Callable>
void intersect(..., Callable callback);
template <typename Payload, typename Callable>
void intersect(...,
const thread Payload &payload_in,
Callable callback);
intersection_result_ref ve ray_data payload işaretçisinin yaşam süresi yalnızca callback süresi boyunca geçerlidir. Bunlardan herhangi birini saklamak ve intersect() çağrısı tamamlandıktan sonra kullanmak tanımsız davranışa yol açar.
Callback gövdesi içinde özyineli ışın izleme gerçekleştirilemez. Callback tamamlandıktan sonra shader tekrar ışın kesişimi gerçekleştirebilir.
Örnek:
[[kernel]] void trace_rays_with_payload(...) {
intersector<instancing, max_levels<2>, triangle_data> i;
i.intersect(ray, acceleration_structure, MyPayload{},
[&](intersection_result_ref<instancing, max_levels<2>, triangle_data> result,
const ray_data MyPayload &final_payload)
{
result.get_primitive_id();
// ...
}
);
}
Intersection Function Buffer (Metal 4+)
Metal 4 ve sonraki sürümlerde, intersection_tags içinde intersection_function_buffer bulunduğunda ek intersect aşırı yüklemeleri mevcuttur.
Örnekler:
result_type intersect(
ray ray,
acceleration_structure<> accel_struct,
intersection_function_buffer_ifba) const;
template <typename T>
result_type intersect(
ray ray,
acceleration_structure<> accel_struct,
intersection_function_buffer_ifba,
thread T &payload) const;
Instancing varyantları:
result_type intersect(
ray ray,
acceleration_structure<instancing> accel_struct,
intersection_function_buffer_ifba) const;
template <typename T>
result_type intersect(
ray ray,
acceleration_structure<instancing> accel_struct,
intersection_function_buffer_ifba,
thread T &payload) const;
Maske varyantları:
result_type intersect(
ray ray,
uint mask,
acceleration_structure<instancing> accel_struct,
intersection_function_buffer_ifba) const;
template <typename T>
result_type intersect(
ray ray,
uint mask,
acceleration_structure<instancing> accel_struct,
intersection_function_buffer_ifba,
thread T &payload) const;
Motion Varyantları
Primitive motion:
result_type intersect(
ray ray,
acceleration_structure<instancing, primitive_motion> as,
float time,
intersection_function_buffer_ifba) const;
template <typename T>
result_type intersect(
ray ray,
acceleration_structure<instancing, primitive_motion> as,
float time,
intersection_function_buffer_ifba,
thread T &payload) const;
Instance motion:
result_type intersect(
ray ray,
acceleration_structure<instancing, instance_motion> as,
float time,
intersection_function_buffer_ifba) const;
template <typename T>
result_type intersect(
ray ray,
acceleration_structure<instancing, instance_motion> as,
float time,
intersection_function_buffer_ifba,
thread T &payload) const;
Intersect Fonksiyonları (Metal 4 ve Sonrası)
Aşağıdaki intersect fonksiyon kümesi yalnızca intersection_tags içinde intersection_function_buffer, user_data bulunduğunda ve instancing bulunmadığında kullanılabilir:
result_type intersect(
ray ray,
acceleration_structure<> accel_struct,
intersection_function_buffer_ifba,
const device void *user_data) const;
template <typename T>
result_type intersect(
ray ray,
acceleration_structure<> accel_struct,
intersection_function_buffer_ifba,
const device void *user_data,
thread T &payload) const;
result_type intersect(
ray ray,
acceleration_structure<instancing> accel_struct,
intersection_function_buffer_ifba,
const device void *user_data) const;
template <typename T>
result_type intersect(
ray ray,
acceleration_structure<instancing> accel_struct,
intersection_function_buffer_ifba,
const device void *user_data,
thread T &payload) const;
result_type intersect(
ray ray,
uint mask,
acceleration_structure<instancing> accel_struct,
intersection_function_buffer_ifba,
const device void *user_data) const;
template <typename T>
result_type intersect(
ray ray,
uint mask,
acceleration_structure<instancing> accel_struct,
intersection_function_buffer_ifba,
const device void *user_data,
thread T &payload) const;
Aşağıdaki intersect fonksiyon kümesi yalnızca intersection_tags içinde intersection_function_buffer, user_data ve instancing bulunduğunda kullanılabilir.
primitive_motion ile
Yalnızca intersection_tags içinde intersection_function_buffer, user_data, instancing ve primitive_motion bulunduğunda kullanılabilir:
result_type intersect(
ray ray,
acceleration_structure<instancing, primitive_motion> as,
float time,
intersection_function_buffer_ifba,
const device void *user_data) const;
template <typename T>
result_type intersect(
ray ray,
acceleration_structure<instancing, primitive_motion> as,
float time,
intersection_function_buffer_ifba,
const device void *user_data,
thread T &payload) const;
result_type intersect(
ray ray,
uint mask,
float time,
acceleration_structure<instancing, primitive_motion> as,
intersection_function_buffer_ifba,
const device void *user_data) const;
template <typename T>
result_type intersect(
ray ray,
uint mask,
float time,
acceleration_structure<instancing, primitive_motion> as,
intersection_function_buffer_ifba,
const device void *user_data,
thread T &payload) const;
instance_motion ile
Yalnızca intersection_tags içinde intersection_function_buffer, instancing, user_data ve instance_motion bulunduğunda kullanılabilir:
result_type intersect(
ray ray,
acceleration_structure<instancing, instance_motion> as,
float time,
intersection_function_buffer_ifba,
const device void *user_data) const;
template <typename T>
result_type intersect(
ray ray,
acceleration_structure<instancing, instance_motion> as,
float time,
intersection_function_buffer_ifba,
const device void *user_data,
thread T &payload) const;
result_type intersect(
ray ray,
uint mask,
float time,
acceleration_structure<instancing, instance_motion> as,
intersection_function_buffer_ifba,
const device void *user_data) const;
template <typename T>
result_type intersect(
ray ray,
uint mask,
float time,
acceleration_structure<instancing, instance_motion> as,
intersection_function_buffer_ifba,
const device void *user_data,
thread T &payload) const;
6.18.3 Dolaşım Davranışını Kontrol Etmek İçin Intersector Fonksiyonları
enum class winding {
clockwise,
counterclockwise
};
void set_triangle_front_facing_winding(winding w);
enum class triangle_cull_mode {
none,
front,
back
};
Tüm OS: Metal 3.1 eğriler için destek ekler.
Varsayılan dolaşım davranışını geçersiz kılmak için intersector<intersection_tags...> nesnesinin aşağıdaki üye fonksiyonlarını kullanın.
Dolaşım Davranışını Kontrol Etme Fonksiyonları
void set_triangle_front_facing_winding(winding)void set_geometry_cull_mode(geometry_cull_mode)void set_opacity_cull_mode(opacity_cull_mode)void force_opacity(forced_opacity)void assume_geometry_type(geometry_type)void assume_identity_transforms(bool)void accept_any_intersection(bool)
Üçgen Yüzleri
Üçgenlerin iki tarafı veya yüzü vardır. Ön yüze bakan winding, ışın başlangıç noktasından bakıldığında hangi üçgen yüzünün ön yüz olarak kabul edileceğini belirler.
Köşeler ışın başlangıç noktasından bakıldığında saat yönünde sıralanmış görünüyorsa ve ön yüze bakan winding clockwise ise görünür yüz ön yüzdür. Diğer yüz arka yüzdür.
Ön yüze bakan winding counterclockwise ise bunun tersi geçerlidir.
Varsayılan winding değerini (clockwise) değiştirmek için:
void set_triangle_cull_mode(triangle_cull_mode tcm);
- Cull modu
frontise, ışın başlangıç noktasından bakıldığında ön yüzü görünür olan üçgenler kesişim için dikkate alınmaz. - Cull modu
backise, ışın başlangıç noktasından bakıldığında arka yüzü görünür olan üçgenler kesişim için dikkate alınmaz.
Geometri Ayıklama (Geometry Culling)
Aşağıdaki fonksiyon, intersector'u aday geometri kümesindeki tüm bounding box veya üçgen primitive'lerini ayıklayacak şekilde ayarlar.
Varsayılan geometri ayıklama modu: none.
enum class geometry_cull_mode {
none,
triangle,
bounding_box,
curve // Metal 3.1 ve sonrası
};
void set_geometry_cull_mode(geometry_cull_mode gcm);
Opaklık Denetimi
Varsayılan opaklık eleme modu: none.
enum class opacity_cull_mode {
none,
opaque,
non_opaque
};
void set_opacity_cull_mode(opacity_cull_mode ocm);
Örnek başına ve geometri başına opaklık ayarlarını geçersiz kılma:
enum class forced_opacity {
none,
opaque,
non_opaque
};
void force_opacity(forced_opacity fo);
Opaklık Davranışı
Üçgen ilkel öğeler opaklığa göre elenebilir:
- Opak üçgen: herhangi bir kesişim fonksiyonunu çalıştırmaz.
- Opak olmayan üçgen: çarpışmayı kabul etmek veya reddetmek için kesişim fonksiyonunu çalıştırır.
PrimitiveAccelerationStructure, üçgen opaklığını şu bayrakla kodlar:
MTLAccelerationStructureGeometryFlagOpaque
Opaklık şu şekilde geçersiz kılınabilir:
intersector.force_opacity()
Opaklık öncelik sırası:
intersector.force_opacity()- Örnek başına bayraklar
MTLAccelerationStructureInstanceFlagOpaqueMTLAccelerationStructureInstanceFlagNonOpaque
- Geometri başına opaklık
Özel bounding box ilkel öğeleri için opaklık şu şekilde değerlendirilir:
intersector.set_opacity_cull_mode()- Örnek bayrakları
- Geometri bayrakları
opaque parametresi, bounding box kesişim programına çözülmüş opaklık durumunu bildirir.
Kısıtlamalar
Eğer intersector.force_opacity() değeri opaque veya non_opaque olarak ayarlanmışsa:
intersector.set_opacity_cull_mode()noneolmak zorundadır.
Opaklık geçersiz kılma ve opaklık eleme birlikte kullanılamaz. Geçersiz kombinasyonlar tanımsız davranış üretir.
Geometri Türü Varsayımları
enum class geometry_type {
none,
triangle,
bounding_box,
curve, // Metal 3.1 ve sonrası
all
};
void assume_geometry_type(geometry_type gt);
Varsayılan geometri varsayımı:
geometry_type::triangle | geometry_type::bounding_box
Metal, performansı artırmak için varsayılan olarak hızlandırma yapılarının eğri geometrisi içermediğini varsayar.
Eğri kesişimini etkinleştirmek için assume_geometry_type() fonksiyonunu geometry_type::curve ile çağırın.
Kimlik Dönüşümü Varsayımı
void assume_identity_transforms(bool);
- Varsayılan:
false
Herhangi Bir Kesişimi Kabul Et
void accept_any_intersection(bool);
true olarak ayarlandığında, kesiştirici ilk kesişmeden hemen sonra geri döner.
Tipik kullanım durumları:
- Gölge ışınları
- Ortam perdeleme (ambient occlusion)
Eğri Geçişi İpuçları (Metal 3.1 ve Sonrası)
Belirli eğri türleri için geçişi optimize etmek amacıyla intersector ve intersection_query nesnelerine ipuçları verilebilir.
enum class curve_basis {
bspline,
catmull_rom,
linear,
bezier,
all
};
// Bitmask değildir
enum class curve_type {
round,
flat,
all
};
Fonksiyonlar aşağıdakilerin belirtilmesine izin verir:
- eğri tabanı
- eğri türü
- kontrol noktası sayısı
Varsayılan değerler:
curve_basis::allcurve_type::all- kontrol noktaları:
0(geçerli herhangi bir sayı)
Geçerli kontrol noktası sayıları: eğri tabanına bağlı olarak 2, 3 veya 4.
Intersector Durum Denetimleri (Metal 4 ve Sonrası)
intersection_tags, intersection_function_buffer içerdiğinde her intersector nesnesi için ışın katkısı ve geometri çarpanı belirtilebilir.
Taban indeks ve geometri çarpanı hesaplamalarında alt 4 bit kullanılır.
void set_base_id(uint index);
void set_geometry_multiplier(uint multiplier);
Varsayılan multiplier değeri: 1.
6.18.5 Kesişim Sorgu Fonksiyonları
Geçişi başlatmak ve geçişe özgü bilgileri sorgulamak için, bir varsayılan olmayan kurucu kullanarak intersection_query nesnesi oluşturun veya önce reset(...) çağırın.
Bu sıra izlenmezse davranış tanımsızdır.
Geçiş sırasında, next() fonksiyonu true döndürdükten sonra fonksiyonlar çağrılabilir.
Bazı ilkel öğeler önce aday kesişimler olarak görünür ve açıkça onaylanmaları gerekir.
Onay fonksiyonları:
commit_triangle_intersection()commit_bounding_box_intersection()commit_curve_intersection()
Opak üçgenler kesiştiklerinde otomatik olarak onaylanır.
Temel Kesişim Sorgu Fonksiyonları
void reset(...)bool next()intersection_type get_candidate_intersection_type()intersection_type get_committed_intersection_type()void commit_curve_intersection()(Metal 3.1 ve sonrası)
Örnek sorgu fonksiyonları:
uint get_candidate_instance_count()uint get_candidate_instance_id(uint depth)uint get_committed_instance_count()uint get_committed_instance_id(uint depth)
Ek Sorgu Denetimleri
void abort();
void commit_triangle_intersection();
void commit_bounding_box_intersection(float distance);
Işın Değeri Sorgu Fonksiyonları
float3 get_world_space_ray_origin();
float3 get_world_space_ray_direction();
float get_ray_min_distance();
intersection_params get_intersection_params();
Aday Kesişim Değer Fonksiyonları
float get_candidate_triangle_distance();
uint get_candidate_instance_id();
uint get_candidate_user_instance_id();
uint get_candidate_geometry_id();
uint get_candidate_primitive_id();
float2 get_candidate_triangle_barycentric_coord();
float4x3 get_candidate_object_to_world_transform();
bool is_candidate_non_opaque_bounding_box();
bool is_candidate_triangle_front_facing();
float3 get_candidate_ray_origin();
float3 get_candidate_ray_direction();
Onaylanmış Kesişim Değer Fonksiyonları
float get_committed_distance();
uint get_committed_instance_id();
uint get_committed_user_instance_id();
uint get_committed_geometry_id();
uint get_committed_primitive_id();
float2 get_committed_triangle_barycentric_coord();
float4x3 get_committed_object_to_world_transform();
float4x3 get_committed_world_to_object_transform();
float get_candidate_curve_parameter(); // Metal 3.1+
float get_committed_curve_parameter(); // Metal 3.1+
bool is_committed_triangle_front_facing();
float3 get_committed_ray_origin();
float3 get_committed_ray_direction();
const device void *get_committed_primitive_data();
Tüm işletim sistemleri: Metal 3 ve sonrası.
Ek Sorgu Fonksiyonları (Metal 3.1+, max_levels<Count>)
Bu fonksiyonlar geçiş bilgilerini sorgulamaya izin verir:
- Aday üçgen çarpışmasının mesafesi
- Şu anda onaylanmış çarpışmanın mesafesi
float get_committed_distance();
Örnek ve geometri sorguları:
uint get_candidate_instance_id();
uint get_candidate_user_instance_id();
uint get_candidate_geometry_id();
uint get_candidate_primitive_id();
uint get_committed_instance_id();
Ek onaylanmış sorgular şunları içerir:
- onaylanmış kullanıcı örnek kimliği
- onaylanmış geometri kimliği
- onaylanmış ilkel kimliği
Nesne uzayı ışın sorguları:
aday çarpışma için ışın başlangıcı
aday çarpışma için ışın yönü
Mevcut onaylanmış çarpışma için nesne uzayındaki ışın başlangıcını sorgulamak üzere aşağıdaki fonksiyonu çağırın:
Mevcut onaylanmış çarpışma için nesne uzayındaki ışın yönünü sorgulamak üzere aşağıdaki fonksiyonu çağırın:
Mevcut çarpışma adayının ışın başlangıcı/yönünü nesne uzayından dünya uzayına dönüştüren matrisi sorgulamak için aşağıdaki fonksiyonu çağırın:
float4x3 get_candidate_object_to_world_transform();
- Mevcut aday çarpışmanın ışın başlangıcı/yönünü dünya uzayından nesne uzayına dönüştüren matrisi sorgulamak için aşağıdaki fonksiyonu çağırın:
uint get_committed_geometry_id();
uint get_committed_primitive_id();
float3 get_candidate_ray_origin();
float3 get_candidate_ray_direction();
float3 get_committed_ray_origin();
float3 get_committed_ray_direction();
float4x3 get_committed_object_to_world_transform();
Tepe noktası öznitelikleri v0, v1 ve v2 için, belirtilen barycentric noktadaki değer şöyledir:
v1 * barycentric_coord.x +
v2 * barycentric_coord.y +
v0 * (1.0f - (barycentric_coord.x + barycentric_coord.y))
Mevcut onaylanmış çarpışmanın ışın başlangıcı/yönünü nesne uzayından dünya uzayına dönüştüren matrisi sorgulamak için aşağıdaki fonksiyonu çağırın.
Mevcut onaylanmış çarpışmanın ışın başlangıcı/yönünü dünya uzayından nesne uzayına dönüştüren matrisi sorgulamak için aşağıdaki fonksiyonu çağırın:
float4x3 get_committed_world_to_object_transform();
- Aday çarpışma konumunun barycentric koordinatlarını sorgulamak için aşağıdaki fonksiyonu çağırın.
get_candidate_intersection_type()triangledöndürdüğünde geçerlidir:
float2 get_candidate_triangle_barycentric_coord();
Onaylanmış çarpışma konumunun barycentric koordinatlarını sorgulamak için aşağıdaki fonksiyonu çağırın.
get_committed_intersection_type()triangledöndürdüğünde geçerlidir.Çarpışan üçgen adayının ön yüz mü arka yüz mü olduğunu sorgulamak için aşağıdaki fonksiyonu çağırın. Ön yüz ise
true, arka yüz isefalsedöndürür.get_candidate_intersection_type()triangledöndürdüğünde geçerlidir.Onaylanmış çarpışmanın ön yüz mü arka yüz mü olduğunu sorgulamak için aşağıdaki fonksiyonu çağırın. Ön yüz ise
true, arka yüz isefalsedöndürür.get_committed_intersection_type()triangledöndürdüğünde geçerlidir.Mevcut aday ilkel öğe için ilkel başına veriyi sorgulamak üzere aşağıdaki fonksiyonu çağırın:
const device void *get_candidate_primitive_data();
float get_committed_curve_parameter();
Aday kesişimdeki örnek sayısını sorgulamak için aşağıdaki fonksiyonu çağırın.
Onaylanmış kesişimdeki örnek sayısını sorgulamak için aşağıdaki fonksiyonu çağırın.
Mevcut onaylanmış çarpışma için ilkel başına veriyi sorgulamak üzere aşağıdaki fonksiyonu çağırın:
const device void *get_committed_primitive_data();
Metal 3.1 ve sonrasında, aşağıdaki iki fonksiyon get_candidate_intersection_type() curve döndürdüğünde ve kesişim etiketi curve_data içerdiğinde çağrılabilir:
- Mevcut aday eğri için eğri parametresini sorgulamak üzere aşağıdaki fonksiyonu çağırın:
float get_candidate_curve_parameter();
- Mevcut onaylanmış kesişim için eğri parametresini sorgulamak üzere aşağıdaki fonksiyonu çağırın.
get_candidate_intersection_type()curvedöndürdüğünde geçerlidir.
Metal 3.1 ve sonrasında, bu bölümdeki geri kalan fonksiyonlar kesişim etiketi max_levels<Count> içerdiğinde çağrılabilir:
uint get_candidate_instance_count();
- Aday kesişimde seviye derinliğindeki örnek kimliğini sorgulamak için aşağıdaki fonksiyonu çağırın.
uint get_candidate_instance_id(uint depth);
- Aday kesişimde seviye derinliğindeki kullanıcı örnek kimliğini sorgulamak için aşağıdaki fonksiyonu çağırın:
uint get_candidate_user_instance_id(uint depth);
- Onaylanmış kesişimde seviye derinliğindeki örnek kimliğini sorgulamak için aşağıdaki fonksiyonu çağırın:
uint get_committed_instance_id(uint depth);
- Onaylanmış kesişimde seviye derinliğindeki kullanıcı örnek kimliğini sorgulamak için aşağıdaki fonksiyonu çağırın.
Dolaylı Instance Tanımlayıcıları (Metal 3.1+)
Metal 3.1 ve sonrasında dolaylı instance tanımlayıcılarını GPU üzerinden doldurabilirsiniz. Metal aşağıdaki tür tanımlarını sağlar:
enum MTLAccelerationStructureInstanceOptions : uint
{
MTLAccelerationStructureInstanceOptionNone = 0,
MTLAccelerationStructureInstanceOptionDisableTriangleCulling = (1 << 0),
MTLAccelerationStructureInstanceOptionTriangleFrontFacingWindingCounterClockwise = (1 << 1),
MTLAccelerationStructureInstanceOptionOpaque = (1 << 2),
MTLAccelerationStructureInstanceOptionNonOpaque = (1 << 3),
};
typedef packed_float3 MTLPackedFloat3;
typedef packed_float3 MTLPackedFloat4x3[4];
struct MTLAccelerationStructureInstanceDescriptor
{
MTLPackedFloat4x3 transformationMatrix;
MTLAccelerationStructureInstanceOptions options;
uint mask;
uint intersectionFunctionTableOffset;
uint accelerationStructureIndex;
};
struct MTLAccelerationStructureUserIDInstanceDescriptor
{
MTLPackedFloat4x3 transformationMatrix;
MTLAccelerationStructureInstanceOptions options;
uint mask;
uint intersectionFunctionTableOffset;
uint accelerationStructureIndex;
uint userID;
};
Tanımlayıcıyı doldurmayı kolaylaştırmak için Metal, acceleration_structure<intersection_tags…> türünden MTLResourceID türüne örtük bir dönüşüm sağlar.
acceleration_structure<primitive_motion> primitiveAStruct = …;
MTLResourceID resource_id = primitiveAStruct;
6.18.7 Eğri Yardımcı Fonksiyonları
Metal 3.1 ve sonrasında, aşağıdaki başlık dosyasında tanımlanan bir dizi eğri yardımcı fonksiyonu sağlanır:
<metal_curves>
Kısaltmalar
- Ps
floatveyahalfanlamına gelir. - P,
Pstüründe bir skaler veya vektördür. EğerPsfloatise,Pfloat4olur.
Fonksiyonlar, eğri parametresi t ve kontrol noktaları p0, p1 vb. verildiğinde bir eğri üzerindeki konumu veya birinci ya da ikinci türevi döndürür.
Desteklenen eğri türleri şunlardır:
- Quadratic Bézier
- Cubic Bézier
- Quadratic B‑Spline
- Cubic B‑Spline
- Cubic Hermite
- Catmull‑Rom
Fonksiyon İmzaları
P bezier(Ps t, P p0, P p1, P p2)
P bezier_derivative(Ps t, P p0, P p1, P p2)
P bezier_second_derivative(Ps t, P p0, P p1, P p2)
P bezier(Ps t, P p0, P p1, P p2, P p3)
P bezier_derivative(Ps t, P p0, P p1, P p2, P p3)
P bezier_second_derivative(Ps t, P p0, P p1, P p2, P p3)
P bspline(Ps t, P p0, P p1, P p2)
P bspline_derivative(Ps t, P p0, P p1, P p2)
P bspline_second_derivative(Ps t, P p0, P p1, P p2)
P bspline(Ps t, P p0, P p1, P p2, P p3)
P bspline_derivative(Ps t, P p0, P p1, P p2, P p3)
P bspline_second_derivative(Ps t, P p0, P p1, P p2, P p3)
P hermite(Ps t, P p0, P p1, P m0, P m1)
P hermite_derivative(Ps t, P p0, P p1, P m0, P m1)
P hermite_second_derivative(Ps t, P p0, P p1, P m0, P m1)
P catmull_rom(Ps t, P p0, P p1, P p2, P p3)
P catmull_rom_derivative(Ps t, P p0, P p1, P p2, P p3)
P catmull_rom_second_derivative(Ps t, P p0, P p1, P p2, P p3)
Dolaylı Fonksiyon Arabellekleri (Metal 4+)
Metal 4 ve sonrasında, bir sahnedeki geometriyi hızlandırma yapısındaki bu geometri üzerinde çalışan shader kümesiyle ilişkilendirmek için dolaylı fonksiyon arabelleklerini kullanabilirsiniz.
Kullanıcı, intersection_function_buffer_arguments içeren bir arabellek sağlar.
struct intersection_function_buffer_arguments
{
// 8 bayta hizalanmış talimat fonksiyon işleyicilerini içeren arabellek.
const device void *intersection_function_buffer;
// Bayt cinsinden maksimum aralık
size_t intersection_function_buffer_size;
// Kesişim fonksiyon girişleri arasındaki adım
size_t intersection_function_stride;
};
intersection_function_stride, 8 bayt artışlarla [0, 4096] aralığını destekler.
Kolaylık sağlamak amacıyla başlık dosyası, intersection_function_buffer_arguments türüne dönüştürülebilen Metal türü MTLIntersectionFunctionBufferArguments sağlar.
6.19 Günlükleme Fonksiyonları
Tüm işletim sistemleri: Metal 3.2 ve sonrası Apple silicon için günlükleme desteği sağlar.
Metal, günlükleme fonksiyonlarını ve türlerini şu başlık dosyasında tanımlar:
<metal_logging>
Günlüklemeyi etkinleştirmek için şu ayarı yapın:
-fmetal-enable-logging
Günlük Türleri
enum log_type
{
log_type_debug,
log_type_info,
log_type_default,
log_type_error,
log_type_fault
};
os_log Yapısı
struct os_log
{
os_log(constant char *subsystem, constant char *category) constant;
void log_with_type(log_type type, constant char *format, ...) constant;
void log_debug(constant char *format, ...) constant;
void log_info(constant char *format, ...) constant;
void log(constant char *format, ...) constant;
void log_error(constant char *format, ...) constant;
void log_fault(constant char *format, ...) constant;
};
Notlar
std::printfbiçim belirteçlerinin çoğunu destekler.%@,%.*Pveya özel CPUos_logbelirteçlerini desteklemez.- Vektörler şu biçim kullanılarak yazdırılabilir:
%v[num_elements][length_modifier][conversion_specifier]
Örnekler:
float4→%v4hlfuchar2→%v2hhu
Diğer kurallar:
- Varsayılan argüman yükseltmesi
halfiçin uygulanır ve bu türdoubletürüne yükseltilir. - Yükseltme vektörlere uygulanmaz.
- Metal, vektörleri yazdırırken
intvefloatgibi 4 baytlık türler içinhluzunluk belirtecini destekler. - Biçim dizgesi bir string literal olmalıdır.
Örnek
constant metal::os_log custom_log("com.custom_log.subsystem", "custom category");
void test_log(float x) {
if (x < M.PI.F)
custom_log.log("custom message %f", x);
}
void test_log(float x) {
if (x < M.PI.F)
os_log_default.log("custom message %f", x);
}
Shader'lardan gelen mesajlar, boyutu MTLLogState tarafından belirlenen bir günlük tamponuna yerleştirilir. Bir komut tamponu içindeki tüm draw/dispatch çağrıları aynı günlük tamponunu paylaşır. Sistem, komut tamponu tamamlanana kadar mesajları kaldırmaz.
Günlük tamponu dolarsa, CPU tamponu boşaltana kadar sonraki mesajlar atılır.
CPU tarafından okunan mesajlar, ilgili alt sistem, kategori ve seviye ile birlikte birleşik günlükleme sistemine gönderilir. os_log_default tarafından gönderilen mesajlar, alt sistem ve kategori nil olacak şekilde karşılık gelen seviyeyi kullanır.
Shader günlükleme mesajlarının özel olarak işlenmesi için Metal API addLogHandler bölümüne bakın.
BELİRTİMDE KOMŞULARI
BELİRTİMİN TAMAMI
- 1. Giriş2.921
- 2. Veri Türleri10.929
- 3. Operatörler1.638
- 4. Adres Alanları1.217
- 5. Fonksiyon ve Değişken Bildirimleri14.679
- 6. Metal Standart Kütüphanesi18.632
- 7. Metal Performans Primitifleri2.111
- 8. Sayısal Uyumluluk2.554
- 9. Ek273
BU SAYFA HAKKINDA
- Çeviri hatası bildire-posta