CUDA Kernel Performansını Geliştirme: Paylaşılan Bellek ve Kayıt Taşması İle İpuçları

CUDA çekirdekleri, kullanılabilir donanım kayıtlarından daha fazla kayıt gerektirdiğinde, derleyici fazla değişkenleri yerel belleğe taşımak zorunda kalır; bu işleme kayıt taşma denir. Kayıt taşması performansı etkiler çünkü çekirdek, taşınan verilere erişmek için yerel bellek ile iletişim kurmak zorundadır; yerel bellek ise fiziksel olarak global bellekte yer alır.

CUDA Toolkit 13.0, NVIDIA tarafından CUDA çekirdekleri için yeni bir optimizasyon özelliği tanıtıldı: paylaşılan bellek kayıt taşması. Bu yazı, yeni özelliği, eklenme motivasyonunu ve nasıl etkinleştirileceğini açıklamaktadır. Ayrıca, bu özelliği kullanmak için ne zaman düşünmeniz gerektiği ve potansiyel etkisini nasıl değerlendirebileceğiniz konusunda rehberlik sunmaktadır.

Paylaşılan Bellek Kayıt Taşması Nasıl Performansı Optimize Eder?

CUDA 13.0 ile birlikte, PTXAS derleyicisinde paylaşılan belleğe kayıt taşması desteği eklenmiştir. Bu özellik etkinleştirildiğinde, derleyici kayıtları paylaşılan belleğe taşımayı öncelikli olarak tercih eder. Eğer yeterli paylaşılan bellek yoksa, kalan taşmalar yerel belleğe yönlendirilir; böylece önceki davranışla tutarlılık sağlanır. Bu değişiklik, taşınan değerleri depolamak için daha düşük gecikmeli, iç bellek avantajından faydalanarak performansı artırır.

Sorunun Genel Görünümü ve Örnek

CUDA 13.0’dan önceki araçlarda, tüm kayıt taşmaları yerel belleğe yerleştiriliyordu; bu, off-chip, yani cihazların global belleklerinde yer alıyordu. Daha büyük L1 önbellek boyutları, birçok uygulamanın taşma maliyetini azaltmaya yardımcı oluyordu. Ancak, taşınan verilerin L2 belleğine yazılabilmesi önemli önbellek satırlarının yerinden edilmesine sebep olabiliyordu ve bu durum, genel performansı olumsuz etkiliyordu. Etki, özellikle yüksek kayıt basıncı olan, performans açısından kritik bölgelerde, döngülerde ve sık yürütülen kod bölümlerinde daha belirgin hale geliyordu.

Birçok iş yükünde, çalıştırma zamanı boyunca paylaşılan belleğin önemli bir kısmı genelde kullanılmıyordu. Bu durum, paylaşılan bellek gereksinimlerinin düşük olması ya da çekirdeğin doluluğu maksimize edecek şekilde tasarlanmamış olması gibi nedenlerle ortaya çıkabiliyordu. Örneğin, bir SM başına iş parçacığı bloğu sayısı, paylaşılan bellek kullanımı yerine yükleme sınırları ya da kayıt basıncı ile sınırlı olduğunda, her blok gerekli olandan daha fazla paylaşılan bellek tahsis edilmiş olabiliyordu. Bu durumda, fazla tahsis edilen paylaşılan bellek boşa gidebiliyordu.

Aşağıdaki kod örneği, yeterli sayıda kayıt kullanan bir çekirdek tasarımıdır; her satırın anlaşılmasını sağlamak amacıyla yüksek kayıt kullanımı sağlanmıştır.

/-- main.cu --
#include <cuda_runtime.h>
#include <stdio.h>

extern "C" __launch_bounds__(256)
__global__ void foo(float *output_tensor, int num_elements) {

    int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
    if (thread_id >= num_elements) return;

    volatile float input_feature[89], weight_scaled[89], bias_added[89], pre_activation[89];
    volatile float activation_sin[89], activation_cos[89], output_accum[89];

    #pragma unroll
    for (int i = 0; i < 89; ++i) {
        input_feature[i] = (float)thread_id + i;
        weight_scaled[i] = input_feature[i] * 2.0f;
        bias_added[i] = 5 + weight_scaled[i];
        activation_sin[i] = __sinf(bias_added[i] * pre_activation[i]);
        activation_cos[i] = __cosf(activation_sin[i % 2] + pre_activation[i]);
        float product = input_feature[i] * weight_scaled[i];
        float squared = product * product;
        float biased = squared + bias_added[i % 4];
        float shifted_sin = __sinf(biased * 0.5f);
        float shifted_cos = __cosf(shifted_sin + 1.0f);
        float amplified = shifted_cos * bias_added[i % 5];
        float combined = amplified + activation_cos[i];
        output_accum[i] = combined;
    }
    volatile float sum = 0.0f;
    #pragma unroll
    for (int i = 0; i < 89; ++i) {
        sum += input_feature[i] + weight_scaled[i] + bias_added[i] + pre_activation[i]
             + activation_sin[i] + activation_cos[i] + output_accum[i];
    }

    output_tensor[thread_id] = sum;
}

