Fix error when raytracing on empty scene (#55)

* Fix error when raytracing on empty scene * Fix NaNs in non-hit XYZ * Update src/gpu/optixPrograms.cu Co-authored-by: Mateusz Szczygielski <[email protected]> Co-authored-by: Mateusz Szczygielski <[email protected]>
RobotecAI · Nov 22, 2022 · 2d4055b · 2d4055b
1 parent 45ba410
commit 2d4055b
Show file tree

Hide file tree

Showing 5 changed files with 35 additions and 30 deletions.
diff --git a/src/DeviceBuffer.hpp b/src/DeviceBuffer.hpp
@@ -116,17 +116,16 @@ struct DeviceBuffer
 
 private:
 	bool ensureDeviceCanFit(std::size_t newElemCount) {
-		if (newElemCount == 0) {
-			auto msg = fmt::format("Attempted to allocate {} bytes of memory", newElemCount);
-			throw std::logic_error(msg);
-		}
 		if (elemCapacity >= newElemCount) {
 			return false;
 		}
 		if (data != nullptr) {
 			CHECK_CUDA(cudaFree(data));
+			data = nullptr;
+		}
+		if (newElemCount > 0) {
+			CHECK_CUDA(cudaMalloc(reinterpret_cast<void**>(&data), newElemCount * sizeof(T)));
 		}
-		CHECK_CUDA(cudaMalloc(reinterpret_cast<void**>(&data), newElemCount * sizeof(T)));
 		elemCapacity = newElemCount;
 		return true;
 	}

diff --git a/src/gpu/RaytraceRequestContext.hpp b/src/gpu/RaytraceRequestContext.hpp
@@ -36,6 +36,7 @@ struct RaytraceRequestContext
 	Field<IS_HIT_I32>::type* isHit;
 	Field<RAY_IDX_U32>::type* rayIdx;
 	Field<RING_ID_U16>::type* ringIdx;
-	Field<DISTANCE_F32>::type* distanceIdx;
-	Field<INTENSITY_F32>::type* intensityIdx;
+	Field<DISTANCE_F32>::type* distance;
+	Field<INTENSITY_F32>::type* intensity;
 };
+static_assert(std::is_trivially_copyable<RaytraceRequestContext>::value);
diff --git a/src/gpu/optixPrograms.cu b/src/gpu/optixPrograms.cu
@@ -23,7 +23,6 @@
 #include <gpu/RaytraceRequestContext.hpp>
 #include <gpu/ShaderBindingTableTypes.h>
 
-
 extern "C" static __constant__ RaytraceRequestContext ctx;
 
 struct Vec3fPayload
@@ -53,21 +52,14 @@ Vec3f decodePayloadVec3f(const Vec3fPayload& src)
 	};
 }
 
-
 template<bool isFinite>
 __forceinline__ __device__
