“CUDA C++ Derleyici Güncellemeleri: ELF Görünürlüğü ve Bağlantısı Üzerindeki Etkileri”

NVIDIA, CUDA’nın bir sonraki büyük sürümü olan CUDA 13.0 ile birlikte NVIDIA CUDA Compiler Driver (NVCC)‘da iki önemli değişiklik tanıtıyor. Bu değişiklikler, __global__ işlevleri ve cihaz değişkenleri için ELF görünürlüğü ve bağlantısını etkileyecek. Güncellemelerin amacı, tespit ve hata ayıklaması zor olan ince çalışma zamanı hatalarını önlemektir. Ancak bu değişiklikler, bazı mevcut CUDA C++ programlarını etkileyebilir.

Değişikliklerin Kapsamı

Bu yazı; kullanıcıları olası kesintiler konusunda uyarmak, değişikliklerin nedenini açıklamak ve eski davranışı geri yüklemek için kullanılabilecek NVCC bayrakları hakkında rehberlik sağlamak amacıyla kaleme alınmıştır. Değişiklikler, Tablo 1’de özetlenmiştir.

NVCC Değişikliği #1: ELF Görünürlüğü

CUDA 13.0’dan önceki sürümlerde, NVCC derleyicisi, __global__ işlevlerinin ve __managed__/__device__/__constant__ değişkenlerin ELF görünürlüğünü değiştirmiyordu. Eğer bu kod, bir paylaşılan kütüphane olarak paketlenirse, bu semboller paylaşılan kütüphaneyi kullanan kişiler için görünür hale gelir.

Problemin Genel Görünümü ve Örnekler

NVCC, varsayılan olarak statik CUDA Runtime Library (CUDART) sürümünü bağlar. Bu durum, hem paylaşılan kütüphaneye hem de ana programa iki ayrı CUDART kütüphanesi bağlanmasına neden olabilir. Eğer __global__ işlev veya __device__/__managed__/__constant__ değişken paylaşılan kütüphane sınırları üzerinden erişiliyorsa, bu durum ince çalışma zamanı sorunlarına yol açabilir.

Örnek 1

foo.cu dosyası şu şekilde:

//-- foo.cu --
#include 
__global__ void foo() {   
printf("n merhaba!"); 
}

main.cu dosyası şu şekilde:

//-- main.cu --
#include 
extern __global__ void foo();
int main() {
  foo>>();
  cudaDeviceSynchronize();
  auto err = cudaGetLastError();
  printf("n cudaGetLastError() = %sn", cudaGetErrorString(err));
}

Bu programda, foo.cu paylaşılan kütüphane libfoo.so olarak derlenir ve ana program tarafından referans gösterilir:

$nvcc foo.cu -shared -o libfoo.so -Xcompiler -fPIC -rdc=true
$nvcc main.cu libfoo.so -o main -rdc=true

Program çalıştığında, beklenen satır (“merhaba”) yazılmayacak, ancak CUDA Runtime’dan hiçbir hata rapor edilmeyecek:

LD_LIBRARY_PATH=. ./main                                                                   
 cudaGetLastError() = hata yok

Altta yatan sorun, __global__ işlev çağrısı sırasının, hem main.cu‘da (şu şekilde: foo>>) hem de foo.cu‘daki ev sahibi kod stub işlevinde CUDA Runtime’a yapılan çağrıları içermesidir. Ancak, libfoo.so ve main programı için iki farklı CUDART kütüphanesi bağlı olduğu için, kernel çağrısı beklenildiği gibi çalışmayacaktır.

Örnek 2

foo.cu dosyası şu şekildedir:

//foo.cu
__managed__ int result = 20;

main.cu dosyası ise:

//main.cu
#include 
extern __managed__ int result;
int main() {
  printf("n sonuç = %d", result);
}

Bu örnekte de, foo.cu paylaşılan kütüphaneye libfoo.so olarak derlenecek ve ana programda referans gösterilecektir:

$nvcc foo.cu -shared -o libfoo.so -Xcompiler -fPIC -rdc=true
$nvcc main.cu libfoo.so -o main -rdc=true -g

Program çalıştırıldığında, result değerine erişmeye çalışırken bir segment hatasıyla karşılaşacaktır:

$LD_LIBRARY_PATH=. gdb -ex=r ./main
..
Thread 1 "main" received signal SIGSEGV, Segmentation fault.
0x000055555555cdaf in main () at main.cu:4
4         printf("n sonuç = %d", result);

Yine altta yatan sorun, libfoo.so ve main programında farklı CUDART kütüphanelerinin statik olarak bağlı olması ve result değişkeninin doğru bir şekilde başlatılmamasıdır.

Etkilenen Platformlar

CUDA 13.0 NVCC değişiklikleri bu bölümde açıklanan şekilde, Windows haricindeki tüm platformları etkilemektedir. Windows’taki ana derleyici aracı (şu şekilde: cl.exe) varsayılan olarak paylaşılan kütüphanelerden dışarı semboller çıkarmayı desteklemediği için, bu bölümde tanımlanan problemler gerçekleşmez.

