 }
 
 __global__
-void computeSVVertexQuadric(SV_Header* headers, SV_Data* datas, Quadric_GPU* faceQuadric, Quadric_GPU* vertexQuadric)
+void computeSVVertexQuadric(SV_Header* headers, SV_Data* datas, Quadric_GPU* faceQuadric, Quadric_GPU* vertexQuadric, size_t vertexCount)
 {
 	int vertexId = blockIdx.x * blockDim.x + threadIdx.x;
+	if (vertexCount <= vertexId)
+		return;
 	
 	SV_Header header = headers[vertexId];
 	SV_Data data = datas[vertexId];

	}

	vertexQuadric[vertexId] = q;
}

int ZEPHYR_GPU_API Zephyr::GPU::decimateSuperVertex(Common::OpenMeshMesh & mesh, unsigned int targetFaceCount, unsigned int binSize)

	std::cout << "Block Needed: " << blockNeeded << std::endl;

	// Setup the header and data
	setupHeaderAndData<<<blockNeeded, maxThreadPerBlock>>>(d_indices_ptr, d_header_ptr, d_datas_ptr, MAX_FACE);
	//CudaCheckError();

	// compute the face quadrics
	thrust::device_vector<Quadric_GPU> FacesQuadric(omesh.n_faces());
	auto d_FaceQuadric_ptr = thrust::raw_pointer_cast(&FacesQuadric[0]);
	computeSVQuadricAllFace<<<blockNeeded, maxThreadPerBlock>>>(d_indices_ptr, d_vertices_ptr, d_FaceQuadric_ptr);
	//CudaCheckError();

	thrust::device_vector<Quadric_GPU> vertexQuadric(vertexCount);
	auto d_vertexQuadric_ptr = thrust::raw_pointer_cast(&vertexQuadric[0]);
	
	blockNeeded = (omesh.n_vertices() + (maxThreadPerBlock - 1)) / maxThreadPerBlock;
	computeSVVertexQuadric<<<blockNeeded, maxThreadPerBlock>>>(d_header_ptr, d_datas_ptr, d_FaceQuadric_ptr, d_vertexQuadric_ptr, vertexCount);
	//CudaCheckError();
	int currentFaceCount = omesh.n_faces();
	//while (targetFaceCount < currentFaceCount)
	{
		thrust::device_vector<bool> d_vertexUsed(vertexCount);
		auto d_vertexUsed_ptr = thrust::raw_pointer_cast(&d_vertexUsed[0]);
		thrust::device_vector<bool> d_independentVertex(vertexCount);
		auto d_independentVertex_ptr = thrust::raw_pointer_cast(&d_independentVertex[0]);

		//blockNeeded = (binSize + (maxThreadPerBlock - 1)) / maxThreadPerBlock;

		selectIndependentVertex<<<blockNeeded, maxThreadPerBlock>>>(d_header_ptr, d_datas_ptr, d_indices_ptr, d_vertexUsed_ptr, d_independentVertex_ptr, vertexCount);
		//CudaCheckError();

		thrust::device_vector<bool> d_checkedvertexUsed(vertexCount);
		auto d_checkedvertexUsed_ptr = thrust::raw_pointer_cast(&d_checkedvertexUsed[0]);
		checkIndependentVertex<<<blockNeeded, maxThreadPerBlock>>>(d_header_ptr, d_datas_ptr, d_indices_ptr, d_independentVertex_ptr, d_checkedvertexUsed_ptr, vertexCount);
		//CudaCheckError();

		thrust::device_vector<BestEdge> d_bestedges(vertexCount);
		auto d_bestEdges_ptr = thrust::raw_pointer_cast(&d_bestedges[0]);
		getBestEdge<<<blockNeeded, maxThreadPerBlock>>>(d_header_ptr, d_datas_ptr, d_indices_ptr, d_vertices_ptr, d_checkedvertexUsed_ptr, d_vertexQuadric_ptr, d_bestEdges_ptr, vertexCount);
		//CudaCheckError();


	}

	time.reportTime();

