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

SCM Repository

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

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

Parent Directory Parent Directory | Revision Log Revision Log | View Patch Patch

revision 2038, Tue Oct 16 11:31:02 2012 UTC revision 2039, Wed Oct 17 16:10:37 2012 UTC
# Line 119  Line 119 
119     sInfo->numInBuffer = count;     sInfo->numInBuffer = count;
120     sInfo->offsets.s0 = row;     sInfo->offsets.s0 = row;
121  }  }
 // 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;  
 }  

Legend:
Removed from v.2038  
changed lines
  Added in v.2039

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