int main() {
    const int num_elements = 896;
    const int ARRAY_BYTES = num_elements * sizeof(float);
    float host_output[num_elements];
    float *device_output;

    cudaMalloc(&device_output, ARRAY_BYTES);

    const int blockSize = 256;
    const int gridSize = (num_elements + blockSize - 1) / blockSize;

    foo<<<gridSize, blockSize>>>(device_output, num_elements);
    cudaDeviceSynchronize();

    cudaMemcpy(host_output, device_output, ARRAY_BYTES,   cudaMemcpyDeviceToHost);

    for (int i = 0; i < num_elements; ++i) {
        printf("host_output[%d] = %fn", i, host_output[i]);
    }

    cudaFree(device_output);
    return 0;
}

nvcc -arch=sm_90 -Xptxas -v main.cu

Bu program normal bir şekilde derlendiğinde (paylaşılan bellek kayıt taşması belirtmeden), çıktı aşağıdaki gibidir:

ptxas info    : Compiling entry function 'foo' for 'sm_90'
ptxas info    : Function properties for foo
    176 bytes stack frame, 176 bytes spill stores, 176 bytes spill loads
ptxas info    : Used 255 registers, used 0 barriers, 176 bytes cumulative stack size

Çıktıda “taşmaların” depolamaları ve yüklemeleri gösterilmektedir. Bu, kayıtların yerel belleğe taşınacağını belirtmektedir.

Ayrıca bu örnekte, derlenmiş çekirdek herhangi bir paylaşılan belleği kullanmamaktadır; bu durum, blok başına paylaşılan bellek tahsisinin tamamen boşa gittiğini göstermektedir.

CUDA 13.0’da Tanıtılan Ek Bir Çözüm Nedir?

Yüksek kayıt sınırına sahip çekirdeklerde performansı artırmak için, CUDA 13.0, kayıt taşmalarının yerel belleğe yönlendirilmesi yerine paylaşılan belleğe yönlendirilmesini sağlayan yeni bir optimizasyon tanıtmıştır. On-chip paylaşılan bellek kullanarak, derleyici taşınan veriyi streaming multiprocessorın yakınında tutar ve bu da erişim gecikmesini önemli ölçüde azaltır, L2 bellek üzerindeki baskıyı hafifletir. Bu gelişme, yerel belleğe taşma yapmanın genellikle bir darboğazı oluşturduğu senaryolar için önemli performans artışları sağlar.

Optimizasyon etkinleştirildiğinde, derleyici önce mevcut paylaşılan bellek alanına taşmalar yapmaya çalışır, yeterli alan yoksa yerel belleğe düşer; bu, programların doğruluğunu korumaktadır.

Önceki çekirdek, paylaşılan bellek kayıt taşmasını etkinleştirerek derlendiğinde, çıktı şu şekildedir:

ptxas info    : Compiling entry function 'foo' for 'sm_90'
ptxas info    : Function properties for foo
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 255 registers, used 0 barriers, 46080 bytes smem

Önceki örnekle karşılaştırıldığında, burada herhangi bir taşma olmadığı görülmektedir. Paylaşılan belleğin kullanımı 46080 bytes smem ifadesi ile belirtilmektedir.

Nsight Compute Sonuçları

Bu optimizasyonun değerini göstermek amacıyla, yukarıda gösterilen CUDA çekirdeği, paylaşılan bellek taşma özelliği etkinleştirilmeden ve etkinleştirildiğinde benchmark edilmiştir. Tablo 2, paylaşılan bellek kayıt taşması optimizasyonu etkinleştirildiğinde çekirdek performansı karşılaştırmasını ve süre, geçen döngüler ve SM aktif döngüleri gibi üç ana ölçütün değişimini göstermektedir.

Paylaşılan Bellek Kayıt Taşmasına Nasıl Geçiş Yapılır?

Paylaşılan bellek kayıt taşması özelliği, CUDA 13.0’da tanıtılmıştır ve önceki sürümlerde mevcut değildir. CUDA 13.0 ve daha sonraki sürümleri hedefleyen geliştiricilerin, enable_smem_spilling PTX pragma’sını montaj yoluyla açıkça etkinleştirmesi gerekmektedir. Bunun için, işlev tanımının hemen ardından inline montajda belirtebilirsiniz:

