Merge branch 'develop' of gitlab.lrz.de:vadere/vadere into develop

VadereUtils/resources/Particles.cl

@@ -15,6 +15,7 @@
 // Common definitions
 ////////////////////////////////////////////////////////////////////////////////
 #define UMAD(a, b, c)  ( (a) * (b) + (c) )
+//#define LOCAL_SIZE_LIMIT 512U
 
 typedef struct{
     float x;
@@ -321,7 +322,7 @@ __kernel void collide(
 ////////////////////////////////////////////////////////////////////////////////
 // Monolithic bitonic sort kernel for short arrays fitting into local memory
 ////////////////////////////////////////////////////////////////////////////////
-__kernel void bitonicSortLocal(
+/*__kernel void bitonicSortLocal1(
     __global uint *d_DstKey,
     __global uint *d_DstVal,
     __global uint *d_SrcKey,
@@ -332,8 +333,7 @@ __kernel void bitonicSortLocal(
     __local uint *l_val
 ){
     uint LOCAL_SIZE_LIMIT = get_local_size(0) * 2;
-
-    //Offset to the beginning of subbatch and load data
+    //Offset to the beginning of subarray and load data
     d_SrcKey += get_group_id(0) * LOCAL_SIZE_LIMIT + get_local_id(0);
     d_SrcVal += get_group_id(0) * LOCAL_SIZE_LIMIT + get_local_id(0);
     d_DstKey += get_group_id(0) * LOCAL_SIZE_LIMIT + get_local_id(0);
@@ -343,6 +343,66 @@ __kernel void bitonicSortLocal(
     l_key[get_local_id(0) + (LOCAL_SIZE_LIMIT / 2)] = d_SrcKey[(LOCAL_SIZE_LIMIT / 2)];
     l_val[get_local_id(0) + (LOCAL_SIZE_LIMIT / 2)] = d_SrcVal[(LOCAL_SIZE_LIMIT / 2)];
 
+    uint comparatorI = get_global_id(0) & ((LOCAL_SIZE_LIMIT / 2) - 1);
+
+    for(uint size = 2; size < LOCAL_SIZE_LIMIT; size <<= 1){
+        //Bitonic merge
+        uint ddd = (comparatorI & (size / 2)) != 0;
+        for(uint stride = size / 2; stride > 0; stride >>= 1){
+            barrier(CLK_LOCAL_MEM_FENCE);
+            uint pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
+            ComparatorLocal(
+                &l_key[pos +      0], &l_val[pos +      0],
+                &l_key[pos + stride], &l_val[pos + stride],
+                ddd
+            );
+        }
+    }
+
+    //Odd / even arrays of LOCAL_SIZE_LIMIT elements
+    //sorted in opposite directions
+    {
+        uint ddd = (get_group_id(0) & 1);
+        for(uint stride = LOCAL_SIZE_LIMIT / 2; stride > 0; stride >>= 1){
+            barrier(CLK_LOCAL_MEM_FENCE);
+            uint pos = 2 * get_local_id(0) - (get_local_id(0) & (stride - 1));
+            ComparatorLocal(
+                &l_key[pos +      0], &l_val[pos +      0],
+                &l_key[pos + stride], &l_val[pos + stride],
+               ddd
+            );
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+    d_DstKey[                     0] = l_key[get_local_id(0) +                      0];
+    d_DstVal[                     0] = l_val[get_local_id(0) +                      0];
+    d_DstKey[(LOCAL_SIZE_LIMIT / 2)] = l_key[get_local_id(0) + (LOCAL_SIZE_LIMIT / 2)];
+    d_DstVal[(LOCAL_SIZE_LIMIT / 2)] = l_val[get_local_id(0) + (LOCAL_SIZE_LIMIT / 2)];
+}*/
+
+__kernel void bitonicSortLocal(
+    __global uint *d_DstKey,
+    __global uint *d_DstVal,
+    __global uint *d_SrcKey,
+    __global uint *d_SrcVal,
+    uint arrayLength,
+    uint dir,
+    __local uint *l_key,
+    __local uint *l_val
+){
+    uint LOCAL_SIZE_LIMIT = get_local_size(0) * 2;
+
+    //Offset to the beginning of subbatch and load data
+    d_SrcKey += get_local_id(0);
+    d_SrcVal += get_local_id(0);
+    d_DstKey += get_local_id(0);
+    d_DstVal += get_local_id(0);
+    l_key[get_local_id(0) +                      0] = d_SrcKey[                     0];
+    l_val[get_local_id(0) +                      0] = d_SrcVal[                     0];
+    l_key[get_local_id(0) + (LOCAL_SIZE_LIMIT / 2)] = d_SrcKey[(LOCAL_SIZE_LIMIT / 2)];
+    l_val[get_local_id(0) + (LOCAL_SIZE_LIMIT / 2)] = d_SrcVal[(LOCAL_SIZE_LIMIT / 2)];
+
     for(uint size = 2; size < arrayLength; size <<= 1){
         //Bitonic merge
         uint ddd = dir ^ ( (get_local_id(0) & (size / 2)) != 0 );

VadereUtils/src/org/vadere/util/opencl/CLLinkedCell.java

@@ -424,8 +424,8 @@ public class CLLinkedCell {
 		    CLInfo.checkCLError(clSetKernelArg(clFindCellBoundsAndReorder, 6, (max_work_group_size+1) * 4)); // local memory
 		    CLInfo.checkCLError(clSetKernelArg1i(clFindCellBoundsAndReorder, 7, numberOfElements));
 
-		    clGlobalWorkSize.put(0, numberOfElements);
-		    clLocalWorkSize.put(0, max_work_group_size);
+		    clGlobalWorkSize.put(0, Math.min(max_work_group_size, numberOfElements));
+		    clLocalWorkSize.put(0, Math.min(max_work_group_size, numberOfElements));
 		    //TODO: local work size? + check 2^n constrain!
 		    CLInfo.checkCLError(clEnqueueNDRangeKernel(clQueue, clFindCellBoundsAndReorder, 1, null, clGlobalWorkSize, clLocalWorkSize, null, null));
 	    }
@@ -451,10 +451,9 @@ public class CLLinkedCell {
 			    CLInfo.checkCLError(clSetKernelArg1p(clBitonicSortLocal, 2, clKeysIn));
 			    CLInfo.checkCLError(clSetKernelArg1p(clBitonicSortLocal, 3, clValuesIn));
 			    CLInfo.checkCLError(clSetKernelArg1i(clBitonicSortLocal, 4, numberOfElements));
-			    //TODO: check the hard coded 1, and the waiting of the queue
-			    CLInfo.checkCLError(clSetKernelArg1i(clBitonicSortLocal, 5, 1));
-			    CLInfo.checkCLError(clSetKernelArg(clBitonicSortLocal, 6, keys.length * 4)); // local memory
-			    CLInfo.checkCLError(clSetKernelArg(clBitonicSortLocal, 7, keys.length * 4)); // local memory
+			    CLInfo.checkCLError(clSetKernelArg1i(clBitonicSortLocal, 5, dir));
+			    CLInfo.checkCLError(clSetKernelArg(clBitonicSortLocal, 6, numberOfElements * 4)); // local memory
+			    CLInfo.checkCLError(clSetKernelArg(clBitonicSortLocal, 7, numberOfElements * 4)); // local memory
 			    clGlobalWorkSize.put(0, numberOfElements / 2);
 			    clLocalWorkSize.put(0, numberOfElements / 2);
 

VadereUtils/tests/org/vadere/util/math/TestCLLinkedList.java

@@ -29,29 +29,51 @@ public class TestCLLinkedList {
 	public void setUp() throws Exception {}
 
 	@Test
-	public void testCalcHash() throws IOException, OpenCLException {
-		CLLinkedCell clUniformHashedGrid = new CLLinkedCell(1024, new VRectangle(0, 0, 10, 10), 1);
+	public void testCalcHashSmall() throws IOException, OpenCLException {
+		int size = 8;
+		CLLinkedCell clUniformHashedGrid = new CLLinkedCell(size, new VRectangle(0, 0, 10, 10), 1);
 		ArrayList<VPoint> positions = new ArrayList<>();
-		for(int i = 0; i < 1024; i++) {
+		for(int i = 0; i < size; i++) {
 			positions.add(new VPoint(random.nextDouble() * 10,random.nextDouble() * 10));
 		}
 		int[] hasehs = clUniformHashedGrid.calcHashes(positions);
 
 		assertEquals(hasehs.length, positions.size());
 
-		logger.info("number of cells = " + clUniformHashedGrid.getGridSize()[0] * clUniformHashedGrid.getGridSize()[1]);
+		//logger.info("number of cells = " + clUniformHashedGrid.getGridSize()[0] * clUniformHashedGrid.getGridSize()[1]);
 		for(int i = 0; i < hasehs.length; i++) {
 			int hash = getGridHash(getGridPosition(positions.get(i), clUniformHashedGrid.getCellSize(), clUniformHashedGrid.getWorldOrign()), clUniformHashedGrid.getGridSize());
 			assertEquals(hasehs[i], hash);
-			logger.info("hash = " + hash);
+			//logger.info("hash = " + hash);
 		}
 	}
 
 	@Test
-	public void testCalcAndSortHash() throws IOException, OpenCLException {
-		CLLinkedCell clUniformHashedGrid = new CLLinkedCell(1024, new VRectangle(0, 0, 10, 10), 1);
+	public void testCalcHashLarge() throws IOException, OpenCLException {
+		int size = 32768; // 2^15
+		CLLinkedCell clUniformHashedGrid = new CLLinkedCell(size, new VRectangle(0, 0, 10, 10), 1);
 		ArrayList<VPoint> positions = new ArrayList<>();
-		for(int i = 0; i < 1024; i++) {
+		for(int i = 0; i < size; i++) {
+			positions.add(new VPoint(random.nextDouble() * 10,random.nextDouble() * 10));
+		}
+		int[] hasehs = clUniformHashedGrid.calcHashes(positions);
+
+		assertEquals(hasehs.length, positions.size());
+
+		//logger.info("number of cells = " + clUniformHashedGrid.getGridSize()[0] * clUniformHashedGrid.getGridSize()[1]);
+		for(int i = 0; i < hasehs.length; i++) {
+			int hash = getGridHash(getGridPosition(positions.get(i), clUniformHashedGrid.getCellSize(), clUniformHashedGrid.getWorldOrign()), clUniformHashedGrid.getGridSize());
+			assertEquals(hasehs[i], hash);
+			//logger.info("hash = " + hash);
+		}
+	}
+
+	@Test
+	public void testCalcAndSortHashSmall() throws IOException, OpenCLException {
+		int size = 8;
+		CLLinkedCell clUniformHashedGrid = new CLLinkedCell(size, new VRectangle(0, 0, 10, 10), 1);
+		ArrayList<VPoint> positions = new ArrayList<>();
+		for(int i = 0; i < size; i++) {
 			positions.add(new VPoint(random.nextDouble() * 10,random.nextDouble() * 10));
 		}
 		int[] hasehs = clUniformHashedGrid.calcSortedHashes(positions);
@@ -71,10 +93,58 @@ public class TestCLLinkedList {
 	}
 
 	@Test
-	public void testGridCell() throws IOException, OpenCLException {
-		CLLinkedCell clUniformHashedGrid = new CLLinkedCell(1024, new VRectangle(0, 0, 10, 10), 1);
+	public void testCalcAndSortHashLarge() throws IOException, OpenCLException {
+		int size = 32768; // 2^15
+		CLLinkedCell clUniformHashedGrid = new CLLinkedCell(size, new VRectangle(0, 0, 10, 10), 1);
+		ArrayList<VPoint> positions = new ArrayList<>();
+		for(int i = 0; i < size; i++) {
+			positions.add(new VPoint(random.nextDouble() * 10,random.nextDouble() * 10));
+		}
+		int[] hasehs = clUniformHashedGrid.calcSortedHashes(positions);
+
+		assertEquals(hasehs.length, positions.size());
+
+		int[] expectedHashes = new int[positions.size()];
+		for(int i = 0; i < hasehs.length; i++) {
+			int hash = getGridHash(getGridPosition(positions.get(i), clUniformHashedGrid.getCellSize(), clUniformHashedGrid.getWorldOrign()), clUniformHashedGrid.getGridSize());
+			expectedHashes[i] = hash;
+		}
+		Arrays.sort(expectedHashes);
+
+		for(int i = 0; i < hasehs.length; i++) {
+			assertEquals(hasehs[i], expectedHashes[i]);
+		}
+	}
+
+	@Test
+	public void testGridCellSmall() throws IOException, OpenCLException {
+		int size = 8;
+		CLLinkedCell clUniformHashedGrid = new CLLinkedCell(size, new VRectangle(0, 0, 10, 10), 1);
+		ArrayList<VPoint> positions = new ArrayList<>();
+		for(int i = 0; i < size; i++) {
+			positions.add(new VPoint(random.nextDouble() * 10,random.nextDouble() * 10));
+		}
+		CLLinkedCell.LinkedCell gridCells = clUniformHashedGrid.calcLinkedCell(positions);
+		int numberOfCells = clUniformHashedGrid.getGridSize()[0] * clUniformHashedGrid.getGridSize()[1];
+		for(int cell = 0; cell < numberOfCells; cell++) {
+			int cellStart = gridCells.cellStarts[cell];
+			int cellEnd = gridCells.cellEnds[cell];
+
+			for(int i = cellStart; i < cellEnd; i++) {
+				VPoint point = new VPoint(gridCells.reorderedPositions[i*2], gridCells.reorderedPositions[i*2+1]);
+				int[] gridPosition = getGridPosition(point, clUniformHashedGrid.getCellSize(), clUniformHashedGrid.getWorldOrign());
+				int gridHash = getGridHash(gridPosition, clUniformHashedGrid.getGridSize());
+				assertEquals(gridHash, cell);
+			}
+		}
+	}
+
+	@Test
+	public void testGridCellLarge() throws IOException, OpenCLException {
+		int size = 32768;
+		CLLinkedCell clUniformHashedGrid = new CLLinkedCell(size, new VRectangle(0, 0, 10, 10), 1);
 		ArrayList<VPoint> positions = new ArrayList<>();
-		for(int i = 0; i < 1024; i++) {
+		for(int i = 0; i < size; i++) {
 			positions.add(new VPoint(random.nextDouble() * 10,random.nextDouble() * 10));
 		}
 		CLLinkedCell.LinkedCell gridCells = clUniformHashedGrid.calcLinkedCell(positions);