WMMA Parçacık API'si

Tensor Core'ları doğrudan programlamak için nvcuda::wmma ad alanındaki warp matris çarpma-birikim API'si olan WMMA'yı kullanırsınız. 🧩

WMMA warp genelinde çalışır: bir warp içindeki 32 iş parçacığının tümü tek bir karo üzerinde birlikte çalışır. Tek tek iş parçacıklarıyla değil, karolarla düşünürsünüz.

Parça, bir warp'ın matris karosundaki bölümünü tutan WMMA veri türüdür. Her iş parçacığı, bu parçanın birkaç elemanına görünmeden sahip olur.

matrix_a, matrix_b veya accumulator etiketiyle parçalar bildirirsiniz. Etiket, WMMA'ya karonun çarpma-birikime nasıl katılacağını söyler.

Parçalar 16'ya 16'ya 16 gibi belirli karo boyutlarını kullanır. Desteklenen bir biçim seçersiniz; donanım yalnızca bu birleşimleri kabul eder.

Önce load_matrix_sync çağrısıyla bellekteki bir karoyu parçaya alırsınız. Bu yükleme, verileri sizin için warp'a dağıtır.

Toplama işleminden önce fill_fragment ile sonuç karosunu sıfırlarsınız. Temiz bir biriktirici, toplamlarınızın bilinen bir değerden başlamasını sağlar.

Temel çağrı mma_sync'tir. Yüklenen parçalar üzerinde D = A çarpı B artı C işlemini doğrudan Tensor Core'ları kullanarak yürütür.

Son olarak store_matrix_sync, biriktirici parçasını belleğe geri yazar. Bu saklama, her iş parçacığının parçalarını tek bir karoda birleştirir.

_sync son eki bir hatırlatmadır: her WMMA çağrısı warp eşzamanlıdır. 32 hattın tümü ona birlikte ulaşmalıdır; aksi durumda davranış tanımsızdır.

Yükleme ve saklama işlemleri, satırlar arasındaki adım olan başlangıç boyutuna ve belleği doğru okumak için satır veya sütun öncelikli bir düzene ihtiyaç duyar.

WMMA akışını izlediniz: bir parça bildirin, karoları yükleyin, biriktiriciyi temizleyin, mma_sync çağrısını yapın ve ardından saklayın. Her adım warp eşzamanlıdır. 🙌