     float3 gridSizeInv = make_float3(1.0f / float(gridSize.x), 1.0f / float(gridSize.y), 1.0f / float(gridSize.z));
     float3 bBoxInv = make_float3(1.0f / bBox.x, 1.0f / bBox.y, 1.0f / bBox.z);
     if (this->useColor) {
-      switch (texformat) {
+        switch (texformat) {
         case RGB3F:
-          generateTriangleColorNormal<float3, float3, float3><<<grid3, threads>>>(vertOut, colOut, normOut, (float3 *) d_colors, this->gridSize, gridSizeInv, bBoxInv, totalVerts);
-          break;
+            generateTriangleColorNormal<float3, float3, float3><<<grid3, threads>>>(
+                vertOut, colOut, normOut, (float3*)d_colors, this->gridSize, gridSizeInv, bBoxInv, totalVerts);
+            break;
 
         case RGB4U:
-          generateTriangleColorNormal<float3, uchar4, float3><<<grid3, threads>>>(vertOut, colOut, normOut, (uchar4 *) d_colors, this->gridSize, gridSizeInv, bBoxInv, totalVerts);
-          break;
-      }
+            generateTriangleColorNormal<float3, uchar4, float3><<<grid3, threads>>>(
+                vertOut, colOut, normOut, (uchar4*)d_colors, this->gridSize, gridSizeInv, bBoxInv, totalVerts);
+            break;
+        }
     } else {
-      generateTriangleNormals<float3><<<grid3, threads>>>(vertOut, normOut, gridSizeInv, bBoxInv, totalVerts);
+        generateTriangleNormals<float3><<<grid3, threads>>>(vertOut, normOut, gridSizeInv, bBoxInv, totalVerts);
     }
 
     offsetTriangleVertices<<<grid3, threads>>>(vertOut, this->origin, totalVerts - 1);

 ***************************************************************************/
/**
 * \file CUDAMarchingCubes.cu
 * \brief CUDA implementation of Marching Cubes 
 *
 * This class computes an isosurface for a given density grid
 * using a CUDA Marching Cubes (MC) alorithm. 
 * The implementation is loosely modeled after the MC demo from 
 * the Nvidia GPU Computing SDK, but the design has been improved 
 * and extended in several ways.  
 * This implementation achieves higher performance by reducing the 
 * number of temporary memory buffers, reduces the number of scan 
 * calls by using vector integer types, and allows extraction of 
 * per-vertex normals and optionally computes per-vertex colors 
 * if a volumetric texture map is provided by the caller.
 *
 * This work is described in the following papers:
 *
 *  "Evaluation of Emerging Energy-Efficient Heterogeneous Computing Platforms 
 *  for Biomolecular and Cellular Simulation Workloads"
 *  John E. Stone, Michael J. Hallock, James C. Phillips, Joseph R. Peterson, 
 *  Zaida Luthey-Schulten, and Klaus Schulten.
 *  25th International Heterogeneity in Computing Workshop, 
 *  2016 IEEE International Parallel and Distributed Processing Symposium 
 *  Workshops (IPDPSW), pp. 89-100, 2016.
 *  http://dx.doi.org/10.1109/IPDPSW.2016.130
 *



// compact voxel array
__global__ void
// __launch_bounds__ ( KERNTHREADS, 1 )
compactVoxels(unsigned int* RESTRICT compactedVoxelArray, const uint2* RESTRICT voxelOccupied, unsigned int lastVoxel,
    unsigned int numVoxels, unsigned int numVoxelsm1) {
    unsigned int blockId = (blockIdx.y * gridDim.x) + blockIdx.x;
    unsigned int i = (blockId * blockDim.x) + threadIdx.x;

    if ((i < numVoxels) && ((i < numVoxelsm1) ? voxelOccupied[i].y < voxelOccupied[i + 1].y : lastVoxel)) {
        compactedVoxelArray[voxelOccupied[i].y] = i;
    }
}


// version that calculates no surface normal or color,  only triangle vertices
__global__ void
// __launch_bounds__ ( NTHREADS, 1 )
generateTriangleVerticesSMEM(float3* RESTRICT pos, const unsigned int* RESTRICT compactedVoxelArray,
    const uint2* RESTRICT numVertsScanned, const float* RESTRICT volume, uint3 gridSize, float3 voxelSize,
    float isoValue, unsigned int activeVoxels, unsigned int maxVertsM3) {
    unsigned int zOffset = blockIdx.z * gridDim.x * gridDim.y;
    unsigned int blockId = (blockIdx.y * gridDim.x) + blockIdx.x;
    unsigned int i = zOffset * (blockDim.x * blockDim.y) + (blockId * blockDim.x) + threadIdx.x;
