CUDA çekirdeklerinin Python dilinde yazılmasını mümkün kılarak, C++ ile uygulanan performansı Python ekosistemine taşıyan Numba, önemli bir köprü işlevi görüyor.
Ancak, CUDA C++ geliştiricileri, şu anda Python ile entegre olmayan pek çok kütüphaneye erişim sağlamaktadır. Bu kütüphaneler arasında CUDA Core Compute Libraries (CCCL), cuRAND ve bfloat16 gibi sayısal türlerin başlık tabanlı uygulamaları gibi birçok örnek bulunmaktadır.
Her CUDA C++ kütüphanesi, Python’a kendi yöntemleriyle tanıtılabilir. Ancak, her kütüphane için bağlayıcı (binding) oluşturmak, hem zahmetli hem de tekrarlayıcı bir süreçtir. Örneğin, float16 ve bfloat16 veri türleri için 60’tan fazla bağımsız işlev tanımı bulunmaktadır ve bu türlerin her biri için benzer bağlayıcılar oluşturulması gerekecektir.
Ayrıca, manuel olarak oluşturulan bağlayıcılar, temel CUDA C++ kütüphanesi yeni özellikler eklediğinde senkronizasyon sorunları yaşayabilir.
Çözüm: Numbast
Numbast, CUDA C/C++ API’lerini otomatik şekilde Numba bağlayıcılarına dönüştüren bir pipeline oluşturur.
Genel hatlarıyla, üst seviye tanımlar, CUDA C++ başlık dosyalarından okunur, serileştirilir ve Python API’lerine aktarılır. Numba bağlayıcı üreticileri, bu tanımlamalar üzerinde döngü oluşturarak her API için Numba uzantıları üretir.
Örnek: C++ Öykünmesi
Aşağıdaki örnekte, myfloat16 adındaki bir tür için Numba bağlayıcılarının nasıl oluşturulduğu gösterilmektedir. Bu C++ tanımlamaları, CUDAfloat16 başlığındakilerden esinlenerek basitleştirilmiştir.
C++ Tanımı
Bu örnek, C++ söz diziminde aşağıdaki unsurları göstermektedir:
// demo.cuh
struct __attribute__((aligned(2))) myfloat16
{public:
half data;
__host__ __device__ myfloat16();
__host__ __device__ myfloat16(double val);
__host__ __device__ operator float() const;
};
__host__ __device__ myfloat16 operator+(const myfloat16 &lh, const myfloat16 &rh);
__host__ __device__ myfloat16 hsqrt(const myfloat16 a);
- Bir yapı (struct) tanımı, şu unsurları içerir:
- Cihaz (device) yapıcıları
- Bazı cihaz yöntemleri, içinde dönüşüm ve aritmetik operatörler yer alır
- İki işlev tanımı: aritmetik operatör aşırı yüklemesi ve bir karekök işlevi.
Daha fazla bilgi için Numbast’ın desteklediği dil özellikleri listesini buradan görebilirsiniz.
Numbast Kurulum Scripti
Numbast kullanımı genellikle iki adım içerir:
- Başlık dosyalarını
AST_Canopyile çözümlemek. - Çözümleme sonuçlarına dayanarak Numba bağlayıcıları oluşturmak.
Aşağıdaki örnek, bu iki adımı uygulayarak Numba bağlayıcılarını kurmayı göstermektedir:
import os
from ast_canopy import parse_declarations_from_source
from numbast import bind_cxx_struct, bind_cxx_function, MemoryShimWriter
from numba import types, cuda
from numba.core.datamodel.models import PrimitiveModel
import numpy as np
# 1. Adım:
# demo.cuh dosyasını AST olarak çözümleyin, tüm bildirimleri okuyun.
source = os.path.join(os.path.dirname(__file__), "demo.cuh")
# "sm_80" uyumlu bir makine için bağlayıcılar oluşturmak istediğinizi varsayalım.
structs, functions, *_ = parse_declarations_from_source(source, [source], "sm_80")
shim_writer = MemoryShimWriter(f'#include "{source}"')
# 2. Adım:
# Açıklamalardan Numba bağlayıcıları oluşturun.
# Yeni tür "myfloat16" bir Sayı (Number) türüdür, veri modeli `PrimitiveModel`dir.
myfloat16 = bind_cxx_struct(shim_writer, structs[0], types.Number, PrimitiveModel)
# İlk işlev, bir işlemi bağlar ve `operator.add` ile bağlanır.
# Kernel'lerde doğrudan kullanabilirsiniz: `myfloat16 + myfloat16`.
bind_cxx_function(shim_writer, functions[0])
# İkinci işlev `hsqrt` olup, Numbast yeni bir Python referansı oluşturur ve sonucu döndürür.
hsqrt = bind_cxx_function(shim_writer, functions[1])
Veri modelleri, Numba’nın temel verileri temsil etme biçimleridir. PrimitiveModel modeli, myfloat16 için uygun olan bir modeldir. StructModel modeli (burada kullanılmıyor), sınıflar ve yapılar için faydalıdır. Diğer veri modelleri daha az yaygın olarak kullanılmaktadır.
En Doğal Yolda Kullanım
CUDA C++’ta, bir myfloat16 nesnesi oluşturulabilir ve aşağıdaki gibi kullanılabilir:
__global__ void kernel()
{
auto one = myfloat16(1.0);
auto two = myfloat16(2.0);
auto three = one + two;
auto sqrt3 = hsqrt(three);
}
Numba kernel’lerinde, aşağıdaki gibi kullanılabilir:
@cuda.jit(link=shim_writer.links())
def kernel():
one = myfloat16(1.0)
two = myfloat16(2.0)
three = one + two
sqrt3 = hsqrt(three)
Numba’daki tür çıkarımı sayesinde, kod, orijinal C++’tan daha temiz hale gelir.
İlk Desteklenen Bağlayıcılar: bfloat16 Veri Türü
Numba aracılığıyla desteklenen ilk bağlayıcı, yeni bfloat16 veri türüdür. Bu tür, PyTorch’un torch.bfloat16 veri türü ile birlikte kullanılabilir, böylece bu yeni veri türü ile özel hesaplama çekirdekleri geliştirmek daha verimli hale gelir.
Aşağıdaki kod örneği, yeni bfloat16 veri türünü kullanarak bir Numba çekirdeği geliştirme sürecini göstermektedir. Bu çekirdek, torch.bfloat16 türünde bir PyTorch dizisini Numba hesaplama çekirdeğine geçirir ve CUDA içsel işlemleri kullanarak matematiksel işlemler yapar.
from numba import float32
import numba.cuda as cuda
import torch
from numbast_extensions.bf16 import get_shims, hsin, nv_bfloat16
@cuda.jit(link=get_shims())
def torch_add_sin(a, b, out):
i, j = cuda.grid(2)
if i < out.shape[0] and j < out.shape[1]:
# bfloat16 türünün aritmetiği
sum = a[i, j] + b[i, j]
# Bfloat16 yerel intrinsics
sin_of_sum = hsin(sum)
# bf16'dan f32'ye yukarı dönüşüm
f32 = float32(sin_of_sum)
# f32'den bf16'ya aşağı dönüşüm
bf16 = nv_bfloat16(f32)
# Harici diziye atama
out[i, j] = bf16
a = torch.ones([2, 2], device=torch.device("cuda:0"), dtype=torch.bfloat16)
b = torch.ones([2, 2], device=torch.device("cuda:0"), dtype=torch.bfloat16)
expected = torch.sin(a + b)
out = torch.zeros([2, 2], device=torch.device("cuda:0"), dtype=torch.bfloat16)
threadsperblock = (16, 16)
blockspergrid = (1, 1)
torch_add_sin[blockspergrid, threadsperblock](a, b, out)
assert torch.equal(expected, out)
Numbast ve bfloat16 Numba bağlayıcılarını conda-forge üzerinden indirebilirsiniz:
conda install -c nvidia -c rapidsai -c conda-forge ml_dtypes numbast-extensions
Yapı: Numbast’ın Temelleri
Numbast, iki bileşenden oluşmaktadır:
AST_Canopy: C++ başlıklarını çözümleyen ve serileştiren temel bir katmandır.- Numbast: Çözümleme sonuçlarını tüketerek dinamik olarak Numba bağlayıcılarını oluşturan kullanıcıya yönelik bir katmandır.
AST_Canopy: Tanım Çözücüsü
Doğa biliminin bir dalı olan kanopi, bir ormanın üst soyunu ifade eder. AST_Canopy, üst seviye tanımları inceleyip, bu bilgileri Python katmanına geçiren bir araçtır.
Şekil 1, AST_Canopy mimarisini göstermektedir:
clangTooling: Numabast gibi bağımsız araçlar yazımını destekleyen bir Clang kütüphanesidir.libastcanopy:clangToolingkullanarak, tanım çözümleme mantığını uygulamaktadır.pylibastcanopy:libastcanopyAPI’lerini Python’da doğrudan sunan bir bağlayıcıdır.- AST_Canopy: Kullanıcı dostu bir Python deneyimi sunan
pylibastcanopyüzerine inşa edilmiş bir katmandır.
Başlık çözümleme ve serileştirme işlemlerinin yanı sıra AST_Canopy ayrıca aşağıdaki özellikleri sunmaktadır:
- Runtime ortam tespiti: Conda paketleri aracılığıyla yüklü
libstdcxxve CUDA başlıklarını otomatik olarak tespit eder veclangderleyicilerini uygun şekilde ayarlar. - Hesaplama yetenekleri çözümleme esnekliği: Farklı hesaplama yeteneklerine dayalı AST çözümleme yapılandırması sağlar. Bazı başlıklar, hesaplama yeteneğine bağlı olarak kodu koşullu olarak sunar.
Numbast: Bağlayıcı Üretici
Numbast, AST_Canopy’nin bir ardılıdır ve çözümleme bilgilerini tüketerek otomatik Numba bağlayıcıları oluşturur. Genel hatlarıyla, Numbast, C++ ve Python söz dizimi arasında bir çeviri katmanı sağlar.
Numba’nın tür sistemi, C ve C++ benzeri diller ile birçok benzerliğe sahiptir. Elbette, Python’da bulunmayan gösterimler gibi bazı özellikler de mevcuttur. Numbast, C++ ve Python arasındaki benzerlikleri ve farklılıkları kapsayan orta katmandır.
Alt Yapının Genel Görünümü
Numbast kullanılarak üretilen bağlayıcılar, Numbada yabancı işlev çağrısı (FFI) özelliğiyle düşürülmektedir (lowering). Vazgeçilebilir, CUDA işlem çağrısını sarmalayan Numba ABI uyumlu shim işlevleri oluşturularak NVRTC ile derlenmektedir. Böylece, bir CUDA C++ geliştiricisinin elde ettiği optimize edilmiş performansı bekleyebilirsiniz; ancak FFI’nin performansı göz önünde bulundurulmalıdır.
Gelecek sürümlerde Numba-cuda, bağlantı zamanı optimizasyonu (LTO) desteği sunarak, hızlandırılmış Numba çekirdekleri ve yerel CUDA C++ performansı arasındaki açığı azaltacaktır.
Önemli Noktalar ve Sonuç
Hem AST_Canopy hem de Numbast üzerinde dikkat edilmesi gereken hususlar bulunmaktadır. AST_Canopy, clangTooling kütüphanesine bağımlıdır. Yeni CUDA dil özellikleri henüz clangTooling tarafından desteklenmiyor olabilir ve bu nedenle yeni dil özelliklerine bağımlı kütüphaneler doğru bir şekilde çözümlemeyebilir. Ancak, çoğu kütüphane, clangTooling tarafından desteklenen özellikleri kullanmaktadır.
Bu yazıda, yeni bir Numba bağlayıcı üretim aracı olan Numbast‘ı tanıttık. Numbast sayesinde, sürekli büyüyen CUDA C++ özellik setinden hızlı bir şekilde faydalanabilirsiniz. Numbast v0.1.0, Numba için yeni bfloat16 veri türünü sunmaktadır. Gelecekte Numbast ile daha fazla bağlayıcının, yeni veri türlerinin, NVSHMEM bağlayıcılarının ve CCCL bağlayıcılarının gelişmesini bekleyebilirsiniz.