Nvdia

CUDA Runtime’da Dinamik Yükleme: Performansı Artırmanın Yolları

0
Paylaşımlar

Dönemsel GPU Aygıt Kodu Yükleme Zamanı

Geçmişte, GPU aygıt kodu, nvcc gibi çevrimdışı araçlar ile uygulama ile birlikte derleniyordu. Bu durumda, GPU aygıt kodu CUDA çalışma zamanında dahili olarak yönetiliyordu. Ardından, çekirdekleri <<<>>> operatörüyle başlatarak CUDA çalışma zamanı, çağrılan çekirdeğin başlatılmasını sağlıyordu.

Ancak bazı durumlarda, GPU aygıt kodunun dinamik olarak derlenmesi ve yüklenmesi gerekebilir. Bu yazıda, CUDA çalışma zamanı aracılığıyla bunu başarmanın bir yolunu gösterecek ve CUDA sürücüsü ile CUDA çalışma zamanı çekirdek eylem işleyicileri arasındaki etkileşimi de sergileyeceğiz.

Dinamizmin Faydaları

Dinamik GPU aygıt kodu yüklemenin sağladığı çeşitli avantajlar vardır:

  • Yüklenen GPU aygıt kodu üzerinde açık kontrol: Yükleme birimiyle ayrı bir şekilde değiştirilse bile bu kontrol sağlanır.
  • Yükleme zamanını kontrol etme: Yükleme API seçenekleriyle nasıl yükleneceğini belirleme imkanı sunar.
  • anlık derleme:NVRTC gibi diğer CUDA Toolkit bileşenleri kullanılarak GPU aygıt kodu modülleri anlık olarak oluşturulabilir.
  • seçici GPU aygıt kodu bağlantısı:nvJitLink gibi diğer CUDA Toolkit bileşenleri kullanarak bağlantı zamanında optimizasyon sağlanabilir.
  • Dinamik GPU aygıt kodu yükleme gerektiren başlık dosyaları:nvcc ile derlendiğinde, bu değişikliklerle birlikte CUDA çalışma zamanı ile bağlanabilir.

CUDA Çalışma Zamanında Statik Yükleme

CUDA çalışma zamanı, başlangıçta yüklenen GPU aygıt kodunun durumunu korur. GPU aygıt kodu modülleri, nvcc gibi derleme araçlarıyla derlenip bağlandıkları dosyalara göre belirlenir. Başlangıçta, CUDA çalışma zamanı bu modülleri yükler ve onlarla dolaylı olarak etkileşimde bulunursunuz. Aşağıda basit bir örnek bulunmaktadır:

main.cu:
#include <stdio.h>
__global__ void helloWorld() { printf(“Hello from the GPU!n”); }

int main(int argc, char *argv[]) {
    cudaSetDevice(0);
    helloWorld<<<1,1,1>>>();
    return cudaDeviceSynchronize();
}

Bu basitleştirilmiş örnek, nvcc ile derlendiğinde, CUDA çalışma zamanının GPU üzerindeki helloWorld çekirdeğini çalıştırabilmesi için uygun GPU aygıt kodu modülünü oluşturan bir yürütülebilir dosya oluşturur.

CUDA Sürücüsünde Dinamik Yükleme

CUDA sürücüsü, çalıştırılacak GPU aygıt kodunun dinamik olarak yüklenmesini gerektirir ve CUDA bağlamları gibi daha fazla durumu yönetir; bu durum, CUDA çalışma zamanının otomatik olarak yaptığı bir işlemdir. Benzer bir örnek, iki dosyaya ayrı ayrı derleme yollarıyla bölünebilir. Daha fazla bilgi için, NVIDIA CUDA Derleyici Sürücü NVCC sayfasına göz atabilirsiniz.

GPU kodu, bağımsız bir GPU aygıt kodu modülü olan bir .fatbin, .cubin veya bağımsız PTX dosyası (bu örnekte, device.fatbin) olarak nvcc kullanılarak derlenecektir.

Sonrasında, ana kaynak dosyasının bir bölümü, bu .fatbin dosyasını kullanıp yönetecek şekilde aşağıdaki gibi olacaktır.

main.c:
#include <cuda.h>

int main(int argc, char *argv[]) {
    …
    cuDeviceGet(&dev, 0);
    cuDevicePrimaryCtxRetain(&ctx, dev);
    cuCtxPushCurrent(ctx);
    cuLibraryLoadFromFile(&library, “device.fatbin”, NULL, NULL, 0, NULL, NULL, 0);
    cuLibraryGetKernel(&kernel, library, “helloWorld”);
    cuLaunchKernel((CUfunction)kernel, 1, 1, 1, 1, 1, 1, 0, NULL, NULL, NULL);
    cuCtxSynchronize();
    cuLibraryUnload(library);
    cuDevicePrimaryCtxRelease(dev);
    return 0;
}