CUDA 13.0’da Sunulan Çözüm

Kullanıcıları yukarıda tanımlanan sorunlardan korumak amacıyla, CUDA 13.0’dan itibaren NVCC, __global__ işlevlerinin ve __managed__/__device__/__constant__ değişkenlerin varsayılan görünürlüğünü gizli olarak değiştirecek ve bu semboller dışa aktarılmayacaktır.

Bu durum, yukarıdaki programların derlenmesini başarısız kılacak ve bu, bir çalışma zamanı hatası veya yanlış davranış yerine daha iyi bir senaryodur. Örneğin, __managed__ değişken örneğinde bir bağlantı hatası ortaya çıkacaktır:

$nvcc foo.cu -shared -o libfoo.so -Xcompiler -fPIC -rdc=true 
$nvcc main.cu libfoo.so -o main -rdc=true   -g 
...
/usr/bin/ld: /tmp/tmpxft_0032ad7a_00000000-11_main.o: in function `main':
/work/bugs/blogexamples/sharedlibrary/managed_var/main.cu:4: undefined reference to `result'

Doğruluk sorunları, paylaşım kütüphanesi sürümünü kullanarak (şu şekilde: -cudart=shared) hem ana programı hem de paylaşılan kütüphaneyi derleyerek önlenebilir. Bu yöntemde, programın tüm bölümlerinin aynı CUDART kullanması sağlanır. Ayrıca NVCC’nin gizli görünürlüğü zorlamasını devre dışı bırakmak için (şu şekilde: --device-entity-has-hidden-visibility=false) ve aşağıdaki gibi derleme işlemini yapmak gerekir:

$nvcc foo.cu -shared -o libfoo.so -Xcompiler -fPIC -rdc=true -cudart=shared -device-entity-has-hidden-visibility=false
$nvcc main.cu libfoo.so -o main -rdc=true   -g -cudart=shared -device-entity-has-hidden-visibility=false

CUDA 13.0’da Opt-Out (Üzerinden Geçme)

Bu değişikliğin NVCC varsayılan davranışındaki bu değişiklik bazı mevcut iş akışlarını bozabilir. Özellikle, sadece bir dinamik (paylaşılan) CUDART kütüphanesi kullanan ve bu kütüphaneyi hem paylaşılan kütüphanede hem de ana programda kullanan iş akışlarını etkilemesi muhtemeldir.

CUDA 13.0 ve sonrasında bu iş akışlarını desteklemeye devam etmek için, değiştirilen NVCC davranışından opt-out yapmak mümkündür. Opt-out yapmak için:

• NVCC komut satırına --device-entity-has-hidden-visibility=false ekleyin.
  • Bu, NVCC davranışını CUDA 13.0’dan önceki araç setlerininkine benzer hale getirir. Bu bayrak, CUDA 12.8’den bu yana mevcuttur, ancak varsayılan değeri CUDA 13.0’a kadar yanlışdı.
• __attribute__((visibility(“default”))) direktifini __global__ işlevinin veya __managed__/__device__/__constant__ değişkenin tanımı üzerine, çevreleyen namespace içinde veya #pragma GCC visibility kullanarak açıkça ekleyin.

Aşağıdaki örneğe göz atabilirsiniz:

__global__ __attribute__((visibility("default"))) void foo1() { }
namespace __attribute__((visibility("default"))) N1 { 
void foo2() { }
}
#pragma GCC visibility push(default)
__global__ void foo3() { }
#pragma GCC visibility pop

