Kendin yaz
Hiç yayınlamadığımız op
Diyelim ki kayan pencere medyanına ihtiyacınız var. Önünüzde tek boyutlu bir seri: yaklaşık bir milyon örnek, arada tek örneklik sivri uçlar. Ortalama filtresi orada tam ters işi yapar — sivri ucu ortadan kaldırmak yerine komşularına bulaştırır. Sağlam, sıkıcı ve doğru araç, kayan bir pencere üzerinden medyandır. Catalog'u açıp arıyorsunuz — hiçbir şey yok. Öyle bir yerleşik hiç olmamış.
Adım adım örnek
Bu yazı bir op'u boş arama kutusundan yayımlanmış içeriğe kadar, platformun bunu yapmak üzere kurulduğu biçimde yürütüyor. Scellis henüz kullanıcılara açılmadı (bugün neredeyiz); dolayısıyla burayı, birinin yaşadığı bir oturumun raporu olarak değil, uçtan uca savunulmuş tasarım olarak okuyun. Çekirdek gerçek WGSL'dir, tolerans muhakemesi denetleyicinin fiilî kuralıdır ve aşağıdaki raporlar düzeneğin bastığı çıktının biçimini gösterir — iddia ettiğimiz bir koşunun çıktısını değil.
Reddetmeye değer kısayol#
Bakımcı bunu bakımcı yolundan düzeltebilir. Bir tanım dosyası, bir referans uygulama, yerleşik kümede bir yer — akşama kalmaz windowed_median yerleşik olur ve eksik olduğunu kimse asla bilmez. Ama bu hiçbir şey kanıtlamaz. Scellis'in genişletilebilirlik üzerine en derin iddiası şu: Bakımcı yalnızca üç grameri elinde tutar — şekil-imza dili, tampon bağlama tanımlayıcısı, geriye doğru gösterim — artı denetleyiciyi ve referans rejimini; her örneği kullanıcılar yazar. Böyle iddialar, dışarıdan sınanmadıkça çürür. O hâlde op'u herhangi bir kullanıcının yazacağı gibi yazın: Studio'daki yazarlık formlarından, içerik olarak, kendi cihazınızda — motor yaması yok, bakımcı anahtarı yok. Yazarlık rehberi tam olarak bu yolu anlatıyor.
Önce imza: op ne vaat ediyor#
Burada bir op, bir fonksiyon gövdesi değildir; gövdesi iliştirilmiş bir sözleşmedir. O yüzden önce imza gelir: girdi, n uzunluğunda 1-D f32 bir tensör; tek parametre w, tek sayı olması zorunlu ve 31 ile sınırlı bir pencere genişliği; çıktı, n − w + 1 uzunluğunda. Çıktı şekli sınıfı static — yalnızca şekillerden ve parametrelerden hesaplanabilir — böylece planlayıcı, daha hiçbir şey çalışmadan belleği bütçeleyebilir. Sonra işin en keyifli yeri: tolerans sınıfı. exact bildirirsiniz — ve bunu gerekçelendirmeniz istenir. Tek sayılı pencereyle medyan, değerler üzerinde hiç aritmetik yapmaz: karşılaştırır ve seçer; çıktı, girdilerden biridir — bit bit aynı. Denetleyici, bildirilen sınıfın op'un anlamsal kategorisiyle örtüşmesini şart koşar: Bir matris çarpımı, uygunluktan kaçmak için gevşek tolerans ilan edemez — aynı kuralla, saf seçim de dürüstçe kesinlik iddia edebilir. Çift sayılı bir pencere iki orta elemanın ortalamasını alırdı — bir yuvarlama adımı — ve farklı bildirmek zorunda kalırdı; imzayı tek sayılı pencerelere kısıtlamak temiz iddiayı korur.
Referans, sözleşmenin ta kendisi#
Kullanıcı hesaplaması için kural, yerleşiklerinkiyle aynı: bir referansa karşı denetlenir, asla öz beyanla yetinilmez. Bu yüzden referans uygulama çekirdekten önce gelir — her çıktı indeksi için: w değeri topla, sırala, ortadakini al — ve motorun bütün sayısal doğruluğu bağladığı CPU referans yolunda çalışır. Çekirdeğiniz siz öyle dediğiniz için doğru değildir; bildirdiğiniz tolerans içinde, seçme şansınızın olmadığı durumlarda o referansı tutturduğu ölçüde doğrudur. Üstelik bu eşleşmenin değeri sizin oturumunuzu aşar: referans uygulama ve uygunluk korpusu serbest lisanslı takas halkasında durur — “referansını tutturuyor” iddiasını isteyen herkes kendisi yeniden türetebilir. Bir satıcının “bana güvenin”i değil.
Çekirdek — ve testin yakaladığı şey#
WGSL'in kendisi gösterişsizdir; bunu bir özellik saymak gerekir. Çıktı elemanı başına bir çağrı; pencereyi sabit boyutlu yerel bir diziye topla (31 sınırı bunu dürüst tutan şey); insertion sort; ortadakini al. Bağlama düzeni — iki depolama tamponu ve bir uniform — platformun sahip olduğu tampon bağlama gramerinde yazılıdır. Süslü parantezlerin içindeki her şeyse sizindir.
// user/kemal/windowed_median@1 -- odd window w <= 31, one thread per output
@group(0) @binding(0) var<storage, read> x : array<f32>;
@group(0) @binding(1) var<storage, read_write> dst : array<f32>;
struct Params { n: u32, w: u32 }
@group(0) @binding(2) var<uniform> p: Params;
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
let i = gid.x;
if (i >= p.n - p.w + 1u) { return; } // the "+ 1u" is the fix the suite forces
var win: array<f32, 31>;
for (var k = 0u; k < p.w; k++) { win[k] = x[i + k]; }
// insertion sort: pure selection -- no arithmetic ever touches the values
for (var a = 1u; a < p.w; a++) {
let v = win[a];
var b = a;
while (b > 0u && win[b - 1u] > v) { win[b] = win[b - 1u]; b--; }
win[b] = v;
}
dst[i] = win[p.w / 2u];
}Sonra uygunluk. Test durumlarını takım kendisi üretir — şekiller, pencereler, uç girdiler — çünkü kendi sınavına not veren yazar, dolambaçlı bir öz beyandan ibarettir. Klasik ilk-taslak hatasını yakalayan durum da egzotik değildir: en küçük geçerli girdi. n = 7 ve w = 7 iken tam olarak bir çıktı elemanı vardır; koşulu i >= n − w + 1 yerine i >= n − w yazarsanız çekirdek hiçbir şey yazmaz. Yukarıdaki + 1u, o düzeltmenin ta kendisi. Çekirdek, üretilen her durumda referansı tutturduğunda — exact'in anlamı “yaklaşık” değil, “eşit”tir — motor bir uygunluk makbuzu üretir: çekirdeğe hash ile bağlı ve kanıtın koştuğu GPU profiliyle kapsamlı. Bu kapsam, işin dürüst tarafıdır: Güven, kanıtı çalıştıran zarfta geçerlidir; başka bir makine sizinkini ithal etmek yerine yerelde yeniden uygunluk koşusu yapar. Ve bunların hiçbiri sizinle kendi GPU'nuzun arasına girmez — uygunluk kazanmamış bir çekirdek, ilk derlemeden itibaren kendi cihazınızda çalışır. Makbuz, op'un yolculuk etmesini sağlayan şeydir.
Altgradyanlar, dürüstçe#
Şimdi sivri uç temizliğini boru hattının içine koyun; üstelik öncesinde eğitilebilir parçalar var — yani gradyanların medyandan geçmesi gerekiyor. Bunun için bir merdiven var ve basamakları dürüst. Bir op, mevcut türevlenebilir op'ların bileşimiyse, autodiff geriye doğru hesabı bedavaya türetir — ve dürüst olalım: Kayan pencere medyanı böyle de kurulabilir; pencereleri aç, her birini medyanla indir. Ama o açma işlemi n çarpı w boyutunda bir ara sonuç üretir — w katı bellek — ve bu op'un varlık sebebi zaten düz bir bütçe. Dolayısıyla: kendine ait bir çekirdek ve merdivenin orta basamağı — bildirilmiş bir geriye doğru hesap; denetleyici onu, sizin kendi ileri hesabınıza karşı sonlu farklarla doğrular. En alt basamaksa dürüst sınırdır: Op'u türevlenemez ilan edersiniz ve motor, -mış gibi yapmak yerine onun içinden eğitmeyi reddeder.
VJP, max-pooling'in yakın akrabasıdır: Medyan girdilerden biri olduğundan, gelen kotanjant argmedyan indeksine yönlenir; başka her yerde sıfırdır. Sıra değişimleri arasında fonksiyon parçalı doğrusaldır ve bu gradyan kesindir — aşağı yukarı değil, kesin. Eşitlik noktasındaysa fonksiyon hiç türevlenemez ve kaydın bunu açıkça söylemesi gerekir: Bildirdiğiniz şey bir altgradyan uzlaşımıdır — seçilen elemana yönlendir — ve op'un üstverisinde yazılıdır, bir yorum satırına gömülü değil. Gradcheck aynı dürüstlüğü kendine de uygular. Adım boyu içinde bir sıra değişimine denk gelen sonda, türevin var olmadığı bir kırılmanın üzerine oturur; düzenek o sondayı, kırışığın üzerinden ortalama almak yerine dışarıda bırakır. Raporu şu biçimdedir:
gradcheck user/kemal/windowed_median@1 (w = 7, reference path, f64)
forward : author reference (gather -> sort -> middle)
vjp : declared backward -- cotangent routed to the argmedian index
probes : sampled (input, output) pairs, central differences, h = 1e-3
skipped : probes within h of an order crossing (no derivative exists there)
bound : max |analytic - numeric| must stay under the declared tolerance
negative control : a corrupted vjp must be REJECTED -- a check that cannot
fail certifies nothingBu dökümde iki satır yük taşır. Parçalı doğrusal bir fonksiyonda merkezî farkların kesme hatası yoktur; kırılmalardan uzakta analitik ve sayısal gradyan, f64 yuvarlamasına kadar örtüşür — bu örtüşme şans değil, fonksiyonun biçimidir. Negatif kontrol de süs değildir: Düzenek, VJP'yi kasıtlı olarak bozup denetimin kalmasını şart koşar — çünkü kalamayan bir denetim hiçbir şey belgelemez.
Dürüst sınır
Sonlu farklarla gradcheck güçlü bir gerek koşuldur, biçimsel bir kanıt değildir. Fonksiyonu örnekler; sembolik olarak doğrulamaz. Platform bunu, denetimin göründüğü her yerde açıkça söyler — denetimi anlatan bir yazı da söylemeli.
Yayınlandı — ve ayırt edilemez#
Yayınlamak en az dramatik adımdır; bütün tasarımın amacı da bu. Op'u bir çalışma alanıyla paylaştığınızda, önemli olan tek kapıyı zaten geçmiş olur: Hiçbir şey, SPDX lisansı iliştirilmeden private'ten çıkmaz — çünkü paylaştığınız şey, bir meslektaşınızın çatallayabileceği şeydir ve lisanssız paylaşılan iş, hukuken bütün hakları saklı sayılır; bu da platformun üzerine kurulduğu çatallama çarkını sessizce geçersiz kılardı. Kamuya açmak bunun üstüne fazlasını ister: eksiksiz üstveri çıtası — yerleşik içeriğin de aştığı çıtanın aynısı. Kart, kökenini özür dilemeden taşır: kaynak user, güven hak edilene dek unverified, exact tolerans rozeti ve makbuz.
Sonra op'u iş akışında iki yerleşik adımın arasına koyup çalıştırın. İşte can alıcı nokta — ve bu bir söz sanatı değil: Çalıştırma anında motor, bir kullanıcının op'unu bir yerleşikten ayırt edemez. “Adil davranır” değil — ayırt edemez. Tek arama, tek doğrulama, tek yürütme yolu; koşunun köken kaydı, yazılan op'u içerik hash'iyle, iki yanındaki yerleşik op'ları sabitlediği gibi sabitler. Sivri uçlar gitti. Boru hattı, gradyanlarıyla birlikte medyanın içinden eğitiliyor.
Hiç kimsenin yayınlamadığı op yine de çıkar — yazılmış, referansa karşı denetlenmiş, makbuzu elinde, bir kullanıcı tarafından yayımlanmış. Bu ilk kullanıcının bakımcı olması ise motorun yapısal olarak fark edemeyeceği bir ayrıntıdır. Bütün mesele zaten bu.