Forum |  HardWare.fr | News | Articles | PC | S'identifier | S'inscrire | Shop Recherche
2774 connectés 

  FORUM HardWare.fr
  Programmation
  Divers

  CUDA : Boucle for qui prend de plus en plus de temps

 
 


 Mot :   Pseudo :  
 
Bas de page
Auteur Sujet :

CUDA : Boucle for qui prend de plus en plus de temps

n°2073347
babyromu
Posté le 02-05-2011 à 21:09:28  profilanswer
 

Bonjour,
 
Je suis actuellement en train d'écrire un programme sous CUDA.
 
J'ai des fichiers .cu .c et .cpp tous reliés via un makefile.
 
Seulement je rencontre un gros pépin pour le moment.
 
Mon programme est une modélisation scientifique d'un gas dont on fait bouger des particules à chaque pas de temps.
 
Cependant je remarque que mon algorithme met de plus en plus de temps pour faire une itération. Pour faire les 5 premieres par exemple il met 5ms tandis que pour en faire 50 il met 1500 ms.
 
Auriez vous une idée d'a quoi pourrait etre du cet accroissement important de temps.
 
Sachant que chaque iteration regroupe une phase de deplacement de particules, un tri et un phénomène de tirage de nombre aléatoire (c'est une Direct simulation Monte Carlo que je fais au cas ou il y aurait un connaisseur).
 
Merci pour votre aide
 
Romuald

mood
Publicité
Posté le 02-05-2011 à 21:09:28  profilanswer
 

n°2073359
gilou
Modérateur
Modzilla
Posté le 02-05-2011 à 21:36:01  profilanswer
 

Peut être que le code nous permettrait de répondre, parce que sinon, bonjour la boule de cristal.
A+,


---------------
There's more than what can be linked! --    Iyashikei Anime Forever!    --  AngularJS c'est un framework d'engulé!  --
n°2073362
babyromu
Posté le 02-05-2011 à 22:17:29  profilanswer
 

o


Message édité par babyromu le 02-05-2011 à 22:32:12
n°2073363
babyromu
Posté le 02-05-2011 à 22:18:42  profilanswer
 

vu la taille du code, j'etais un peu perplexe à le mettre ici...

n°2073390
Joel F
Real men use unique_ptr
Posté le 03-05-2011 à 00:45:23  profilanswer
 

ca sent les alloc memoires foireuses ou des acces non coalescents

n°2073459
babyromu
Posté le 03-05-2011 à 11:09:21  profilanswer
 

Je vous met ici une partie du code. Le seul kernel que j'utilise pour le moment dans la grosse boucle est advection (car deja elle seule ralentit fortement), les autres sont commentés et n'agissent pas pour le moment. Par exemple, pour 1000 iteration elle met 13ms, our 1500 elle en met 280ms.
 

Code :
  1. extern "C" void molecule_set_CollisionsGPU(molecule_set_t molecule_set)//, histogram_t *histograms)
  2. {
  3.     unsigned int timer = 0;
  4.     cutilCheckError( cutCreateTimer( &timer));
  5.  
  6.     /*Définition des variables utilisées*/
  7.     int TAILLE;
  8.     TAILLE = molecule_set->number*11/10;
  9.     printf("On fixe la taille des tableaux à : %d\n",TAILLE); 
  10.  
  11.    /////////////////Definition des variables utilisées
  12.  
  13. int i,k,j,m;
  14. v3 *tmp,*tmp1;
  15. mlcl *h_mol;
  16. mlcl *d_mol;
  17.     float *collisions,*d_collisions, *samples,*d_samples,*statsize,*d_statsize;
  18.     int *appcell,*d_appcell,*nbrecumu, *d_nbrecumu,*nbrecell,*d_nbrecell;
  19.     molecule_t mol;
  20. cells_table_item_t cell;
  22. ////////////////Allocation mémoire sur CPU
  24. h_mol = (mlcl*)malloc(sizeof(mlcl)*TAILLE);
  25. appcell = (int*)malloc(sizeof(int)*TAILLE);
  26. nbrecumu = (int*)malloc(sizeof(int)*(molecule_set->table->size+1));
  27. nbrecumu[0]=0;
  28. collisions = (float*)malloc(sizeof(float)*molecule_set->table->size);
  29. samples = (float*)malloc(sizeof(float)*molecule_set->table->size);
  30. statsize = (float*)malloc(sizeof(float)*molecule_set->table->size);
  31. nbrecell = (int*)malloc(sizeof(int)*molecule_set->table->size);
  33. static float f_max_1, f_max_2;
  34. static float normalization_1, normalization_2;
  35.     cuda_init_inflow_distribution(&normalization_1,&normalization_2,&f_max_1,&f_max_2);
  36.     float u_max_at_boundary_1 = mean_vel_1 + width_cutoff;
  37.    
  38.    
  39.     static float flux_rate_1 = (2 *sqrt(2) *sigma) / ( exp(-(mean_vel_1*mean_vel_1))/sqrt(PI) + mean_vel_1*(1.0 + erf(mean_vel_1)) );
  41. /*Remplissage des tableaux*/
  43. for (k=0;k<molecule_set->table->size;k++) {
  44.  cell = &(molecule_set->table->array[k]);
  45.  if(cell->number>1)
  46.  {
  47.  for (i=0; i<cell->number; i++) {
  48.   mol = (molecule_t)cells_table_item_Get_Data(cell,i);
  49.   for (j=0;j<3;j++)
  50.   {
  51.    tmp = molecule_GetVelocity(mol);
  52.    tmp1= molecule_GetPosition(mol);
  53.    h_mol[i+nbrecumu[k]].vit[j]=(*tmp)[j];
  54.    h_mol[i+nbrecumu[k]].pos[j]=(*tmp1)[j];
  55.   }
  56.   appcell[i+nbrecumu[k]] =k;
  58.  }
  59.   nbrecumu[k+1]=cell->number + nbrecumu[k];
  60.   nbrecell[k]=0;
  61.   samples[k]=cell->statistics.samples;
  62.   collisions[k]=cell->collision_delay;
  63.   statsize[k]=cell->statistics.size;
  64.  }
  65. }
  67. for(k=nbrecumu[molecule_set->table->size];k<TAILLE;k++)
  68. {
  69.  appcell[k]=10000;
  70.  }
  72. ////////////////// Tableau pour les sorties de donnés : histo + statistiques
  74. int jj;
  76. statistic Stat[molecule_set->table->size];
  78. for (i=0;i<molecule_set->table->size;i++) Stat[i]=((molecule_set->table->array)[i]).statistics;
  80. statistic *d_Stat;
  81. HANDLE_ERROR( cudaMalloc((void **)&d_Stat, sizeof(statistic)*(molecule_set->table->size)));
  82. HANDLE_ERROR( cudaMemcpy(d_Stat, Stat, molecule_set->table->size* sizeof(statistic), cudaMemcpyHostToDevice));
  84. /*
  85. histoGPU Histo[n_histos];
  86. for (i=0;i<n_histos;i++)
  87. {
  88. for (jj=0;jj<3;jj++) for (j=0;j<2;j++) (Histo[i].range[j])[jj]=(histograms[i]->range[j])[jj];
  89. Histo[i].sample_size=histograms[i]->sample_size;
  90. for (j=0;j<3;j++) Histo[i].n_bins[j]=histograms[i]->n_bins[j];
  91. Histo[i].cell[0]=histo_cells[i];
  92. }
  93. */
  96. ///////////////////////////////////////////////////////////////////////////////////////////////////  
  97. //GENERATION DE LA STRUCTURE INITIALE OUR LES NOMBRES ALEATOIRES
  100. mt_struct_stripped* ds_MTConfigs;
  101. unsigned int* ds_MTStates;
  102. int* ds_MTiStates;
  104. mt_struct_stripped h_MTConfigs[N_GPU_RNG_THREADS_TOTAL];
  105.  const char* fname = "./MersenneTwister.dat";
  106.  FILE* fd = fopen(fname, "rb" );
  107.  if(!fd){
  108.   printf("initMTGPU(): failed to open %s\n", fname);
  109.   printf("TEST FAILED\n" );
  110.   exit(0);
  111.  }
  112.  if( !fread(h_MTConfigs, sizeof(h_MTConfigs), 1, fd) ){
  113.   printf("initMTGPU(): failed to load %s\n", fname);
  114.   printf("TEST FAILED\n" );
  115.   exit(0);
  116.  }
  117. fclose(fd);
  119. cudaMalloc((void **)&ds_MTConfigs,  sizeof(h_MTConfigs));
  120. cudaMemcpy(ds_MTConfigs, h_MTConfigs, sizeof(h_MTConfigs), cudaMemcpyHostToDevice);
  121. cudaMalloc((void **)&ds_MTStates,  N_GPU_RNG_THREADS_TOTAL*MT_NN*sizeof(unsigned int));
  122. cudaMalloc((void **)&ds_MTiStates,  N_GPU_RNG_THREADS_TOTAL*sizeof(int));
  123. initialize_CUDA_UniformMT(ds_MTConfigs, ds_MTStates, ds_MTiStates);
  125. //printf("N_GPU_RNG_THREADS_TOTAL %d\n",N_GPU_RNG_THREADS_TOTAL );
  126. //printf("%10lu\n",sizeof(mt_struct_stripped));
  127. //printf("%10lu\n",sizeof(h_MTConfigs));
  128. ///////////////////////////////////////////////////////////////////////////////////////////////////////
  130. dim3 block;
  131.     dim3 grid;
  132.    
  133.     /*Allocation et envoi mémoire vers GPU*/
  134.  
  135.     HANDLE_ERROR( cudaMalloc((void **)&d_mol, sizeof(mlcl)*TAILLE));
  136. HANDLE_ERROR( cudaMalloc((void **)&d_appcell, sizeof(int)*TAILLE));
  137. HANDLE_ERROR( cudaMalloc((void **)&d_nbrecumu, sizeof(int)*(molecule_set->table->size+1)));
  138. HANDLE_ERROR( cudaMalloc((void **)&d_nbrecell, sizeof(int)*(molecule_set->table->size)));
  139. HANDLE_ERROR( cudaMalloc((void **)&d_collisions, sizeof(float)*molecule_set->table->size));
  140. HANDLE_ERROR( cudaMalloc((void **)&d_samples, sizeof(float)*molecule_set->table->size));
  141. HANDLE_ERROR( cudaMalloc((void **)&d_statsize, sizeof(float)*molecule_set->table->size));
  144. HANDLE_ERROR( cudaMemcpy(d_mol, h_mol, TAILLE* sizeof(mlcl), cudaMemcpyHostToDevice));
  145. HANDLE_ERROR( cudaMemcpy(d_nbrecumu, nbrecumu, (molecule_set->table->size+1)* sizeof(int), cudaMemcpyHostToDevice));
  146. HANDLE_ERROR( cudaMemcpy(d_nbrecell, nbrecell, (molecule_set->table->size)* sizeof(int), cudaMemcpyHostToDevice));
  147. HANDLE_ERROR( cudaMemcpy(d_samples, samples, molecule_set->table->size* sizeof(float), cudaMemcpyHostToDevice));
  148. HANDLE_ERROR( cudaMemcpy(d_collisions, collisions, molecule_set->table->size* sizeof(float), cudaMemcpyHostToDevice));
  149. HANDLE_ERROR( cudaMemcpy(d_statsize, statsize, molecule_set->table->size* sizeof(float), cudaMemcpyHostToDevice));
  150. HANDLE_ERROR( cudaMemcpy(d_appcell, appcell, TAILLE *sizeof(int), cudaMemcpyHostToDevice));
  153. //////////////////////////////////////////////////////////////////////////////////  
  154. ///////Boucle de travail//////////////////////////////////////////////////////////  
  155. //////////////////////////////////////////////////////////////////////////////////
  157. float* dtsum;
  158. dtsum = (float*)malloc(sizeof(float));
  159. dtsum[0]=0;
  160. float* d_dtsum;
  161. HANDLE_ERROR( cudaMalloc((void **)&d_dtsum, sizeof(float)));
  162. cudaMemcpy(d_dtsum, dtsum,sizeof(float), cudaMemcpyHostToDevice);
  165. int* increm;
  166. increm = (int*)malloc(sizeof(int));
  167. dtsum[0]=0;
  168. int* d_increm;
  169. HANDLE_ERROR( cudaMalloc((void **)&d_increm, sizeof(int)));
  170. cudaMemcpy(d_increm, increm,sizeof(int), cudaMemcpyHostToDevice);
  172. int* col;
  173. col = (int*)malloc(molecule_set->table->size*sizeof(int));
  174. int* d_col;
  175. HANDLE_ERROR( cudaMalloc((void **)&d_col, molecule_set->table->size*sizeof(int)));
  176. cudaMemcpy(d_col, col,molecule_set->table->size*sizeof(int), cudaMemcpyHostToDevice);
  178. int sommegpu=0;
  180. cutilCheckError( cutStartTimer( timer));
  182. for(m=1;m<1500;m++)
  183. {
  185. //////////////////                     ADVECTION  
  187. block.x = 32;
  188.     grid.x  = TAILLE/32;
  190. AdvectionGPU<<<grid,block>>>(d_mol,d_appcell,dt,x_max,x_min,y_min,y_max,z_min,z_max,molecule_set->table->dx);
  191. //cutilCheckMsg("Kernel execution failed" );
  194. /////////////////////                         TRI
  196. //grid.x = molecule_set->table->size;
  197.     //block.x  = 1;
  198.    
  199.     //tri<<<block,grid>>>(d_mol,d_nbrecumu,d_nbrecell,d_appcell);
  200.     //cutilCheckMsg("Kernel execution failed" );
  201.    
  202. // cudaMemcpy(nbrecumu, d_nbrecumu,sizeof(int)*201, cudaMemcpyDeviceToHost);
  203. // cudaMemcpy(nbrecell, d_nbrecell,sizeof(int)*200, cudaMemcpyDeviceToHost);
  204. // for (i=199;i<200;i++) printf("nombre de part dans la cellule %d vaut %d ajout de %d parti: total parti = %d et 1ere cell : %d\n",i,nbrecumu[i],nbrecell[i],nbrecell[199],nbrecell[0]);
  205. //cudaMemcpy(col, d_col,sizeof(int)*200, cudaMemcpyDeviceToHost);
  206. //for (i=199;i<200;i++) printf("le nombre de coll dans la cell %d vaut %d collisions \n",i,col[i]);
  208.  // sommegpu =0;
  209. //  for (i=0;i<200;i++) sommegpu+=col[i];
  210.  // printf("%d\n",sommegpu/m);
  212. //////////////////                          COLLISIONS
  214. // grid.x = 1;
  215.   //  block.x  = molecule_set->table->size;
  216. // CollGPU<<<grid,block>>>(d_mol,d_nbrecumu,d_nbrecell,d_collisions,d_statsize,d_samples,c_r_max,dt,n_per_cell_ref, ds_MTStates, ds_MTConfigs, ds_MTiStates,d_appcell,d_col);
  217. // cutilCheckMsg("Kernel execution failed" );
  219. /*
  220. cudaMemcpy(appcell, d_appcell,sizeof(int)*TAILLE, cudaMemcpyDeviceToHost);
  221. cudaMemcpy(h_mol, d_mol,sizeof(mlcl)*TAILLE, cudaMemcpyDeviceToHost);
  222. for (i=19100;i<19300;i++) printf("la particule %d avec une vitesse de %f est dans la cellule %d\n",i,h_mol[i].vit[0],appcell[i]);
  223. */
  226.     ///////////////////                             STAT
  227.     //if (m%25==0) UpdateStat<<<grid,block>>>(d_Stat,d_mol,d_nbrecumu);
  228.     if (m%500==0)
  229.     {
  230.      /*HANDLE_ERROR( cudaMemcpy(Stat, d_Stat, molecule_set->table->size* sizeof(statistic), cudaMemcpyDeviceToHost));
  231.  for (i=0;i<molecule_set->table->size;i++) ((molecule_set->table->array)[i]).statistics=Stat[i];
  232.      molecule_set_Output_Cell_Stats_File(molecule_set);*/
  233.      }
  234.    
  235.    
  236. }
  237.  
  238.   cutilCheckError( cutStopTimer( timer));
  239.   float gpu_time = cutGetTimerValue( timer);
  240.   printf( "GPU Processing time: %f (ms)\n", gpu_time);
  242. HANDLE_ERROR( cudaMemcpy(Stat, d_Stat, molecule_set->table->size* sizeof(statistic), cudaMemcpyDeviceToHost));
  243. cudaMemcpy(nbrecumu, d_nbrecumu,sizeof(int)*201, cudaMemcpyDeviceToHost);
  244. cudaMemcpy(nbrecell, d_nbrecell,sizeof(int)*200, cudaMemcpyDeviceToHost);
  245. for (i=0;i<molecule_set->table->size;i++) ((molecule_set->table->array)[i]).statistics=Stat[i];
  246.     cells_table_item_t myptr;
  247.     cells_table_t mytable;
  249.     mytable = molecule_set->table;
  251.     for (i=0;i<mytable->size;i++) {
  252.     myptr = &(mytable->array[i]);
  253.     myptr->number = nbrecumu[i+1]-nbrecumu[i];
  254.     }
  256. //for (i=0;i<200;i++) printf("nombre de part dans la cellule %d vaut %d ajout de %d parti\n",i,nbrecumu[i],nbrecell[i]);
  258. ///////////////////Renvoi des données sur le processeur
  259. cudaMemcpy(h_mol, d_mol, TAILLE* sizeof(mlcl), cudaMemcpyDeviceToHost);
  260. cudaMemcpy(nbrecumu, d_nbrecumu, (molecule_set->table->size+1)* sizeof(int), cudaMemcpyDeviceToHost);
  261. cudaMemcpy(samples, d_samples, molecule_set->table->size* sizeof(float), cudaMemcpyDeviceToHost);
  262. cudaMemcpy(collisions, d_collisions, molecule_set->table->size* sizeof(float), cudaMemcpyDeviceToHost);
  263. cudaMemcpy(statsize, d_statsize, molecule_set->table->size* sizeof(float), cudaMemcpyDeviceToHost);
  265. cudaFree(d_mol);
  266. cudaFree(d_nbrecumu);
  267. cudaFree(d_collisions);
  268. cudaFree(d_statsize);
  269. cudaFree(d_samples);
  270. cudaFree(d_nbrecell);
  271. cudaFree(d_appcell);
  272. cudaFree(ds_MTConfigs);
  273. cudaFree(ds_MTStates);
  274. cudaFree(ds_MTiStates);
  275. cudaFree(d_dtsum);
  277.     free(h_mol);
  278. free(nbrecumu);
  279. free(collisions);
  280. free(statsize);
  281. free(samples);
  282. free(appcell);
  283. free(nbrecell);
  284. free(dtsum);
  285. }
  287. ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
  293. void initialize_CUDA_UniformMT(mt_struct_stripped* ds_MTConfigs, unsigned int* ds_MTStates, int* ds_MTiStates)
  294. {
  295. c_initialize_CUDA_UniformMT(ds_MTConfigs, ds_MTStates, ds_MTiStates);
  296. }
  299. __device__ Real generateUniform(unsigned int *m_mtstate, mt_struct_stripped *m_config, int *m_istate)
  300. {
  301. int     iState = *m_istate;
  302. unsigned int mti1 = m_mtstate[iState];
  303. int iState1 = iState + 1;
  304. int iStateM = iState + MT_MM;
  305. if(iState1 >= MT_NN) iState1 -= MT_NN;
  306. if(iStateM >= MT_NN) iStateM -= MT_NN;
  308. unsigned int mti  = mti1;
  309.      mti1 = m_mtstate[iState1];
  310. unsigned int mtiM = m_mtstate[iStateM];
  312. unsigned int x;
  313. x = (mti & MT_UMASK) | (mti1 & MT_LMASK);
  314. x =  mtiM ^ (x >> 1) ^ ((x & 1) ? (*m_config).matrix_a : 0);
  316. *m_istate    = iState1;
  317. m_mtstate[iState] = x;
  319. //Tempering transformation
  320. x ^= (x >> MT_SHIFT0);
  321. x ^= (x << MT_SHIFTB) & (*m_config).mask_b;
  322. x ^= (x << MT_SHIFTC) & (*m_config).mask_c;
  323. x ^= (x >> MT_SHIFT1);
  325. //Convert to (0, 1] Real and write to global memory
  326. return ((Real)x + 1.0f) / (double)4294967296.0f;
  327. }
  330. __global__ void AdvectionGPU(mlcl *mol,int* cellule, float dt, int x_max, int x_min,int y_min,int y_max,int z_min,int z_max,float dx)
  331. {
  333.  int tid;
  334.  int i;
  335.  int cell;
  336.  float pos_collision, t_collision;
  337.  tid = blockIdx.x*blockDim.x + threadIdx.x;
  339.  if (tid<19283)
  340.  {
  342.  for (i=0;i<3;++i)
  343.  {
  344.   mol[tid].pos[i]+= dt*mol[tid].vit[i];
  345.  }
  346.  if (mol[tid].pos[0]<x_min) mol[tid].pos[0] = x_max-x_min+mol[tid].pos[0];
  347.  if (mol[tid].pos[0]>x_max) mol[tid].pos[0] = x_min-x_max+mol[tid].pos[0];
  348.  if (mol[tid].pos[1]<y_min) mol[tid].pos[1] = y_max-y_min+mol[tid].pos[1];
  349.  if (mol[tid].pos[1]>y_max) mol[tid].pos[1] = y_min-y_max+mol[tid].pos[1];
  350.  if (mol[tid].pos[2]<z_min) mol[tid].pos[2] = z_max-z_min+mol[tid].pos[2];
  351.  if (mol[tid].pos[2]>z_max) mol[tid].pos[2] = z_min-z_max+mol[tid].pos[2];
  353.  cell=(mol[tid].pos[0]-x_min)/dx;
  354.  cellule[tid] = cell;
  355.  }
  356. }
  359. __global__ void tri(mlcl *mol, int* nbrecumule,int* nbrecell,int *appcell)
  360. {
  361. int tid =0;
  362. tid = blockIdx.x*blockDim.x + threadIdx.x;
  363. int tic =0;
  364. int i;
  365. mlcl retrouve[250];
  367.  //On les rerange dans la bonne cellule
  369.  if (tid!=0 && tid!= 199)
  370.  {
  371.   for (i=nbrecumule[tid-1];i<nbrecumule[tid+2];i++)
  372.   {
  373.    if(appcell[i]==tid)
  374.    {
  375.     retrouve[tic]=mol[i];
  376.     tic++;
  377.    }
  378.   }
  379.   nbrecell[tid]=tic;
  380.  }
  381.  else if (tid==0)/////////Pas oublier les Boundary Left
  382.  {
  383.   for (i=nbrecumule[tid];i<nbrecumule[tid+2];i++)
  384.   {
  385.    if(appcell[i]==0)
  386.    {
  387.     retrouve[tic]=mol[i];
  388.     tic++;
  389.    }
  390.   }
  391.   for (i=nbrecumule[200];i<nbrecumule[200]+100;i++)
  392.   {
  393.    if(appcell[i]==0)
  394.    {
  395.     retrouve[tic]=mol[i];
  396.     tic++;
  397.    }
  398.   }
  400.   nbrecell[tid]=tic;
  401.  }
  402.  else
  403.  {
  404.   for (i=nbrecumule[tid-1];i<nbrecumule[tid+1];i++)
  405.   {
  406.    if(appcell[i]==tid)
  407.    {
  408.     retrouve[tic]=mol[i];
  409.     tic++;
  410.    }
  411.   }
  412.   nbrecell[tid]=tic;
  413.  }
  414.  nbrecumule[tid+1]=0;
  415.  __syncthreads();
  417.  for(i=0;i<tid+1;i++) nbrecumule[tid+1]+=nbrecell[i];
  418.  __syncthreads();
  419.  for(i=0;i<nbrecell[tid];i++)
  420.  {
  421.   mol[nbrecumule[tid]+i]=retrouve[i];
  422.   appcell[nbrecumule[tid]+i]=tid;
  423.  }
  425. }
  428. __global__ void CollGPU(mlcl *mol, int* nbrecumule,int* nbrecell, float* delai,float* taille, float* sample, float c_r_max,int dt, int n_per_cell_ref,unsigned int *m_mtstate, mt_struct_stripped *m_config, int *m_istate, int *appcell, int* col)
  429. {
  430.  int tid,i,j,index1,index2;
  431.  tid = blockIdx.x*blockDim.x + threadIdx.x;
  432.  float epsilon,chi,seps,schi,ceps,cchi,norm_c_r;
  433.  float B;
  434.  float n_aver_for_collision_interval, delta_t_coll;
  435.  v3 c_m,c_r,c_r_star;
  436.  float arret;
  437.  arret = float(1.0)/100;
  438.  col[tid]=0;
  439.  if(nbrecell[tid]>1)
  440.  {
  441.  ///////////////Collisions
  442.  while (delai[tid]<=arret) {
  443.   do {
  444.    index1 = nbrecumule[tid]+(nbrecumule[tid+1]-nbrecumule[tid])*generateUniform(&m_mtstate[tid],&m_config[tid],&m_istate[tid]);
  445.    index2 = nbrecumule[tid]+(nbrecumule[tid+1]-nbrecumule[tid])*generateUniform(&m_mtstate[tid],&m_config[tid],&m_istate[tid]);
  446.       index2 = (index2>=index1) ? (index2+1) : index2;
  447.    
  448.       for (i=0;i<3;i++) {
  449.        c_m[i] = (mol[index1].vit[i] + mol[index2].vit[i])/2;
  450.        c_r[i] = mol[index1].vit[i] - mol[index2].vit[i];
  451.      }
  453.      norm_c_r = sqrt(c_r[0]*c_r[0] + c_r[1]*c_r[1] + c_r[2]*c_r[2]);
  454.    } while(norm_c_r<c_r_max*generateUniform(&m_mtstate[tid],&m_config[tid],&m_istate[tid]));
  455.  col[tid]++;
  456.  epsilon = 2*generateUniform(&m_mtstate[tid],&m_config[tid],&m_istate[tid])*float(M_PI);
  457.  chi = float(M_PI)*generateUniform(&m_mtstate[tid],&m_config[tid],&m_istate[tid]);
  458.  seps = sin(epsilon);
  459.  schi=sin(chi);
  460.  cchi=cos(chi);
  461.  ceps=cos(epsilon);
  463.  B = sqrt(c_r[1]*c_r[1] + c_r[2]*c_r[2]);
  464.  c_r_star[0] = (cchi * c_r[0] + schi * seps * B);
  465.  c_r_star[1] = cchi * c_r[1] + schi * (norm_c_r * c_r[2] * ceps - c_r[0] * c_r[1] * seps)/B;
  466.  c_r_star[2] = cchi * c_r[2] - schi * (norm_c_r * c_r[1] * ceps + c_r[0] * c_r[2] * seps)/B;
  468.  for (j=0;j<3;++j) {
  469.     mol[index1].vit[j] = (c_r_star[j] + 2*c_m[j])/2;
  470.     mol[index2].vit[j] = (-c_r_star[j] + 2*c_m[j])/2;
  471.  }
  473.  n_aver_for_collision_interval = ((taille[tid] == 0) ? (nbrecumule[tid+1]-nbrecumule[tid]) : ((taille[tid]) / (sample[tid])));
  474.     delta_t_coll = (2*float(M_SQRT2)/((nbrecumule[tid+1]-nbrecumule[tid]))) * (n_per_cell_ref)/(n_aver_for_collision_interval * norm_c_r);
  475.     delai[tid] += delta_t_coll; 
  476. }
  477. delai[tid] = delai[tid] - float(1.0)/100;
  478.  }
  479. nbrecell[tid]=0;
  480. }
  485. void cuda_init_inflow_distribution(float *normalization_1,float *normalization_2,float *f_max_1,float *f_max_2)
  486.      /* Initializes the static variables used repeatedly in inflow_distribution_left  
  487. and inflow_distribution_right  
  488.      */
  489. {
  490.   float u_f_max_1, u_f_max_2;
  492.   *normalization_1 = 0.5 * exp(- beta2_o_beta1 * beta2_o_beta1 * mean_vel_1 * mean_vel_1) / (beta2_o_beta1 * beta2_o_beta1)
  493.     + mean_vel_1 * 0.5 * sqrt(M_PI) * (1.0 + erf( beta2_o_beta1 * mean_vel_1)) / beta2_o_beta1;
  494.  
  495.   *normalization_2 = 0.5 * exp(- beta2_o_beta1 * beta2_o_beta1 * mean_vel_2 * mean_vel_2) / (beta2_o_beta1 * beta2_o_beta1)
  496.     + mean_vel_2 * 0.5 * sqrt(M_PI) * (1.0 + erf( beta2_o_beta1 * mean_vel_2)) / beta2_o_beta1;
  497.  
  498.   u_f_max_1 = 0.5 * (mean_vel_1 + sqrt(mean_vel_1 * mean_vel_1 + 2.0 / (beta2_o_beta1 * beta2_o_beta1)));
  499.  
  500.   *f_max_1 = u_f_max_1 * exp(- beta2_o_beta1 * beta2_o_beta1 * pow((u_f_max_1 - mean_vel_1),2)) / *normalization_1;
  501.  
  502.   u_f_max_2 = 0.5 * (mean_vel_2 + sqrt(mean_vel_2 * mean_vel_2 + 2.0 / (beta2_o_beta1 * beta2_o_beta1)));
  503.  
  504.   *f_max_2 = u_f_max_2 * exp(- beta2_o_beta1 * beta2_o_beta1 * pow((u_f_max_2 - mean_vel_2),2)) / *normalization_2;
  505. }
  507. __global__ void UpdateStat(statistic *stat,mlcl *mol,int *nbrecumu)
  508. {
  510. int tid = blockIdx.x*blockDim.x + threadIdx.x;
  511. int i;
  512. stat[tid].size = stat[tid].size + nbrecumu[tid+1]-nbrecumu[tid];
  513.     stat[tid].samples ++;
  514.    
  515.     for (i=nbrecumu[tid];i<nbrecumu[tid+1];i++) {
  516.     stat[tid].sum_v[0] += mol[i].vit[0];
  517.     stat[tid].sum_v[1] += mol[i].vit[1];
  518.     stat[tid].sum_v[2] += mol[i].vit[2];
  519.     stat[tid].sum_v2[0] += mol[i].vit[0]*mol[i].vit[0];
  520.     stat[tid].sum_v2[1] += mol[i].vit[1]*mol[i].vit[1];
  521.     stat[tid].sum_v2[2] += mol[i].vit[2]*mol[i].vit[2];
  522.   }
  525. }


Message édité par gilou le 03-05-2011 à 11:11:10
n°2073757
GrosBocdel
Posté le 04-05-2011 à 13:21:23  profilanswer
 

cudaThreadSynchronize(); après ton appel de kernel?


Aller à :
Ajouter une réponse
  FORUM HardWare.fr
  Programmation
  Divers

  CUDA : Boucle for qui prend de plus en plus de temps

 

Sujets relatifs
[PHP] Utiliser 2 requêtes SQL dans 1 boucle ?Quel language, et combien de temps ?
temps nécessaire au développement d'une couche securiséeTemps d'exécution en millisecondes (linux/windows) + multithreads
Excel : boucle copie de plages sous conditionBoucle vba qui ouvre des fichiers avec noms variables
Mesurer le temps d'exec d'une fonctionvariable dans une boucle
[Résolu] Update multiple dans une boucle whileEquivalence boucle while(- && - && -) et while (- || - || -)
Plus de sujets relatifs à : CUDA : Boucle for qui prend de plus en plus de temps

Forum MesDiscussions.Net, Version 2010.2
(c) 2000-2011 Doctissimo
Page générée en 0.206 secondes

Copyright © 1997-2022 Hardware.fr SARL (Signaler un contenu illicite / Données personnelles) / Groupe LDLC / Shop HFR