Edit 2: S'il vous plaît jeter un oeil à ce crosspost pour TLDR.
modifier: Étant donné que les particules sont segmentées en cellules de grille (disons 16^3 grid), est-ce une meilleure idée de laisser exécuter un groupe de travail pour chaque cellule de la grille et autant d'éléments de travail dans un groupe de travail qu'il peut y avoir un nombre maximal de particules par cellule de la grille?
Dans ce cas, je pourrais charger toutes les particules decellules voisines dans la mémoire locale et itérer à travers elles en calculant certaines propriétés. Ensuite, je pourrais écrire une valeur spécifique dans chaque particule de la cellule de grille actuelle.
Cette approche serait-elle avantageuse par rapport à l'exécution du noyau pour toutes les particules et pour chaque itération sur (la plupart du temps les mêmes) voisins?
Aussi, quel est le ratio idéal de number of particles/number of grid cells?
J'essaye de réimplémenter (et de modifier) Particules CUDA pour OpenCL et utilisez-le pour interroger les voisins les plus proches pour chaque particule. J'ai créé les structures suivantes:
- Tampon
Pmaintenant toutes les particules "positions 3D (float3) Tampon
Spstockerint2paires d'identifiants de particules et leurs hachages spatiaux.Spest trié en fonction du hachage. (Le hachage n'est qu'un simple mappage linéaire de la 3D vers 1D - pas encore d'indexation Z.)Tampon
Lstockerint2paires de positions de début et de fin de hachages spatiaux particuliers dans le tamponSp. Exemple:L[12] = (int2)(0, 50).L[12].xest l'index (dansSp) du premier particule avec hachage spatial12.L[12].yest l'index (dansSp) du dernier particule avec hachage spatial12.
Maintenant que j'ai tous ces tampons, je veux parcourir toutes les particules de P et pour chaque particule, parcourir ses voisins les plus proches. Actuellement, j'ai un noyau qui ressemble à ceci (pseudocode):
__kernel process_particles(float3* P, int2* Sp, int2* L, int* Out) {
size_t gid = get_global_id(0);
float3 curr_particle = P[gid];
int processed_value = 0;
for(int x=-1; x<=1; x++)
for(int y=-1; y<=1; y++)
for(int z=-1; z<=1; z++) {
float3 neigh_position = curr_particle + (float3)(x,y,z)*GRID_CELL_SIDE;
// ugly boundary checking
if ( dot(neigh_position<0, (float3)(1)) +
dot(neigh_position>BOUNDARY, (float3)(1)) != 0)
continue;
int neigh_hash = spatial_hash( neigh_position );
int2 particles_range = L[ neigh_hash ];
for(int p=particles_range.x; p<particles_range.y; p++)
processed_value += heavy_computation( P[ Sp[p].y ] );
}
Out[gid] = processed_value;
}
Le problème avec ce code est qu'il est lent. Je soupçonne l'accès à la mémoire GPU non linéaire (en particulier P[Sp[p].y] dans le plus intérieur for loop) à l'origine de la lenteur.
Ce que je veux faire, c'est utiliser Courbe d'ordre Z comme hachage spatial. De cette façon, je pourrais avoir seulement 1 for boucle itérant à travers une plage continue de mémoire lors de l'interrogation des voisins. Le seul problème est que je ne sais pas quelles devraient être les valeurs Z-index de début et de fin.
Le Saint Graal que je veux atteindre:
__kernel process_particles(float3* P, int2* Sp, int2* L, int* Out) {
size_t gid = get_global_id(0);
float3 curr_particle = P[gid];
int processed_value = 0;
// How to accomplish this??
// `get_neighbors_range()` returns start and end Z-index values
// representing the start and end near neighbors cells range
int2 nearest_neighboring_cells_range = get_neighbors_range(curr_particle);
int first_particle_id = L[ nearest_neighboring_cells_range.x ].x;
int last_particle_id = L[ nearest_neighboring_cells_range.y ].y;
for(int p=first_particle_id; p<=last_particle_id; p++) {
processed_value += heavy_computation( P[ Sp[p].y ] );
}
Out[gid] = processed_value;
}
Réponses:
-1 pour la réponse № 1Vous devriez étudier attentivement les algorithmes de Morton Code. La détection de collision en temps réel d'Ericsons l'explique très bien.
Ericson - Détection de collision en temps réel
Voici une autre explication intéressante comprenant quelques tests:
Encodage / décodage Morton par entrelacement de bits: implémentations
Les algorithmes Z-Order ne définissent que les chemins ducoordonnées dans lesquelles vous pouvez hacher des coordonnées 2 ou 3D en un seul entier. Bien que l'algorithme approfondisse chaque itération, vous devez définir vous-même les limites. Habituellement, l'index d'arrêt est indiqué par une sentinelle. Laisser la sentinelle s'arrêter vous dira à quel niveau la particule est placée. Ainsi, le niveau maximum que vous souhaitez définir vous indiquera le nombre de cellules par dimension. Par exemple, avec un niveau maximum à 6, vous avez 2 ^ 6 = 64. Vous aurez 64x64x64 cellules dans votre système (3D). Cela signifie également que vous devez utiliser des coordonnées basées sur des nombres entiers. Si vous utilisez des flottants, vous devez convertir comme coord.x = 64*float_x etc.
Si vous savez combien de cellules vous avez dans votre système, vous pouvez définir vos limites. Essayez-vous d'utiliser un octree binaire?
Puisque les particules sont en mouvement (dans cet exemple CUDA), vous devriez essayer de paralléliser le nombre de particules au lieu de cellules.
Si vous souhaitez créer des listes de voisins les plus prochesvous devez mapper les particules sur les cellules. Cela se fait à travers un tableau qui est ensuite trié par cellules en particules. Vous devez tout de même parcourir les particules et accéder à ses voisins.
À propos de votre code:
Le problème avec ce code est qu'il est lent. Je soupçonne que l'accès à la mémoire GPU non linéaire (en particulier P [Sp [p] .y] dans la boucle for la plus interne) est à l'origine de la lenteur.
Souvenez-vous de Donald Knuth. Vous devez mesurer où se trouve le goulot de la bouteille. Vous pouvez utiliser NVCC Profiler et rechercher les goulots d'étranglement. Je ne sais pas ce qu'OpenCL a en tant que profileur.
// ugly boundary checking if ( dot(neigh_position<0, (float3)(1)) + dot(neigh_position>BOUNDARY, (float3)(1)) != 0) continue;
Je pense que vous ne devriez pas le brancher de cette façon, que diriez-vous de renvoyer zéro lorsque vous appelez heavy_computation. Pas sûr, mais peut-être avez-vous une sorte de prédiction de branche ici. Essayez de supprimer cela d'une manière ou d'une autre.
Courir en parallèle sur les cellules est une bonne idéeseulement si vous n'avez pas d'accès en écriture aux données de particules, sinon vous devrez utiliser atomics. Si vous dépassez la plage de particules à la place, vous lisez les accès aux cellules et aux voisins, mais vous créez votre somme en parallèle et vous n'êtes pas obligé de recourir à un paradigme de condiction raciale.
En outre, quel est le rapport idéal du nombre de particules / nombre de cellules de la grille?
Dépend vraiment de vos algorithmes et de laempaquetage de particules dans votre domaine, mais dans votre cas, je définirais la taille de cellule équivalente au diamètre de particule et utiliserais simplement le nombre de cellules que vous obtenez.
Donc, si vous voulez utiliser l'ordre Z et atteindre votre Saint Graal, essayez d'utiliser des coordonnées entières et de les hacher.
Essayez également d'utiliser de plus grandes quantités de particules.Vous devriez prendre en compte environ 65000 particules comme les exemples de CUDA, car de cette façon, la parallélisation est principalement efficace; les unités de traitement en cours d'exécution sont exploitées (moins de threads inactifs).