Derleyici ile AMD OpenCL çekirdeğinde döngüler açmanın bir yolu var mı?

Question

AMD ve Nvidia GPU'lar için OpenCL arasındaki performans farklılıklarını değerlendirmeye çalışıyorum. Matris-vektör çarpımını gerçekleştiren bir çekirdeğim var. Çekirdeği iki farklı sistemde çalıştırıyorum, dizüstü bilgisayarım Ubuntu 12.04 ve CUDA 4.0 (OpenCL kitaplıklarını ve başlıklarını içeren) bir NVidia GT525m ve diğeri ise Ubuntu ile tekrar AMD Radeon HD7970 ile bir masaüstü. 12.04 ve en son Catalyst sürücüleri.

Nvidia OpenCL uygulaması için (~ 6x) büyük bir hız artışı sağlayan iki #pragma unroll deyiminde çekirdeğim var. Ancak AMD OpenCL sürümü herhangi bir hızlanma üretmiyor. AMD APP çekirdek analizörü ile çekirdeğe bakıldığında, açma sayısının bilinmediği için açma işleminin kullanılmadığı hatası verilir. Yani benim sorum şu ki, #pragma unroll AMD OpenCL ile çalışıyor mu yoksa bir alternatif var mı (belki de habersiz olduğum bir derleyici bayrağı). Ben (onlara yorum yaparak kontrol)

__kernel void mvKernel(__global float* a, const __global float* x, __global float* y, int m, int n) 
{ 
    float sum = 0.0f; 
    __global float* A; 
    int i; 
    int j = 0; 
    int indx = get_global_id(0); 
    __local float xs[12000]; 
#pragma unroll 
    for(i = get_local_id(0); i < n; i+= get_local_size(0)) { 
     xs[i] = x[i]; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 
    A = &a[indx]; 
#pragma unroll 256 
    for(i = 0; i < n; i++) { 
     sum += xs[i] * A[j]; 
     j += m; 
    } 
    y[indx] = sum; 
} 

Bu aynı çekirdek hem uygulamalarda doğru sonuçlar üretir ama #pragma göz önüne sermek komutları AMD için bir şey yapmayın aşağıda çekirdek yer verdik.

Answer 1

Bu belge belgelenmemiş, ancak aslında #pragma unroll ile çalışmalıdır. Kayıt uygulamasının uygulanıp uygulanmadığını görmek için derleyici günlüğünü kontrol edebilir misiniz? Çekirdek analizörünün OpenCL çalışma zamanı ile aynı derleyiciyi kullanıp kullanmadığını kontrol edemiyorum, kontrol etmek isteyebilirsiniz.

Aksi takdirde, n'un 256 parça kümesine geldiğini biliyorsanız, 256 öğeden oluşan blokların üzerinde bir döngü ve 256 bir sabit boyutta bir tane daha içeri girerek el ile açabilirsiniz; bu da daha kolay olabilir. Bu kesinlikle, yolculuk sayımının statik olarak bilinmediği problemi çözecektir. Bununla birlikte, bir döngüyü açmanızın genellikle bir kazanın büyüklüğü olmadığı göz önünde bulundurulduğundan, hesaplamayı önbelleğe almak için çok fazla yazmaçınız olmadığı için, unutmayın. Döngünün açılmasından kaynaklanan artan kayıt basıncı, daha da yavaş olan döküntünün dökülmesine yol açabilir. Çekirdeğin gerçekte AMD kartında ne kadar hızlı olduğunu kontrol etmelisiniz. Daha yeni bir NVIDIA OpenCL derleyicisi de artık pragma'dan faydalanmayabilir.