11.3.2021, 9:00 - 11:00: Due to updates GitLab may be unavailable for some minutes between 09:00 and 11:00.

Commit a9edc936 authored by Benedikt Zoennchen's avatar Benedikt Zoennchen

implement the linked cell on the GPU. JUnit tests included.

parent a58b9135
...@@ -109,7 +109,7 @@ uint getGridHash(uint2 gridPos, __constant uint2* gridSize){ ...@@ -109,7 +109,7 @@ uint getGridHash(uint2 gridPos, __constant uint2* gridSize){
//Wrap addressing, assume power-of-two grid dimensions //Wrap addressing, assume power-of-two grid dimensions
gridPos.x = gridPos.x & ((*gridSize).x - 1); gridPos.x = gridPos.x & ((*gridSize).x - 1);
gridPos.y = gridPos.y & ((*gridSize).y - 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( ...@@ -156,7 +156,6 @@ __kernel void findCellBoundsAndReorder(
__global uint *d_CellStart, //output: cell start index __global uint *d_CellStart, //output: cell start index
__global uint *d_CellEnd, //output: cell end index __global uint *d_CellEnd, //output: cell end index
__global float2 *d_ReorderedPos, //output: reordered by cell hash positions __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_Hash, //input: sorted grid hashes
__global const uint *d_Index, //input: particle indices sorted by hash __global const uint *d_Index, //input: particle indices sorted by hash
......
...@@ -41,12 +41,14 @@ import static org.lwjgl.opencl.CL10.clEnqueueWaitForEvents; ...@@ -41,12 +41,14 @@ import static org.lwjgl.opencl.CL10.clEnqueueWaitForEvents;
import static org.lwjgl.opencl.CL10.clEnqueueWriteBuffer; import static org.lwjgl.opencl.CL10.clEnqueueWriteBuffer;
import static org.lwjgl.opencl.CL10.clFinish; import static org.lwjgl.opencl.CL10.clFinish;
import static org.lwjgl.opencl.CL10.clGetDeviceIDs; 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.clGetPlatformIDs;
import static org.lwjgl.opencl.CL10.clReleaseCommandQueue; import static org.lwjgl.opencl.CL10.clReleaseCommandQueue;
import static org.lwjgl.opencl.CL10.clReleaseContext; import static org.lwjgl.opencl.CL10.clReleaseContext;
import static org.lwjgl.opencl.CL10.clReleaseKernel; import static org.lwjgl.opencl.CL10.clReleaseKernel;
import static org.lwjgl.opencl.CL10.clReleaseMemObject; import static org.lwjgl.opencl.CL10.clReleaseMemObject;
import static org.lwjgl.opencl.CL10.clReleaseProgram; 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.clSetKernelArg1i;
import static org.lwjgl.opencl.CL10.clSetKernelArg1p; import static org.lwjgl.opencl.CL10.clSetKernelArg1p;
import static org.lwjgl.system.MemoryStack.stackPush; import static org.lwjgl.system.MemoryStack.stackPush;
...@@ -125,7 +127,7 @@ public class CLUniformHashedGrid { ...@@ -125,7 +127,7 @@ public class CLUniformHashedGrid {
private static final Logger logger = LogManager.getLogger(CLUniformHashedGrid.class); 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; private boolean debug = false;
...@@ -170,30 +172,30 @@ public class CLUniformHashedGrid { ...@@ -170,30 +172,30 @@ public class CLUniformHashedGrid {
clBitonicSort(clHashes, clIndices, clHashes, clIndices, numberOfElements, 1); clBitonicSort(clHashes, clIndices, clHashes, clIndices, numberOfElements, 1);
clFindCellBoundsAndReorder(clCellStarts, clCellEnds, clReorderedPositions, clHashes, clIndices, clPositions, numberOfElements); 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, clCellEnds, true, 0, cellEnds, null, null);
clEnqueueReadBuffer(clQueue, clReorderedPositions, true, 0, reorderedPositions, null, null); clEnqueueReadBuffer(clQueue, clReorderedPositions, true, 0, reorderedPositions, null, null);
clEnqueueReadBuffer(clQueue, clIndices, true, 0, indices, null, null); clEnqueueReadBuffer(clQueue, clIndices, true, 0, indices, null, null);
clEnqueueReadBuffer(clQueue, clHashes, true, 0, hashes, 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[] aCellStarts = CLUtils.toIntArray(cellStarts, numberOfGridCells);
int[] aCellEnds = CLUtils.toIntArray(cellEnds, numberOfElements); int[] aCellEnds = CLUtils.toIntArray(cellEnds, numberOfGridCells);
float[] aReorderedPositions = CLUtils.toFloatArray(reorderedPositions, numberOfElements * 2); float[] aReorderedPositions = CLUtils.toFloatArray(reorderedPositions, numberOfElements * 2);
int[] aIndices = CLUtils.toIntArray(indices, numberOfElements); int[] aIndices = CLUtils.toIntArray(indices, numberOfElements);
int[] aHashes = CLUtils.toIntArray(hashes, 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 gridCells = new GridCells();
/*gridCells.cellEnds = aCellEnds; gridCells.cellEnds = aCellEnds;
gridCells.cellStarts = aCellStarts; gridCells.cellStarts = aCellStarts;
gridCells.reorderedPositions = aReorderedPositions; gridCells.reorderedPositions = aReorderedPositions;
gridCells.indices = aIndices; gridCells.indices = aIndices;
gridCells.hashes = aHashes; gridCells.hashes = aHashes;
gridCells.positions = aPositions;*/ gridCells.positions = aPositions;
//clearMemory(); clearMemory();
//clearCL(); clearCL();
return gridCells; return gridCells;
//clBitonicSort(clHashes, clIndices, clHashes, clIndices, numberOfElements, 1); //clBitonicSort(clHashes, clIndices, clHashes, clIndices, numberOfElements, 1);
...@@ -267,8 +269,8 @@ public class CLUniformHashedGrid { ...@@ -267,8 +269,8 @@ public class CLUniformHashedGrid {
this.gridSize = CLUtils.toIntBuffer(iGridSize, CLUtils.toIntBuffer(iGridSize)); this.gridSize = CLUtils.toIntBuffer(iGridSize, CLUtils.toIntBuffer(iGridSize));
this.cellStarts = MemoryUtil.memAllocInt(numberOfElements); this.cellStarts = MemoryUtil.memAllocInt(numberOfGridCells);
this.cellEnds = MemoryUtil.memAllocInt(numberOfElements); this.cellEnds = MemoryUtil.memAllocInt(numberOfGridCells);
this.indices = MemoryUtil.memAllocInt(numberOfElements); this.indices = MemoryUtil.memAllocInt(numberOfElements);
this.reorderedPositions = MemoryUtil.memAllocFloat(numberOfElements * 2); this.reorderedPositions = MemoryUtil.memAllocFloat(numberOfElements * 2);
} }
...@@ -303,15 +305,6 @@ public class CLUniformHashedGrid { ...@@ -303,15 +305,6 @@ public class CLUniformHashedGrid {
buildProgram(); 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( private void clCalcHash(
final long clHashes, final long clHashes,
final long clIndices, final long clIndices,
...@@ -356,8 +349,8 @@ public class CLUniformHashedGrid { ...@@ -356,8 +349,8 @@ public class CLUniformHashedGrid {
CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 3, clHashes)); CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 3, clHashes));
CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 4, clIndices)); CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 4, clIndices));
CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 5, clPositions)); CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 5, clPositions));
CLInfo.checkCLError(clSetKernelArg1p(clFindCellBoundsAndReorder, 5, (max_work_group_size+1) * 4)); // local memory CLInfo.checkCLError(clSetKernelArg(clFindCellBoundsAndReorder, 6, (max_work_group_size+1) * 4)); // local memory
CLInfo.checkCLError(clSetKernelArg1i(clFindCellBoundsAndReorder, 6, numberOfElements)); CLInfo.checkCLError(clSetKernelArg1i(clFindCellBoundsAndReorder, 7, numberOfElements));
clGlobalWorkSize.put(0, numberOfElements); clGlobalWorkSize.put(0, numberOfElements);
clLocalWorkSize.put(0, max_work_group_size); clLocalWorkSize.put(0, max_work_group_size);
...@@ -599,7 +592,10 @@ public class CLUniformHashedGrid { ...@@ -599,7 +592,10 @@ public class CLUniformHashedGrid {
clFindCellBoundsAndReorder = clCreateKernel(clProgram, "findCellBoundsAndReorder", errcode_ret); clFindCellBoundsAndReorder = clCreateKernel(clProgram, "findCellBoundsAndReorder", errcode_ret);
CLInfo.checkCLError(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); logger.info("CL_DEVICE_MAX_WORK_GROUP_SIZE = " + max_work_group_size);
} }
......
...@@ -39,9 +39,11 @@ public class TestCellGridSort { ...@@ -39,9 +39,11 @@ public class TestCellGridSort {
assertEquals(hasehs.length, positions.size()); assertEquals(hasehs.length, positions.size());
logger.info("number of cells = " + clUniformHashedGrid.getGridSize()[0] * clUniformHashedGrid.getGridSize()[1]);
for(int i = 0; i < hasehs.length; i++) { for(int i = 0; i < hasehs.length; i++) {
int hash = getGridHash(getGridPosition(positions.get(i), clUniformHashedGrid.getCellSize(), clUniformHashedGrid.getWorldOrign()), clUniformHashedGrid.getGridSize()); int hash = getGridHash(getGridPosition(positions.get(i), clUniformHashedGrid.getCellSize(), clUniformHashedGrid.getWorldOrign()), clUniformHashedGrid.getGridSize());
assertEquals(hasehs[i], hash); assertEquals(hasehs[i], hash);
logger.info("hash = " + hash);
} }
} }
...@@ -76,7 +78,18 @@ public class TestCellGridSort { ...@@ -76,7 +78,18 @@ public class TestCellGridSort {
positions.add(new VPoint(random.nextDouble() * 10,random.nextDouble() * 10)); positions.add(new VPoint(random.nextDouble() * 10,random.nextDouble() * 10));
} }
CLUniformHashedGrid.GridCells gridCells = clUniformHashedGrid.calcPositionsInCell(positions); 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 { ...@@ -102,6 +115,6 @@ public class TestCellGridSort {
private static int getGridHash(final int[] gridPos, int[] gridSize) { private static int getGridHash(final int[] gridPos, int[] gridSize) {
gridPos[0] = gridPos[0] & (gridSize[0] - 1); gridPos[0] = gridPos[0] & (gridSize[0] - 1);
gridPos[1] = gridPos[1] & (gridSize[1] - 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]);
} }
} }
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment