Fix bounds computation on GPU

szellmann · Jan 20, 2025 · 84fc2a4 · 84fc2a4
1 parent abe6ea5
commit 84fc2a4
Show file tree

Hide file tree

Showing 4 changed files with 89 additions and 3 deletions.
diff --git a/scene/VisionarayScene.cpp b/scene/VisionarayScene.cpp
@@ -316,6 +316,9 @@ aabb VisionaraySceneImpl::getBounds() const
 void VisionaraySceneImpl::attachInstance(
     dco::Instance inst, unsigned instID, unsigned userID)
 {
+#if defined(WITH_CUDA) || defined(WITH_HIP)
+  m_gpuScene->attachInstance(inst, instID, userID);
+#else
   m_bounds[boundsID].insert(get_bounds(inst));
 
   m_instances.set(instID, inst.instID);
@@ -326,20 +329,21 @@ void VisionaraySceneImpl::attachInstance(
 
   // Upload/set accessible pointers
   deviceState()->onDevice.instances = deviceState()->dcos.instances.devicePtr();
+#endif
 }
 
 void VisionaraySceneImpl::attachGeometry(
     dco::Geometry geom, unsigned geomID, unsigned userID)
 {
-  m_bounds[boundsID].insert(get_bounds(geom));
-
 #if defined(WITH_CUDA) || defined(WITH_HIP)
   m_gpuScene->attachGeometry(geom, geomID, userID);
 #else
 
   if (geom.primitives.len == 0)
     return;
 
+  m_bounds[boundsID].insert(get_bounds(geom));
+
   m_geometries.set(geomID, geom.geomID);
   m_objIds.set(geomID, userID);
 
@@ -392,7 +396,8 @@ void VisionaraySceneImpl::attachGeometry(
 void VisionaraySceneImpl::attachVolume(
     dco::Volume vol, unsigned volID, unsigned userID)
 {
-  m_bounds[boundsID].insert(get_bounds(vol));
+  // use bounds member, that way we don't need to reach for the GPU:
+  m_bounds[boundsID].insert(vol.bounds);
 
   // Patch volID into scene primitives:
   vol.volID = volID;

diff --git a/scene/VisionaraySceneGPU.cu b/scene/VisionaraySceneGPU.cu
@@ -4,6 +4,15 @@
 
 namespace visionaray {
 
+template <typename Obj>
+__global__ static void getBoundsGPU(Obj obj, aabb *bounds)
+{
+  if (blockIdx.x != 0 || threadIdx.x != 0)
+    return;
+
+  *bounds = get_bounds(obj);
+}
+
 struct VisionaraySceneGPU::Impl
 {
   typedef cuda_index_bvh<basic_triangle<3,float>> TriangleBVH;
@@ -402,12 +411,43 @@ void VisionaraySceneGPU::dispatch()
   }
 }
 
+void VisionaraySceneGPU::attachInstance(
+    dco::Instance inst, unsigned instID, unsigned userID)
+{
+  aabb *bounds;
+  CUDA_SAFE_CALL(cudaMalloc(&bounds, sizeof(aabb)));
+  getBoundsGPU<<<1,1>>>(inst, bounds);
+  aabb hostBounds;
+  CUDA_SAFE_CALL(
+      cudaMemcpy(&hostBounds, bounds, sizeof(aabb), cudaMemcpyDeviceToHost));
+  m_impl->parent->m_bounds[m_impl->parent->boundsID].insert(hostBounds);
+  CUDA_SAFE_CALL(cudaFree(bounds));
+
+  m_impl->parent->m_instances.set(instID, inst.instID);
+  m_impl->parent->m_objIds.set(instID, userID); // TODO: separate inst/geom
+
+  m_impl->parent->m_instances.set(instID, inst.instID);
+  deviceState()->dcos.instances.update(inst.instID, inst);
+
+  // Upload/set accessible pointers
+  deviceState()->onDevice.instances = deviceState()->dcos.instances.devicePtr();
+}
+
 void VisionaraySceneGPU::attachGeometry(
     dco::Geometry geom, unsigned geomID, unsigned userID)
 {
   if (geom.primitives.len == 0)
     return;
 
+  aabb *bounds;
+  CUDA_SAFE_CALL(cudaMalloc(&bounds, sizeof(aabb)));
+  getBoundsGPU<<<1,1>>>(geom, bounds);
+  aabb hostBounds;
+  CUDA_SAFE_CALL(
+      cudaMemcpy(&hostBounds, bounds, sizeof(aabb), cudaMemcpyDeviceToHost));
+  m_impl->parent->m_bounds[m_impl->parent->boundsID].insert(hostBounds);
+  CUDA_SAFE_CALL(cudaFree(bounds));
+
   m_impl->parent->m_geometries.set(geomID, geom.geomID);
   m_impl->parent->m_objIds.set(geomID, userID);
 

diff --git a/scene/VisionaraySceneGPU.h b/scene/VisionaraySceneGPU.h
@@ -17,6 +17,7 @@ struct VisionaraySceneGPU
   aabb getBounds() const;
   void commit();
   void dispatch();
+  void attachInstance(dco::Instance inst, unsigned instID, unsigned userID=~0u);
   void attachGeometry(dco::Geometry geom, unsigned geomID, unsigned userID=~0u);
 
 #ifdef WITH_CUDA

diff --git a/scene/VisionaraySceneGPU.hip b/scene/VisionaraySceneGPU.hip
@@ -17,6 +17,15 @@ void split_primitive(aabb& L, aabb& R, float plane, int axis, dco::Instance cons
   // empty
 }
 
+template <typename Obj>
+__global__ static void getBoundsGPU(Obj obj, aabb *bounds)
+{
+  if (blockIdx.x != 0 || threadIdx.x != 0)
+    return;
+
+  *bounds = get_bounds(obj);
+}
+
 struct VisionaraySceneGPU::Impl
 {
   typedef hip_index_bvh<basic_triangle<3,float>> TriangleBVH;
@@ -419,12 +428,43 @@ void VisionaraySceneGPU::dispatch()
   }
 }
 
+void VisionaraySceneGPU::attachInstance(
+    dco::Instance inst, unsigned instID, unsigned userID)
+{
+  aabb *bounds;
+  HIP_SAFE_CALL(hipMalloc(&bounds, sizeof(aabb)));
+  getBoundsGPU<<<1,1>>>(inst, bounds);
+  aabb hostBounds;
+  HIP_SAFE_CALL(
+      hipMemcpy(&hostBounds, bounds, sizeof(aabb), hipMemcpyDeviceToHost));
+  m_impl->parent->m_bounds[m_impl->parent->boundsID].insert(hostBounds);
+  HIP_SAFE_CALL(hipFree(bounds));
+
+  m_impl->parent->m_instances.set(instID, inst.instID);
+  m_impl->parent->m_objIds.set(instID, userID); // TODO: separate inst/geom
+
+  m_impl->parent->m_instances.set(instID, inst.instID);
+  deviceState()->dcos.instances.update(inst.instID, inst);
+
+  // Upload/set accessible pointers
+  deviceState()->onDevice.instances = deviceState()->dcos.instances.devicePtr();
+}
+
 void VisionaraySceneGPU::attachGeometry(
     dco::Geometry geom, unsigned geomID, unsigned userID)
 {
   if (geom.primitives.len == 0)
     return;
 
+  aabb *bounds;
+  HIP_SAFE_CALL(hipMalloc(&bounds, sizeof(aabb)));
+  getBoundsGPU<<<1,1>>>(geom, bounds);
+  aabb hostBounds;
+  HIP_SAFE_CALL(
+      hipMemcpy(&hostBounds, bounds, sizeof(aabb), hipMemcpyDeviceToHost));
+  m_impl->parent->m_bounds[m_impl->parent->boundsID].insert(hostBounds);
+  HIP_SAFE_CALL(hipFree(bounds));
+
   m_impl->parent->m_geometries.set(geomID, geom.geomID);
   m_impl->parent->m_objIds.set(geomID, userID);