Geliştirici yapay zeka çağında, GPU’ları maksimum potansiyeli ile kullanmak, daha iyi modeller eğitmek ve geniş ölçekli kullanıcı hizmetleri sunmak için kritik öneme sahiptir. Sıklıkla bu modeller, ince modifikasyonlar nedeniyle hazır kütüphane işlemleri ile ifade edilemeyen katmanlara sahip olabilir. Derin öğrenme derleyicileri genellikle dağıtımın uygulanabilirliğini sağlamak adına son birkaç optimizasyon aşamasından feragat eder.
NVIDIA CUDA geliştiricilerine DL ve HPC çekirdeklerinin performansını en üst düzeye çıkarmak için gereken güç ve kontrolü sağlamak amacıyla 2017’den bu yana CUTLASS üzerinde çalışıyoruz.
Artık yeni Python arayüzü ile geliştirmelerin bir sonraki aşamasına geçiyor. CUTLASS 4.0 ile CUTLASS 3.x yeniden tasarımında tanıtılan temel soyutlamalar, Python’da doğrudan erişilebilir hale geliyor. Bu yazıda, CUTLASS 3.x’in arkasındaki tasarım ilkelerini, temel arka uç kütüphanesini, CUDA Tensörlerini ve Mekansal Mikro Çekirdekleri (CuTe) aktardık ve CuTe’nin ana özelliklerini kullanarak yapılan optimizasyon örneklerine değindik.
CUTLASS 3.x Döneminde Öne Çıkanlar
CUTLASS 3, CuTe adı verilen, işleri daha anlamlı hale getiren yeni bir kütüphane sunmuştur. Bu kütüphane, birim ve bileşenleri tanımlamak için bir düzenleme konsepti üzerinden çalışmaktadır. CuTe’nin düzeni, thread-veri organizasyonunu oldukça basit hale getirir. Geliştiricilere daha iyi bir anlayış ve statik olarak kontrol edilebilir bir yapı sunan CuTe, aynı zamanda CUTLASS 2.x ile karşılaştırıldığında yüksek performansı ve Tensor Core işlem kapsamını korur.
CUTLASS 3, bu daha anlamlı yaklaşımın ötesinde, bütün versiyonların hedeflediği amaçları paylaşmaktadır: CUDA geliştiricilerine, en son donanım özellikleri etrafında yüksek performanslı GPU lineer cebir çekirdekleri yazmaları için sezgisel bir programlama modeli geliştirmek. Bu yeni büyük iterasyonda, aşağıdakilere vurgu yapıldı:
- Geliştirici verimliliği ve bileşenlerin temiz bir şekilde ayrılması için kütüphane tasarımının her katmanını özelleştirme yeteneği.
- Kernel yapılarının doğruluğunu sağlamaya yönelik derleme zamanında kontroller. Bu, eğer derlenirse, doğru çalışacaktır garantisi verir ve aksi takdirde yapılabilir statik hata mesajları sağlar.
- Daha az tanımlı tür ile daha az API yüzey alanı azaltımı ve özelleştirme kıvrım noktalarıyla daha düz bir öğrenme eğrisi.
- NVIDIA Hopper H100 ve NVIDIA Blackwell B200 üzerinde büyük performans, WGMMA (Hopper için) veya UMMA (Blackwell için), Tensor Bellek Hızlandırıcısı (Hopper için) TMA ve threadblock kümeleri gibi özelliklerden yararlanarak.
CuTe Kütüphanesi
CUTLASS 3.x’in kalbinde, CuTe bulunmaktadır. CuTe, thread ve veri tensörlerini tanımlamak ve işlemek için yeni bir kütüphanedir. CuTe, güçlü bir düzenleme temsili ve bu düzenlemeler üzerinde işlem yapan bir bilgi işlem cebiri içerir.
CuTe’nin düzen temsili, doğal olarak hiyerarşiktir, statik ve dinamik bilgileri destekler ve çok boyutlu tensörleri temsil etmek için kullanılır. Aynı düzen temsili, veri tensörleri ve thread tensörlerini tanımlamak için de kullanılır. Bu ortak terminolojinin kullanımı, CuTe Düzeni kavramının geniş uygulanabilirliğini gösterir.
Bu temsil yeteneğinin üzerine inşa edilen CuTe, kullanıcıların basit bilinen düzenlerden karmaşık düzenler oluşturmalarını veya bir düzeni diğerinin üzerinde bölmelerini sağlayan bir düzenler cebiri sunar. Tüketiciler, algoritmalarının mantıksal tanımlarına odaklanabilirken, CuTe mekanik kaydını onların yerine yapar. Bu araçlarla, kullanıcılar yoğun lineer cebir algoritmalarını hızlıca tasarlayıp, uygulayıp ve değiştirebilirler.
CuTe Düzeni ve Tensörleri
Daha fazla CuTe belgesi ve düzenler ile tensörler hakkında bilgiye belirtilen doküman dizininde erişebilirsiniz.
CuTe, veri tipinin, şeklinin, bellek alanının ve düzeninin derli toplu bir şekilde paketlendiği Layout ve Tensor nesneleri sağlar. Bu nesneler, kullanıcı için karmaşık indekslemeleri gerçekleştirir.
Layout<Shape, Stride>,Shapeiçindeki mantıksal koordinatlar ileStrideile hesaplanan indeksler arasında bir harita sağlar.Shape, bir veya daha fazla koordinat boşluğunu tanımlar ve bunlar arasında haritalama yapılır.Stride, koordinatların indekslere dönüştürülmesi için indeks haritasını tanımlar.
Tensor<Engine, Layout>birLayoutile bir iteratoru bir araya getirir. Bu iterator, küresel bellek, paylaşılan bellek, kayıt belleği veya rastgele erişim ve derecelendirme sağlayan başka herhangi bir şeyin işaretçisi olabilir.
CuTe’deki düzenlerin hiyerarşik olması ve tensor cebirinde tensor işlemlerinin katlanması ilhamı aldığı dikkat çekicidir. Şekil, hiyerarşik Shape ve Stride düzenlerini gösterir ve bu, basit satır-bazlı ve sütun-bazlı düzenlerin çok ötesine geçmektedir. Aynı zamanda, hiyerarşik düzenler halen normal tensörler gibi erişilebilir, böylece daha gelişmiş veri düzenleri algoritmik geliştirme aşamasında soyutlanmış olur.
CuTe Düzeni ile Dönüştürme ve Bölme
CuTe Düzenleri, temel bir işlem olarak işlevsel bileşimi destekler. İşlevsel bileşim, bir düzenin şekli ve sırasını dönüştürmek için kullanılabilir. Eğer elimizde (m,n) koordinatına sahip bir veri düzeni varsa ve bu düzen ile (thread_idx,value_idx) koordinatlarını kullanmak istiyorsak, veri düzenini bir düzen ile bileştiririz. Bu durum, her iş parçacığına ve değere erişmemizi kolaylaştırır.
Örneğin, bir 4×8 veri düzenini ele alalım. Bu düzeni her bir koordinat için iş parçacıkları ve değerler atamak istiyoruz. Bu durumu derleyen bir “TV düzeni” yazarak veri düzeninin bir düzeni ile yapıyoruz.
Yukarıda belirtilen bileşimin sonucunda, her iş parçacığının değerleri sonuçtaki her satıra düzenlenir. Bu yapıyı daha iyi anlamak için, TV düzeninin tersini incelediğimizde daha sezgisel bir görünüm elde ederiz.
Belirtilen düzende, 4×8 veri düzenindeki her koordinatın iş parçacığına ve değere nasıl harita göre düzenlendiği gösterilmektedir. İstediğimiz bölme düzeni ile yazılıp, istenilen veri düzenlerine uygulanabilir.
CuTe Düzen Cebiri hakkında daha fazla belgelere GitHub üzerinde erişebilirsiniz.
CuTe Matriks Çarpma-Toplama Atomları
Bir atom, donanım hızlandırılmış bir matematiksel işlem veya kopyalama işlemine birlikte katılması gereken en küçük thread ve veri topluluğudur.
Bir atom, bir PTX talimatını ve bu talimata katılması gereken thread ve değerlerin şekli ile düzenine dair meta verilerle birleştirir. Bu meta veriler, veri ve çıktının bölünmesi için kullanılan CuTe TV düzenleri şeklinde ifade edilir. Kullanıcılar genellikle bu katmanı uzatmamalıdır, çünkü yeni mimariler için CuTe atomları sağlanacaktır.
Yukarıda gösterilen görüntü, SM70_8x8x4_F32F16F16F32_NT talimatı ve ilişkili MMA_Traits meta verileri listesini gösterir. Solda, TV düzenleri (thread_id,value_id) -> coord şeklinde kaydedilirken, sağda ise meta verilerin görselleştirildiği inverse coord -> (thread_id,value_id) haritası bulunmaktadır.
Bu görüntüyü print_latex(make_tiled_mma(cute::SM70_8x8x4_F32F16F16F32_NT{})) komutuyla oluşturulabilir.
CuTe Karo MMAları
Karo MMA ve karo kopya, MMA atomlarının ve kopya atomlarının karolarıdır. Karo kelimesini, atomları bir araya getirerek iki boyutlu bir matris gibi yapı inşa etmek için kullanıyoruz. Karo düzeyi, bireysel atomların tekrar etmesini ve veri ile thread’lerin kombinasyonunu içermesini sağlar.
Bu katman, CUTLASS 2.x’teki MMA talimatlarının warp-level döşenmesine benzer; ancak döşemeyi işlemlerin katılımıyla genel bir görünümde sunar. İşlevsellik amacıyla, çok sayıda donanım hızlandırılmış matematiksel ve veri taşıma işlemi ile bileşen oluşturan GPU mikro-kernelleri oluşturmak için kullanılır. Karo MMA ve Karo Kopya türleri, bu çeşitli donanım hızlandırıcı CuTe atomları için açıkça tanımlanmış bir API sunar.
CuTe GEMM ve Ana Döngüler
Mimariden bağımsız karo API ile, kullanıcılar GEMM dış döngülerine tutarlı bir arayüz oluştururken, iç döngüler atom katmanından gelir.
Örneğin:
Tensor gA = . . . // 64x16 gmem için A
Tensor gB = . . . // 96x16 gmem için B
Tensor gC = . . . // 64x96 gmem için C
// 64x16 statik düzenli row-major smem için A
Tensor sA = make_tensor(make_smem_ptr<TA>(smemAptr),
Layout<Shape < _64,_16>,
Stride<Int<17>, _1>>{});
// 96x16 statik düzenli interleaved col-major smem için B
Tensor sB = make_tensor(make_smem_ptr<TB>(smemBptr),
Layout<Shape <Shape <_32, _3>,_16>,
Stride<Stride< _1,_512>,_32>>{});
// Karo MMA'ya göre thread'ler arasında verilen tensorları böl.
ThrMMA thr_mma = tiled_mma.get_slice(thread_idx);
Tensor tCsA = thr_mma.partition_A(sA); // (MMA, MMA_M, MMA_K) smem
Tensor tCsB = thr_mma.partition_B(sB); // (MMA, MMA_N, MMA_K) smem
Tensor tCgC = thr_mma.partition_C(gC); // (MMA, MMA_M, MMA_N) gmem
// Kayıt tensörlerinin yukarıda belirtilen biçimlerle aynı şekil/düzenini oluştur.
Tensor tCrA = thr_mma.make_fragment_A(tCsA); // (MMA, MMA_M, MMA_K) rmem
Tensor tCrB = thr_mma.make_fragment_B(tCsB); // (MMA, MMA_N, MMA_K) rmem
Tensor tCrC = thr_mma.make_fragment_C(tCgC); // (MMA, MMA_M, MMA_N) rmem
// Kopyalama işlemi sonrasında rmem'den thread-düzeyinde bölmeleri temizle.
cute::copy(tCsA, tCrA);
cute::copy(tCsB, tCrB);
// Thread düzeyindeki toplama (accumulators) için rmem'i temizle.
cute::clear(tCrC);
// Rmem’de GEMM işlemini gerçekleştir.
// Eşdeğer olarak:
// for(int k = 0; k < size<2>(tCrA); ++k)
// for(int m = 0; m < size<1>(tCrC); ++m)
// for(int n = 0; n < size<2>(tCrC); ++n)
// tiled_mma.call(tCrA(_,m,k), tCrB(_,n,k), tCrC(_,m,n));
// Rmem'den gmem'e AXPBY yapmak için geri döndürme işlemi
cute::axpby(alpha, tCrC, beta, tCgC);
// Eşdeğer olarak:
// for(int i = 0; i < size(tCrC); ++i)
// tCgC(i) = alpha * tCrC(i) + beta * tCgC(i)Yukarıdaki kodda, bellek bölme ve işlem talimatlarının zamanlamalarının kararlarını almak için birçok karar mevcuttur.
- rmem olarak sadece
A: (MMA,MMA_M),B: (MMA,MMA_N)veC: (MMA,MMA_M,MMA_N)Tensörlerini tahsis et ve her k-blok döngüsünde kopyala. - k-tile’ları için gmem üzerinde birden fazla k-tile’yi hesaba katın ve smem’e kopyalayın.
- Yukarıdaki kopyalama aşamalarını hesaplama aşamalarıyla asenkron bir şekilde örtüşmek.
- s mem’den rmem’e kopyalama işleminde erişim desenlerini artıracak daha iyi düzenlemeler bulmak.
- gmem’den smem’ye kopyalama işlemi için verimli Karo Kopya bölme düzenleri bulmak.
Bu sorunlar, CuTe Tensörleri üzerinde talimatların zamanlaması ve uygulanması süreçleri “mekansal mikro-kerneller” yönünde değerlendirilmiştir. Genel olarak, CuTe’yi kullanarak talimatların sıralanması ve yürütülmesi konuları CUTLASS düzeyine bırakılmıştır ve bu konuların detayları bir sonraki kısmı oluşturacaktır.
Özet
Sonuç olarak, CuTe, geliştiricilerin CUDA kodunu daha okunabilir, sürdürülebilir ve yüksek performanslı bir şekilde yazmalarını sağlarken, tensör düzeni ve thread haritalama gibi düşük düzey detaylardan uzaklaştırır. Aynı zamanda modern NVIDIA GPU’lar üzerinde yoğun lineer cebir için bütünleşik ve cebirsel bir arayüz sunmaktadır.
Daha fazla bilgi için yazılımı GitHub üzerinden indirebilir, belgelere göz atabilir veya derin tartışmalar için Geliştirici Forumumuza katılabilirsiniz.
Teşekkürler
Bu yazıya katkılarından dolayı Jack Kosaian, Mark Hoemmen, Haicheng Wu ve Matt Nicely’e teşekkür ederiz. Özel teşekkürlerimizi Colfax International ekibine, Jay Shah, Paul VanKoughnett ve Ryo Asai’ya iletiyoruz.
