GPU osm copy potential field onto the gpu.

VadereModelTests/TestOSM/scenarios/rimea_09_publicRoom_osm1_2.scenario

@@ -211,6 +211,7 @@
         "spawnAtRandomPositions" : false,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       } ],
       "dynamicElements" : [ ],

VadereModelTests/TestOSM/scenarios/rimea_09_publicRoom_osm1_4.scenario

@@ -202,6 +202,7 @@
         "spawnAtRandomPositions" : false,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       } ],
       "dynamicElements" : [ ],

VadereModelTests/TestOSM/scenarios/rimea_10_pathfinding_osm1.scenario

@@ -552,6 +552,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -571,6 +572,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -590,6 +592,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -609,6 +612,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -628,6 +632,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -647,6 +652,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -666,6 +672,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 2 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -685,6 +692,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 2 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -704,6 +712,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -723,6 +732,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 2 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -742,6 +752,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 2 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       }, {
         "id" : -1,
@@ -761,6 +772,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       } ],
       "dynamicElements" : [ ],

VadereModelTests/TestOSM/scenarios/rimea_11_exitSelection_osm1.scenario

@@ -190,6 +190,7 @@
         "spawnAtRandomPositions" : false,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       } ],
       "dynamicElements" : [ ],

VadereModelTests/TestOSM/scenarios/rimea_12_evacuation_osm1.scenario

@@ -154,6 +154,7 @@
         "spawnAtRandomPositions" : true,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       } ],
       "dynamicElements" : [ ],

VadereModelTests/TestOSM/scenarios/rimea_13_stairs_osm1_a.scenario

@@ -150,6 +150,7 @@
         "spawnAtRandomPositions" : false,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       } ],
       "dynamicElements" : [ ],

VadereModelTests/TestOSM/scenarios/rimea_13_stairs_osm1_b.scenario

@@ -150,6 +150,7 @@
         "spawnAtRandomPositions" : false,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       } ],
       "dynamicElements" : [ ],

VadereModelTests/TestOSM/scenarios/rimea_14_selectRoute_osm1.scenario

@@ -200,6 +200,7 @@
         "spawnAtRandomPositions" : false,
         "useFreeSpaceOnly" : false,
         "targetIds" : [ 1 ],
+        "groupSizeDistribution" : [ 0.0, 0.0, 1.0 ],
         "dynamicElementType" : "PEDESTRIAN"
       } ],
       "dynamicElements" : [ ],

VadereSimulator/src/org/vadere/simulator/models/osm/optimization/ParticleSwarmOptimizer.java

