Sabit argümanları bir CUDA çekirdeğine geçirmenin en hızlı (veya en zarif) yolu

Question

Bir çok şey yapması gereken bir CUDA çekirdeği olmasını istediğimi söyleyebilirim, ancak tüm çekirdekler için sabit olan kubbe parametreleri var. Bu argümanlar bir giriş olarak ana programa geçirilir, böylece #DEFINE'da tanımlanamazlar.

Çekirdek, birden çok kez (yaklaşık 65 K) çalışır ve matematiğini yapmak için bu parametrelere (ve bazı diğer girdilere) gereksinim duyar.

Sorum şu: Bu sabitlerin çekirdeklere aktarılmasının en hızlı (veya en zarif) yolu nedir?

Sabitler 2 veya 3 öğe uzunluğu float* veya int* dizileridir. Bunların yaklaşık 5 ~ 10'u olacaklar.

---

oyuncak Örnek: 2 sabitleri const1 ve const2

__global__ void kernelToyExample(int inputdata, ?????){ 
     value=inputdata*const1[0]+const2[1]/const1[2]; 
} 

daha iyi

__global__ void kernelToyExample(int inputdata, float* const1, float* const2){ 
     value=inputdata*const1[0]+const2[1]/const1[2]; 
} 

veya

__global__ void kernelToyExample(int inputdata, float const1x, float const1y, float const1z, float const2x, float const2y){ 
     value=inputdata*const1x+const2y/const1z; 
} 

olanya da belki bazı küresel salt okunur bellekte bunları ilan edip çekirdekleri oradan okuyalım mı? Öyleyse, L1, L2, global? Hangisi?

Bilmediğim daha iyi bir yolu var mı?

Bir Tesla K40 üzerinde çalışıyor.

Answer 1

Sadece bunlara değer verin. Derleyici, her bloktaki tüm iş parçacıklarına önbelleğe alınmış yayını kolaylaştırmak için bunları otomatik olarak en uygun yere yerleştirir - bilgi işlem biriminde paylaşımlı bellek 1.x aygıtları veya sabit bellek/sabit önbellek> 100 aygıtlarda.

struct arglist { 
    float magicfloat_1; 
    float magicfloat_2; 
    //...... 
    float magicfloat_19; 
    int magicint1; 
    //...... 
}; 

__global__ void kernel(...., const arglist args) 
{ 
    // you get the idea 
} 

[standart uyarı:: tarayıcıda yazılı, gerçek değil çekirdekten verilecek argüman uzun bir liste vardı Örneğin

, değer geçirilecek bir yapı gitmek için temiz bir yoldur

template<int magiconstant1> 
__global__ void kernel(....) 
{ 
    for(int i=0; i < magconstant1; ++i) { 
     // ..... 
    } 
} 

template kernel<3>(....); 
template kernel<4>(....); 
template kernel<5>(....); 

: aslında sizin magicint birini ortaya çıktı ise kod, sorumluluğun alıcıya]

sadece önceden bilmek değerlerin az sayıda birini aldı, daha sonra templating son derece güçlü bir araçtır

Derleyici, magconstant tanımayı tanımak için yeterince akıllıdır, döngü çağrısı derleme zamanında bilinir ve otomatik olarak döngüyü sizin için açacaktır. Templating, hızlı, esnek kodlar oluşturmak için bir very powerful technique'dur ve daha önce yapmadıysanız kendinize alışmanız tavsiye edilir.