-void saveRayResult(Vec3f* xyz=nullptr)
+void saveRayResult(const Vec3f* xyz=nullptr, const Vec3f* origin=nullptr)
 {
 	const int rayIdx = optixGetLaunchIndex().x;
-	Vec3f origin = decodePayloadVec3f({
-		optixGetPayload_0(),
-		optixGetPayload_1(),
-		optixGetPayload_2()
-	});
-
 	if (ctx.xyz != nullptr) {
-		// Return actual XYZ of the hit point or infinity vector with signs of the ray.
-		ctx.xyz[rayIdx] = isFinite ? *xyz : ctx.rays[rayIdx] * Vec3f{CUDART_INF_F, CUDART_INF_F, CUDART_INF_F};
+		// Return actual XYZ of the hit point or infinity vector.
+		ctx.xyz[rayIdx] = isFinite ? *xyz : Vec3f{CUDART_INF_F, CUDART_INF_F, CUDART_INF_F};
 	}
 	if (ctx.isHit != nullptr) {
 		ctx.isHit[rayIdx] = isFinite;
@@ -78,21 +70,26 @@ void saveRayResult(Vec3f* xyz=nullptr)
 	if (ctx.ringIdx != nullptr && ctx.ringIds != nullptr) {
 		ctx.ringIdx[rayIdx] = ctx.ringIds[rayIdx % ctx.ringIdsCount];
 	}
-	if (ctx.distanceIdx != nullptr) {
-		ctx.distanceIdx[rayIdx] = isFinite
-		                        ? sqrt(pow(
-		                            (*xyz)[0] - origin[0], 2) +
-		                            pow((*xyz)[1] - origin[1], 2) +
-		                            pow((*xyz)[2] - origin[2], 2))
+	if (ctx.distance != nullptr) {
+		ctx.distance[rayIdx] = isFinite
+		                        ? sqrt(
+		                            pow((*xyz)[0] - (*origin)[0], 2) +
+		                            pow((*xyz)[1] - (*origin)[1], 2) +
+		                            pow((*xyz)[2] - (*origin)[2], 2))
 		                        : CUDART_INF_F;
 	}
-	if (ctx.intensityIdx != nullptr) {
-		ctx.intensityIdx[rayIdx] = 100;
+	if (ctx.intensity != nullptr) {
+		ctx.intensity[rayIdx] = 100;
 	}
 }
 
 extern "C" __global__ void __raygen__()
 {
+	if (ctx.scene == 0) {
+		saveRayResult<false>();
+		return;
+	}
+
 	Mat3x4f ray = ctx.rays[optixGetLaunchIndex().x];
 
 	Vec3f origin = ray * Vec3f{0, 0, 0};
@@ -124,7 +121,12 @@ extern "C" __global__ void __closesthit__()
 	Vec3f hitObject = Vec3f((1 - u - v) * A + u * B + v * C);
 	Vec3f hitWorld = optixTransformPointFromObjectToWorldSpace(hitObject);
 
-	saveRayResult<true>(&hitWorld);
+	Vec3f origin = decodePayloadVec3f({
+		optixGetPayload_0(),
+		optixGetPayload_1(),
+		optixGetPayload_2()
+	});
+	saveRayResult<true>(&hitWorld, &origin);
 }
 
 extern "C" __global__ void __miss__()

diff --git a/src/graph/RaytraceNode.cpp b/src/graph/RaytraceNode.cpp
@@ -57,8 +57,8 @@ void RaytraceNode::schedule(cudaStream_t stream)
 		.isHit = getPtrTo<IS_HIT_I32>(),
 		.rayIdx = getPtrTo<RAY_IDX_U32>(),
 		.ringIdx = getPtrTo<RING_ID_U16>(),
-		.distanceIdx = getPtrTo<DISTANCE_F32>(),
-		.intensityIdx = getPtrTo<INTENSITY_F32>(),
+		.distance = getPtrTo<DISTANCE_F32>(),
+		.intensity = getPtrTo<INTENSITY_F32>(),
 	};
 
 	CUdeviceptr pipelineArgsPtr = requestCtx->getCUdeviceptr();

diff --git a/src/scene/Scene.cpp b/src/scene/Scene.cpp
@@ -102,14 +102,17 @@ OptixShaderBindingTable Scene::buildSBT()
 		.missRecordBase = dMissRecords.readDeviceRaw(),
 		.missRecordStrideInBytes = sizeof(MissRecord),
 		.missRecordCount = 1U,
-		.hitgroupRecordBase = dHitgroupRecords.readDeviceRaw(),
+		.hitgroupRecordBase = getObjectCount() > 0 ? dHitgroupRecords.readDeviceRaw() : static_cast<CUdeviceptr>(0),
 		.hitgroupRecordStrideInBytes = sizeof(HitgroupRecord),
 		.hitgroupRecordCount = static_cast<unsigned>(dHitgroupRecords.getElemCount()),
 	};
 }
 
 OptixTraversableHandle Scene::buildAS()
 {
+	if (getObjectCount() == 0) {
+		return static_cast<OptixTraversableHandle>(0);
+	}
 	std::vector<OptixInstance> instances;
 	for (auto&& entity : entities) {
 		// TODO(prybicki): this is somewhat inefficient, because most of the time only transform changes.