GPU并行优化

apps/TextureMesh/TextureMesh.cpp

@@ -1049,6 +1049,7 @@ bool checkFileExistence(const std::string& filename) {
 
 int main(int argc, LPCTSTR* argv)
 {
+	TD_TIMER_START();
 	#ifdef _DEBUGINFO
 	// set _crtBreakAlloc index to stop in <dbgheap.c> at allocation
 	_CrtSetDbgFlag(_CRTDBG_ALLOC_MEM_DF | _CRTDBG_LEAK_CHECK_DF);// | _CRTDBG_CHECK_ALWAYS_DF);
@@ -1128,7 +1129,6 @@ int main(int argc, LPCTSTR* argv)
 	printf("MAKE_PATH_SAFE(OPT::strViewsFileName): %s\n", MAKE_PATH_SAFE(OPT::strViewsFileName).c_str());
 	if (!OPT::strViewsFileName.empty())
 		views = ParseViewsFile(MAKE_PATH_SAFE(OPT::strViewsFileName), scene);
-	printf("Second\n");
 	// compute mesh texture
 	TD_TIMER_START();
 	if (!scene.TextureMesh(OPT::nResolutionLevel, OPT::nMinResolution, OPT::minCommonCameras, OPT::fOutlierThreshold, OPT::fRatioDataSmoothness,
@@ -1164,6 +1164,7 @@ int main(int argc, LPCTSTR* argv)
 		sml.Save(baseFileName+_T("_orthomap.txt"));
 	}
 
+	VERBOSE("All TextureMesh completed (%s)!", TD_TIMER_GET_FMT().c_str());
 	return EXIT_SUCCESS;
 }
 /*----------------------------------------------------------------*/

libs/MVS/CMakeLists.txt

@@ -12,6 +12,37 @@ if(VCG_FOUND)
 	add_definitions(${VCG_DEFINITIONS})
 endif()
 
+# 在find_package(CUDA)之前设置全局PIC
+SET(CMAKE_POSITION_INDEPENDENT_CODE ON)
+
+find_package(CUDA REQUIRED)
+
+SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} 
+    -Xcompiler=-fPIC 
+    --relocatable-device-code=true
+    --ptxas-options=-v
+    -gencode arch=compute_75,code=sm_75
+")
+
+# 启用CUDA语言支持（现代CMake方式）
+IF(CMAKE_VERSION VERSION_GREATER_EQUAL 3.18)
+    ENABLE_LANGUAGE(CUDA)
+ENDIF()
+
+cuda_add_library(MeshTextureCUDA 
+    cuda/MeshTextureCUDA.cu
+)
+
+SET_TARGET_PROPERTIES(MeshTextureCUDA PROPERTIES
+    CUDA_SEPARABLE_COMPILATION ON
+    CUDA_RESOLVE_DEVICE_SYMBOLS ON  # 关键：解析设备符号
+    POSITION_INDEPENDENT_CODE ON
+)
+# 为CUDA目标显式添加编译标志
+TARGET_COMPILE_OPTIONS(MeshTextureCUDA PRIVATE
+    $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-fPIC>
+)
+
 set(CERES_LIBS "")
 if(OpenMVS_USE_CERES)
 	FIND_PACKAGE(Ceres)
@@ -44,11 +75,19 @@ IF(CMAKE_VERSION VERSION_GREATER_EQUAL 3.16.0)
 	TARGET_PRECOMPILE_HEADERS(MVS PRIVATE "Common.h")
 endif()
 
+TARGET_LINK_LIBRARIES(MeshTextureCUDA
+    ${CUDA_LIBRARIES}
+    ${OpenCV_LIBRARIES}
+	cudart_static
+)
+
 # Link its dependencies
 if(_USE_CUDA)
 	SET_TARGET_PROPERTIES(MVS PROPERTIES CUDA_ARCHITECTURES "50;72;75")
 endif()
-TARGET_LINK_LIBRARIES(MVS PRIVATE Common Math IO CGAL::CGAL ${CERES_LIBRARIES} ${CUDA_CUDA_LIBRARY})
+TARGET_LINK_LIBRARIES(MVS PRIVATE Common Math IO CGAL::CGAL ${CERES_LIBRARIES} ${CUDA_CUDA_LIBRARY}
+    MeshTextureCUDA
+    ${CUDA_LIBRARIES})
 
 if(OpenMVS_USE_PYTHON)
 	# Create the Python wrapper

libs/MVS/SceneTexture.cpp

@@ -40,6 +40,7 @@
 #include <pybind11/embed.h>
 #include <pybind11/stl.h>
 #include "ConfigEnv.h"
+#include "cuda/MeshTextureCUDA.h"
 
 namespace py = pybind11;
 
@@ -73,6 +74,7 @@ using namespace MVS;
 #define MASK_FACE_OCCLUSION
 // #define DISPLAY_DEMO
 #define CACHE_MASK
+#define USE_CUDA
 
 // inference algorithm
 #if TEXOPT_INFERENCE == TEXOPT_INFERENCE_LBP
@@ -6078,7 +6080,7 @@ bool MeshTexture::FaceViewSelection3( unsigned minCommonCameras, float fOutlierT
 		LabelArr labels;
 
 		// 为每个面片单独决定是否使用虚拟面算法
-        Mesh::FaceIdxArr virtualFacesI; // 使用虚拟面算法的面片
+        // Mesh::FaceIdxArr virtualFacesI; // 使用虚拟面算法的面片
         Mesh::FaceIdxArr perFaceFaces; // 使用逐面算法的面片
 
 		/*
@@ -6104,12 +6106,9 @@ bool MeshTexture::FaceViewSelection3( unsigned minCommonCameras, float fOutlierT
 		// construct and use virtual faces for patch creation instead of actual mesh faces;
 		// the virtual faces are composed of coplanar triangles sharing same views
 		if (bUseVirtualFaces) {
-
 			Mesh::FaceIdxArr mapFaceToVirtualFace(faces.size()); // for each mesh face ID, store the virtual face ID witch contains it
-
 			 // 标记使用虚拟面算法的面片
 			std::vector<bool> isVirtualFace(faces.size(), true);
-
 			// 1) create FaceToVirtualFaceMap
 			FaceDataViewArr virtualFacesDatas;
 			VirtualFaceIdxsArr virtualFaces; // stores each virtual face as an array of mesh face ID
@@ -6117,7 +6116,9 @@ bool MeshTexture::FaceViewSelection3( unsigned minCommonCameras, float fOutlierT
 			// CreateVirtualFaces3(facesDatas, virtualFacesDatas, virtualFaces, minCommonCameras);
 			// CreateVirtualFaces4(facesDatas, virtualFacesDatas, virtualFaces, mapFaceToVirtualFace, minCommonCameras);
 			// CreateVirtualFaces6(facesDatas, virtualFacesDatas, virtualFaces, isVirtualFace, minCommonCameras);
+			TD_TIMER_STARTD();
 			CreateVirtualFaces7(facesDatas, virtualFacesDatas, virtualFaces, isVirtualFace, minCommonCameras);
+			DEBUG_EXTRA("CreateVirtualFaces7 completed: %s", TD_TIMER_GET_FMT().c_str());
 
 			size_t controlCounter(0);
 			FOREACH(idxVF, virtualFaces) {
@@ -6137,7 +6138,7 @@ bool MeshTexture::FaceViewSelection3( unsigned minCommonCameras, float fOutlierT
 				}
 			}
 
-			printf("perFaceFaces.size = %d\n", perFaceFaces.size());
+			// printf("perFaceFaces.size = %d\n", perFaceFaces.size());
 
 			ASSERT(controlCounter == faces.size());
 			// 2) create function to find virtual faces neighbors
@@ -8154,7 +8155,20 @@ void MeshTexture::LocalSeamLeveling3()
 		// extract image
 		const Image8U3& image0(images[texturePatch.label].image);
 		Image32F3 image, imageOrg;
+#ifdef USE_CUDA		
+		if (MeshTextureCUDA::ConvertToCUDA(image0(texturePatch.rect), image, 1.0/255.0))
+		{
+			// printf("ConvertToCUDA Successed!\n");
+		}
+		else
+		{
+			image0(texturePatch.rect).convertTo(image, CV_32FC3, 1.0/255.0);
+			// printf("ConvertToCUDA Failed!\n");
+		}
+#else
 		image0(texturePatch.rect).convertTo(image, CV_32FC3, 1.0/255.0);
+#endif
+
 		image.copyTo(imageOrg);
 		// render patch coverage
 		Image8U mask(image.size()); {
@@ -8252,11 +8266,43 @@ void MeshTexture::LocalSeamLeveling3()
 		}
 		// make sure the border is continuous and
 		// keep only the exterior tripe of the given size
+#ifdef USE_CUDA		
+		if (MeshTextureCUDA::ProcessMaskCUDA(mask, 20))
+		{
+			// printf("Success ProcessMaskCUDA!\n");
+            // 成功使用CUDA加速
+        } 
+		else 
+		{
+            // 回退到CPU版本
+			// printf("Failed ProcessMaskCUDA!\n");
+			ProcessMask(mask, 20);
+        }
+#else
 		ProcessMask(mask, 20);
+#endif
+
 		// compute texture patch blending
+#ifdef USE_CUDA
+		if (MeshTextureCUDA::PoissonBlendCUDA(image, imageOrg, mask, 1.0f))
+		{
+			// printf("Success PoissonBlendCUDA!");
+            // 成功使用CUDA加速
+        } 
+		else 
+		{
+            // 回退到CPU版本
+			// printf("Failed PoissonBlendCUDA!");
+            PoissonBlending(imageOrg, image, mask, 1.0f);
+        }
+#else
 		PoissonBlending(imageOrg, image, mask);
+#endif
 		// apply color correction to the patch image
 		cv::Mat imagePatch(image0(texturePatch.rect));
+		#ifdef TEXOPT_USE_OPENMP
+		#pragma omp parallel for collapse(2)
+		#endif
 		for (int r=0; r<image.rows; ++r) {
 			for (int c=0; c<image.cols; ++c) {
 				if (mask(r,c) == empty)
@@ -8372,6 +8418,8 @@ void MeshTexture::GenerateTexture(bool bGlobalSeamLeveling, bool bLocalSeamLevel
 
 
 	LOG_OUT() << "First loop completed" << std::endl;
+
+	TD_TIMER_STARTD();
 	// perform seam leveling
 	if (texturePatches.size() > 2 && (bGlobalSeamLeveling || bLocalSeamLeveling)) {
 		// create seam vertices and edges
@@ -8381,16 +8429,18 @@ void MeshTexture::GenerateTexture(bool bGlobalSeamLeveling, bool bLocalSeamLevel
 		if (bGlobalSeamLeveling) {
 			TD_TIMER_STARTD();
 			GlobalSeamLeveling3();
-			DEBUG_ULTIMATE("\tglobal seam leveling completed (%s)", TD_TIMER_GET_FMT().c_str());
+			DEBUG_EXTRA("\tglobal seam leveling completed (%s)", TD_TIMER_GET_FMT().c_str());
 		}
 
 		// perform local seam leveling
 		if (bLocalSeamLeveling) {
 			TD_TIMER_STARTD();
-			LocalSeamLeveling();
+			// LocalSeamLeveling();
-			DEBUG_ULTIMATE("\tlocal seam leveling completed (%s)", TD_TIMER_GET_FMT().c_str());
+			LocalSeamLeveling3();
+			DEBUG_EXTRA("\tlocal seam leveling completed (%s)", TD_TIMER_GET_FMT().c_str());
 		}
 	}
+	DEBUG_EXTRA("seam (%s)", TD_TIMER_GET_FMT().c_str());
 
 	// merge texture patches with overlapping rectangles
 	for (unsigned i=0; i<texturePatches.size()-1; ++i) {
@@ -9635,12 +9685,12 @@ bool Scene::TextureMesh(unsigned nResolutionLevel, unsigned nMinResolution, unsi
 				dict3 = result["result3"].cast<py::dict>();
 			}
 
-			printf("dict1 size=%d, dict2 size=%d, dict3 size=%d\n", dict1.size(), dict2.size(), dict3.size());
+			// printf("dict1 size=%d, dict2 size=%d, dict3 size=%d\n", dict1.size(), dict2.size(), dict3.size());
 
 			// 处理返回的可见面字典
 			for (auto item : dict1) {
 				std::string image_name = item.first.cast<std::string>();
-				printf("dict1 mask image name=%s\n", image_name.c_str());
+				// printf("dict1 mask image name=%s\n", image_name.c_str());
 				py::list visible_faces = item.second.cast<py::list>();
 				
 				std::unordered_set<int> face_set;
@@ -9657,7 +9707,7 @@ bool Scene::TextureMesh(unsigned nResolutionLevel, unsigned nMinResolution, unsi
 
 			for (auto item : dict2) {
 				std::string image_name = item.first.cast<std::string>();
-				printf("dict2 mask image name=%s\n", image_name.c_str());
+				// printf("dict2 mask image name=%s\n", image_name.c_str());
 				py::list edge_faces = item.second.cast<py::list>();
 				
 				std::unordered_set<int> face_set;
@@ -9669,7 +9719,7 @@ bool Scene::TextureMesh(unsigned nResolutionLevel, unsigned nMinResolution, unsi
 
 			for (auto item : dict3) {
 				std::string image_name = item.first.cast<std::string>();
-				printf("dict3 mask image name=%s\n", image_name.c_str());
+				// printf("dict3 mask image name=%s\n", image_name.c_str());
 				py::list delete_edge_faces = item.second.cast<py::list>();
 				
 				std::unordered_set<int> face_set;
@@ -9762,7 +9812,6 @@ bool Scene::TextureMesh(unsigned nResolutionLevel, unsigned nMinResolution, unsi
 		texture.GenerateTexture(bGlobalSeamLeveling, bLocalSeamLeveling, nTextureSizeMultiple, nRectPackingHeuristic, colEmpty, fSharpnessWeight, maxTextureSize, baseFileName, bOriginFaceview);
 		DEBUG_EXTRA("Generating texture atlas and image completed: %u patches, %u image size, %u textures (%s)", texture.texturePatches.size(), mesh.texturesDiffuse[0].width(), mesh.texturesDiffuse.size(), TD_TIMER_GET_FMT().c_str());
 	}		
-
 	#ifdef DISPLAY_DEMO
 	try {
 		py::scoped_interpreter guard{};  // 自动管理解释器生命周期