#include <cuda_runtime.h>
#include <stdio.h>

extern "C" __launch_bounds__(256)
__global__ void foo(float *output_tensor, int num_elements) {
    asm volatile (".pragma "enable_smem_spilling";");

    int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
    if (thread_id >= num_elements) return;
    volatile float input_feature[89], weight_scaled[89], bias_added[89], pre_activation[89];
    volatile float activation_sin[89], activation_cos[89], output_accum[89];
    #pragma unroll
    for (int i = 0; i < 89; ++i) {
        input_feature[i] = (float)thread_id + i;
        weight_scaled[i] = input_feature[i] * 2.0f;
        bias_added[i] = 5 + weight_scaled[i];
        activation_sin[i] = __sinf(bias_added[i] * pre_activation[i]);
        activation_cos[i] = __cosf(activation_sin[i % 2] + pre_activation[i]);

        float product = input_feature[i] * weight_scaled[i];
        float squared = product * product;
        float biased = squared + bias_added[i % 4];
        float shifted_sin = __sinf(biased * 0.5f);
        float shifted_cos = __cosf(shifted_sin + 1.0f);
        float amplified = shifted_cos * bias_added[i % 5];
        float combined = amplified + activation_cos[i];
        output_accum[i] = combined;
    }
    volatile float sum = 0.0f;

    #pragma unroll
    for (int i = 0; i < 89; ++i) {
        sum += input_feature[i] + weight_scaled[i] + bias_added[i] + pre_activation[i]
             + activation_sin[i] + activation_cos[i] + output_accum[i];
    }
    output_tensor[thread_id] = sum;
}

Paylaşılan Bellek Kayıt Taşmasının Sınırları

Bu optimizasyon, cihaz kodu ve PTXAS derleyicisi için bir performans fırsatı sunuyor; ancak önemli kısıtlamaları da beraberinde getiriyor. Sadece işlev kapsamı içinde geçerlidir ve aşağıdaki durumlarda kullanılmamalıdır, aksi halde derleme hataları oluşabilir:

• Per-function derleme modu, örneğin nvcc -rdc=true ya da ptxas -c, nvcc -G veya ptxas -g, nvcc -ewp. Unutmayın ki cihaz-debug derleme modları (nvcc -G ya da ptxas -g) de per-function derlemeyi ifade eder.
• Dynamik paylaşılan bellek kullanan çekirdekler.
• Warplar arasında dinamik kayıt yeniden tahsisi yapan çekirdekler.

Eğer yükleme sınırları açıkça belirtilmezse, PTXAS tahminlerde, mümkün olan maksimum iş parçacığı sayısını dikkate alır. Eğer çekirdek, tahmin edilenden daha az iş parçacığı ile başlatılırsa, blok başına tahsis edilen gerçek paylaşılan bellek gereksinimlerden fazla olabilir; bu da birlikte çalışan iş parçacığı bloklarının sayısını azaltabilir. Bu durum, doluluğu azaltabilir ve performans geri dönüşlerine yol açabilir. Daha öngörülebilir bir davranış ve daha iyi performans için, bu özelliğin sadece yükleme sınırları açıkça tanımlandığında kullanılması önerilir.

Gerçek İş Yüklerinde Olası Performans Artışları

Bu optimizasyon, CUDA ile Lattice QCD hesaplamaları için kullanılan QUDA kütüphanesinden çeşitli CUDA çekirdekleri üzerinde değerlendirildi. Optimizasyon, genellikle 5-10% arasında performans artışları sağladı. Bu iyileşmeler, paylaşılan bellek yerine yerel belleğe kayıt taşmalarının azaltılmasından ya da tamamen ortadan kaldırılmasından kaynaklanmaktadır.

Paylaşılan Bellek Kayıt Taşması Optimizasyonu ile Başlayın

CUDA 13.0, yüksek kayıt baskısı olan çekirdeklerde kayıt taşmalarının etkili bir şekilde yönetilmesini sağlayan bir PTXAS optimizasyonu sunmaktadır. Eğer CUDA çekirdeğinizin iyi tanımlanmış yükleme sınırları ve tutarlı bir paylaşılan bellek kullanımı varsa, iç montajda enable_smem_spilling kullanarak paylaşılan bellek taşmasını denemenizi tavsiye ederiz.

Teşekkürler

Aşağıdaki NVIDIA katkıda bulunanlarına teşekkür ederiz: Jerry Zheng, Kate Clark, Howard Chen, Neumann Hon, Jaewook Shin, Abhishek Patwardhan ve Yufan Cheng.