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 2039 - (download) (as text) (annotate)
Wed Oct 17 16:10:37 2012 UTC (7 years ago) by lamonts
File size: 6619 byte(s)
Added Query Code to tree-to-cl
#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;
}

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