NVIDIA HPC SDK v25.7, GPU hızlandırmalı yüksek performans hesaplama (HPC) uygulamaları geliştiren yazılımcılar için önemli bir ilerleme sunuyor. Bu sürüm, yaklaşık iki yıllık sürekli geliştirmelerin bir sonucu olarak, bütünleşik bellek programlaması üzerine odaklandı ve CPU ile GPU arasında veri transferini otomatikleştiren kapsamlı bir araç seti oluşturdu. Geleneksel olarak gereken manuel veri yönetimini ortadan kaldırarak, GPU geliştirmelerini kolaylaştırıyor, porting sürelerini kısaltıyor, hataları azaltıyor ve bilimsel yükler üzerinde optimizasyon yapma esnekliğini artırıyor.
Yeni Özellikler ve Gelişmeler
CPU ve GPU mimarilerinin sıkı entegrasyonu ile koheren NVIDIA platformları, HPC pazarında giderek daha fazla ilgi görüyor. NVIDIA GH200 Grace Hopper Süper Çipi ve NVIDIA GB200 NVL72 sistemleri, bu yeni nesil mimarinin öne çıkan örneklerinden. Bu sistemler, İsviçre Ulusal Süper Bilgi Merkezi’ndeki ALPS süper bilgisayarında ve Jülich Süper Bilgi Merkezi’ndeki JUPITER’de kullanılıyor.
Bunlar sadece yüksek performanslarıyla değil, aynı zamanda geliştirici verimliliğini artırmalarıyla da dikkat çekiyor. Paylaşılan adres alanı ile birlikte, geliştiriciler artık CPU ve GPU arasında veri transferlerini manuel olarak yönetmek zorunda kalmıyor. NVIDIA CUDA sürücüsü, veri transferini otomatik olarak yürütüyor ve bu, gerçek dünya projelerinde değerli faydalar sağlıyor.
“Bütünleşik bellek programlamasından faydalanmak, NEMO okyanus modelinin GPU’lara geçirilmesi sürecinde hızlanmamızı sağladı. Ayrıca, geleneksel yöntemlere kıyasla daha fazla yükü GPU’da çalıştırabilme esnekliği sunuyor.”— Alexey Medvedev, Kıdemli Araştırma Mühendisi, Barcelona Süper Bilgi Merkezi
Veri Yönetimini Kolaylaştıran Yöntemler
NVIDIA Grace Hopper Süper Çip Mimarisi’nin tanıtılmasından bu yana, NVIDIA HPC SDK, bütünleşik bellek programlamasına yönelik özellikler eklemeye devam ediyor. NVIDIA HPC SDK 25.7 sürümü, CPU ile GPU arasında manuel veri hareketini basitleştiren ve birçok durumda ortadan kaldıran kapsamlı bir araç seti sunuyor. Bu, bilimsel uygulama geliştiricileri için büyük bir verimlilik artışı sağlıyor.
Veri yönetimi, GPU programlamanın en zorlayıcı yönlerinden biri olarak kabul ediliyor. CPU ve GPU alt programları arasında doğru ve etkili veri akışını sağlamak, genellikle birden fazla hata ayıklama ve optimizasyon döngüsü gerektiriyor. Birçok HPC uygulamasında, bu karmaşıklık, dinamik olarak tahsis edilmiş verilerin ve bileşik yapının kullanımıyla daha da artıyor. Aşağıda verilmiş olan Fortran örneği, “derin kopya” olarak adlandırılan yaygın bir deseni gösteriyor; burada, bir türetilmiş türün tahsis edilebilir dizi üyesini paraleleliz etmek için ek bir döngü gerekiyor.
type mytype
integer, allocatable::arrmem(:)
integer :: scalarmem
end type mytype
! Türetilmiş dizi tanımı.
type(mytype)::base(:)
…
! Bu döngü, GPU'ya veri kopyalamak için.
!$acc enter data copyin(base(:))
do i=1,N
!$acc enter data copyin(base(i)%arrmem(:))
end do
! Türetilmiş tür mytype'ın üyelerine erişim sağlayan paralel OpenACC döngüsü.
!$acc parallel loop collapse(2) default(present)
do i=1,N
do j=1,M
base(i)%arrmem(j)=base(i)%scalarmem
end do
end doGerçek dünya HPC uygulamalarında, döngüler genellikle birden fazla diziye erişir; bunlardan bazıları türetilmiş türler içinde, diğerleri modüllerde veya ortak bloklarda tanımlanmıştır. Veri hareketini yönetmek, çoğu gerçekçi kodda oldukça karmaşık hale gelebilir ve çoğu zaman paralel döngülerin kendilerini belirlemek ve düzenlemekten daha zor olabilir. Veri bağımlılıklarını anlamak, doğru transferleri sağlamak ve bellek sızıntıları veya yarış durumlarından kaçınmak gibi sorunlar, GPU programlamasına önemli bir yük bindirir.
Bu karmaşıklığın büyük bir kısmı, Grace Hopper ve benzeri GPU mimarileri üzerinde sunulan bütünleşik bellek modeli ile ortadan kaldırılabilir. Hem CPU hem de GPU bir tek adres alanını paylaştığı için, açık veri yönetimi sıklıkla gerekli değildir. Daha önceki örnekteki türetilmiş tür dizisi artık ek veri hareketi direktifleri olmadan doğrudan paralelleştirilebilir.
type mytype
integer, allocatable::arrmem(:)
integer :: scalarmem
end type mytype
! Türetilmiş dizi tanımı.
type(mytype)::base(:)
…
! Türetilmiş tür mytype'ın üyelerine erişim sağlayan paralel OpenACC döngüsü.
!$acc parallel loop collapse(2)
do i=1,N
do j=1,M
base(i)%arrmem(j)=base(i)%scalarmem
end do
end doC++ Kodlarının Zorlukları
C++ kodlarıyla GPU’lara aktarım yapıldığında, nesne yönelimli soyutlamanın ve veri kapsüllemenin yoğun kullanımı, geliştiricilerin veri kopyalamak için uygulama içi detaylara erişimini engelleyebiliyor. Örneğin, std::vector ile yazılmış bir OpenACC kodu, bütünleşik bellek olmadan düzgün bir şekilde çalışmaz. Döngüde kullanılan std::vector işaretçi operatörü, hem std::vector sınıfındaki veriye hem de öğeleri için tahsis edilen verilere erişim gerektirir; bu veriler, std::vector sınıfının kendisiyle ardışık olarak yerleştirilmemiştir.
std::vector<int> v(N);
#pragma acc kernels
for (i = 0; i < v.size(); i++)
v[i] = i;
Bütünleşik bellek olmadan, örneğin copyout(v) ifadesinin eklenmesi, sadece std::vector nesnesinin kopyalanmasıyla sonuçlanır; fakat içeriğindeki öğeler kopyalanmaz. Bu nedenle, bu tür kodlar genellikle ham işaretçi üzerinde doğrudan çalışacak şekilde yeniden yazılır, bu da nesne yönelimli programlama tarzına dönülmesine sebep olur ve diğer STL konteynırlarından verileri GPU’ya kopyalamak mümkün olmayabilir. Bunun nedeni, bu standart C++ konteynırlarının elemanlarına erişim sağlayan bir arayüz eksikliğidir.
std::vector<int> v(N);
auto ptr = v.data();
#pragma acc kernels copyout(ptr[0:N])
for (i = 0; i < v.size(); i++)
ptr[i] = i;
Okyanus ve İklim Bilimleri için kullanılan gelişmiş bir modelleme çerçevesi olan NEMO, araştırma faaliyetleri ve tahmin hizmetleri için kullanılır.
BSC, NEMO’yu GPU’lara aktarma sürecine başlamadan önce, mevcut kod tabanının mevcut faydalarını keşfetmek amacıyla iç değerlendirme yaptı.
NVIDIA GTC oturumlarında NVIDIA Grace Hopper Platformu ile Bilimsel İş Akışlarını Hızlandırma ve NVIDIA HPC Derleyicileri ile Bilim Hızlandırma, bu gerçek dünya kodunu bir vaka çalışması olarak kullandık. Geliştirici verimliliğinin, Grace Hopper gibi koheren sistemlerde önemli ölçüde artırıldığını gösterdik.
NEMO İçin GPU Portlama Stratejisi
Başarıyla gerçekleştirdiğimiz portlama sürecinde, bütünleşik bellek, açık veri yönetimi kodlarına gerek duymadan sadece paralelleşmeye odaklanmamıza imkan tanıdı. Daha az kod ile, geliştiriciler GPU portlama sürecinin daha erken aşamalarında hız kazandılar. Bu gösterim için, NEMO v4.2.0’daki GYRE_PISCES benchmark’ını kullandık.
Bu bellek bant genişliği ile sınırlandırılmış bir benchmark’tır ve başlangıçta çok çekirdekli CPU’lara eş zamanlı olarak MPI kullanılarak paraleleştirilmiştir. Ana noktalar, aktif ve pasif izleyicilerin difüzyon ve advecting süreçleridir. Bu bölümler üzerinde, OpenACC kullanarak döngüleri sadece paralelleştirip, bellek yönetimini CUDA sürücüsü ve donanıma bıraktık.
Başlangıçta GPU’ya portlama stratejimiz, hiçbir GPU veri yönetimi kodu eklenmeden şu şekildeydi:
- Tam paralel, sıkı iç içe geçmiş döngüler
!$acc parallel loop gang vector collapse()kullanılarak paralelleştirildi. - Çapraz etkileşim bağımlılıkları olan döngüler
!$acc loop seqolarak belirtildi. - Dizi notasyonu içindeki işlemler
!$acc kernelsiçinde sarıldı. - Paralel döngüler içindeki dış rutinler
!$acc routine seqile işaretlendi.
NEMO’nun yapısı gereği, birçok fonksiyon birbirini takip eden birden fazla paralel bölge içeriyor ve bu durum, her OpenACC paralel yapısının sonunda gizli senkronizasyonlara sebep oluyor. Bu senkronizasyonlar, programlamayı basit hale getirmek ve paralelize edilmeyen ve sıralı kod arasında bir yürütme sırası korumak için eklenmiştir.
Bu gizli engellerden kaçınmak için, paralel ve kernel yapılarında async ekleri eklendi ve böylece daha iyi bir eşzamanlılık sağlandı. Asenkron modda paralel bölgeler çalıştırıldığında, !$acc ile senkronizasyon, GPU’da hesaplanan verilerin, sonraki MPI çağrılarından önce mevcut olmasını sağlamak için veya yerel değişkenlerin, alt rutinlerin sonuna kadar kapsam dışına çıkmasını önlemek için eklendi.
Aşağıda, kamuya açık açık kaynak NEMO deposunda bulunan trazdf.f90 içinden alınmış bir kod parçası, daha önce açıklanan OpenACC paralelleştirme stratejisini göstermektedir.
SUBROUTINE tra_zdf_imp(...)
...
REAL(wp), DIMENSION(ntsi-(nn_hls):ntei+(nn_hls),ntsj-(nn_hls):ntej+(nn_hls),jpk) :: &
& zwi, zwt, zwd, zws
...
DO jn = 1, kjpt
...
!* 3D tekrarı: Xk = (Zk - Sk Xk+1) / Tk (sonuç izleyicinin arkasındaki)
!$acc parallel loop gang vector collapse(2) async(1)
DO jj = ntsj, ntej
DO ji = ntsi, ntei
pt(ji,jj,jpkm1,jn,Kaa) = pt(ji,jj,jpkm1,jn,Kaa)/zwt(ji,jj,jpkm1)*tmask(ji,jj,jpkm1)
END DO
END DO
!$acc end parallel
!$acc parallel async(1)
!$acc loop seq
DO jk = jpk-2, 1, -1
!$acc loop gang vector collapse(2)
DO jj = ntsj,
Kaynak
