From aa8f03e29dbf5ea8e61039d49c5cea798a0a9c3c Mon Sep 17 00:00:00 2001 From: Sameer Sheorey Date: Wed, 9 Oct 2024 16:42:54 -0700 Subject: [PATCH] Fix no tbb target with BUILD_SYCL_MODULE Use sycl_target_sources for files with sycl code in cmake. style fix --- 3rdparty/embree/embree.cmake | 7 +- 3rdparty/find_dependencies.cmake | 6 +- 3rdparty/mkl/tbb.cmake | 1 + cpp/open3d/t/geometry/CMakeLists.txt | 11 +- cpp/open3d/t/geometry/RaycastingScene.cpp | 626 ++++++++++-------- cpp/open3d/t/geometry/RaycastingScene.h | 3 +- cpp/pybind/CMakeLists.txt | 2 +- cpp/pybind/core/sycl_utils.cpp | 2 +- cpp/pybind/t/geometry/raycasting_scene.cpp | 3 +- docker/Dockerfile.ci | 24 +- .../test/t/geometry/test_raycasting_scene.py | 46 +- 11 files changed, 414 insertions(+), 317 deletions(-) diff --git a/3rdparty/embree/embree.cmake b/3rdparty/embree/embree.cmake index 1bd89b2ba58..6d55198ec69 100644 --- a/3rdparty/embree/embree.cmake +++ b/3rdparty/embree/embree.cmake @@ -68,7 +68,9 @@ if(BUILD_SYCL_MODULE) set(ISA_ARGS ${ISA_ARGS} -DCMAKE_CXX_COMPILER=icpx) set(ISA_ARGS ${ISA_ARGS} -DCMAKE_C_COMPILER=icx) set(ISA_ARGS ${ISA_ARGS} -DEMBREE_SYCL_SUPPORT=ON) - #list(APPEND ISA_LIBS embree4_sycl) + list(APPEND ISA_LIBS embree4_sycl ze_wrapper) + list(APPEND ISA_BUILD_BYPRODUCTS "/${Open3D_INSTALL_LIB_DIR}/${CMAKE_STATIC_LIBRARY_PREFIX}embree4_sycl${CMAKE_STATIC_LIBRARY_SUFFIX}" + "/${Open3D_INSTALL_LIB_DIR}/${CMAKE_STATIC_LIBRARY_PREFIX}ze_wrapper${CMAKE_STATIC_LIBRARY_SUFFIX}") endif() @@ -101,11 +103,10 @@ ExternalProject_Add( /${Open3D_INSTALL_LIB_DIR}/${CMAKE_STATIC_LIBRARY_PREFIX}math${CMAKE_STATIC_LIBRARY_SUFFIX} /${Open3D_INSTALL_LIB_DIR}/${CMAKE_STATIC_LIBRARY_PREFIX}tasking${CMAKE_STATIC_LIBRARY_SUFFIX} /${Open3D_INSTALL_LIB_DIR}/${CMAKE_STATIC_LIBRARY_PREFIX}ze_wrapper${CMAKE_STATIC_LIBRARY_SUFFIX} - /${Open3D_INSTALL_LIB_DIR}/${CMAKE_STATIC_LIBRARY_PREFIX}embree_rthwif${CMAKE_STATIC_LIBRARY_SUFFIX} ${ISA_BUILD_BYPRODUCTS} ) ExternalProject_Get_Property(ext_embree INSTALL_DIR) set(EMBREE_INCLUDE_DIRS ${INSTALL_DIR}/include/ ${INSTALL_DIR}/src/ext_embree/) # "/" is critical. set(EMBREE_LIB_DIR ${INSTALL_DIR}/${Open3D_INSTALL_LIB_DIR}) -set(EMBREE_LIBRARIES embree4_sycl embree4 simd lexers sys math tasking ze_wrapper embree_rthwif ${ISA_LIBS}) +set(EMBREE_LIBRARIES embree4 simd lexers sys math tasking ${ISA_LIBS}) diff --git a/3rdparty/find_dependencies.cmake b/3rdparty/find_dependencies.cmake index c3dc3885565..d2afbe405c9 100644 --- a/3rdparty/find_dependencies.cmake +++ b/3rdparty/find_dependencies.cmake @@ -1384,7 +1384,7 @@ if(BUILD_GUI) ${CPP_LIBRARY} ${CPPABI_LIBRARY}) message(STATUS "Filament C++ libraries: ${CPP_LIBRARY} ${CPPABI_LIBRARY}") if (LIBCPP_VERSION GREATER 11) - message(WARNING "libc++ (LLVM) version ${LIBCPP_VERSION} > 11 includes libunwind that " + message(WARNING "libc++ (LLVM) version ${LIBCPP_VERSION} > 11 includes libunwind that " "interferes with the system libunwind.so.8 and may crash Python code when exceptions " "are used. Please consider using libc++ (LLVM) v11.") endif() @@ -1697,7 +1697,7 @@ else(OPEN3D_USE_ONEAPI_PACKAGES) INCLUDE_DIRS ${STATIC_MKL_INCLUDE_DIR} LIB_DIR ${STATIC_MKL_LIB_DIR} LIBRARIES ${STATIC_MKL_LIBRARIES} - DEPENDS ext_tbb ext_mkl_include ext_mkl + DEPENDS Open3D::3rdparty_tbb ext_mkl_include ext_mkl ) if(UNIX) target_compile_options(3rdparty_blas INTERFACE "$<$:-m64>") @@ -1719,7 +1719,7 @@ else(OPEN3D_USE_ONEAPI_PACKAGES) endif() if(NOT USE_SYSTEM_TBB) include(${Open3D_3RDPARTY_DIR}/mkl/tbb.cmake) - list(APPEND Open3D_3RDPARTY_PRIVATE_TARGETS_FROM_CUSTOM TBB::tbb) + list(APPEND Open3D_3RDPARTY_PRIVATE_TARGETS_FROM_CUSTOM Open3D::3rdparty_tbb) else() list(APPEND Open3D_3RDPARTY_PRIVATE_TARGETS_FROM_SYSTEM Open3D::3rdparty_tbb) endif() diff --git a/3rdparty/mkl/tbb.cmake b/3rdparty/mkl/tbb.cmake index 385af4b644b..a36435ee1b4 100644 --- a/3rdparty/mkl/tbb.cmake +++ b/3rdparty/mkl/tbb.cmake @@ -38,3 +38,4 @@ install(TARGETS tbb EXPORT ${PROJECT_NAME}Targets RUNTIME DESTINATION ${Open3D_INSTALL_BIN_DIR} COMPONENT tbb ) +add_library(${PROJECT_NAME}::3rdparty_tbb ALIAS tbb) diff --git a/cpp/open3d/t/geometry/CMakeLists.txt b/cpp/open3d/t/geometry/CMakeLists.txt index 1b1ba5760cf..e345c819b6f 100644 --- a/cpp/open3d/t/geometry/CMakeLists.txt +++ b/cpp/open3d/t/geometry/CMakeLists.txt @@ -7,7 +7,6 @@ target_sources(tgeometry PRIVATE LineSet.cpp BoundingVolume.cpp PointCloud.cpp - RaycastingScene.cpp RGBDImage.cpp TensorMap.cpp TriangleMesh.cpp @@ -16,6 +15,16 @@ target_sources(tgeometry PRIVATE VtkUtils.cpp ) +if (BUILD_SYCL_MODULE) + open3d_sycl_target_sources(tgeometry PRIVATE + RaycastingScene.cpp + ) +else() + target_sources(tgeometry PRIVATE + RaycastingScene.cpp + ) +endif() + open3d_show_and_abort_on_warning(tgeometry) open3d_set_global_properties(tgeometry) open3d_set_open3d_lib_properties(tgeometry) diff --git a/cpp/open3d/t/geometry/RaycastingScene.cpp b/cpp/open3d/t/geometry/RaycastingScene.cpp index aa0b200bc76..2691d9d6fe9 100644 --- a/cpp/open3d/t/geometry/RaycastingScene.cpp +++ b/cpp/open3d/t/geometry/RaycastingScene.cpp @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -27,7 +28,6 @@ #include "open3d/utility/Helper.h" #include "open3d/utility/Logging.h" - namespace callbacks { struct GeomPrimID { @@ -43,7 +43,8 @@ struct CountIntersectionsContext { }; #ifdef BUILD_SYCL_MODULE -RTC_SYCL_INDIRECTLY_CALLABLE void CountIntersectionsFunc(const RTCFilterFunctionNArguments* args) { +RTC_SYCL_INDIRECTLY_CALLABLE void CountIntersectionsFunc( + const RTCFilterFunctionNArguments* args) { #else void CountIntersectionsFunc(const RTCFilterFunctionNArguments* args) { #endif @@ -57,7 +58,8 @@ void CountIntersectionsFunc(const RTCFilterFunctionNArguments* args) { // Avoid crashing when debug visualizations are used. if (context == nullptr) return; - GeomPrimID *previous_geom_prim_ID_tfar = context->previous_geom_prim_ID_tfar; + GeomPrimID* previous_geom_prim_ID_tfar = + context->previous_geom_prim_ID_tfar; int* intersections = context->intersections; // Iterate over all rays in ray packet. @@ -100,7 +102,8 @@ struct ListIntersectionsContext { }; #ifdef BUILD_SYCL_MODULE -RTC_SYCL_INDIRECTLY_CALLABLE void ListIntersectionsFunc(const RTCFilterFunctionNArguments* args) { +RTC_SYCL_INDIRECTLY_CALLABLE void ListIntersectionsFunc( + const RTCFilterFunctionNArguments* args) { #else void ListIntersectionsFunc(const RTCFilterFunctionNArguments* args) { #endif @@ -114,7 +117,8 @@ void ListIntersectionsFunc(const RTCFilterFunctionNArguments* args) { // Avoid crashing when debug visualizations are used. if (context == nullptr) return; - GeomPrimID *previous_geom_prim_ID_tfar = context->previous_geom_prim_ID_tfar; + GeomPrimID* previous_geom_prim_ID_tfar = + context->previous_geom_prim_ID_tfar; unsigned int* ray_ids = context->ray_ids; unsigned int* geometry_ids = context->geometry_ids; unsigned int* primitive_ids = context->primitive_ids; @@ -157,7 +161,7 @@ void ListIntersectionsFunc(const RTCFilterFunctionNArguments* args) { } } -} // namespace callbacks +} // namespace callbacks namespace { @@ -166,16 +170,17 @@ typedef Eigen::AlignedVector3 Vec3fa; typedef Eigen::Matrix Vec2f; typedef Eigen::Vector3f Vec3f; -void enablePersistentJITCache() -{ +#ifdef BUILD_SYCL_MODULE +void enablePersistentJITCache() { #if defined(_WIN32) - _putenv_s("SYCL_CACHE_PERSISTENT","1"); - _putenv_s("SYCL_CACHE_DIR","cache"); + _putenv_s("SYCL_CACHE_PERSISTENT", "1"); + _putenv_s("SYCL_CACHE_DIR", "cache"); #else - setenv("SYCL_CACHE_PERSISTENT","1",1); - setenv("SYCL_CACHE_DIR","cache",1); + setenv("SYCL_CACHE_PERSISTENT", "1", 1); + setenv("SYCL_CACHE_DIR", "cache", 1); #endif } +#endif // Error function called by embree. void ErrorFunction(void* userPtr, enum RTCError error, const char* str) { @@ -378,47 +383,47 @@ struct RaycastingScene::Impl { } virtual void CastRays(const float* const rays, - const size_t num_rays, - float* t_hit, - unsigned int* geometry_ids, - unsigned int* primitive_ids, - float* primitive_uvs, - float* primitive_normals, - const int nthreads, - const bool line_intersection) = 0; - + const size_t num_rays, + float* t_hit, + unsigned int* geometry_ids, + unsigned int* primitive_ids, + float* primitive_uvs, + float* primitive_normals, + const int nthreads, + const bool line_intersection) = 0; + virtual void TestOcclusions(const float* const rays, - const size_t num_rays, - const float tnear, - const float tfar, - int8_t* occluded, - const int nthreads) = 0; - + const size_t num_rays, + const float tnear, + const float tfar, + int8_t* occluded, + const int nthreads) = 0; + virtual void CountIntersections(const float* const rays, - const size_t num_rays, - int* intersections, - const int nthreads) = 0; - + const size_t num_rays, + int* intersections, + const int nthreads) = 0; + virtual void ListIntersections(const float* const rays, - const size_t num_rays, - const size_t num_intersections, - int* cumsum, - unsigned int* track_intersections, - unsigned int* ray_ids, - unsigned int* geometry_ids, - unsigned int* primitive_ids, - float* primitive_uvs, - float* t_hit, - const int nthreads) = 0; - + const size_t num_rays, + const size_t num_intersections, + int* cumsum, + unsigned int* track_intersections, + unsigned int* ray_ids, + unsigned int* geometry_ids, + unsigned int* primitive_ids, + float* primitive_uvs, + float* t_hit, + const int nthreads) = 0; + virtual void ComputeClosestPoints(const float* const query_points, - const size_t num_query_points, - float* closest_points, - unsigned int* geometry_ids, - unsigned int* primitive_ids, - float* primitive_uvs, - float* primitive_normals, - const int nthreads) = 0; + const size_t num_query_points, + float* closest_points, + unsigned int* geometry_ids, + unsigned int* primitive_ids, + float* primitive_uvs, + float* primitive_normals, + const int nthreads) = 0; }; #ifdef BUILD_SYCL_MODULE @@ -433,8 +438,9 @@ struct RaycastingScene::SYCLImpl : public RaycastingScene::Impl { try { sycl_device_ = sycl::device(rtcSYCLDeviceSelector); - } catch(std::exception& e) { - utility::LogError("Caught exception creating sycl::device: {}", e.what()); + } catch (std::exception& e) { + utility::LogError("Caught exception creating sycl::device: {}", + e.what()); return; } @@ -445,9 +451,8 @@ struct RaycastingScene::SYCLImpl : public RaycastingScene::Impl { rtcSetDeviceSYCLDevice(device_, sycl_device_); if (!device_) { - utility::LogError( - "Error %d: cannot create device\n", - rtcGetDeviceError(NULL)); + utility::LogError("Error %d: cannot create device\n", + rtcGetDeviceError(NULL)); } } @@ -462,65 +467,72 @@ struct RaycastingScene::SYCLImpl : public RaycastingScene::Impl { const bool line_intersection) override { CommitScene(); - auto scene = this->scene_; - queue_.submit([=](sycl::handler& cgh) { - cgh.parallel_for(sycl::range<1>(num_rays),[=](sycl::item<1> item, sycl::kernel_handler kh) { - const size_t i = item.get_id(0); - - struct RTCRayHit rh; - const float* r = &rays[i * 6]; - rh.ray.org_x = r[0]; - rh.ray.org_y = r[1]; - rh.ray.org_z = r[2]; - if (line_intersection) { - rh.ray.dir_x = r[3] - r[0]; - rh.ray.dir_y = r[4] - r[1]; - rh.ray.dir_z = r[5] - r[2]; - } else { - rh.ray.dir_x = r[3]; - rh.ray.dir_y = r[4]; - rh.ray.dir_z = r[5]; - } - rh.ray.tnear = 0; - if (line_intersection) { - rh.ray.tfar = 1.f; - } else { - rh.ray.tfar = std::numeric_limits::infinity(); - } - rh.ray.mask = -1; - rh.ray.id = i; - rh.ray.flags = 0; - rh.hit.geomID = RTC_INVALID_GEOMETRY_ID; - rh.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; - - rtcIntersect1(scene, &rh); - - t_hit[i] = rh.ray.tfar; - if (rh.hit.geomID != RTC_INVALID_GEOMETRY_ID) { - geometry_ids[i] = rh.hit.geomID; - primitive_ids[i] = rh.hit.primID; - primitive_uvs[i * 2 + 0] = rh.hit.u; - primitive_uvs[i * 2 + 1] = rh.hit.v; - float inv_norm = 1.f / std::sqrt(rh.hit.Ng_x * rh.hit.Ng_x + - rh.hit.Ng_y * rh.hit.Ng_y + - rh.hit.Ng_z * rh.hit.Ng_z); - primitive_normals[i * 3 + 0] = rh.hit.Ng_x * inv_norm; - primitive_normals[i * 3 + 1] = rh.hit.Ng_y * inv_norm; - primitive_normals[i * 3 + 2] = rh.hit.Ng_z * inv_norm; - } else { - geometry_ids[i] = RTC_INVALID_GEOMETRY_ID; - primitive_ids[i] = RTC_INVALID_GEOMETRY_ID; - primitive_uvs[i * 2 + 0] = 0; - primitive_uvs[i * 2 + 1] = 0; - primitive_normals[i * 3 + 0] = 0; - primitive_normals[i * 3 + 1] = 0; - primitive_normals[i * 3 + 2] = 0; - } - }); + auto scene = this->scene_; + queue_.submit([=](sycl::handler& cgh) { + cgh.parallel_for( + sycl::range<1>(num_rays), + [=](sycl::item<1> item, sycl::kernel_handler kh) { + const size_t i = item.get_id(0); + + struct RTCRayHit rh; + const float* r = &rays[i * 6]; + rh.ray.org_x = r[0]; + rh.ray.org_y = r[1]; + rh.ray.org_z = r[2]; + if (line_intersection) { + rh.ray.dir_x = r[3] - r[0]; + rh.ray.dir_y = r[4] - r[1]; + rh.ray.dir_z = r[5] - r[2]; + } else { + rh.ray.dir_x = r[3]; + rh.ray.dir_y = r[4]; + rh.ray.dir_z = r[5]; + } + rh.ray.tnear = 0; + if (line_intersection) { + rh.ray.tfar = 1.f; + } else { + rh.ray.tfar = + std::numeric_limits::infinity(); + } + rh.ray.mask = -1; + rh.ray.id = i; + rh.ray.flags = 0; + rh.hit.geomID = RTC_INVALID_GEOMETRY_ID; + rh.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; + + rtcIntersect1(scene, &rh); + + t_hit[i] = rh.ray.tfar; + if (rh.hit.geomID != RTC_INVALID_GEOMETRY_ID) { + geometry_ids[i] = rh.hit.geomID; + primitive_ids[i] = rh.hit.primID; + primitive_uvs[i * 2 + 0] = rh.hit.u; + primitive_uvs[i * 2 + 1] = rh.hit.v; + float inv_norm = + 1.f / std::sqrt(rh.hit.Ng_x * rh.hit.Ng_x + + rh.hit.Ng_y * rh.hit.Ng_y + + rh.hit.Ng_z * rh.hit.Ng_z); + primitive_normals[i * 3 + 0] = + rh.hit.Ng_x * inv_norm; + primitive_normals[i * 3 + 1] = + rh.hit.Ng_y * inv_norm; + primitive_normals[i * 3 + 2] = + rh.hit.Ng_z * inv_norm; + } else { + geometry_ids[i] = RTC_INVALID_GEOMETRY_ID; + primitive_ids[i] = RTC_INVALID_GEOMETRY_ID; + primitive_uvs[i * 2 + 0] = 0; + primitive_uvs[i * 2 + 1] = 0; + primitive_normals[i * 3 + 0] = 0; + primitive_normals[i * 3 + 1] = 0; + primitive_normals[i * 3 + 2] = 0; + } + }); }); queue_.wait_and_throw(); } - + void TestOcclusions(const float* const rays, const size_t num_rays, const float tnear, @@ -529,41 +541,44 @@ struct RaycastingScene::SYCLImpl : public RaycastingScene::Impl { const int nthreads) override { CommitScene(); - auto scene = this->scene_; - queue_.submit([=](sycl::handler& cgh) { - cgh.parallel_for(sycl::range<1>(num_rays),[=](sycl::item<1> item, sycl::kernel_handler kh) { - struct RTCRayQueryContext context; - rtcInitRayQueryContext(&context); - - RTCOccludedArguments args; - rtcInitOccludedArguments(&args); - args.context = &context; - - const size_t i = item.get_id(0); - - struct RTCRay ray; - const float* r = &rays[i * 6]; - ray.org_x = r[0]; - ray.org_y = r[1]; - ray.org_z = r[2]; - ray.dir_x = r[3]; - ray.dir_y = r[4]; - ray.dir_z = r[5]; - ray.tnear = tnear; - ray.tfar = tfar; - ray.mask = -1; - ray.id = i; - ray.flags = 0; - - rtcOccluded1(scene, &ray, &args); - - occluded[i] = int8_t( - -std::numeric_limits::infinity() == ray.tfar); - }); + auto scene = this->scene_; + queue_.submit([=](sycl::handler& cgh) { + cgh.parallel_for( + sycl::range<1>(num_rays), + [=](sycl::item<1> item, sycl::kernel_handler kh) { + struct RTCRayQueryContext context; + rtcInitRayQueryContext(&context); + + RTCOccludedArguments args; + rtcInitOccludedArguments(&args); + args.context = &context; + + const size_t i = item.get_id(0); + + struct RTCRay ray; + const float* r = &rays[i * 6]; + ray.org_x = r[0]; + ray.org_y = r[1]; + ray.org_z = r[2]; + ray.dir_x = r[3]; + ray.dir_y = r[4]; + ray.dir_z = r[5]; + ray.tnear = tnear; + ray.tfar = tfar; + ray.mask = -1; + ray.id = i; + ray.flags = 0; + + rtcOccluded1(scene, &ray, &args); + + occluded[i] = int8_t( + -std::numeric_limits::infinity() == + ray.tfar); + }); }); queue_.wait_and_throw(); } - + void CountIntersections(const float* const rays, const size_t num_rays, int* intersections, @@ -571,63 +586,74 @@ struct RaycastingScene::SYCLImpl : public RaycastingScene::Impl { CommitScene(); queue_.memset(intersections, 0, sizeof(int) * num_rays).wait(); - - callbacks::GeomPrimID* previous_geom_prim_ID_tfar = sycl::malloc_device(num_rays, queue_); + + callbacks::GeomPrimID* previous_geom_prim_ID_tfar = + sycl::malloc_device(num_rays, queue_); // Check if allocation was successful if (!previous_geom_prim_ID_tfar) { throw std::runtime_error("Failed to allocate device memory"); } - - auto host_previous_geom_prim_ID_tfar = std::unique_ptr>(new callbacks::GeomPrimID[num_rays]); + + auto host_previous_geom_prim_ID_tfar = + std::unique_ptr>( + new callbacks::GeomPrimID[num_rays]); for (size_t i = 0; i < num_rays; ++i) { - host_previous_geom_prim_ID_tfar[i] = {uint32_t(RTC_INVALID_GEOMETRY_ID), - uint32_t(RTC_INVALID_GEOMETRY_ID), 0.f}; + host_previous_geom_prim_ID_tfar[i] = { + uint32_t(RTC_INVALID_GEOMETRY_ID), + uint32_t(RTC_INVALID_GEOMETRY_ID), 0.f}; } // Copy the initialized data to the device - queue_.memcpy(host_previous_geom_prim_ID_tfar.get(), previous_geom_prim_ID_tfar, num_rays * sizeof(callbacks::GeomPrimID)).wait(); - - auto scene = this->scene_; - queue_.submit([=](sycl::handler& cgh) { - cgh.parallel_for(sycl::range<1>(num_rays),[=](sycl::item<1> item, sycl::kernel_handler kh) { - callbacks::CountIntersectionsContext context; - rtcInitRayQueryContext(&context.context); - context.previous_geom_prim_ID_tfar = previous_geom_prim_ID_tfar; - context.intersections = intersections; - - RTCIntersectArguments args; - rtcInitIntersectArguments(&args); - args.filter = callbacks::CountIntersectionsFunc; - args.context = &context.context; - - const size_t i = item.get_id(0); - - struct RTCRayHit rh; - const float* r = &rays[i * 6]; - rh.ray.org_x = r[0]; - rh.ray.org_y = r[1]; - rh.ray.org_z = r[2]; - rh.ray.dir_x = r[3]; - rh.ray.dir_y = r[4]; - rh.ray.dir_z = r[5]; - rh.ray.tnear = 0; - rh.ray.tfar = std::numeric_limits::infinity(); - rh.ray.mask = -1; - rh.ray.flags = 0; - rh.ray.id = i; - rh.hit.geomID = RTC_INVALID_GEOMETRY_ID; - rh.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; - - rtcIntersect1(scene, &rh, &args); - }); + queue_.memcpy(host_previous_geom_prim_ID_tfar.get(), + previous_geom_prim_ID_tfar, + num_rays * sizeof(callbacks::GeomPrimID)) + .wait(); + + auto scene = this->scene_; + queue_.submit([=](sycl::handler& cgh) { + cgh.parallel_for( + sycl::range<1>(num_rays), + [=](sycl::item<1> item, sycl::kernel_handler kh) { + callbacks::CountIntersectionsContext context; + rtcInitRayQueryContext(&context.context); + context.previous_geom_prim_ID_tfar = + previous_geom_prim_ID_tfar; + context.intersections = intersections; + + RTCIntersectArguments args; + rtcInitIntersectArguments(&args); + args.filter = callbacks::CountIntersectionsFunc; + args.context = &context.context; + + const size_t i = item.get_id(0); + + struct RTCRayHit rh; + const float* r = &rays[i * 6]; + rh.ray.org_x = r[0]; + rh.ray.org_y = r[1]; + rh.ray.org_z = r[2]; + rh.ray.dir_x = r[3]; + rh.ray.dir_y = r[4]; + rh.ray.dir_z = r[5]; + rh.ray.tnear = 0; + rh.ray.tfar = std::numeric_limits::infinity(); + rh.ray.mask = -1; + rh.ray.flags = 0; + rh.ray.id = i; + rh.hit.geomID = RTC_INVALID_GEOMETRY_ID; + rh.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; + + rtcIntersect1(scene, &rh, &args); + }); }); queue_.wait_and_throw(); // Free the allocated memory sycl::free(previous_geom_prim_ID_tfar, queue_); } - + void ListIntersections(const float* const rays, const size_t num_rays, const size_t num_intersections, @@ -641,72 +667,87 @@ struct RaycastingScene::SYCLImpl : public RaycastingScene::Impl { const int nthreads) override { CommitScene(); - queue_.memset(track_intersections, 0, sizeof(uint32_t) * num_rays).wait(); + queue_.memset(track_intersections, 0, sizeof(uint32_t) * num_rays) + .wait(); queue_.memset(ray_ids, 0, sizeof(uint32_t) * num_intersections).wait(); - queue_.memset(geometry_ids, 0, sizeof(uint32_t) * num_intersections).wait(); - queue_.memset(primitive_ids, 0, sizeof(uint32_t) * num_intersections).wait(); - queue_.memset(primitive_uvs, 0, sizeof(float) * num_intersections * 2).wait(); + queue_.memset(geometry_ids, 0, sizeof(uint32_t) * num_intersections) + .wait(); + queue_.memset(primitive_ids, 0, sizeof(uint32_t) * num_intersections) + .wait(); + queue_.memset(primitive_uvs, 0, sizeof(float) * num_intersections * 2) + .wait(); queue_.memset(t_hit, 0, sizeof(float) * num_intersections).wait(); - callbacks::GeomPrimID* previous_geom_prim_ID_tfar = sycl::malloc_device(num_rays, queue_); + callbacks::GeomPrimID* previous_geom_prim_ID_tfar = + sycl::malloc_device(num_rays, queue_); // Check if allocation was successful if (!previous_geom_prim_ID_tfar) { throw std::runtime_error("Failed to allocate device memory"); } - - auto host_previous_geom_prim_ID_tfar = std::unique_ptr>(new callbacks::GeomPrimID[num_rays]); + + auto host_previous_geom_prim_ID_tfar = + std::unique_ptr>( + new callbacks::GeomPrimID[num_rays]); for (size_t i = 0; i < num_rays; ++i) { - host_previous_geom_prim_ID_tfar[i] = {uint32_t(RTC_INVALID_GEOMETRY_ID), - uint32_t(RTC_INVALID_GEOMETRY_ID), 0.f}; + host_previous_geom_prim_ID_tfar[i] = { + uint32_t(RTC_INVALID_GEOMETRY_ID), + uint32_t(RTC_INVALID_GEOMETRY_ID), 0.f}; } // Copy the initialized data to the device - queue_.memcpy(host_previous_geom_prim_ID_tfar.get(), previous_geom_prim_ID_tfar, num_rays * sizeof(callbacks::GeomPrimID)).wait(); + queue_.memcpy(host_previous_geom_prim_ID_tfar.get(), + previous_geom_prim_ID_tfar, + num_rays * sizeof(callbacks::GeomPrimID)) + .wait(); // cumsum int* cumsum_ = sycl::malloc_device(num_rays, queue_); queue_.memcpy(cumsum, cumsum_, num_rays * sizeof(int)).wait(); - auto scene = this->scene_; - queue_.submit([=](sycl::handler& cgh) { - cgh.parallel_for(sycl::range<1>(num_rays),[=](sycl::item<1> item, sycl::kernel_handler kh) { - callbacks::ListIntersectionsContext context; - rtcInitRayQueryContext(&context.context); - context.previous_geom_prim_ID_tfar = previous_geom_prim_ID_tfar; - context.ray_ids = ray_ids; - context.geometry_ids = geometry_ids; - context.primitive_ids = primitive_ids; - context.primitive_uvs = primitive_uvs; - context.t_hit = t_hit; - context.cumsum = cumsum_; - context.track_intersections = track_intersections; - - RTCIntersectArguments args; - rtcInitIntersectArguments(&args); - //args.filter = callbacks::ListIntersectionsFunc; - args.context = &context.context; - - const size_t i = item.get_id(0); - - struct RTCRayHit rh; - const float* r = &rays[i * 6]; - rh.ray.org_x = r[0]; - rh.ray.org_y = r[1]; - rh.ray.org_z = r[2]; - rh.ray.dir_x = r[3]; - rh.ray.dir_y = r[4]; - rh.ray.dir_z = r[5]; - rh.ray.tnear = 0; - rh.ray.tfar = std::numeric_limits::infinity(); - rh.ray.mask = -1; - rh.ray.flags = 0; - rh.ray.id = i; - rh.hit.geomID = RTC_INVALID_GEOMETRY_ID; - rh.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; - - rtcIntersect1(scene, &rh, &args); - }); + auto scene = this->scene_; + queue_.submit([=](sycl::handler& cgh) { + cgh.parallel_for( + sycl::range<1>(num_rays), + [=](sycl::item<1> item, sycl::kernel_handler kh) { + callbacks::ListIntersectionsContext context; + rtcInitRayQueryContext(&context.context); + context.previous_geom_prim_ID_tfar = + previous_geom_prim_ID_tfar; + context.ray_ids = ray_ids; + context.geometry_ids = geometry_ids; + context.primitive_ids = primitive_ids; + context.primitive_uvs = primitive_uvs; + context.t_hit = t_hit; + context.cumsum = cumsum_; + context.track_intersections = track_intersections; + + RTCIntersectArguments args; + rtcInitIntersectArguments(&args); + // args.filter = callbacks::ListIntersectionsFunc; + args.context = &context.context; + + const size_t i = item.get_id(0); + + struct RTCRayHit rh; + const float* r = &rays[i * 6]; + rh.ray.org_x = r[0]; + rh.ray.org_y = r[1]; + rh.ray.org_z = r[2]; + rh.ray.dir_x = r[3]; + rh.ray.dir_y = r[4]; + rh.ray.dir_z = r[5]; + rh.ray.tnear = 0; + rh.ray.tfar = std::numeric_limits::infinity(); + rh.ray.mask = -1; + rh.ray.flags = 0; + rh.ray.id = i; + rh.hit.geomID = RTC_INVALID_GEOMETRY_ID; + rh.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; + + rtcIntersect1(scene, &rh, &args); + }); }); queue_.wait_and_throw(); @@ -714,7 +755,7 @@ struct RaycastingScene::SYCLImpl : public RaycastingScene::Impl { sycl::free(previous_geom_prim_ID_tfar, queue_); sycl::free(cumsum_, queue_); } - + void ComputeClosestPoints(const float* const query_points, const size_t num_query_points, float* closest_points, @@ -880,12 +921,16 @@ struct RaycastingScene::CPUImpl : public RaycastingScene::Impl { const int nthreads) override { CommitScene(); - memset(intersections, 0, sizeof(int) * num_rays); + std::memset(intersections, 0, sizeof(int) * num_rays); - auto previous_geom_prim_ID_tfar = std::unique_ptr>(new callbacks::GeomPrimID[num_rays]); + auto previous_geom_prim_ID_tfar = + std::unique_ptr>( + new callbacks::GeomPrimID[num_rays]); for (size_t i = 0; i < num_rays; ++i) { previous_geom_prim_ID_tfar[i] = {uint32_t(RTC_INVALID_GEOMETRY_ID), - uint32_t(RTC_INVALID_GEOMETRY_ID), 0.f}; + uint32_t(RTC_INVALID_GEOMETRY_ID), + 0.f}; } callbacks::CountIntersectionsContext context; @@ -949,17 +994,21 @@ struct RaycastingScene::CPUImpl : public RaycastingScene::Impl { const int nthreads) override { CommitScene(); - memset(track_intersections, 0, sizeof(uint32_t) * num_rays); - memset(ray_ids, 0, sizeof(uint32_t) * num_intersections); - memset(geometry_ids, 0, sizeof(uint32_t) * num_intersections); - memset(primitive_ids, 0, sizeof(uint32_t) * num_intersections); - memset(primitive_uvs, 0, sizeof(float) * num_intersections * 2); - memset(t_hit, 0, sizeof(float) * num_intersections); - - auto previous_geom_prim_ID_tfar = std::unique_ptr>(new callbacks::GeomPrimID[num_rays]); + std::memset(track_intersections, 0, sizeof(uint32_t) * num_rays); + std::memset(ray_ids, 0, sizeof(uint32_t) * num_intersections); + std::memset(geometry_ids, 0, sizeof(uint32_t) * num_intersections); + std::memset(primitive_ids, 0, sizeof(uint32_t) * num_intersections); + std::memset(primitive_uvs, 0, sizeof(float) * num_intersections * 2); + std::memset(t_hit, 0, sizeof(float) * num_intersections); + + auto previous_geom_prim_ID_tfar = + std::unique_ptr>( + new callbacks::GeomPrimID[num_rays]); for (size_t i = 0; i < num_rays; ++i) { previous_geom_prim_ID_tfar[i] = {uint32_t(RTC_INVALID_GEOMETRY_ID), - uint32_t(RTC_INVALID_GEOMETRY_ID), 0.f}; + uint32_t(RTC_INVALID_GEOMETRY_ID), + 0.f}; } callbacks::ListIntersectionsContext context; @@ -1073,14 +1122,17 @@ struct RaycastingScene::CPUImpl : public RaycastingScene::Impl { RaycastingScene::RaycastingScene(int64_t nthreads #ifdef BUILD_SYCL_MODULE - , const core::Device& device + , + const core::Device& device #endif - ) { +) { #ifdef BUILD_SYCL_MODULE if (device.IsSYCL()) { impl_ = std::make_unique(); - dynamic_cast(impl_.get())->InitializeDevice(); + dynamic_cast(impl_.get()) + ->InitializeDevice(); + impl_->tensor_device_ = device; } else { #endif impl_ = std::make_unique(); @@ -1095,7 +1147,6 @@ RaycastingScene::RaycastingScene(int64_t nthreads } #endif - impl_->tensor_device_ = device; rtcSetDeviceErrorFunction(impl_->device_, ErrorFunction, NULL); impl_->scene_ = rtcNewScene(impl_->device_); @@ -1142,14 +1193,18 @@ uint32_t RaycastingScene::AddTriangles(const core::Tensor& vertex_positions, 3 * sizeof(uint32_t), num_triangles); { -#ifdef BUILD_SYCL_MODULE auto data = vertex_positions.Contiguous(); +#ifdef BUILD_SYCL_MODULE if (impl_->tensor_device_.IsSYCL()) { - dynamic_cast(impl_.get())->queue_.memcpy(vertex_buffer, data.GetDataPtr(), sizeof(float) * 3 * num_vertices).wait(); + dynamic_cast(impl_.get()) + ->queue_ + .memcpy(vertex_buffer, data.GetDataPtr(), + sizeof(float) * 3 * num_vertices) + .wait(); } else { #endif - memcpy(vertex_buffer, data.GetDataPtr(), - sizeof(float) * 3 * num_vertices); + std::memcpy(vertex_buffer, data.GetDataPtr(), + sizeof(float) * 3 * num_vertices); #ifdef BUILD_SYCL_MODULE } #endif @@ -1158,11 +1213,15 @@ uint32_t RaycastingScene::AddTriangles(const core::Tensor& vertex_positions, auto data = triangle_indices.Contiguous(); #ifdef BUILD_SYCL_MODULE if (impl_->tensor_device_.IsSYCL()) { - dynamic_cast(impl_.get())->queue_.memcpy(index_buffer, data.GetDataPtr(), sizeof(uint32_t) * 3 * num_triangles).wait(); + dynamic_cast(impl_.get()) + ->queue_ + .memcpy(index_buffer, data.GetDataPtr(), + sizeof(uint32_t) * 3 * num_triangles) + .wait(); } else { #endif - memcpy(index_buffer, data.GetDataPtr(), - sizeof(uint32_t) * 3 * num_triangles); + std::memcpy(index_buffer, data.GetDataPtr(), + sizeof(uint32_t) * 3 * num_triangles); #ifdef BUILD_SYCL_MODULE } #endif @@ -1201,22 +1260,25 @@ std::unordered_map RaycastingScene::CastRays( std::unordered_map result; result["t_hit"] = core::Tensor(shape, core::Float32, rays.GetDevice()); - result["geometry_ids"] = core::Tensor(shape, core::UInt32, rays.GetDevice()); - result["primitive_ids"] = core::Tensor(shape, core::UInt32, rays.GetDevice()); + result["geometry_ids"] = + core::Tensor(shape, core::UInt32, rays.GetDevice()); + result["primitive_ids"] = + core::Tensor(shape, core::UInt32, rays.GetDevice()); shape.push_back(2); - result["primitive_uvs"] = core::Tensor(shape, core::Float32, rays.GetDevice()); + result["primitive_uvs"] = + core::Tensor(shape, core::Float32, rays.GetDevice()); shape.back() = 3; - result["primitive_normals"] = core::Tensor(shape, core::Float32, rays.GetDevice()); + result["primitive_normals"] = + core::Tensor(shape, core::Float32, rays.GetDevice()); auto data = rays.Contiguous(); impl_->CastRays(data.GetDataPtr(), num_rays, - result["t_hit"].GetDataPtr(), - result["geometry_ids"].GetDataPtr(), - result["primitive_ids"].GetDataPtr(), - result["primitive_uvs"].GetDataPtr(), - result["primitive_normals"].GetDataPtr(), - nthreads, - false); + result["t_hit"].GetDataPtr(), + result["geometry_ids"].GetDataPtr(), + result["primitive_ids"].GetDataPtr(), + result["primitive_uvs"].GetDataPtr(), + result["primitive_normals"].GetDataPtr(), nthreads, + false); return result; } @@ -1251,7 +1313,8 @@ core::Tensor RaycastingScene::CountIntersections(const core::Tensor& rays, // results. size_t num_rays = shape.NumElements(); - core::Tensor intersections(shape, core::Dtype::FromType(), impl_->tensor_device_); + core::Tensor intersections(shape, core::Dtype::FromType(), + impl_->tensor_device_); auto data = rays.Contiguous(); @@ -1272,8 +1335,10 @@ RaycastingScene::ListIntersections(const core::Tensor& rays, size_t num_rays = shape.NumElements(); // determine total number of intersections - core::Tensor intersections(shape, core::Dtype::FromType(), impl_->tensor_device_); - core::Tensor track_intersections(shape, core::Dtype::FromType(), impl_->tensor_device_); + core::Tensor intersections(shape, core::Dtype::FromType(), + impl_->tensor_device_); + core::Tensor track_intersections(shape, core::Dtype::FromType(), + impl_->tensor_device_); auto data = rays.Contiguous(); impl_->CountIntersections(data.GetDataPtr(), num_rays, intersections.GetDataPtr(), nthreads); @@ -1286,7 +1351,8 @@ RaycastingScene::ListIntersections(const core::Tensor& rays, } // prepare ray allocations (cumsum) - core::Tensor cumsum_tensor = core::Tensor::Zeros(shape, core::Dtype::FromType(), impl_->tensor_device_); + core::Tensor cumsum_tensor = core::Tensor::Zeros( + shape, core::Dtype::FromType(), impl_->tensor_device_); int* cumsum_ptr = cumsum_tensor.GetDataPtr(); cumsum_ptr[0] = 0; @@ -1300,17 +1366,21 @@ RaycastingScene::ListIntersections(const core::Tensor& rays, shape.push_back(num_rays + 1); result["ray_splits"] = core::Tensor(shape, core::UInt32); uint32_t* ptr = result["ray_splits"].GetDataPtr(); - for (int i = 0; i < num_rays; ++i) { + for (size_t i = 0; i < num_rays; ++i) { ptr[i] = cumsum_ptr[i]; } ptr[num_rays] = num_intersections; shape[0] = num_intersections; - result["ray_ids"] = core::Tensor(shape, core::UInt32, impl_->tensor_device_); - result["geometry_ids"] = core::Tensor(shape, core::UInt32, impl_->tensor_device_); - result["primitive_ids"] = core::Tensor(shape, core::UInt32, impl_->tensor_device_); + result["ray_ids"] = + core::Tensor(shape, core::UInt32, impl_->tensor_device_); + result["geometry_ids"] = + core::Tensor(shape, core::UInt32, impl_->tensor_device_); + result["primitive_ids"] = + core::Tensor(shape, core::UInt32, impl_->tensor_device_); result["t_hit"] = core::Tensor(shape, core::Float32, impl_->tensor_device_); shape.push_back(2); - result["primitive_uvs"] = core::Tensor(shape, core::Float32, impl_->tensor_device_); + result["primitive_uvs"] = + core::Tensor(shape, core::Float32, impl_->tensor_device_); impl_->ListIntersections(data.GetDataPtr(), num_rays, num_intersections, cumsum_ptr, @@ -1320,7 +1390,7 @@ RaycastingScene::ListIntersections(const core::Tensor& rays, result["primitive_ids"].GetDataPtr(), result["primitive_uvs"].GetDataPtr(), result["t_hit"].GetDataPtr(), nthreads); - + return result; } diff --git a/cpp/open3d/t/geometry/RaycastingScene.h b/cpp/open3d/t/geometry/RaycastingScene.h index 20237344279..3d636ba012c 100644 --- a/cpp/open3d/t/geometry/RaycastingScene.h +++ b/cpp/open3d/t/geometry/RaycastingScene.h @@ -32,7 +32,8 @@ class RaycastingScene { /// \brief Default Constructor. RaycastingScene(int64_t nthreads = 0 #ifdef BUILD_SYCL_MODULE - , const core::Device& device = core::Device("CPU:0") + , + const core::Device &device = core::Device("CPU:0") #endif ); diff --git a/cpp/pybind/CMakeLists.txt b/cpp/pybind/CMakeLists.txt index c79bbd96719..6efae9a17fd 100644 --- a/cpp/pybind/CMakeLists.txt +++ b/cpp/pybind/CMakeLists.txt @@ -106,7 +106,7 @@ endif() # Include additional libraries that may be absent from the user system # eg: libc++.so, libc++abi.so (needed by filament) for Linux. # libc++.so is a linker script including libc++.so.1 and libc++abi.so, so append 1 to libc++.so -set(PYTHON_EXTRA_LIBRARIES $) +set(PYTHON_EXTRA_LIBRARIES $) if (BUILD_GUI AND CMAKE_SYSTEM_NAME STREQUAL "Linux") list(APPEND PYTHON_EXTRA_LIBRARIES ${CPP_LIBRARY}.1 ${CPPABI_LIBRARY}) endif() diff --git a/cpp/pybind/core/sycl_utils.cpp b/cpp/pybind/core/sycl_utils.cpp index ba90cb6fc51..7db39b7acbc 100644 --- a/cpp/pybind/core/sycl_utils.cpp +++ b/cpp/pybind/core/sycl_utils.cpp @@ -12,7 +12,7 @@ namespace open3d { namespace core { -void pybind_sycl_utils_definitions(py::module& m) { +void pybind_sycl_utils_definitions(py::module& m) { m.def("sycl_demo", &sy::SYCLDemo); py::module m_sycl = m.def_submodule("sycl"); diff --git a/cpp/pybind/t/geometry/raycasting_scene.cpp b/cpp/pybind/t/geometry/raycasting_scene.cpp index a9156b212e5..732abf6cdd2 100644 --- a/cpp/pybind/t/geometry/raycasting_scene.cpp +++ b/cpp/pybind/t/geometry/raycasting_scene.cpp @@ -59,7 +59,8 @@ void pybind_raycasting_scene_definitions(py::module& m) { static_cast>(m.attr("RaycastingScene")); // Constructors. #ifdef BUILD_SYCL_MODULE - raycasting_scene.def(py::init(), "nthreads"_a = 0, "device"_a = core::Device("CPU:0"), R"doc( + raycasting_scene.def(py::init(), "nthreads"_a = 0, + "device"_a = core::Device("CPU:0"), R"doc( Create a RaycastingScene. Args: diff --git a/docker/Dockerfile.ci b/docker/Dockerfile.ci index 6629b129fc1..a54cbc0320e 100755 --- a/docker/Dockerfile.ci +++ b/docker/Dockerfile.ci @@ -3,7 +3,7 @@ ARG BASE_IMAGE FROM ${BASE_IMAGE} # For bash-specific commands -SHELL ["/bin/bash", "-c"] +SHELL ["/bin/bash", "-c", "-o", "pipefail"] # Required build args, should be specified in docker_build.sh ARG DEVELOPER_BUILD @@ -71,7 +71,6 @@ RUN if [ "${BUILD_SYCL_MODULE}" = "ON" ]; then \ RUN apt-get update && apt-get install -y \ git \ wget \ - curl \ build-essential \ pkg-config \ zlib1g \ @@ -89,7 +88,7 @@ RUN apt-get update && apt-get install -y \ liblzma-dev \ && rm -rf /var/lib/apt/lists/* -# pyenv or Intel Python +# pyenv # The pyenv python paths are used during docker run, in this way docker run # does not need to activate the environment again. # The soft link from the python patch level version to the python mino version @@ -97,16 +96,13 @@ RUN apt-get update && apt-get install -y \ # which patch level pyenv will install (latest). ENV PYENV_ROOT=/root/.pyenv ENV PATH="$PYENV_ROOT/shims:$PYENV_ROOT/bin:$PYENV_ROOT/versions/$PYTHON_VERSION/bin:$PATH" -ENV PATH="/opt/intel/oneapi/intelpython/latest/bin:${PATH}" -RUN if [ "${BUILD_SYCL_MODULE}" = "OFF" ]; then \ - curl https://pyenv.run | bash \ - && pyenv update \ - && pyenv install $PYTHON_VERSION \ - && pyenv global $PYTHON_VERSION \ - && pyenv rehash \ - && ln -s $PYENV_ROOT/versions/${PYTHON_VERSION}* $PYENV_ROOT/versions/${PYTHON_VERSION}; \ - fi -RUN python --version && pip --version +RUN wget -q https://pyenv.run -O- | bash \ + && pyenv update \ + && pyenv install $PYTHON_VERSION \ + && pyenv global $PYTHON_VERSION \ + && pyenv rehash \ + && ln -s $PYENV_ROOT/versions/${PYTHON_VERSION}* $PYENV_ROOT/versions/${PYTHON_VERSION} \ + && python --version && pip --version SHELL ["/bin/bash", "-o", "pipefail", "-c"] @@ -174,7 +170,7 @@ RUN source util/ci_utils.sh \ # Open3D Jupyter dependencies RUN mkdir -p /etc/apt/keyrings \ - && curl -fsSL https://deb.nodesource.com/gpgkey/nodesource-repo.gpg.key \ + && wget -qO- https://deb.nodesource.com/gpgkey/nodesource-repo.gpg.key \ | gpg --dearmor -o /etc/apt/keyrings/nodesource.gpg \ && echo "deb [signed-by=/etc/apt/keyrings/nodesource.gpg] https://deb.nodesource.com/node_16.x nodistro main" \ | tee /etc/apt/sources.list.d/nodesource.list \ diff --git a/python/test/t/geometry/test_raycasting_scene.py b/python/test/t/geometry/test_raycasting_scene.py index 4dceb41e7eb..13e3e8cce51 100755 --- a/python/test/t/geometry/test_raycasting_scene.py +++ b/python/test/t/geometry/test_raycasting_scene.py @@ -19,14 +19,18 @@ @pytest.mark.parametrize("device", list_devices()) def test_cast_rays(device): vertices = o3d.core.Tensor([[0, 0, 0], [1, 0, 0], [1, 1, 0]], - dtype=o3d.core.float32, device=device) - triangles = o3d.core.Tensor([[0, 1, 2]], dtype=o3d.core.uint32, device=device) + dtype=o3d.core.float32, + device=device) + triangles = o3d.core.Tensor([[0, 1, 2]], + dtype=o3d.core.uint32, + device=device) scene = o3d.t.geometry.RaycastingScene(device=device) geom_id = scene.add_triangles(vertices, triangles) rays = o3d.core.Tensor([[0.2, 0.1, 1, 0, 0, -1], [10, 10, 10, 1, 0, 0]], - dtype=o3d.core.float32, device=device) + dtype=o3d.core.float32, + device=device) ans = scene.cast_rays(rays) # first ray hits the triangle @@ -34,7 +38,8 @@ def test_cast_rays(device): assert np.isclose(ans['t_hit'][0].item(), 1.0) # second ray misses - assert o3d.t.geometry.RaycastingScene.INVALID_ID == ans['geometry_ids'][1].cpu() + assert o3d.t.geometry.RaycastingScene.INVALID_ID == ans['geometry_ids'][ + 1].cpu() assert np.isinf(ans['t_hit'][1].item()) @@ -43,8 +48,11 @@ def test_cast_rays(device): @pytest.mark.parametrize("device", list_devices()) def test_cast_lots_of_rays(device): vertices = o3d.core.Tensor([[0, 0, 0], [1, 0, 0], [1, 1, 0]], - dtype=o3d.core.float32, device=device) - triangles = o3d.core.Tensor([[0, 1, 2]], dtype=o3d.core.uint32, device=device) + dtype=o3d.core.float32, + device=device) + triangles = o3d.core.Tensor([[0, 1, 2]], + dtype=o3d.core.uint32, + device=device) scene = o3d.t.geometry.RaycastingScene(device=device) scene.add_triangles(vertices, triangles) @@ -60,14 +68,18 @@ def test_cast_lots_of_rays(device): @pytest.mark.parametrize("device", list_devices()) def test_test_occlusions(device): vertices = o3d.core.Tensor([[0, 0, 0], [1, 0, 0], [1, 1, 0]], - dtype=o3d.core.float32, device=device) - triangles = o3d.core.Tensor([[0, 1, 2]], dtype=o3d.core.uint32, device=device) + dtype=o3d.core.float32, + device=device) + triangles = o3d.core.Tensor([[0, 1, 2]], + dtype=o3d.core.uint32, + device=device) scene = o3d.t.geometry.RaycastingScene(device=device) scene.add_triangles(vertices, triangles) rays = o3d.core.Tensor([[0.2, 0.1, 1, 0, 0, -1], [10, 10, 10, 1, 0, 0]], - dtype=o3d.core.float32, device=device) + dtype=o3d.core.float32, + device=device) ans = scene.test_occlusions(rays).cpu() # first ray is occluded by the triangle @@ -90,8 +102,11 @@ def test_test_occlusions(device): @pytest.mark.parametrize("device", list_devices()) def test_test_lots_of_occlusions(device): vertices = o3d.core.Tensor([[0, 0, 0], [1, 0, 0], [1, 1, 0]], - dtype=o3d.core.float32, device=device) - triangles = o3d.core.Tensor([[0, 1, 2]], dtype=o3d.core.uint32, device=device) + dtype=o3d.core.float32, + device=device) + triangles = o3d.core.Tensor([[0, 1, 2]], + dtype=o3d.core.uint32, + device=device) scene = o3d.t.geometry.RaycastingScene(device=device) scene.add_triangles(vertices, triangles) @@ -118,7 +133,8 @@ def test_add_triangle_mesh(device): rays = o3d.core.Tensor([[0.5, 0.5, -1, 0, 0, 1], [0.5, 0.5, 0.5, 0, 0, 1], [10, 10, 10, 1, 0, 0]], - dtype=o3d.core.float32, device=device) + dtype=o3d.core.float32, + device=device) ans = scene.count_intersections(rays) np.testing.assert_equal(ans.cpu().numpy(), [2, 1, 0]) @@ -139,7 +155,8 @@ def test_count_intersections(device): rays = o3d.core.Tensor([[0.5, 0.5, -1, 0, 0, 1], [0.5, 0.5, 0.5, 0, 0, 1], [10, 10, 10, 1, 0, 0]], - dtype=o3d.core.float32, device=device) + dtype=o3d.core.float32, + device=device) ans = scene.count_intersections(rays) np.testing.assert_equal(ans.cpu().numpy(), [2, 1, 0]) @@ -182,7 +199,8 @@ def test_list_intersections(device): rays = o3d.core.Tensor([[0.5, 0.5, -1, 0, 0, 1], [0.5, 0.5, 0.5, 0, 0, 1], [10, 10, 10, 1, 0, 0]], - dtype=o3d.core.float32, device=device) + dtype=o3d.core.float32, + device=device) print("PYTHON TEST 1", device) ans = scene.list_intersections(rays) print("PYTHON TEST 2")