OpenCL πρόβλημα

[computeras13]

Καλησπέρα. Για άλλη μια φορά έχω ερώτηση σε θέματα HPC

 

Χθες ξεκίνησε η πρώτη μου επαφή με OpenCL και αντιμετωπίζω ένα πρόβλημα όταν προσπαθώ αν τρέξω το παρακάτω kernel. Στην ουσία η μόνη μεταβλητή που θέλω να διαβάσω μετά είναι η tmp_cells. Όμως κατά την επιστροφή όλες οι τιμές του πίνακα speeds κάθε στοιχείου του struct είναι 0.

Κάτω από το kernel παραθέτω και τον κώδικα που δημιουργώ τα buffers, τρέχω το kernel και επιστρέφω το struct. Το compile γίνεται κανονικά οπότε δεν το παραθέτω. Τι ακριβώς κάνω λάθος και οι τιμές που επιστρέφονται δεν είναι σωστές;

 

Kernel

>
#pragma OPENCL EXTENSION cl_khr_fp64 : enable                                 
typedef struct                                                                
{                                                                             
 int    nx;                                                                  
 int    ny;                                                                  
 int    maxIters;                                                            
 int    reynolds_dim;                                                        
 double density;                                                             
 double accel;                                                               
 double omega;                                                               
} t_param;                                                                    
                                                                             
typedef struct                                                                
{                                                                             
 double speeds[9];                                                           
} t_speed;                                                                    
	                                                                              
__kernel void propagate(__global const t_param* params,                       
                       __global const t_speed* cells,                        
                       __global       t_speed* tmp_cells,                    
                       __global const int*     obstacles,                    
                                const unsigned int count)                    
{                                                                             
 int pos = get_global_id(0);                                                 
 if (pos < count)                                                            
 {                                                                           
   int temp, x_e, x_w, y_n, y_s, ii, jj;                                     
   ii = (int) pos / params->nx;                                              
   jj = pos % params->nx;                                                    
   y_n = (ii + 1) % params->ny;                                              
   x_e = (jj + 1) % params->nx;                                              
   y_s = (ii == 0) ? (ii + params->ny -1) : (ii -1);                         
   x_w = (jj == 0) ? (jj + params->nx -1) : (jj -1);                         
	                                                                              
   temp = ii * params->nx;                                                   
   tmp_cells[pos].speeds[0] = cells[pos].speeds[0];                          
   tmp_cells[pos].speeds[1] = cells[temp + x_w].speeds[1];                   
   tmp_cells[pos].speeds[3] = cells[temp + x_e].speeds[3];                   
   temp = y_n * params->nx;                                                  
   tmp_cells[pos].speeds[4] = cells[temp + jj].speeds[4];                    
   tmp_cells[pos].speeds[8] = cells[temp + x_w].speeds[8];                   
   tmp_cells[pos].speeds[7] = cells[temp + x_e].speeds[7];                   
   temp = y_s * params->nx;                                                  
   tmp_cells[pos].speeds[2] = cells[temp + jj].speeds[2];                    
   tmp_cells[pos].speeds[5] = cells[temp + x_w].speeds[5];                   
   tmp_cells[pos].speeds[6] = cells[temp + x_e].speeds[6];                   
 }                                                                           
}                                                                             

 

c code

>
m_params = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(t_param), NULL, NULL);
	m_cells = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(t_speed) * count, NULL, NULL);
	m_tmp_cells = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(t_speed) * count, NULL, NULL);
	m_obstacles = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int) * count, NULL, NULL);
if (!m_params || !m_obstacles || !m_cells || !m_tmp_cells)
{
 printf("Error: Failed to allocate device memory collision!\n");
 exit(1);
}
err = clEnqueueWriteBuffer(commands, m_params, CL_TRUE, 0, sizeof(t_param), &params, 0, NULL, NULL);
err |= clEnqueueWriteBuffer(commands, m_cells, CL_TRUE, 0, sizeof(t_speed) * count, cells, 0, NULL, NULL);
err |= clEnqueueWriteBuffer(commands, m_tmp_cells, CL_TRUE, 0, sizeof(t_speed) * count, tmp_cells, 0, NULL, NULL);
err |= clEnqueueWriteBuffer(commands, m_obstacles, CL_TRUE, 0, sizeof(int) * count, obstacles, 0, NULL, NULL);
if (err != CL_SUCCESS)
{
 printf("Error: Failed to write data to source array!\n");
 exit(1);
}
err = 0;
err  = clSetKernelArg(propagate_kernel, 0, sizeof(cl_mem), &m_params);
err |= clSetKernelArg(propagate_kernel, 1, sizeof(cl_mem), &m_cells);
err |= clSetKernelArg(propagate_kernel, 2, sizeof(cl_mem), &m_tmp_cells);
err |= clSetKernelArg(propagate_kernel, 3, sizeof(cl_mem), &m_obstacles);
err |= clSetKernelArg(propagate_kernel, 4, sizeof(unsigned int), &count);
if (err != CL_SUCCESS)
{
 printf("Error: Failed to set accelarate flow kernel arguments! %d\n", err);
 exit(1);
}
err = clGetKernelWorkGroupInfo(propagate_kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL);
if (err != CL_SUCCESS)
{
 printf("Error: Failed to retrieve kernel work group info! %d\n", err);
exit(1);
}
err = clEnqueueNDRangeKernel(commands, propagate_kernel, count, 0, &global, &local, 0, NULL, NULL);
clFinish(commands);
err = clEnqueueReadBuffer(commands, m_tmp_cells, CL_TRUE, 0, sizeof(t_speed) * count, tmp_cells, 0, NULL, NULL);
if (err != CL_SUCCESS)
{
 printf("Error: Failed to read output array! %d\n", err);
 exit(1);
}

[computeras13]

Συμπληρώνω οτι παρατήρησα πρόβλημα και σε ένα άλλο kernel. Όταν τρέχει ο kernel παίρνω σφάλμα "Floating point exception" όπου λογικά όπως το καταλαβαίνω εγώ η διαίρεση που πάει να γίνει γίνεται με 0 το οποίο σημαίνει οτι ούτε ο kernel βλέπει σωστά τις τιμές. Μήπως διαβάζω εγώ λάθος τα δεδομένα από τον kernel;! Ο τρόπος που περνάω δεδομένα είναι ίδιος με το από πάνω post.

 

>
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
typedef struct
{
double speeds[9];
} t_speed;
__kernel void prepare(__global const t_speed* cells,
                     __global const int*     obstacles,
                     __global       double*  results,
                              const unsigned int count)
{
 int pos = get_global_id(0);
 if(pos >= count) return;
 if(obstacles[pos] == 1) results[pos] = 0.00;
 else
 {
   double local_density = 0.00;
   for(int kk = 0; kk < 9; kk++)
     local_density += cells[pos].speeds[kk];
   results[pos] = (cells[pos].speeds[1] + cells[pos].speeds[5] +
                   cells[pos].speeds[8] - (cells[pos].speeds[3] +
                   cells[pos].speeds[6] + cells[pos].speeds[7])) /
                   local_density;
 }
}

[computeras13]

Άκυρο δεν είναι στον kernel το πρόβλημα. Είχε να κάνει με το πως καλούσα να εκτελεστεί ο kernel. Υπήρχε εκεί ένα θέμα οτι ο αριθμός των global "εργατών" δεν ήταν διαιρέσιμος με αυτόν των local.