Home My Page Projects Code Snippets Project Openings diderot
Summary Activity Tracker Tasks SCM

SCM Repository

[diderot] View of /branches/lamont_dev/src/include/Diderot/cl-spatial.h
ViewVC logotype

View of /branches/lamont_dev/src/include/Diderot/cl-spatial.h

Parent Directory Parent Directory | Revision Log Revision Log


Revision 2032 - (download) (as text) (annotate)
Fri Oct 12 01:55:56 2012 UTC (9 years, 3 months ago) by lamonts
File size: 10107 byte(s)
Added commented to spatial.h
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable

//GridContext -> Meta information about the grid 
typedef struct { 
   int2 gridSize;         // the size(s) of the spatial grid structure
   int numberOfStrands;   // The nuber of strands in execution
   int2 cellSizes;        // The size of each cell in the grid
} GridContext;

//SpatialInfo -> this keeps track of information about query being performed for a certain Strand
typedef struct {   
    int numInBuffer;    // the number of particles currently in the buffer 
    int maxBufferSize;  // the maximum number of particles that can be in the buffer 
    int8 offsets;       // the offsets into where next particle should come from.
                        // s0 -> current x grid position
                        // s1 -> stopping x grid position
                        // s2 -> current y grid position
                        // s3 -> stopping y grid position
                        // s4 -> current strand within a grid cell
    int hasOffsets;     // determines whether the offsets have been calulated already
}SpatialInfo; 

#define MAX_NEIGHBORS 8            //This is just for this example of finding the 7 closestest neighbors
#define MAX_PARTICLE_PER_CELL 4    //This is the number of particles that can be in a cell at one time without overlapping.
#define MAX_SPATIAL_BUFFER_SIZE 2  //The size of the spatial buffer for each strand.

//worldToGrid - converts the world position into it's grid position
//@param worldCoord: the world coordinates
//@param ctx: the grid context
inline int2 worldToGrid(float2 worldCoord, __global GridContext * ctx) {
   int2 gridPos;
   gridPos.x = ((int)worldCoord.x / ctx->cellSizes.x) % ctx->gridSize.x;
   gridPos.y = ((int)worldCoord.y / ctx->cellSizes.y) % ctx->gridSize.y;

    return gridPos;
}

//isWithinRadius - checks to see if the two positions are within each other by the given radius 
inline bool isWithinRadius(float2 pos, float2 pos2, int radius) {
    return distance(pos,pos2) < radius;
} 

