 		fprintf(stderr, "CUDA Runtime API error : %s.\n", cudaGetErrorString(err));
 		return;
 	}
+}
+
+__device__
+double MyAtomicAdd(double* address, double val)
+{
+	unsigned long long int* address_as_ull = (unsigned long long int*)address;
+
+	unsigned long long int old = *address_as_ull, assumed;
+
+	do{
+		assumed = old;
+		old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
+	} while (assumed != old);
+
+	return __longlong_as_double(old);
 }
 
 /*

				total_weight = 0.0;
				centriodTimesWeight = { 0.0, 0.0, 0.0 };
			}
			
			current_posTimesWeight.x /= current_weight;
			current_posTimesWeight.y /= current_weight;
			current_posTimesWeight.z /= current_weight;

			current_weight /= triangle_nb;
			
			double3 temp_pos;
			temp_pos.x = SeedsInformation[current_seed * 4 + 0];
			temp_pos.y = SeedsInformation[current_seed * 4 + 1];
			temp_pos.z = SeedsInformation[current_seed * 4 + 2];
			double temp_weight = SeedsInformation[current_seed * 4 + 3] ;

			temp_weight += current_weight;
			temp_pos.x += current_posTimesWeight.x * current_weight;
			temp_pos.y += current_posTimesWeight.x * current_weight;
			temp_pos.z += current_posTimesWeight.x * current_weight;

			atomicAdd(&SeedsInformation[current_seed * 4 + 0], temp_pos.x);
			atomicAdd(&SeedsInformation[current_seed * 4 + 1], temp_pos.y);
			atomicAdd(&SeedsInformation[current_seed * 4 + 2], temp_pos.z);
			atomicAdd(&SeedsInformation[current_seed * 4 + 3], temp_weight);

			/*------------------------ test part -------------------------*/
			if (tid == 0 && i == 3)
			{
				test_index[20] = current_seed;
				test_index[88] = vertex_nb;
				/*for (int j = 0; j < vertex_nb * 4; ++j)
				{
					test_dis[j] = polygon[j];

		free(polygon);
	}
}

extern "C" void runCuda(double* host_seeds_pointer, double* host_mesh_vertex_pointer,
	int* host_facet_index, int points_nb, int mesh_vertex_nb, int mesh_facet_nb)