@@ -37,10 +37,6 @@ public class ParticleSwarmOptimizer implements StepCircleOptimizer {
 		VCircle circle = ((VCircle) reachableArea);
 		double stepSize = circle.getRadius();
 
-		final PotentialEvaluationFunction potentialEvaluationFunction = new PotentialEvaluationFunction(pedestrian);
-		potentialEvaluationFunction.setStepSize(stepSize);
-
-
 		List<VPoint> positions = StepCircleOptimizerDiscrete.getReachablePositions(pedestrian, random);
 		// maximum possible angle of movement relative to ankerAngle
 		double angle;

VadereSimulator/src/org/vadere/simulator/models/potential/fields/ObstacleDistancePotential.java

@@ -67,4 +67,8 @@ public class ObstacleDistancePotential implements IPotentialField {
 	public double getPotential(@NotNull VPoint pos, @Nullable Agent agent) {
 		return eikonalSolver.getPotential(pos, 0.0, 1.0);
 	}
+
+	public EikonalSolver getEikonalSolver() {
+		return eikonalSolver;
+	}
 }

VadereSimulator/src/org/vadere/simulator/models/potential/fields/PotentialFieldSingleTargetGrid.java

@@ -14,6 +14,7 @@ import org.vadere.state.scenario.Topography;
 import org.vadere.util.geometry.Vector2D;
 import org.vadere.util.geometry.shapes.VPoint;
 import org.vadere.util.geometry.shapes.VShape;
+import org.vadere.util.potential.calculators.EikonalSolver;
 
 /**
  * A IPotentialTargetGrid, that creates for a single target the a floor field
@@ -27,6 +28,7 @@ public class PotentialFieldSingleTargetGrid extends AbstractPotentialFieldTarget
 	private AttributesFloorField attributesFloorField;
 	private AttributesAgent attributesPedestrian;
 	private final int targetId;
+	private EikonalSolver eikonalSolver;
 
 	public PotentialFieldSingleTargetGrid(final Topography topography,
 			final AttributesAgent attributesPedestrian,
@@ -71,9 +73,14 @@ public class PotentialFieldSingleTargetGrid extends AbstractPotentialFieldTarget
 			PotentialFieldAndInitializer potentialFieldAndInitializer = PotentialFieldAndInitializer.create(topography,
 					targetId, shapes, this.attributesPedestrian, this.attributesFloorField);
 			targetPotentialFields.put(targetId, potentialFieldAndInitializer);
+			eikonalSolver = potentialFieldAndInitializer.eikonalSolver;
 		}
 	}
 
+	public EikonalSolver getEikonalSolver() {
+		return eikonalSolver;
+	}
+
 	@Override
 	public void postLoop(double simTimeInSec) {}
 

VadereSimulator/tests/org/vadere/simulator/models/seating/osm/opencl/TestCLOptimalStepsModel.java

+package org.vadere.simulator.models.seating.osm.opencl;
+
+import org.junit.Before;
+import org.junit.Test;
+import org.vadere.simulator.models.osm.PedestrianOSM;
+import org.vadere.simulator.models.osm.opencl.CLOptimalStepsModel;
+import org.vadere.simulator.models.potential.fields.ObstacleDistancePotential;
+import org.vadere.simulator.models.potential.fields.PotentialFieldSingleTargetGrid;
+import org.vadere.simulator.projects.io.IOVadere;
+import org.vadere.state.attributes.models.AttributesFloorField;
+import org.vadere.state.attributes.models.AttributesOSM;
+import org.vadere.state.attributes.models.AttributesPotentialCompact;
+import org.vadere.state.attributes.scenario.AttributesAgent;
+import org.vadere.state.scenario.Topography;
+import org.vadere.state.util.StateJsonConverter;
+import org.vadere.state.util.TextOutOfNodeException;
+import org.vadere.util.geometry.shapes.VPoint;
+import org.vadere.util.geometry.shapes.VRectangle;
+import org.vadere.util.opencl.OpenCLException;
+
+import java.io.IOException;
+import java.util.ArrayList;
+import java.util.List;
+import java.util.Random;
+import java.util.stream.Collectors;
+
+import static org.junit.Assert.assertEquals;
+
+/**
+ * @author Benedikt Zoennchen
+ */
+public class TestCLOptimalStepsModel {
+
+	private String topographyStringChicken = "{\n" +
+			"  \"attributes\" : {\n" +
+			"    \"bounds\" : {\n" +
+			"      \"x\" : 0.0,\n" +
+			"      \"y\" : 0.0,\n" +
+			"      \"width\" : 35.0,\n" +
+			"      \"height\" : 60.0\n" +
+			"    },\n" +
+			"    \"boundingBoxWidth\" : 0.5,\n" +
+			"    \"bounded\" : true\n" +
+			"  },\n" +
+			"  \"obstacles\" : [ {\n" +
+			"    \"shape\" : {\n" +
+			"      \"x\" : 9.0,\n" +
+			"      \"y\" : 21.0,\n" +
+			"      \"width\" : 1.0,\n" +
+			"      \"height\" : 20.0,\n" +
+			"      \"type\" : \"RECTANGLE\"\n" +
+			"    },\n" +
+			"    \"id\" : -1\n" +
+			"  }, {\n" +
+			"    \"shape\" : {\n" +
+			"      \"x\" : 25.0,\n" +
+			"      \"y\" : 21.0,\n" +
+			"      \"width\" : 1.0,\n" +
+			"      \"height\" : 20.0,\n" +
+			"      \"type\" : \"RECTANGLE\"\n" +
+			"    },\n" +
+			"    \"id\" : -1\n" +
+			"  }, {\n" +
+			"    \"shape\" : {\n" +
+			"      \"x\" : 10.0,\n" +
+			"      \"y\" : 40.0,\n" +
+			"      \"width\" : 15.0,\n" +
+			"      \"height\" : 1.0,\n" +
+			"      \"type\" : \"RECTANGLE\"\n" +
+			"    },\n" +
+			"    \"id\" : -1\n" +
+			"  } ],\n" +
+			"  \"stairs\" : [ ],\n" +
+			"  \"targets\" : [ {\n" +
+			"    \"id\" : 1,\n" +
+			"    \"absorbing\" : true,\n" +
+			"    \"shape\" : {\n" +
+			"      \"x\" : 10.0,\n" +
+			"      \"y\" : 51.0,\n" +
+			"      \"width\" : 15.0,\n" +
+			"      \"height\" : 5.0,\n" +
+			"      \"type\" : \"RECTANGLE\"\n" +
+			"    },\n" +
+			"    \"waitingTime\" : 0.0,\n" +
+			"    \"waitingTimeYellowPhase\" : 0.0,\n" +
+			"    \"parallelWaiters\" : 0,\n" +
+			"    \"individualWaiting\" : true,\n" +
+			"    \"deletionDistance\" : 0.1,\n" +
+			"    \"startingWithRedLight\" : false,\n" +
+			"    \"nextSpeed\" : -1.0\n" +
+			"  } ],\n" +
+			"  \"sources\" : [ {\n" +
+			"    \"id\" : -1,\n" +
+			"    \"shape\" : {\n" +
+			"      \"x\" : 10.0,\n" +
+			"      \"y\" : 6.0,\n" +
+			"      \"width\" : 15.0,\n" +
+			"      \"height\" : 5.0,\n" +
+			"      \"type\" : \"RECTANGLE\"\n" +
+			"    },\n" +
+			"    \"interSpawnTimeDistribution\" : \"org.vadere.state.scenario.ConstantDistribution\",\n" +
+			"    \"distributionParameters\" : [ 1.0 ],\n" +
+			"    \"spawnNumber\" : 300,\n" +
+			"    \"maxSpawnNumberTotal\" : -1,\n" +
+			"    \"startTime\" : 0.0,\n" +
+			"    \"endTime\" : 0.0,\n" +
+			"    \"spawnAtRandomPositions\" : true,\n" +
+			"    \"useFreeSpaceOnly\" : false,\n" +
+			"    \"targetIds\" : [ 1 ],\n" +
+			"    \"groupSizeDistribution\" : [ 0.0, 0.0, 1.0 ],\n" +
+			"    \"dynamicElementType\" : \"PEDESTRIAN\"\n" +
+			"  } ],\n" +
+			"  \"dynamicElements\" : [ ],\n" +
+			"  \"attributesPedestrian\" : {\n" +
+			"    \"radius\" : 0.195,\n" +
+			"    \"densityDependentSpeed\" : false,\n" +
+			"    \"speedDistributionMean\" : 1.34,\n" +
+			"    \"speedDistributionStandardDeviation\" : 0.26,\n" +
+			"    \"minimumSpeed\" : 0.5,\n" +
+			"    \"maximumSpeed\" : 2.2,\n" +
+			"    \"acceleration\" : 2.0\n" +
+			"  },\n" +
+			"  \"attributesCar\" : null\n" +
+			"}";
+
+	private AttributesFloorField attributesFloorField;
+	private AttributesOSM attributesOSM;
+	private List<CLOptimalStepsModel.PedestrianOpenCL> pedestrians;
+	private Topography topography;
+	private ObstacleDistancePotential obstacleDistancePotential;
+	private PotentialFieldSingleTargetGrid targetPotentialField;
+	private AttributesAgent attributesAgent;
+	private AttributesPotentialCompact attributesPotentialCompact;
+	private Random random;
+	private int numberOfElements;
+	private VRectangle bound;
+
+	/*
+	AttributesOSM attributesOSM,
+				  AttributesAgent attributesPedestrian, Topography topography,
+				  Random random, IPotentialFieldTarget potentialFieldTarget,
+				  PotentialFieldObstacle potentialFieldObstacle,
+				  PotentialFieldAgent potentialFieldPedestrian,
+				  List<SpeedAdjuster> speedAdjusters,
+				  StepCircleOptimizer stepCircleOptimizer
+	 */
+	@Before
+	public void setUp() throws IOException, TextOutOfNodeException {
+		random = new Random();
+		numberOfElements = 256;
+		attributesOSM = new AttributesOSM();
+		attributesFloorField = new AttributesFloorField();
+		attributesAgent = new AttributesAgent();
+		attributesPotentialCompact = new AttributesPotentialCompact();
+		topography = StateJsonConverter.deserializeTopography(topographyStringChicken);
+		bound = new VRectangle(topography.getBounds());
+		obstacleDistancePotential = new ObstacleDistancePotential(
+				topography.getObstacles().stream().map(obs -> obs.getShape()).collect(Collectors.toList()),
+				bound, attributesFloorField);
+		targetPotentialField = new PotentialFieldSingleTargetGrid(topography, attributesAgent, attributesFloorField, 1);
+		targetPotentialField.preLoop(0.4f);
+		pedestrians = new ArrayList<>();
+
+		for(int i = 0; i < numberOfElements; i++) {
+			VPoint randomPosition = new VPoint(
+					(float)(bound.getMinX() + random.nextDouble() * bound.getWidth()),
+					(float)(bound.getMinY() + random.nextDouble() * bound.getHeight()));
+			CLOptimalStepsModel.PedestrianOpenCL pedestrian = new CLOptimalStepsModel.PedestrianOpenCL(randomPosition, 1.9f);
+			pedestrians.add(pedestrian);
+		}
+	}
+
+
+	@Test
+	public void testIdentity() throws OpenCLException {
+		CLOptimalStepsModel clOptimalStepsModel = new CLOptimalStepsModel(
+				pedestrians.size(),
+				new VRectangle(topography.getBounds()),
+				1.5f + attributesPotentialCompact.getPedPotentialWidth(), // max step length + function width
+				attributesFloorField,
+				targetPotentialField.getEikonalSolver(),
+				obstacleDistancePotential.getEikonalSolver());
+
+		List<CLOptimalStepsModel.PedestrianOpenCL> result = clOptimalStepsModel.getNextSteps(pedestrians);
+
+		for(int i = 0; i < numberOfElements; i++) {
+			assertEquals("not equals for index = " + i, result.get(i).position, result.get(i).newPosition);
+		}
+	}
+}

VadereUtils/resources/OSM.cl

+/*
+ * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
+ *
+ * Please refer to the NVIDIA end user license agreement (EULA) associated
+ * with this source code for terms and conditions that govern your use of
+ * this software. Any use, reproduction, disclosure, or distribution of
+ * this software and related documentation outside the terms of the EULA
+ * is strictly prohibited.
+ *
+ */
+
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Common definitions
+////////////////////////////////////////////////////////////////////////////////
+#define UMAD(a, b, c)  ( (a) * (b) + (c) )
+
+typedef struct{
+    float x;
+    float y;
+    float z;
+} Float3;
+
+typedef struct{
+    uint x;
+    uint y;
+    uint z;
+}Uint3;
+
+typedef struct{
+    int x;
+    int y;
+    int z;
+}Int3;
+
+
+typedef struct{
+    Float3 colliderPos;
+    float  colliderRadius;
+
+    Float3 gravity;
+    float globalDamping;
+    float particleRadius;
+
+    Uint3 gridSize;
+    uint numCells;
+    Float3 worldOrigin;
+    Float3 cellSize;
+
+    uint numBodies;
+    uint maxParticlesPerCell;
+
+    float spring;
+    float damping;
+    float shear;
+    float attraction;
+    float boundaryDamping;
+} simParams_t;
+
+typedef struct {
+    float2 position;
+    float stepLength;
+} pedestrian;
+
+inline void ComparatorPrivate(
+    uint *keyA,
+    uint *valA,
+    uint *keyB,
+    uint *valB,
+    uint dir
+){
+    if( (*keyA > *keyB) == dir ){
+        uint t;
+        t = *keyA; *keyA = *keyB; *keyB = t;
+        t = *valA; *valA = *valB; *valB = t;
+    }
+}
+
+inline void ComparatorLocal(
+    __local uint *keyA,
+    __local uint *valA,
+    __local uint *keyB,
+    __local uint *valB,
+    uint dir
+){
+    if( (*keyA > *keyB) == dir ){
+        uint t;
+        t = *keyA; *keyA = *keyB; *keyB = t;
+        t = *valA; *valA = *valB; *valB = t;
+    }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Save particle grid cell hashes and indices
+////////////////////////////////////////////////////////////////////////////////
+uint2 getGridPos(float3 p, __constant float* cellSize, __constant float2* worldOrigin){
+    uint2 gridPos;
+    float2 wordOr = (*worldOrigin);
+    gridPos.x = (int)floor((p.x - wordOr.x) / (*cellSize));
+    gridPos.y = (int)floor((p.y - wordOr.y) / (*cellSize));
+    return gridPos;
+}
+
+//Calculate address in grid from position (clamping to edges)
+uint getGridHash(uint2 gridPos, __constant uint2* gridSize){
+    //Wrap addressing, assume power-of-two grid dimensions
+    gridPos.x = gridPos.x & ((*gridSize).x - 1);
+    gridPos.y = gridPos.y & ((*gridSize).y - 1);
+    return UMAD(  (*gridSize).x, gridPos.y, gridPos.x );
+}
+
+__kernel void nextSteps(
+    __global float        *newPositions,        //output
+    __global const float  *orderedPedestrians,  //input
+    __global const uint   *d_CellStart,         //input: cell boundaries
+    __global const uint   *d_CellEnd,           //input
+    __global const float  *obstaclePotential,   //input
+    __global const float  *targetPotential,     //input
+    __constant float2     *worldOrigin,         //input
+    float                  potentialCellSize    //input
+){
+    const uint index = get_global_id(0);
+    newPositions[index*2] = orderedPedestrians[index*3];
+    newPositions[index*2+1] = orderedPedestrians[index*3+1];
+}
+
+//Calculate grid hash value for each particle
+__kernel void calcHash(
+    __global uint           *d_Hash, //output
+    __global uint           *d_Index, //output
+    __global const float3   *d_Pos, //input: positions
+    __constant float        *cellSize,
+    __constant float2       *worldOrigin,
+    __constant uint2        *gridSize,
+    uint numParticles
+){
+    const uint index = get_global_id(0);
+    if(index >= numParticles)
+        return;
+
+    float3 p = d_Pos[index];
+
+    //Get address in grid
+    uint2  gridPos = getGridPos(p, cellSize, worldOrigin);
+    uint gridHash = getGridHash(gridPos, gridSize);
+
+    //Store grid hash and particle index
+    d_Hash[index] = gridHash;
+    d_Index[index] = index;
+}
+
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Find cell bounds and reorder positions+velocities by sorted indices
+////////////////////////////////////////////////////////////////////////////////
+__kernel void Memset(
+    __global uint *d_Data,
+    uint val,
+    uint N
+){
+    if(get_global_id(0) < N)
+        d_Data[get_global_id(0)] = val;
+}
+
+__kernel void findCellBoundsAndReorder(
+    __global uint   *d_CellStart,     //output: cell start index
+    __global uint   *d_CellEnd,       //output: cell end index
+    __global float3 *d_ReorderedPos,  //output: reordered by cell hash positions
+
+    __global const uint   *d_Hash,    //input: sorted grid hashes
+    __global const uint   *d_Index,   //input: particle indices sorted by hash
+    __global const float3 *d_Pos,     //input: positions array sorted by hash
+    __local uint *localHash,          //get_group_size(0) + 1 elements
+    uint    numParticles
+){
+    uint hash;
+    const uint index = get_global_id(0);
+
+    //Handle case when no. of particles not multiple of block size
+    if(index < numParticles){
+        hash = d_Hash[index];
+
+        //Load hash data into local memory so that we can look
+        //at neighboring particle's hash value without loading
+        //two hash values per thread
+        localHash[get_local_id(0) + 1] = hash;
+
+        //First thread in block must load neighbor particle hash
+        if(index > 0 && get_local_id(0) == 0)
+            localHash[0] = d_Hash[index - 1];
+    }
+