//speherical_query - given 
inline void spherical_query(__global @STRAND_TY@ * q_strand  //The strand that is performing the query
                            __global @STRAND_TY@ * strands,  //All the strands running
                            __global int * grid,             //An array of grid cells
                            __global int * gridCounter,      //the number of strands located within a grid cell
                            __global GridContext * ctx,      //information about the grids
                            SpatialInfo  * sInfo,            //Information about the querying being performed
                            float radius,                    //the radius size for th spherical_query
                            __global int * spatialBuffer)    // the actual spatial buffer for this array that stores the strand id
{ 
    int id = get_global_id(0);
    int count = 0, row, col, k;
    
    // This will probably be removed 
    if(id >= ctx->numberOfStrands) {
        return; 
    }
   
    int2 queryBounds,gridPos;
    
    float2 particlePosition = q_strand.pos;

    /* This function will be called multiple times but I don't want to always calculate the offsets each time 
       so this if statement only gets called on the first call to spherical_query to setup the offsets */
    if(!sInfo->hasOffsets){
      // This sets up the bounds for this query based on the radius we are using
      queryBounds = (int2)(ceil(radius/(float)ctx->cellSizes.x)+1,ceil(radius/(float)ctx->cellSizes.y) + 1);
      // This finds the grid position of strand thats performing the query 
      gridPos = worldToGrid(particlePosition,ctx);
      // This stores the offsets of where we are at for this query (i.e. starting, stopping, x & y grid cells) 
      sInfo->offsets = (int8)(gridPos.x - queryBounds.x, gridPos.x + queryBounds.x, 
                           gridPos.y - queryBounds.y, gridPos.y + queryBounds.y,0,gridPos.y - queryBounds.y,0,0);  
      sInfo->hasOffsets = 1;
    }

   //FOREACH cell defined within the query bounds
   for(row = sInfo->offsets.s0; row <= sInfo->offsets.s1; row++){
        for(col = sInfo->offsets.s2; col <= sInfo->offsets.s3; col++) {
            //CHECK to make sure the query has not gone outside the bounds of the grid
            if((col >=0 && col < ctx->gridSize.y) && (row >= 0 && row < ctx->gridSize.x)){
                //Retrieve the index of the grid cell
                int neighborCellIndx =  (row * ctx->gridSize.x) + col;
                 //Retrieve the number of strands within that current grid cell
                 int neighborCellStrands =  gridCounter[neighborCellIndx];
                 //FOREACH strand within the current cell
                 for(k = sInfo->offsets.s4; k < neighborCellStrands; k++) { 
                      sInfo->offsets.s4 = 0;
                      //Retrieve the index of the current strand within the cell
                      int strandCellLIdx = (row* ctx->gridSize.x * MAX_PARTICLE_PER_CELL) + (col* MAX_PARTICLE_PER_CELL);
                      // Retrieve it's strand id
                      int neighborStrand = grid[strandCellLIdx + k];
                      /* Check to see if the strand performing the query and the strand within the cell are actually defined
                         within the given radius */  
                      if(isWithinRadius(particlePosition,strands[neighborStrand].pos,radius)){
                                //Check to make sure the particle can be placed within the spatial buffer
                               if(count < sInfo->maxBufferSize) {
                                   spatialBuffer[count] = neighborParticle; 
                                   count++;
                                }else {
                                   // If the particle can't be placed in the spatial buffer then save the state of the
                                   // query so the query doesn't start over again when the function is called. 
                                   sInfo->numInBuffer = count; 
                                   sInfo->offsets.s0 = row; 
                                   sInfo->offsets.s2 = col; 
                                   sInfo->offsets.s4 = k;  
                                   return;      
                                }
                      }
                 }
            }

        }

       sInfo->offsets.s2 = sInfo->offsets.s5; 
    }   
   sInfo->numInBuffer = count;
   sInfo->offsets.s0 = row;
}
// constructGrid - contructs the spatial grid
__kernel void constructGrid ( __global @STRAND_TY@ * strands,
                              __global int * grid,
                              __global int * gridCounter, 
                             __global GridContext * ctx) 
{ 
    int id = get_global_id(0);
    
    // This will probably be removed 
    if(id >= ctx->numberOfParticles) {
        return;
    }

    //Retrieve the grid coordinates from the particle's position. 
    int2 gridPos = worldToGrid(strands[id].pos,ctx);
    //Retrieve the index of the grid cell 
    int gridIndx = (gridPos.x * ctx->gridSize.x) + gridPos.y;
    
    //Increment the number of strands at that index
    int count = atom_inc(&gridCounter[gridIndx]);
    
    
    //if(count < MAX_PARTICLE_PER_CELL) {
    //Place the strand id into the list of strands at that grid cell location 
    int gridLoc = (gridPos.x * ctx->gridSize.x * MAX_PARTICLE_PER_CELL) + (gridPos.y * MAX_PARTICLE_PER_CELL);
    grid[gridLoc + count] = id;
    
    //}
}
#define RADIUS_QUERY 100
//FindMaxNeighbor - is an example of how the spherical_query will be used in Diderot. 
__kernel void findMaxNeighbor( __global @STRAND_TY@ * strands,
                               __global int * grid, 
                               __global int * gridCounter,
                               __global GridContext * ctx,
                               __global SpatialInfo * spatialInfo, 
                               __global int * spatialBuffer, 
                               __global int * results) 
{

   int id = get_global_id(0); 
   int neighborCount = 0; 
   float currentMax = 0; 
   int maxId = 0; 
   int count = 0; 

   //Will be removed. 
   if(id >= ctx->numberOfStrands) {
        return;
    }
   
   SpatialInfo sInfo = spatialInfo[id]; 
   sInfo.hasOffsets = 0; 

   //This represents a pointer to the spatial buffer for this particular strand
   //In this case each strand will have a maximum of two neighbors it can process at
   //a time. If there are more strands defined within the speherical query then it will
   // process 2 and then call spherical_query again to place more strands in the buffer. 
   __global int * neighborBuffer =  (spatialBuffer + (id * MAX_SPATIAL_BUFFER_SIZE));
   
    spherical_query(particles,
                  grid,
                  gridCounter, 
                  ctx, 
                  &sInfo,
                  RADIUS_QUERY, 
                  neighborBuffer); 

   //In this example, each strand has a power field. And this loop is trying to figure out
   //of all the strands within the spherical query which one has the greatest power.  
   while(sInfo.numInBuffer > 0) { 
        //NOTE: this code represents the code that will be within the FOREACH block
        for(int i= 0; i < sInfo.numInBuffer; i++){
              @STRAND_TY@ p = strands[neighborBuffer[i]];
              if(currentMax < p.power){
                    maxId = neighborBuffer[i];
                    currentMax = p.power;  
               }
              neighborCount++; 
        }
     //Call speherical_query again to retrieve more strands within the given radius. 
    spherical_query(particles, 
                  grid,
                  gridCounter, 
                  ctx, 
                  &sInfo,
                  RADIUS_QUERY, 
                  neighborBuffer);
  }
  //Return the strand of the greatest power defined within the speherical_query 
  results[id] = maxId;  
} 

root@smlnj-gforge.cs.uchicago.edu
ViewVC Help
Powered by ViewVC 1.0.0