From a9edc936829fc6cfe5621a4c3d3af3f776f5e969 Mon Sep 17 00:00:00 2001 From: Benedikt Zoennchen Date: Thu, 12 Jul 2018 16:21:25 +0200 Subject: [PATCH] implement the linked cell on the GPU. JUnit tests included. --- VadereUtils/resources/Particles.cl | 3 +- .../util/opencl/CLUniformHashedGrid.java | 44 +++++++++---------- .../vadere/util/math/TestCellGridSort.java | 17 ++++++- 3 files changed, 36 insertions(+), 28 deletions(-) diff --git a/VadereUtils/resources/Particles.cl b/VadereUtils/resources/Particles.cl index 3a33ccccd..65c423486 100644 --- a/VadereUtils/resources/Particles.cl +++ b/VadereUtils/resources/Particles.cl @@ -109,7 +109,7 @@ 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( UMAD(1.0, (*gridSize).y, gridPos.y), (*gridSize).x, gridPos.x ); + return UMAD( (*gridSize).x, gridPos.y, gridPos.x ); } @@ -156,7 +156,6 @@ __kernel void findCellBoundsAndReorder( __global uint *d_CellStart, //output: cell start index __global uint *d_CellEnd, //output: cell end index __global float2 *d_ReorderedPos, //output: reordered by cell hash positions - __global float2 *d_ReorderedVel, //output: reordered by cell hash velocities __global const uint *d_Hash, //input: sorted grid hashes __global const uint *d_Index, //input: particle indices sorted by hash diff --git a/VadereUtils/src/org/vadere/util/opencl/CLUniformHashedGrid.java b/VadereUtils/src/org/vadere/util/opencl/CLUniformHashedGrid.java index 6bc0d78d4..942ea39af 100644 --- a/VadereUtils/src/org/vadere/util/opencl/CLUniformHashedGrid.java +++ b/VadereUtils/src/org/vadere/util/opencl/CLUniformHashedGrid.java @@ -41,12 +41,14 @@ import static org.lwjgl.opencl.CL10.clEnqueueWaitForEvents; import static org.lwjgl.opencl.CL10.clEnqueueWriteBuffer; import static org.lwjgl.opencl.CL10.clFinish; import static org.lwjgl.opencl.CL10.clGetDeviceIDs; +import static org.lwjgl.opencl.CL10.clGetDeviceInfo; import static org.lwjgl.opencl.CL10.clGetPlatformIDs; import static org.lwjgl.opencl.CL10.clReleaseCommandQueue; import static org.lwjgl.opencl.CL10.clReleaseContext; import static org.lwjgl.opencl.CL10.clReleaseKernel; import static org.lwjgl.opencl.CL10.clReleaseMemObject; import static org.lwjgl.opencl.CL10.clReleaseProgram; +import static org.lwjgl.opencl.CL10.clSetKernelArg; import static org.lwjgl.opencl.CL10.clSetKernelArg1i; import static org.lwjgl.opencl.CL10.clSetKernelArg1p; import static org.lwjgl.system.MemoryStack.stackPush; @@ -125,7 +127,7 @@ public class CLUniformHashedGrid { private static final Logger logger = LogManager.getLogger(CLUniformHashedGrid.class); - private int max_work_group_size; + private long max_work_group_size; private boolean debug = false; @@ -170,30 +172,30 @@ public class CLUniformHashedGrid { clBitonicSort(clHashes, clIndices, clHashes, clIndices, numberOfElements, 1); clFindCellBoundsAndReorder(clCellStarts, clCellEnds, clReorderedPositions, clHashes, clIndices, clPositions, numberOfElements); - /*clEnqueueReadBuffer(clQueue, clCellStarts, true, 0, cellStarts, null, null); + clEnqueueReadBuffer(clQueue, clCellStarts, true, 0, cellStarts, null, null); clEnqueueReadBuffer(clQueue, clCellEnds, true, 0, cellEnds, null, null); clEnqueueReadBuffer(clQueue, clReorderedPositions, true, 0, reorderedPositions, null, null); clEnqueueReadBuffer(clQueue, clIndices, true, 0, indices, null, null); clEnqueueReadBuffer(clQueue, clHashes, true, 0, hashes, null, null); - clEnqueueReadBuffer(clQueue, clPositions, true, 0, this.positions, null, null);*/ + clEnqueueReadBuffer(clQueue, clPositions, true, 0, this.positions, null, null); - /*int[] aCellStarts = CLUtils.toIntArray(cellStarts, numberOfElements); - int[] aCellEnds = CLUtils.toIntArray(cellEnds, numberOfElements); + int[] aCellStarts = CLUtils.toIntArray(cellStarts, numberOfGridCells); + int[] aCellEnds = CLUtils.toIntArray(cellEnds, numberOfGridCells); float[] aReorderedPositions = CLUtils.toFloatArray(reorderedPositions, numberOfElements * 2); int[] aIndices = CLUtils.toIntArray(indices, numberOfElements); int[] aHashes = CLUtils.toIntArray(hashes, numberOfElements); - float[] aPositions = CLUtils.toFloatArray(this.positions, numberOfElements * 2);*/ + float[] aPositions = CLUtils.toFloatArray(this.positions, numberOfElements * 2); GridCells gridCells = new GridCells(); - /*gridCells.cellEnds = aCellEnds; + gridCells.cellEnds = aCellEnds; gridCells.cellStarts = aCellStarts; gridCells.reorderedPositions = aReorderedPositions; gridCells.indices = aIndices; gridCells.hashes = aHashes; - gridCells.positions = aPositions;*/ + gridCells.positions = aPositions; - //clearMemory(); - //clearCL(); + clearMemory(); + clearCL(); return gridCells; //clBitonicSort(clHashes, clIndices, clHashes, clIndices, numberOfElements, 1); @@ -267,8 +269,8 @@ public class CLUniformHashedGrid { this.gridSize = CLUtils.toIntBuffer(iGridSize, CLUtils.toIntBuffer(iGridSize)); - this.cellStarts = MemoryUtil.memAllocInt(numberOfElements); - this.cellEnds = MemoryUtil.memAllocInt(numberOfElements); + this.cellStarts = MemoryUtil.memAllocInt(numberOfGridCells); + this.cellEnds = MemoryUtil.memAllocInt(numberOfGridCells); this.indices = MemoryUtil.memAllocInt(numberOfElements); this.reorderedPositions = MemoryUtil.memAllocFloat(numberOfElements * 2); } @@ -303,15 +305,6 @@ public class CLUniformHashedGrid { buildProgram(); } - /*__kernel void calcHash( - __global uint *d_Hash, //output - __global uint *d_Index, //output - __global const float2 *d_Pos, //input: positions - __constant float cellSize, - __constant float2 worldOrigin, - __constant uint2 gridSize - uint numParticles*/ - private void clCalcHash( final long clHashes, final long clIndices, @@ -356,8 +349,8 @@ public class CLUniformHashedGrid { CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 3, clHashes)); CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 4, clIndices)); CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 5, clPositions)); - CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 5, (max_work_group_size+1) * 4)); // local memory - CLInfo.checkCLError(clSetKernelArg1i(clFindCellBoundsAndReorder, 6, numberOfElements)); + 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); @@ -599,7 +592,10 @@ public class CLUniformHashedGrid { clFindCellBoundsAndReorder = clCreateKernel(clProgram, "findCellBoundsAndReorder", errcode_ret); CLInfo.checkCLError(errcode_ret); - max_work_group_size = CLInfo.getDeviceInfoInt(clDevice, CL_DEVICE_MAX_WORK_GROUP_SIZE); + PointerBuffer pp = stack.mallocPointer(1); + clGetDeviceInfo(clDevice, CL_DEVICE_MAX_WORK_GROUP_SIZE, pp, null); + max_work_group_size = pp.get(0); + logger.info("CL_DEVICE_MAX_WORK_GROUP_SIZE = " + max_work_group_size); } diff --git a/VadereUtils/tests/org/vadere/util/math/TestCellGridSort.java b/VadereUtils/tests/org/vadere/util/math/TestCellGridSort.java index 8d9445ab5..b51180b36 100644 --- a/VadereUtils/tests/org/vadere/util/math/TestCellGridSort.java +++ b/VadereUtils/tests/org/vadere/util/math/TestCellGridSort.java @@ -39,9 +39,11 @@ public class TestCellGridSort { 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); } } @@ -76,7 +78,18 @@ public class TestCellGridSort { positions.add(new VPoint(random.nextDouble() * 10,random.nextDouble() * 10)); } CLUniformHashedGrid.GridCells gridCells = clUniformHashedGrid.calcPositionsInCell(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); + } + } } /** @@ -102,6 +115,6 @@ public class TestCellGridSort { private static int getGridHash(final int[] gridPos, int[] gridSize) { gridPos[0] = gridPos[0] & (gridSize[0] - 1); gridPos[1] = gridPos[1] & (gridSize[1] - 1); - return umad(umad(1, gridSize[1], gridPos[1]), gridSize[0], gridPos[0]); + return umad(gridSize[0], gridPos[1], gridPos[0]); } } -- GitLab