GCC ve Clang derleyici bayrağı -fvisibility, bu sembolleri etkilemez çünkü CUDA derleyicisi tıpkı yukarıda belirtildiği gibi, ana derleyiciye gönderilen kodda dekorasyonları attribute((visibility(“hidden”)) ile açık bir şekilde gösterir.

Pre-CUDA 13.0’da Opt-In (Katılım Sağlamak)

CUDA 12.8 ve sonrası için --device-entity-has-hidden-visibility=true bayrağı belirlenebilir ve bu, CUDA 13.0 ve sonrasında varsayılan durumdadır. Bu, gizli ELF görünürlüğünü __global__ işlevleri ve __managed__/__device__/__constant__ değişkenler için zorlar; yukarıda belirtilen opt-out mekanizmalarından biri kullanılmadığı sürece.

NVCC Değişikliği #2: Zorlanmış İç Bağlantı

CUDA programlama modelinde, __global__ işlevleri ana koddan başlatılabilir. NVCC, ana derleyiciye gönderilen kodda orijinal __global__ işlevini, CUDA’yı çalıştırmak için gerekli olan stub işleviyle değiştirir. Benzer stublar, __managed__/__device__/__constant__ değişkenler için de oluşturulmaktadır.

Problemin Genel Görünümü ve Örnekler

Birçok CUDA kütüphanesi (örneğin Thrust), header içerisinde bulunur ve içerir. Eğer iki farklı CUDA dosyası (a.cu ve b.cu) aynı başlık dosyalarını içeriyorsa ve tüm program modunda (-rdc=false) derleniyorsa, her çeviri birimi ayrı bir cihaz programı oluşturacak. Ancak, ana bağlayıcı, a.o ve b.o dosyalarının içindeki __global__ stub işlevlerini birleştirir.

Örnek 1

common.h dosyası şu şekildedir:

//common.h
#include 
__managed__ int result;
template 
__global__ void foo() { result = 1; }

a.cu dosyası ise:

//a.cu
#include "common.h"
int first() { 
  foo>>();       // HATA: b.cu'dan gelen
                             // foo işlevini yanlışlıkla başlatabilir!
  cudaDeviceSynchronize(); 
  return result; 
}

b.cu dosyası ise:

//b.cu
#include "common.h"
int first();
int main() {
  int val = first();
  assert(val == 1);           // assert başarısız olabilir!
  foo>>();        // HATA: a.cu'dan gelen
                              // foo işlevini yanlışlıkla başlatabilir!
  cudaDeviceSynchronize();
  assert(result == 1);        // assert başarısız olabilir!
}

Şu şekilde bir derleme yapılır:

$ nvcc  a.cu b.cu -o prog

Burada, hem a.cu hem de b.cufoo işlevini başlatır. Her iki dosya da (çalışma zamanında) result‘in kopyaları olacaktır ve başlatılan foo işlevi beklenmeyen durumlara neden olacaktır, bu nedenle assertion hatası meydana gelebilir.

Etkilenen Platformlar

Bu NVCC davranışının değişikliği, tüm platformları etkilemektedir. Ancak sadece tüm program derleme modu (-rdc=false) kullanıldığında, bu durum kullanılacak olan NVCC varsayılan modudur.

CUDA 13.0’da Sunulan Çözüm

CUDA 13.0’dan itibaren, NVCC __global__ işlev şablonu stub işlevlerinin iç bağlantıya sahip olmasını zorlayacaktır. Böylece çeviri birimi dosyaları ayrı olarak işlenir ve ana bağlantı sırasında birleştirilmez.

CUDA 13.0’da Opt-Out (Üzerinden Geçme)

CUDA 13.0’daki bu değişiklik, bazı yasal mevcut programları bozabilir. Özellikle, bir __global__ işlev şablonunun bir çeviri biriminde açıkça örneklendiği ve başka bir çeviri biriminde referans gösterildiği durumlarda problem çıkabilir.

//first.cu
template 
__global__ void foo() { }

template
__global__ void foo(); // açık örnekleme

// second.cu
template 
__global__ void foo(); // first.cu'da açıktan örnekleme yapılmış.

int main() { foo>>(); cudaDeviceSynchronize(); }

Bu kod, CUDA 13.0 ve sonraki sürümlerde artık bağlantı yapmayacaktır çünkü ev sahibi stub foo için bağlantı içeriği içsel olacaktır ve bu yüzden second.o‘daki foo referansı, ana bağlayıcı tarafından çözülmeyecektir.

$nvcc first.cu second.cu -o prog
/usr/bin/ld: /tmp/tmpxft_0032b262_00000000-18_second.o: in function `main':
tmpxft_0032b262_00000000-10_second.cudafe1.cpp:(.text+0xdb): undefined reference to `void foo()'

Bununla beraber -static-global-template-stub=false bayrağını kullanmak, eski NVCC davranışını geri yükleyecek ve yukarıdaki programın derlenmesine olanak sağlayacaktır. Bu bayrak, CUDA 12.8’den itibaren kullanılabilir, ama varsayılan değerinin CUDA 13.0 ile değiştirileceği düşünülmektedir.

Pre-CUDA 13.0’da Opt-In (Katılım Sağlamak)

Opt-in yapmak için, CUDA 12.8’den itibaren kullanılabilir olan -static-global-template-stub=true bayrağı eklenebilir.

Özet

CUDA 12.8 ve sonrası NVCC bayrakları, bazı uzun süreli ve ince çalışma zamanı hatalarını önlemenize yardımcı olur. Onları kodunuzda kullanmayı göz önünde bulundurmanız önemlidir. CUDA 13.0’daki bayrak varsayılanlarının değişmesi, bazı mevcut CUDA C++ kodları için derleme veya bağlantı hatalarına neden olabilir. Kodunuzda güncellemeler yapmak veya yeni bayrakları kullanarak NVCC davranışından açıkça opt-out yapmak için yolları düşünmelisiniz.

Teşekkürler

Aşağıdaki NVIDIA katkıcılarına teşekkür ederiz: Chu-Cheow Lim, Jonathan Bentz ve Tony Scudiero.