Appeler un noyau à partir d'un noyau

Question

Un suivi Q de: CUDA: Calling a __device__ function from a kernel

Je suis en train d'accélérer une opération de tri. Une version pseudo simplifiée suit:

// some costly swap operation 
__device__ swap(float* ptrA, float* ptrB){ 
    float saveData;   // swap some 
    saveData= *Adata;  // big complex 
    *Adata= *Bdata   //  data chunk 
    *Bdata= saveData; 
} 

// a rather simple sort operation 
__global__ sort(float data[]){ 
    for (i=0; i<limit: i++){ 
    find left swap point 
    find right swap point 
    swap<<<1,1>>>(left, right); 
    } 
} 

(Note:. Cette version simple ne montre pas les techniques de réduction dans les blocs) L'idée est qu'il est facile (rapide) pour identifier les points d'échange. L'opération d'échange est coûteuse (lente). Utilisez donc un bloc pour trouver/identifier les points d'échange. Utilisez d'autres blocs pour effectuer les opérations d'échange. C'est-à-dire que l'échange est effectué en parallèle. Cela ressemble à un plan décent. Mais si le compilateur inline les appels de l'appareil, il n'y a pas d'échange parallèle. Existe-t-il un moyen de dire au compilateur de ne PAS entrer en ligne un appel de périphérique?

Answer 1

Edit (2016):

parallélisme dynamique a été introduit dans la deuxième génération de GPU d'architecture Kepler. Le lancement des noyaux dans le périphérique est pris en charge sur les périphériques de capacité de calcul 3.5 et supérieure.

---

Réponse d'origine:

Vous devrez attendre la fin de l'année où la prochaine génération de matériel est disponible. Aucun périphérique CUDA actuel ne peut lancer des noyaux à partir d'autres noyaux - il n'est actuellement pas supporté.

Answer 2

Je sais que cela fait longtemps que cette question a été posée. Quand j'ai googlé le même problème, je suis arrivé à cette page. On dirait que j'ai la solution.

Solution:

je suis arrivé en quelque sorte here et vu l'approche cool de lancer le noyau à partir d'un autre noyau.

__global__ void kernel_child(float *var1, int N){ 
    //do data operations here 
} 


__global__ void kernel_parent(float *var1, int N) 
{ 
    kernel_child<<<1,2>>>(var1,N); 
} 

Le parallélisme dynamique sur cuda 5.0 et plus a rendu cela possible. Aussi en cours d'exécution assurez-vous que vous utilisez l'architecture de calcul_35 ou ci-dessus.

Chemin du terminal Vous pouvez exécuter le noyau parent ci-dessus (qui finira par exécuter le noyau enfant) à partir de termial. Vérifié sur une machine Linux.

$ nvcc -arch=sm_35 -rdc=true yourFile.cu 
$ ./a.out 

Espérons que ça aide. Je vous remercie!