From d9ea99613fdf0e61d4ec2d174d098c2ca1963f12 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Sun, 3 Nov 2024 11:11:59 -0500 Subject: [PATCH 01/11] Check bvh pointers and rewrite for loop as while --- src/geocel/vg/detail/BVHNavigator.hh | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/src/geocel/vg/detail/BVHNavigator.hh b/src/geocel/vg/detail/BVHNavigator.hh index f2916b4b85..75e26c96a8 100644 --- a/src/geocel/vg/detail/BVHNavigator.hh +++ b/src/geocel/vg/detail/BVHNavigator.hh @@ -46,6 +46,8 @@ class BVHNavigator bool top, vecgeom::VPlacedVolume const* exclude = nullptr) { + CELER_EXPECT(vecgeom::BVHManager::GetBVH(0) != nullptr); + if (top) { assert(vol != nullptr); @@ -58,15 +60,22 @@ class BVHNavigator Vector3D currentpoint(point); Vector3D daughterlocalpoint; - for (auto v = vol; v->GetDaughters().size() > 0;) + while (vol->GetDaughters().size() > 0) { - auto bvh = vecgeom::BVHManager::GetBVH(v->GetLogicalVolume()->id()); + auto* bvh + = vecgeom::BVHManager::GetBVH(vol->GetLogicalVolume()->id()); + CELER_ASSERT(bvh); - if (!bvh->LevelLocate(exclude, currentpoint, v, daughterlocalpoint)) + // Note: vol is updated by this call + if (!bvh->LevelLocate( + exclude, currentpoint, vol, daughterlocalpoint)) + { + // Not inside any daughter break; + } currentpoint = daughterlocalpoint; - path.Push(v); + path.Push(vol); // Only exclude the placed volume once since we could enter it // again via a different volume history. exclude = nullptr; From 9c99c70b9d448898b0673613db264345ef2a7b09 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Sun, 3 Nov 2024 12:05:11 -0500 Subject: [PATCH 02/11] Move surface detail to cuda detail --- src/geocel/CMakeLists.txt | 7 +---- src/geocel/vg/VecgeomParams.cc | 17 +++-------- .../VecgeomSetup.cu} | 18 ++++++++--- .../VecgeomSetup.hh} | 30 +++++++++++++++++-- 4 files changed, 47 insertions(+), 25 deletions(-) rename src/geocel/vg/{VecgeomParams.surface.cu => detail/VecgeomSetup.cu} (79%) rename src/geocel/vg/{VecgeomParams.surface.hh => detail/VecgeomSetup.hh} (53%) diff --git a/src/geocel/CMakeLists.txt b/src/geocel/CMakeLists.txt index 4b05c41e9c..d7fb8241ba 100644 --- a/src/geocel/CMakeLists.txt +++ b/src/geocel/CMakeLists.txt @@ -80,12 +80,6 @@ if(CELERITAS_USE_VecGeom) if(VecGeom_GDML_FOUND) list(APPEND PRIVATE_DEPS VecGeom::vgdml) endif() - if(VecGeom_CUDA_FOUND AND VecGeom_SURF_FOUND) - # Special routines needed for surface - list(APPEND SOURCES - vg/VecgeomParams.surface.cu - ) - endif() if(CELERITAS_USE_CUDA) # This needs to be public because its might be needed # to resolve the symbols generate by the `nvcc -dlink` of @@ -93,6 +87,7 @@ if(CELERITAS_USE_VecGeom) list(APPEND PUBLIC_DEPS VecGeom::vecgeom) list(APPEND SOURCES vg/RaytraceImager.cu + vg/detail/VecgeomSetup.cu ) elseif(VecGeom_CUDA_FOUND) # VecGeom is built with CUDA but Celeritas is not diff --git a/src/geocel/vg/VecgeomParams.cc b/src/geocel/vg/VecgeomParams.cc index 964e417523..a607d4ccc3 100644 --- a/src/geocel/vg/VecgeomParams.cc +++ b/src/geocel/vg/VecgeomParams.cc @@ -49,10 +49,7 @@ #include "VecgeomData.hh" // IWYU pragma: associated #include "detail/VecgeomCompatibility.hh" - -#ifdef VECGEOM_USE_SURF -# include "VecgeomParams.surface.hh" -#endif +#include "detail/VecgeomSetup.hh" static_assert(std::is_same_v, "Celeritas and VecGeom real types do not match"); @@ -82,12 +79,6 @@ namespace } while (0) #endif -#if defined(VECGEOM_ENABLE_CUDA) && defined(VECGEOM_USE_SURF) -# define VG_CUDASURF_CALL(CODE) CODE -#else -# define VG_CUDASURF_CALL(CODE) CELER_UNREACHABLE -#endif - //---------------------------------------------------------------------------// // HELPER FUNCTIONS //---------------------------------------------------------------------------// @@ -197,7 +188,7 @@ VecgeomParams::~VecgeomParams() if (VecgeomParams::use_surface_tracking()) { CELER_LOG(debug) << "Clearing VecGeom surface GPU data"; - VG_CUDASURF_CALL(teardown_surface_tracking_device()); + VG_SURF_CALL(detail::teardown_surface_tracking_device()); } else { @@ -367,8 +358,8 @@ void VecgeomParams::build_surface_tracking() ScopedTimeAndRedirect time_and_output_( "BrepCudaManager::TransferSurfData"); - VG_CUDASURF_CALL( - setup_surface_tracking_device(brep_helper.GetSurfData())); + VG_SURF_CALL( + detail::setup_surface_tracking_device(brep_helper.GetSurfData())); CELER_DEVICE_CHECK_ERROR(); } } diff --git a/src/geocel/vg/VecgeomParams.surface.cu b/src/geocel/vg/detail/VecgeomSetup.cu similarity index 79% rename from src/geocel/vg/VecgeomParams.surface.cu rename to src/geocel/vg/detail/VecgeomSetup.cu index ba85c3131f..36c00660c4 100644 --- a/src/geocel/vg/VecgeomParams.surface.cu +++ b/src/geocel/vg/detail/VecgeomSetup.cu @@ -3,31 +3,41 @@ // See the top-level COPYRIGHT file for details. // SPDX-License-Identifier: (Apache-2.0 OR MIT) //---------------------------------------------------------------------------// -//! \file geocel/vg/VecgeomParams.surface.cu +//! \file geocel/vg/detail/VecgeomSetup.cu //---------------------------------------------------------------------------// -#include "VecgeomParams.hh" +#include "VecgeomSetup.cuda.hh" -#include +#ifdef VECGEOM_USE_SURF +# include +#endif #include "corecel/Assert.hh" +#ifdef VECGEOM_USE_SURF using BrepCudaManager = vgbrep::BrepCudaManager; using SurfData = vgbrep::SurfData; +#endif namespace celeritas { +namespace detail +{ +//---------------------------------------------------------------------------// +// VECGEOM SURFACE //---------------------------------------------------------------------------// +#ifdef VECGEOM_USE_SURF void setup_surface_tracking_device(SurfData const& surf_data) { BrepCudaManager::Instance().TransferSurfData(surf_data); CELER_DEVICE_CALL_PREFIX(DeviceSynchronize()); } -//---------------------------------------------------------------------------// void teardown_surface_tracking_device() { BrepCudaManager::Instance().Cleanup(); } +#endif //---------------------------------------------------------------------------// +} // namespace detail } // namespace celeritas diff --git a/src/geocel/vg/VecgeomParams.surface.hh b/src/geocel/vg/detail/VecgeomSetup.hh similarity index 53% rename from src/geocel/vg/VecgeomParams.surface.hh rename to src/geocel/vg/detail/VecgeomSetup.hh index b80a9f811d..61dd3fead3 100644 --- a/src/geocel/vg/VecgeomParams.surface.hh +++ b/src/geocel/vg/detail/VecgeomSetup.hh @@ -3,20 +3,46 @@ // See the top-level COPYRIGHT file for details. // SPDX-License-Identifier: (Apache-2.0 OR MIT) //---------------------------------------------------------------------------// -//! \file geocel/vg/VecgeomParams.surface.hh +//! \file geocel/vg/detail/VecgeomSetup.hh //---------------------------------------------------------------------------// #pragma once -#include +#include +#ifdef VECGEOM_USE_SURF +# include +#endif namespace celeritas { +namespace detail +{ +#ifdef VECGEOM_USE_SURF //---------------------------------------------------------------------------// // Set up surface tracking void setup_surface_tracking_device(vgbrep::SurfData const&); // Tear down surface tracking void teardown_surface_tracking_device(); +#endif //---------------------------------------------------------------------------// +// INLINE DEFINITIONS +//---------------------------------------------------------------------------// +#ifndef VECGEOM_ENABLE_CUDA +# ifdef VECGEOM_USE_SURF +// Set up surface tracking +inline void +setup_surface_tracking_device(vgbrep::SurfData const&) +{ + CELER_ASSERT_UNREACHABLE(); +} + +inline void teardown_surface_tracking_device() +{ + CELER_ASSERT_UNREACHABLE(); +} +# endif +#endif +//---------------------------------------------------------------------------// +} // namespace detail } // namespace celeritas From fc0cf2583201b2792a07753ac157b103bb44d3c8 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Sun, 3 Nov 2024 13:06:02 -0500 Subject: [PATCH 03/11] Add functions to get and check BVH pointers --- src/corecel/sys/KernelLauncher.device.hh | 4 +- src/geocel/vg/VecgeomParams.cc | 13 ++++++- src/geocel/vg/detail/VecgeomSetup.cu | 47 ++++++++++++++++++++++++ src/geocel/vg/detail/VecgeomSetup.hh | 29 ++++++++++++++- 4 files changed, 88 insertions(+), 5 deletions(-) diff --git a/src/corecel/sys/KernelLauncher.device.hh b/src/corecel/sys/KernelLauncher.device.hh index 8378eae791..cba451e24d 100644 --- a/src/corecel/sys/KernelLauncher.device.hh +++ b/src/corecel/sys/KernelLauncher.device.hh @@ -106,7 +106,9 @@ void KernelLauncher::operator()(Range threads, if (!threads.empty()) { using StreamT = CELER_DEVICE_PREFIX(Stream_t); - StreamT stream = celeritas::device().stream(stream_id).get(); + StreamT stream = stream_id + ? celeritas::device().stream(stream_id).get() + : nullptr; auto config = calc_launch_params_(threads.size()); detail::launch_action_impl <<>>( diff --git a/src/geocel/vg/VecgeomParams.cc b/src/geocel/vg/VecgeomParams.cc index a607d4ccc3..a7f399416e 100644 --- a/src/geocel/vg/VecgeomParams.cc +++ b/src/geocel/vg/VecgeomParams.cc @@ -448,9 +448,18 @@ void VecgeomParams::build_volume_tracking() vecgeom::cxx::BVHManager::DeviceInit(); CELER_DEVICE_CHECK_ERROR(); } -#else - CELER_NOT_CONFIGURED("CUDA"); #endif + + // Check BVH pointers + auto ptrs = detail::bvh_pointers_device(); + CELER_VALIDATE(ptrs.symbol, + << "VecGeom device BVH is not correctly initialized: " + "runtime symbol is null"); + CELER_VALIDATE( + ptrs.kernel == ptrs.symbol, + << "inconsistenct VecGeom BVH device pointers: " + << static_cast(ptrs.kernel) << " from device kernel, " + << static_cast(ptrs.symbol) << " from runtime symbol"); } } diff --git a/src/geocel/vg/detail/VecgeomSetup.cu b/src/geocel/vg/detail/VecgeomSetup.cu index 36c00660c4..9483898c27 100644 --- a/src/geocel/vg/detail/VecgeomSetup.cu +++ b/src/geocel/vg/detail/VecgeomSetup.cu @@ -5,6 +5,8 @@ //---------------------------------------------------------------------------// //! \file geocel/vg/detail/VecgeomSetup.cu //---------------------------------------------------------------------------// +#include + #include "VecgeomSetup.cuda.hh" #ifdef VECGEOM_USE_SURF @@ -12,6 +14,8 @@ #endif #include "corecel/Assert.hh" +#include "corecel/Macros.hh" +#include "corecel/sys/KernelLauncher.device.hh" #ifdef VECGEOM_USE_SURF using BrepCudaManager = vgbrep::BrepCudaManager; @@ -22,6 +26,49 @@ namespace celeritas { namespace detail { +namespace +{ +//---------------------------------------------------------------------------// +//! Access +struct BvhGetter +{ + VecGeom::cuda::BVH const** dest{nullptr}; + + CELER_FUNCTION operator()(ThreadId tid) + { + CELER_EXPECT(tid == ThreadId{0}); + *dest = vecgeom::cuda::BVHManager::GetBVH(0); + } +}; +} // namespace + +//---------------------------------------------------------------------------// +/*! + * Get pointers to the device BVH after setup, for consistency checking. + */ +CudaPointers bvh_pointers_device() +{ + CudaPointers result; + + // Copy from kernel using 1-thread launch + { + DeviceVector bvh_ptr{1, StreamId{}}; + BvhGetter execute_thread{params, state, seeds}; + static KernelLauncher const launch_kernel( + "vecgeom-get-bvhptr"); + launch_kernel(1u, StreamId{}, execute_thread); + CELER_CUDA_CALL(cudaDeviceSynchronize()); + bvh_ptr.copy_to_host({&result.kernel, 1}); + } + + // Copy from symbol using runtime API + CELER_CUDA_CALL(cudaMemcpyFromSymbol( + &result.symbol, vecgeom::cuda::dBVH, 1, 0, cudaMemcpyDeviceToHost)); + + // Return + return result; +} + //---------------------------------------------------------------------------// // VECGEOM SURFACE //---------------------------------------------------------------------------// diff --git a/src/geocel/vg/detail/VecgeomSetup.hh b/src/geocel/vg/detail/VecgeomSetup.hh index 61dd3fead3..2478ce2041 100644 --- a/src/geocel/vg/detail/VecgeomSetup.hh +++ b/src/geocel/vg/detail/VecgeomSetup.hh @@ -8,6 +8,8 @@ #pragma once #include + +#include "corecel/Assert.hh" #ifdef VECGEOM_USE_SURF # include #endif @@ -16,8 +18,27 @@ namespace celeritas { namespace detail { -#ifdef VECGEOM_USE_SURF //---------------------------------------------------------------------------// +/*! + * Pointers to device data, obtained from a kernel launch or from runtime. + * + * The \c kernel data is copied from inside a kernel to global heap memory, and + * thence to this result. The \c symbol data is copied via \c + * cudaMemcpyFromSymbol . + */ +template +struct CudaPointers +{ + T* kernel{nullptr}; + T* symbol{nullptr}; +}; + +//---------------------------------------------------------------------------// +// Get pointers to the device BVH after setup, for consistency checking +CudaPointers bvh_pointers_device(); + +//---------------------------------------------------------------------------// +#ifdef VECGEOM_USE_SURF // Set up surface tracking void setup_surface_tracking_device(vgbrep::SurfData const&); @@ -29,8 +50,12 @@ void teardown_surface_tracking_device(); // INLINE DEFINITIONS //---------------------------------------------------------------------------// #ifndef VECGEOM_ENABLE_CUDA +inline CudaPointers bvh_pointers_device() +{ + CELER_ASSERT_UNREACHABLE(); +} + # ifdef VECGEOM_USE_SURF -// Set up surface tracking inline void setup_surface_tracking_device(vgbrep::SurfData const&) { From 3345ff295ed9e1171289524a19351d822e291663 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Mon, 4 Nov 2024 09:57:23 -0500 Subject: [PATCH 04/11] Fix missing libdl for static build on cuda 11.5 --- src/corecel/CMakeLists.txt | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/src/corecel/CMakeLists.txt b/src/corecel/CMakeLists.txt index 0e5a873543..fd791df8d8 100644 --- a/src/corecel/CMakeLists.txt +++ b/src/corecel/CMakeLists.txt @@ -188,11 +188,13 @@ if(CELERITAS_USE_CUDA OR CELERITAS_USE_HIP) ) endif() if(CELERITAS_USE_CUDA) - # we currently depend on CUDA install including the nvtx header + # Note: we currently assume the CUDA install includes the nvtx header list(APPEND SOURCES sys/ScopedLimitSaver.cuda.cc sys/ScopedProfiling.cuda.cc ) + # NVTX requires `-ldl` + list(APPEND PRIVATE_DEPS ${CMAKE_DL_LIBS}) endif() if(CELERITAS_USE_HIP) From b904594b5445cd939f6520bdb6ac0680ec0b6a8e Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Mon, 4 Nov 2024 10:24:45 -0500 Subject: [PATCH 05/11] Fix missing quote and deprecation notes in .h files --- src/celeritas_cmake_strings.h | 1 + src/celeritas_sys_config.h | 1 + src/corecel/device_runtime_api.h | 5 +++-- 3 files changed, 5 insertions(+), 2 deletions(-) diff --git a/src/celeritas_cmake_strings.h b/src/celeritas_cmake_strings.h index 0c59ce7a86..970e4999b6 100644 --- a/src/celeritas_cmake_strings.h +++ b/src/celeritas_cmake_strings.h @@ -8,6 +8,7 @@ * \brief Detailed CMake configuration information. * \deprecated This file should be replaced by "corecel/Config.hh". */ +// DEPRECATED: remove in Celeritas v1.0 //---------------------------------------------------------------------------// #ifndef celeritas_cmake_strings_h #define celeritas_cmake_strings_h diff --git a/src/celeritas_sys_config.h b/src/celeritas_sys_config.h index b16a4cc412..81f8dbc9d6 100644 --- a/src/celeritas_sys_config.h +++ b/src/celeritas_sys_config.h @@ -8,6 +8,7 @@ * \brief System-specific configuration options for Celeritas. * \deprecated This file should be replaced by "corecel/Config.hh". */ +// DEPRECATED: remove in Celeritas v1.0 //---------------------------------------------------------------------------// #ifndef celeritas_sys_config_h #define celeritas_sys_config_h diff --git a/src/corecel/device_runtime_api.h b/src/corecel/device_runtime_api.h index 9f48e003d4..3e98ce2a32 100644 --- a/src/corecel/device_runtime_api.h +++ b/src/corecel/device_runtime_api.h @@ -8,6 +8,7 @@ * \brief Include CUDA or HIP runtime APIs for compiling with host/cc compiler. * \deprecated This file should be replaced by "corecel/DeviceRuntimeApi.hh". */ +// DEPRECATED: remove in Celeritas v1.0 //---------------------------------------------------------------------------// #ifndef CELERITAS_DEVICE_RUNTIME_API_H #define CELERITAS_DEVICE_RUNTIME_API_H @@ -21,6 +22,6 @@ "corecel/device_runtime_api.h is deprecated and should be replaced by \"corecel/DeviceRuntimeApi.hh\"" #endif -#include "corecel/DeviceRuntimeApi.hh +#include "corecel/DeviceRuntimeApi.hh" -#endif /* CELERITAS_DEVICE_RUNTIME_API_H */ \ No newline at end of file +#endif /* CELERITAS_DEVICE_RUNTIME_API_H */ From d102da16868c326b09caff06e13fa2b613494dd6 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Mon, 4 Nov 2024 10:25:15 -0500 Subject: [PATCH 06/11] Add missing space to message --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 4a718d4636..40ffb83ba9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -408,7 +408,7 @@ if(CELERITAS_USE_VecGeom) AND NOT (VecGeom_CUDA_ARCHITECTURES STREQUAL CMAKE_CUDA_ARCHITECTURES)) message(WARNING "CUDA architecture types between " "VecGeom (${VecGeom_CUDA_ARCHITECTURES}) and " - "Celeritas (${CMAKE_CUDA_ARCHITECTURES}) should probably match:" + "Celeritas (${CMAKE_CUDA_ARCHITECTURES}) should probably match: " "runtime errors may result" ) endif() From 436649bd0b67babedaaf59d8031f3516c5e4577a Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Mon, 4 Nov 2024 10:25:38 -0500 Subject: [PATCH 07/11] fixup! Add functions to get and check BVH pointers --- src/corecel/sys/KernelLauncher.device.hh | 1 - src/geocel/vg/VecgeomParams.cc | 50 ++++++++++++++++++++---- src/geocel/vg/detail/VecgeomSetup.cu | 18 +++++---- src/geocel/vg/detail/VecgeomSetup.hh | 5 ++- 4 files changed, 55 insertions(+), 19 deletions(-) diff --git a/src/corecel/sys/KernelLauncher.device.hh b/src/corecel/sys/KernelLauncher.device.hh index cba451e24d..5d9d5f53e5 100644 --- a/src/corecel/sys/KernelLauncher.device.hh +++ b/src/corecel/sys/KernelLauncher.device.hh @@ -132,7 +132,6 @@ void KernelLauncher::operator()(size_type num_threads, StreamId stream_id, F const& execute_thread) const { - CELER_EXPECT(stream_id); (*this)(range(ThreadId{num_threads}), stream_id, execute_thread); } diff --git a/src/geocel/vg/VecgeomParams.cc b/src/geocel/vg/VecgeomParams.cc index a7f399416e..49b5cf68cf 100644 --- a/src/geocel/vg/VecgeomParams.cc +++ b/src/geocel/vg/VecgeomParams.cc @@ -452,14 +452,48 @@ void VecgeomParams::build_volume_tracking() // Check BVH pointers auto ptrs = detail::bvh_pointers_device(); - CELER_VALIDATE(ptrs.symbol, - << "VecGeom device BVH is not correctly initialized: " - "runtime symbol is null"); - CELER_VALIDATE( - ptrs.kernel == ptrs.symbol, - << "inconsistenct VecGeom BVH device pointers: " - << static_cast(ptrs.kernel) << " from device kernel, " - << static_cast(ptrs.symbol) << " from runtime symbol"); + + vecgeom::cuda::BVH const* bvh_symbol_ptr{nullptr}; +#ifdef VECGEOM_BVHMANAGER_DEVICE + bvh_symbol_ptr = BVHManager::GetDeviceBVH(); +#endif + if (ptrs.kernel == nullptr || ptrs.kernel != ptrs.symbol + || (bvh_symbol_ptr && (ptrs.kernel != bvh_symbol_ptr))) + { + // It's very bad if the kernel-viewed BVH pointer is null or + // inconsistent with the VecGeom-provided BVH pointer (only + // available in very recent VecGeom). It's bad (but not really + // necessary?) if cudaMemcpyFromSymbol fails when accessed from + // Celeritas + auto msg = world_logger()( + CELER_CODE_PROVENANCE, + (ptrs.kernel == nullptr + || (bvh_symbol_ptr && (ptrs.kernel != bvh_symbol_ptr))) + ? LogLevel::error + : LogLevel::debug); + auto msg_pointer = [&msg](auto* p) { + if (p) + { + msg << p; + } + else + { + msg << "nullptr"; + } + }; + msg << "VecGeom CUDA may not be correctly linked or initialized (" + "BVH device pointers are null or inconsistent: "; + msg_pointer(ptrs.kernel); + msg << " from Celeritas device kernel, "; + msg_pointer(ptrs.symbol); + msg << " from Celeritas runtime symbol, "; +#ifdef VECGEOM_BVHMANAGER_DEVICE + msg_pointer(bvh_symbol_ptr); +#else + msg << "unavailable"; +#endif + msg << " from VecGeom runtime symbol)"; + } } } diff --git a/src/geocel/vg/detail/VecgeomSetup.cu b/src/geocel/vg/detail/VecgeomSetup.cu index 9483898c27..dc0610b4e2 100644 --- a/src/geocel/vg/detail/VecgeomSetup.cu +++ b/src/geocel/vg/detail/VecgeomSetup.cu @@ -5,9 +5,11 @@ //---------------------------------------------------------------------------// //! \file geocel/vg/detail/VecgeomSetup.cu //---------------------------------------------------------------------------// +#include "VecgeomSetup.hh" + #include -#include "VecgeomSetup.cuda.hh" +#include "corecel/data/DeviceVector.hh" #ifdef VECGEOM_USE_SURF # include @@ -32,9 +34,9 @@ namespace //! Access struct BvhGetter { - VecGeom::cuda::BVH const** dest{nullptr}; + vecgeom::cuda::BVH const** dest{nullptr}; - CELER_FUNCTION operator()(ThreadId tid) + CELER_FUNCTION void operator()(ThreadId tid) { CELER_EXPECT(tid == ThreadId{0}); *dest = vecgeom::cuda::BVHManager::GetBVH(0); @@ -46,14 +48,14 @@ struct BvhGetter /*! * Get pointers to the device BVH after setup, for consistency checking. */ -CudaPointers bvh_pointers_device() +CudaPointers bvh_pointers_device() { - CudaPointers result; + CudaPointers result; // Copy from kernel using 1-thread launch { - DeviceVector bvh_ptr{1, StreamId{}}; - BvhGetter execute_thread{params, state, seeds}; + DeviceVector bvh_ptr{1, StreamId{}}; + BvhGetter execute_thread{bvh_ptr.data()}; static KernelLauncher const launch_kernel( "vecgeom-get-bvhptr"); launch_kernel(1u, StreamId{}, execute_thread); @@ -64,8 +66,8 @@ CudaPointers bvh_pointers_device() // Copy from symbol using runtime API CELER_CUDA_CALL(cudaMemcpyFromSymbol( &result.symbol, vecgeom::cuda::dBVH, 1, 0, cudaMemcpyDeviceToHost)); + CELER_CUDA_CALL(cudaDeviceSynchronize()); - // Return return result; } diff --git a/src/geocel/vg/detail/VecgeomSetup.hh b/src/geocel/vg/detail/VecgeomSetup.hh index 2478ce2041..56d4ff8687 100644 --- a/src/geocel/vg/detail/VecgeomSetup.hh +++ b/src/geocel/vg/detail/VecgeomSetup.hh @@ -7,6 +7,7 @@ //---------------------------------------------------------------------------// #pragma once +#include #include #include "corecel/Assert.hh" @@ -35,7 +36,7 @@ struct CudaPointers //---------------------------------------------------------------------------// // Get pointers to the device BVH after setup, for consistency checking -CudaPointers bvh_pointers_device(); +CudaPointers bvh_pointers_device(); //---------------------------------------------------------------------------// #ifdef VECGEOM_USE_SURF @@ -50,7 +51,7 @@ void teardown_surface_tracking_device(); // INLINE DEFINITIONS //---------------------------------------------------------------------------// #ifndef VECGEOM_ENABLE_CUDA -inline CudaPointers bvh_pointers_device() +inline CudaPointers bvh_pointers_device() { CELER_ASSERT_UNREACHABLE(); } From 81fedfbe7c471ecd96aa12dd5e0ca85e01a12712 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Mon, 4 Nov 2024 10:27:53 -0500 Subject: [PATCH 08/11] Check BVH consistency for new versions of vecgeom --- src/geocel/vg/VecgeomParams.cc | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/src/geocel/vg/VecgeomParams.cc b/src/geocel/vg/VecgeomParams.cc index 49b5cf68cf..4d8b6210f4 100644 --- a/src/geocel/vg/VecgeomParams.cc +++ b/src/geocel/vg/VecgeomParams.cc @@ -445,7 +445,18 @@ void VecgeomParams::build_volume_tracking() CELER_LOG(debug) << "Initializing BVH on GPU"; ScopedTimeAndRedirect time_and_output_( "vecgeom::BVHManager::DeviceInit"); - vecgeom::cxx::BVHManager::DeviceInit(); +#if defined(VECGEOM_BVHMANAGER_DEVICE) + auto* bvh_ptr = BVHManager::DeviceInit(); +#elif defined(VECGEOM_ENABLE_CUDA) + BVHManager::DeviceInit(); +#endif +#ifdef VECGEOM_BVHMANAGER_DEVICE + auto* bvh_symbol_ptr = BVHManager::GetDeviceBVH(); + CELER_VALIDATE(bvh_ptr && bvh_ptr == bvh_symbol_ptr, + << "inconsistent BVH device pointer: allocated " + << bvh_ptr << " but copy-from-symbol returned " + << bvh_symbol_ptr); +#endif CELER_DEVICE_CHECK_ERROR(); } #endif From 2abb4574e9599e3e75face68469ed67dacf94b7a Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Mon, 4 Nov 2024 10:32:06 -0500 Subject: [PATCH 09/11] fixup! Check bvh pointers and rewrite for loop as while --- src/geocel/vg/detail/BVHNavigator.hh | 2 -- 1 file changed, 2 deletions(-) diff --git a/src/geocel/vg/detail/BVHNavigator.hh b/src/geocel/vg/detail/BVHNavigator.hh index 75e26c96a8..1e25b427c5 100644 --- a/src/geocel/vg/detail/BVHNavigator.hh +++ b/src/geocel/vg/detail/BVHNavigator.hh @@ -46,8 +46,6 @@ class BVHNavigator bool top, vecgeom::VPlacedVolume const* exclude = nullptr) { - CELER_EXPECT(vecgeom::BVHManager::GetBVH(0) != nullptr); - if (top) { assert(vol != nullptr); From 8fad2c43df5f7fc020f62722795f608b34bf01cf Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Mon, 4 Nov 2024 10:32:41 -0500 Subject: [PATCH 10/11] Fix link error (missing corecel_final) with vecgeom+celeritas static build --- test/geocel/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/geocel/CMakeLists.txt b/test/geocel/CMakeLists.txt index 149e75a778..811eb72a2b 100644 --- a/test/geocel/CMakeLists.txt +++ b/test/geocel/CMakeLists.txt @@ -56,7 +56,7 @@ celeritas_target_link_libraries(testcel_geocel #-----------------------------------------------------------------------------# celeritas_setup_tests(SERIAL - LINK_LIBRARIES testcel_geocel testcel_core Celeritas::geocel + LINK_LIBRARIES testcel_geocel testcel_core Celeritas::geocel Celeritas::corecel ) #-----------------------------------------------------------------------------# From 09901ed945eb3f1ae35babb9bc7e242d03285795 Mon Sep 17 00:00:00 2001 From: Seth R Johnson Date: Mon, 4 Nov 2024 10:33:08 -0500 Subject: [PATCH 11/11] Use VG_CUDA_CALL and VECGEOM_ENABLE_CUDA for macros in VecgeomParams --- src/geocel/vg/VecgeomParams.cc | 22 ++++++++++++---------- 1 file changed, 12 insertions(+), 10 deletions(-) diff --git a/src/geocel/vg/VecgeomParams.cc b/src/geocel/vg/VecgeomParams.cc index 4d8b6210f4..3954faa908 100644 --- a/src/geocel/vg/VecgeomParams.cc +++ b/src/geocel/vg/VecgeomParams.cc @@ -17,9 +17,9 @@ #include #include "corecel/Config.hh" -#if CELERITAS_USE_CUDA +#include "corecel/DeviceRuntimeApi.hh" +#ifdef VECGEOM_ENABLE_CUDA # include -# include #endif #ifdef VECGEOM_USE_SURF # include @@ -28,8 +28,6 @@ # include #endif -#include "corecel/DeviceRuntimeApi.hh" - #include "corecel/Assert.hh" #include "corecel/Macros.hh" #include "corecel/cont/Range.hh" @@ -54,6 +52,8 @@ static_assert(std::is_same_v, "Celeritas and VecGeom real types do not match"); +using vecgeom::cxx::BVHManager; + namespace celeritas { namespace @@ -385,7 +385,7 @@ void VecgeomParams::build_volume_tracking() } // Init the bounding volume hierarchy structure - vecgeom::cxx::BVHManager::Init(); + BVHManager::Init(); if (celeritas::device()) { @@ -421,22 +421,25 @@ void VecgeomParams::build_volume_tracking() set_cuda_heap_size(heap_size); } -#if CELERITAS_USE_CUDA +#ifdef VECGEOM_ENABLE_CUDA auto& cuda_manager = vecgeom::cxx::CudaManager::Instance(); cuda_manager.set_verbose(vecgeom_verbosity()); +#endif { CELER_LOG(debug) << "Converting to CUDA geometry"; ScopedTimeAndRedirect time_and_output_( "vecgeom::CudaManager.LoadGeometry"); - cuda_manager.LoadGeometry(); + VG_CUDA_CALL(cuda_manager.LoadGeometry()); CELER_DEVICE_CALL_PREFIX(DeviceSynchronize()); } { CELER_LOG(debug) << "Transferring geometry to GPU"; ScopedTimeAndRedirect time_and_output_( "vecgeom::CudaManager.Synchronize"); - auto world_top_devptr = cuda_manager.Synchronize(); + void const* world_top_devptr{nullptr}; + VG_CUDA_CALL( + world_top_devptr = cuda_manager.Synchronize().GetPtr()); CELER_DEVICE_CHECK_ERROR(); CELER_VALIDATE(world_top_devptr != nullptr, << "VecGeom failed to copy geometry to GPU"); @@ -459,7 +462,6 @@ void VecgeomParams::build_volume_tracking() #endif CELER_DEVICE_CHECK_ERROR(); } -#endif // Check BVH pointers auto ptrs = detail::bvh_pointers_device(); @@ -522,7 +524,7 @@ void VecgeomParams::build_data() if (celeritas::device()) { -#if CELERITAS_USE_CUDA +#ifdef VECGEOM_ENABLE_CUDA auto& cuda_manager = vecgeom::cxx::CudaManager::Instance(); device_ref_.world_volume = cuda_manager.world_gpu(); #endif