CUDA noyau lancement paramètres expliqué droit?

Ici, j'ai essayé de l'auto-expliquer le CUDA de lancement des paramètres du modèle (ou l'exécution de la configuration du modèle) à l'aide de pseudo-codes, mais je ne sais pas si il y avait quelques grand erreurs, Donc j'espère que quelqu'un l'aide à la revoir, et me donner quelques conseils. Grâce avancé.

Ici, il est:

/*
normally, we write kernel function like this.
note, __global__ means this function will be called from host codes,
and executed on device. and a __global__ function could only return void.
if there's any parameter passed into __global__ function, it should be stored
in shared memory on device. so, kernel function is so different from the *normal*
C/C++ functions. if I was the CUDA authore, I should make the kernel function more
different  from a normal C function.
*/
__global__ void
kernel(float *arr_on_device, int n) {
int idx = blockIdx.x * blockDIm.x + threadIdx.x;
if (idx < n) {
arr_on_device[idx] = arr_on_device[idx] * arr_on_device[idx];
}
}
/*
after this definition, we could call this kernel function in our normal C/C++ codes !!
do you feel something wired ? un-consistant ?
normally, when I write C codes, I will think a lot about the execution process down to
the metal in my mind, and this one...it's like some fragile codes. break the sequential
thinking process in my mind.
in order to make things normal, I found a way to explain: I expand the *__global__ * function
to some pseudo codes:
*/
#define __foreach(var, start, end) for (var = start, var < end; ++var)
__device__ int
__indexing() {
const int blockId = blockIdx.x * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
return 
blockId * (blockDim.x * blockDim.y * blockDim.z) +
threadIdx.z * (blockDim.x * blockDim.y) +
threadIdx.x;
}
global_config =:
{
/*
global configuration.
note the default values are all 1, so in the kernel codes,
we could just ignore those dimensions.
*/ 
gridDim.x = gridDim.y = gridDim.z = 1;
blockDim.x = blockDim.y = blockDim.z = 1;
};
kernel =:
{
/*
I thought CUDA did some bad evil-detail-covering things here.
it's said that CUDA C is an extension of C, but in my mind,
CUDA C is more like C++, and the *<<<>>>* part is too tricky.
for example:
kernel<<<10, 32>>>(); means kernel will execute in 10 blocks each have 32 threads.
dim3 dimG(10, 1, 1);
dim3 dimB(32, 1, 1);
kernel<<<dimG, dimB>>>(); this is exactly the same thing with above.
it's not C style, and C++ style ? at first, I thought this could be done by
C++'s constructor stuff, but I checked structure *dim3*, there's no proper
constructor for this. this just brroke the semantics of both C and C++. I thought
force user to use *kernel<<<dim3, dim3>>>* would be better. So I'd like to keep
this rule in my future codes.
*/
gridDim  = dimG;
blockDim = dimB;
__foreach(blockIdx.z,  0, gridDim.z)
__foreach(blockIdx.y,  0, gridDim.y)
__foreach(blockIdx.x,  0, gridDim.x)
__foreach(threadIdx.z, 0, blockDim.z)
__foreach(threadIdx.y, 0, blockDim.y)
__foreach(threadIdx.x, 0, blockDim.x)
{
const int idx = __indexing();        
if (idx < n) {
arr_on_device[idx] = arr_on_device[idx] * arr_on_device[idx];
}
}
};
/*
so, for me, gridDim & blockDim is like some boundaries.
e.g. gridDim.x is the upper bound of blockIdx.x, this is not that obvious for people like me.
*/
/* the declaration of dim3 from vector_types.h of CUDA/include */
struct __device_builtin__ dim3
{
unsigned int x, y, z;
#if defined(__cplusplus)
__host__ __device__ dim3(unsigned int vx = 1, unsigned int vy = 1, unsigned int vz = 1) : x(vx), y(vy), z(vz) {}
__host__ __device__ dim3(uint3 v) : x(v.x), y(v.y), z(v.z) {}
__host__ __device__ operator uint3(void) { uint3 t; t.x = x; t.y = y; t.z = z; return t; }
#endif /* __cplusplus */
};
typedef __device_builtin__ struct dim3 dim3;

Il existe de nombreuses classes disponibles, y compris l'introduction de celles que vous pouvez regarder n'importe quel moment et à seulement 1 heure, ici. Essayez le GPU Computing using CUDA C de la série.
avez-vous lu ce que j'ai expliqué ci-dessous ?
Vous semblez être assez confus sur les grilles et les blocs sans parler de l'architecture sous-jacente qui est sacrément important. J'ai signé pour la ligne gratuit CUDA cours à udacity et a obtenu efficace (si pas vraiment avancé) codage en une semaine. Check it out parce que la programmation sur GPU semble avoir besoin de fondations solides.
J'ai vérifié le cours que vous avez suggéré, en particulier la partie expliquant le noyau de la configuration de lancement udacity.com/course/viewer#!/c-cs344/l-55120467/m-67074291 (commence dans ce clip, et après 3 clips), l'enseignant n'est toujours pas clairement comment le GPU distribue les tâches aux threads. Regardez cette fonction noyau, comment le compilateur élargir les codes ? Donc, beaucoup de gens ont juste dit de vous rappeler la règle, mais la règle a été mise en œuvre dans les codes ? Alors maintenant, je suis en train de lire le conducteur de l'API et OpenCL documents, veulent juste savoir ce qui s'est passé sous le capot.
Un langage de haut niveau sont faits de sorte que vous N'avez PAS à vous soucier de la façon dont le code est intégré à la machine des instructions sur le CPU. Et à chaque nouvelle génération de PROCESSEUR apporte quelque chose de nouveau, même si le jeu d'instruction est la même, ce qui se passe RÉELLEMENT est pas nécessairement la même chose sur des architectures différentes.

OriginalL'auteur blackball | 2013-10-08

c cuda

11

CUDA DRIVER API

Le CUDA Driver API v4.0 et ci-dessus utilise les fonctions suivantes pour le contrôle d'un noyau de lancement:
```
cuFuncSetCacheConfig
cuFuncSetSharedMemConfig
cuLaunchKernel
```
Suivantes CUDA Driver fonctions de l'API ont été utilisés avant l'introduction de cuLaunchKernel en v4.0.
```
cuFuncSetBlockShape()
cuFuncSetSharedSize()
cuParamSet{Size,i,fv}()
cuLaunch
cuLaunchGrid
```
Plus d'informations sur ces fonctions peuvent être trouvés dans cuda.h.
```
CUresult CUDAAPI cuLaunchKernel(CUfunction f,
unsigned int gridDimX,
unsigned int gridDimY,
unsigned int gridDimZ,
unsigned int blockDimX,
unsigned int blockDimY,
unsigned int blockDimZ,
unsigned int sharedMemBytes,
CUstream hStream,
void **kernelParams,
void **extra);
```
cuLaunchKernel prend comme paramètres l'ensemble de la configuration de lancement.

Voir le Pilote NVIDIA API[Contrôle de l'Exécution]Un pour plus de détails.

CUDA NOYAU LANCEMENT

cuLaunchKernel va
1. vérifiez les paramètres de lancement
2. modifier la configuration de la mémoire partagée
3. modifier le local de l'allocation de mémoire
4. pousser un flux de synchronisation jeton dans la mémoire tampon de commande pour s'assurer de deux commandes dans le cours d'eau ne se chevauchent pas
4. poussez le lancement des paramètres dans la mémoire tampon de commande
5. poussez la commande de lancement dans la mémoire tampon de commande
6. envoyez la commande de la mémoire tampon de l'appareil (sur les pilotes wddm cette étape peut être reportée)
7. sur wddm le noyau pilote page tous les de mémoire dans la mémoire de l'appareil

Le GPU sera
1. vérifier la commande
2. envoyer les commandes pour les tâches de calcul distributeur
3. envoi de la configuration de lancement et de thread se bloque au SMs

Lorsque tous les threads blocs ont terminé le travail de distributeur de vider les caches à l'honneur le CUDA de la mémoire de modèle et il marquera le noyau comme achevé pour l'élément suivant dans le flux peut progresser.

L'ordre que le thread se bloque sont distribué diffère entre les architectures.

De calcul de la capacité 1.x dispositifs de stocker les paramètres du noyau dans la mémoire partagée.
Calculer la capacité de 2.0-3.5 dispositifs de stocker les kenrel des paramètres dans la mémoire constante.

CUDA RUNTIME API

Le CUDA Runtime C++ bibliothèque de logiciels et outil de construction de la chaîne sur le dessus de la CUDA Driver API. Le CUDA Runtime utilise les fonctions suivantes pour le contrôle d'un noyau de lancement:

cudaConfigureCall
cudaFuncSetCacheConfig
cudaFuncSetSharedMemConfig
cudaLaunch
cudaSetupArgument

Voir NVIDIA Exécution de l'API[Contrôle de l'Exécution]Deux

La <<<>>> CUDA extension du langage est la méthode la plus utilisée pour le lancement d'un noyau.

Lors de la compilation nvcc permettra de créer un nouveau PROCESSEUR stub fonction pour chaque fonction noyau appelé à l'aide de <<<>>>, et il permettra de remplacer les <<<>>> avec un appel à la stub de la fonction.

Par exemple
```
__global__ void kernel(float* buf, int j)
{
//...
}
kernel<<<blocks,threads,0,myStream>>>(d_buf,j);
```
génère
```
void __device_stub__Z6kernelPfi(float *__par0, int __par1){__cudaSetupArgSimple(__par0, 0U);__cudaSetupArgSimple(__par1, 4U);__cudaLaunch(((char *)((void ( *)(float *, int))kernel)));}
```
Vous pouvez inspecter les fichiers générés par l'ajout de --gardez votre nvcc ligne de commande.

cudaLaunch appels cuLaunchKernel.

CUDA DYNAMIQUE PARALLÉLISME

CUDA CDP fonctionne de manière similaire à l'API CUDA Runtime décrit ci-dessus.

Merci!!!! C'est utile pour moi...c'est presque la seule réponse utile ici.

OriginalL'auteur Greg Smith
3

En utilisant <<<...>>>, vous êtes le lancement d'un certain nombre de fils dans le GPU. Ces fils sont regroupés en blocs et les formes d'une grande grille. Tous les threads d'exécuter la invoquée noyau de la fonction de code.

Dans le noyau de la fonction, de construire des variables comme threadIdx et blockIdx activer le code de savoir quel thread, il s'exécute et ne la date prévue de la partie de l'œuvre.

modifier

Fondamentalement, <<<...>>> simplifie la procédure de configuration pour le lancement d'un noyau. Sans l'utiliser, on peut avoir à appeler 4~5 Api pour un seul noyau de lancement, tout comme la prise en charge d'OpenCL façon, qui n'utilisent que le C99 syntaxe.

En fait, vous pourriez vérifier CUDA driver Api. Il peut fournir toutes les Api de sorte que vous n'avez pas besoin d'utiliser <<<>>>.

J'ai lu que le doc il y a longtemps, et quelques autres. Je pensais que CUDA cassé la sémantique de C et de C++ à la kernel de la fonction de la partie, ce qui le rend difficile à comprendre et faciles à confondre. Donc, vous avez dit En utilisant <<<...>>>, le lancement de..., oui, vous avez raison, comme tous ceux de document dit, mais en expliquant CUDA exécution de la configuration comme ça, c'est tout simplement trop mauvais. Que tu viens de raconter comment, mais pas POURQUOI, oh, non, tu viens de dire ce, comment.
J'ai édité la réponse. Vous souhaiterez peut-être réviser votre question en ajoutant le POURQUOI et comment, pour éviter tout malentendu, de votre question tous le public.

OriginalL'auteur kangshiyin
1

Fondamentalement, le GPU est divisée en "périphérique" Gpu (par exemple la GeForce 690 a 2) -> plusieurs SM (streaming multiprocesseurs) -> plusieurs CUDA cores. Autant que je sache, la dimensionnalité d'un bloc ou d'une grille est juste une logique d'affectation hors de propos de matériel, mais le total de la taille d'un bloc (x*y*z) est très important.

Threads d'un bloc doivent être sur la même CM, à utiliser ses installations de la mémoire partagée et la synchronisation. Si vous ne pouvez pas avoir des blocs avec plus de threads que de CUDA cores sont contenues dans un SM.

Si nous avons un scénario simple où nous avons 16 SMs avec 32 CUDA cores chacun, et nous avons 31x1x1 la taille du bloc, et 20x1x1 taille de la grille, on perdra au moins 1/32 de la puissance de traitement de la carte. Chaque fois qu'un bloc est exécuté, un SM aura seulement 31 de ses 32 cœurs occupé. Les blocs se charge de remplir les SMs, nous avons 16 blocs de terminer à peu près au même moment, et que les 4 premiers SMs gratuit, ils vont commencer à traiter les 4 derniers blocs (PAS nécessairement des blocs #17-20).

Commentaires et corrections sont les bienvenues.

Je pense que la partie que je veux comprendre, c'est: comment le compilateur CUDA analyse le noyau de codes, et distribue les tâches aux threads dans le dispositif. Dans mon post original, je imaginé un moyen d'auto-expliquer le mécanisme.

OriginalL'auteur Boyko Perfanov

Vous devez vous connecter pour publier un commentaire.

modifier