Unverified Commit acb55e38 authored by Ľubomír Jurčišin's avatar Ľubomír Jurčišin
Browse files

add coments and rename kernels

parent 46d6dca4

GPU/src/main/resources/opencl/MeshDistanceCalculations.c

+20 −20
Original line number Diff line number Diff line 
kernel void closestOnSurface(global const MeshTriangle* triangles,

                             global const int* triangleStarts,

                             global const int* numTriangles,

                             global const Point3D* fromPoints,

                             global const Point3D* fromNormals,

                             global double* resultDistances,

                             global Point3D* resultPoints,

                             const int numPoints,

                             const int calcRelativeDist) {

/*

    Kernels calculating distance from point to surface.

    @author Lubomir Jurcisin

*/



/*

    Takes triangles around each point and calculates the closest point on them.

    Each thread works through its own triangleFan.

*/

kernel void closestOnTriangleFan(global const MeshTriangle* triangles, global const int* triangleStarts, global const int* numTriangles,

                                 global const Point3D* fromPoints, global const Point3D* fromNormals,

                                 global double* resultDistances, global Point3D* resultPoints,

                                 const int numPoints, const int calcRelativeDist) {

    const int Gid = get_global_id(0);

    if (Gid >= numPoints) {

        return;
@@ -36,8 +40,11 @@ kernel void closestOnSurface(global const MeshTriangle* triangles, 
    }

}





kernel void oneClosest(global const MeshTriangle* triangles, global const Point3D* fromPoint,

/*

    Every thread has its own triangle which it calculates the closest point on.

    It then uses work-group reduction to find the closest point.

*/

kernel void closestOnTriangles(global const MeshTriangle* triangles, global const Point3D* fromPoint,

                               global Point3D* resultPoint, local double* distanceBuffer, const int numTriangles) {

    const int Gid = get_global_id(0);

    const int Lid = get_local_id(0);
@@ -52,13 +59,6 @@ kernel void oneClosest(global const MeshTriangle* triangles, global const Point3 
        distanceBuffer[Lid] = dist;

    }



    //printf("[%d] pointX: %f, pointY: %f, pointZ: %f, distance: %f \n", Gid, fromPoint[0].x, fromPoint[0].y, fromPoint[0].z, dist);

    //printf("[%d] pointX: %f, pointY: %f, pointZ: %f, distance: %f \n", Gid, closestPoint.x, closestPoint.y, closestPoint.z, dist);

    //distanceBuffer[0] = 42;

    //distanceBuffer[25] = 43;

    //printf("[%d] orig: %f bufferBefore[%d]: %f \n", Gid, dist, Lid, distanceBuffer[Lid]);

    //printf("[%d] orig: %f \n", Gid, dist);



    barrier(CLK_LOCAL_MEM_FENCE);



    int processingBlock = numTriangles / 2;
@@ -78,7 +78,7 @@ kernel void oneClosest(global const MeshTriangle* triangles, global const Point3 
        barrier(CLK_LOCAL_MEM_FENCE);

    }



    if (Gid < numTriangles && dist == distanceBuffer[0]) {  // thread 0 writes closest point to output buffer

    if (Gid < numTriangles && dist == distanceBuffer[0]) {  // thread which has the matching distance writes its point to output

        resultPoint[0] = closestPoint;

    }

}

GPU/src/main/resources/opencl/NNKdTree.c

+4 −0
Original line number Diff line number Diff line 
/*

    Kernel uses stack-free traversal of a left-balanced kd-tree to find closest vertex.

    @author Lubomir Jurcisin

*/

kernel void NNKdTree(global const Point3D* from, global const Point3D* fromNormals, global const Point3D* kdTree,

                        global double* closestDistances, global int* closestPointsIndexes,

                        const int numFrom, const int kdTreeSize, const int calcRelativeDist) {

GeometryEngines/src/main/java/cz/fidentis/analyst/engines/distance/impl/MeshDistanceNNGPU.java

+1 −1
Original line number Diff line number Diff line
@@ -144,7 +144,7 @@ public class MeshDistanceNNGPU extends MeshDistanceVisitorImpl { 
        // * 4 because every point is made of 3 floats and one zeroed alignment float

        pointsOnSurface = clContext.createFloatBuffer(numPoints * 4, WRITE_ONLY);



        CLKernel triangleSearch = triangleProgram.getKernel("closestOnSurface")

        CLKernel triangleSearch = triangleProgram.getKernel("closestOnTriangleFan")

                .putArgs(trianglesBuffer.get(), triangleStartsBuffer, trianglesToReadBuffer,

                        pointsFromBuffer.get(), normalsBuffer.get(),

                        resultDistancesTriangle, pointsOnSurface)

GeometryEngines/src/main/java/cz/fidentis/analyst/engines/distance/impl/MeshDistancePreciseGPU.java

+1 −1
Original line number Diff line number Diff line
@@ -192,7 +192,7 @@ public class MeshDistancePreciseGPU extends MeshDistanceVisitorImpl { 
                workGroupSize

        );



        CLKernel meshSearch = meshProgram.getKernel("oneClosest")

        CLKernel meshSearch = meshProgram.getKernel("closestOnTriangles")

                .putArgs(trianglesBuffer.get(), pointBuffer.get(), resultPoint)

                .putNullArg(triangles.size())

                .putArg(triangles.size());