Dinamik yüklemenin daha önce belirtilen faydaları, CUDA çalışma zamanına da genişletilmiş olup kullanım durumu etkinleştirme bölümünde daha ayrıntılı bir şekilde açıklanacaktır.

CUDA Çalışma Zamanında Dinamik Yükleme

CUDA’daki dinamik yükleme destekleyen değişikliklerle, GPU aygıt kodunu dinamik olarak yükleme esnekliği, artık CUDA çalışma zamanına da sağlanmıştır. Bu, daha önceki örneğin aşağıdaki gibi kısaltılabileceği anlamına gelir. Bu yaklaşım, sürücü örneğinde gerekli olan açık CUDA bağlamı yönetim yükünü ortadan kaldırır. Aşağıda güncellenmiş ana kaynak dosyasının bir bölümü verilmiştir:

main.cu:
#include <cuda_runtime_api.h>

int main(int argc, char *argv[]) {
    …
    cudaLibraryLoadFromFile(&library, “device.fatbin”, NULL, NULL, 0, NULL, NULL, 0);
    cudaLibraryGetKernel(&kernel, library, “helloWorld”);
    cudaLaunchKernel((const void*)kernel, 1, 1, NULL, 0, NULL);
    cudaDeviceSynchronize();
    cudaLibraryUnload(library);
    return 0;
}

Kullanım Durumu Etkinleştirme

Bu yeni özelliklerle hangi kullanım durumlarının mümkün hale geldiğine bakalım:

  • Pür CUDA çalışma zamanı API kullanımı: Artık tüm dinamik GPU aygıt kodu modül yükleme, yalnızca sürücü API’lerini kullanmayı gerektiriyor.
  • CUDA sürücüsü ve CUDA çalışma zamanı arasındaki türlerin karşılıklı kullanabilirliği: Çekirdek işlevi işleyicilerini değiştiren herhangi bir durumda yararlanma olanağına sahip olursunuz.
  • CUDA çalışma zamanı örnekleri arasında durum paylaşımı: Şimdi cudaGetKernel API’si ile çekirdek işlevlerini paylaşmayı mümkün kılıyor.

Pure CUDA Çalışma Zamanı API Kullanımı

Artık, tüm dinamik GPU aygıt kodu modül yükleme yalnızca sürücü API’lerini gerektiriyordu. Diğer kütüphaneler veya uygulamalar, NVRTC kullanarak GPU aygıt kodunu oluşturan veya dinamik yükleyen nvJitLink gibi yükleme işlemlerini gerçekleştirebilirler. Ancak yeni CUDA çalışma zamanı dinamik yükleme API’leri ile bu derleme ve yönetim süreçleri, artık tam anlamıyla CUDA çalışma zamanı üzerinden gerçekleştirilebilir.

NevRTC belgelere göre güncellenmiş bir SAXPY örneği şu şekilde olmuştur:

Güncel NVRTC SAXPY Örneği

// Üretilmiş PTX'i yükleyin ve SAXPY çekirdeği için bir tanım alın.
CUdevice cuDevice;
CUcontext context;
CUmodule module;
CUfunction kernel;
CUDA_SAFE_CALL(cuInit(0));
CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));
CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, 0, 0));
CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, “saxpy”));
…
//SAXPY'yi çalıştırın
void *args[] = {&a, &dX, &dY, &dOut, &n};
CUDA_SAFE_CALL(
cuLaunchKernel(kernel,
    NUM_BLOCKS, 1, 1,   // grid dim
    NUM_THREADS, 1, 1,  // block dim
    0,                  // shmem
    NULL,               // stream
    args, 0));          // arguments

Güncellenmiş NVRTC SAXPY Örneği

// Üretilmiş PTX'i yükleyin ve SAXPY çekirdeği için bir tanım alın.
cudaLibrary_t library;
cudaKernel_t kernel;
CUDART_SAFE_CALL(cudaLibraryLoadData(&library, ptx, 0,0,0,0,0,0));
CUDART_SAFE_CALL(cudaLibraryGetKernel(&kernel, library, “saxpy”));
…
//SAXPY'yi çalıştırın
void *args[] = {&a, &dX, &dY, &dOut, &n};
CUDART_SAFE_CALL(
cudaLaunchKernel((void*)kernel,
    NUM_BLOCKS,    // grid dim
    NUM_THREADS,   // block dim
    args,          // arguments
    0,             // shmem
    NULL));        // stream

Ayrıca, daha önce, dinamik GPU aygıt kodu yüklemesi gerektiren başlık dosyaları nvcc ile derlendiğinde, kullanıcınınCUDA sürücüsüne bağlantı yapması gerekecekti. Artık bu, CUDA çalışma zamanı ile yüklenebilen başlık dosyalarının kendi başlarına bağlanma gereğini ortadan kaldırdığı anlamına geliyor.

CUDA Sürücüsü ve CUDA Çalışma Zamanı Arasında Türlerin Karşılıklı Kullanılabilirliği

Önceleri, çekirdek işlevi işleyicileri CUDA çalışma zamanı ve CUDA sürücüsü arasında değişmeyen şekilde kullanılmıyordu; oysa CUDA akışları ve CUDA olayları gibi birçok diğer işleyici değiştirilebilir durumdaydı.

Artık cudaKernel_t ve CUkernel (aynı zamanda cudaLibrary_t ve CUlibrary) türleri değiştirilebilir hale geldi. CUDA çalışma zamanı API’lerini kullanarak yüklemek ancak CUDA sürücüsü API’sini kullanarak çekirdek veya özellik ayarlarını başlatmak için türleri dönüştürebilirsiniz.

CUDA Çalışma Zamanı Örnekleri Arasında Durum Paylaşma

İki teorik kütüphane, kütüphane A ve kütüphane B olduğunu varsayalım; her biri kendi statik CUDA çalışma zamanına bağlıdır.

Tarihsel CUDA çalışma zamanının otomatik yüklemesi, CUDA çekirdek işlevi işleyicilerinin birden çok CUDA çalışma zamanı örneği arasında paylaşılmasını mümkün kılmamıştır. Bu bağlamda, bu kütüphanelerin çekirdek işlevlerini paylaşmanın bir yolu olamazdı.

Artık cudaGetKernel API’si ile çekirdek işlevleri paylaşılarak elde edilebilir ve diğer bir CUDA çalışma zamanı örneğine iletilebilir. Böylece, iki kütüphane arasındaki CUDA çekirdek işlevinin paylaşılmasına ihtiyaç duyulduğunda, kütüphane A cudaGetKernel çağrısını yapabilir ve işleyiciyi kütüphane B’ye iletebilir. Bu durum, kütüphaneler arasında kod paylaşımını artırma ve her bir kütüphanenin kendi çekirdek uygulamalarını içermesi gerekliliğini azaltma potansiyeline sahiptir.

// matrix_mul.cu - dinamik paylaşılan işleyici kullanımı
void matrix_mul() {
  cudaLibrary_t lib;
  cudaKernel_t kern;
  cudaLibraryLoadData(&lib, ptx, …); // ptx den nvrtc
  cudaLibraryGetKernel(&kern, lib, “matrixMul”);
  libcommon.foo(kern);
}
// vector_add.cu  - örtük paylaşılan işleyici kullanımı
__global__ void vectorAdd() { … }
void vector_add() {
  cudaGetKernel(&kern, vectorAdd);
  libcommon.foo(kern);
}
// libcommon.cu - paylaşılabilir çekirdek işlevi alır
void foo(cudaKernel_t kern) {
  cudaLaunchKernel(kern, ...);
}

Bu örnek basit görünse de, kütüphanelerin, aralarında gereken çekirdeklerin kopyalarını azaltarak bellek alanından tasarruf sağlayabileceğini göstermektedir.

CUDA Çalışma Zamanında Dinamik Yüklemeye Başlayın

Bu yazıda, GPU aygıt kodunu yükleme kabiliyeti sağlayan yeni CUDA çalışma zamanı API’lerini tanıttık. Bu, yalnızca CUDA çalışma zamanı API’sinin gerekli olduğu durumlarda cihaz kodunu yüklemek ve çalıştırmak için daha basit bir yol sunmaktadır.

Bu API’leri kullanmaya başlamak için, CUDA Toolkit sürüm 12.8 veya üzerini CUDA Toolkit adresinden indirmeniz gerekmektedir. cudaLibrary* ve cudaKernel* API’leri hakkında daha fazla bilgi için, CUDA çalışma zamanı API belgelerine göz atabilirsiniz.

Kaynak

Nvdia Blog

0
Paylaşımlar
Yorumlar

Benzer Haberler

Exit mobile version