Add MR test

include/rgl/api/core.h

@@ -387,8 +387,9 @@ typedef enum : int32_t
  */
 typedef enum : int32_t
 {
-	RGL_RETURN_TYPE_FIRST = 0,
-	RGL_RETURN_TYPE_LAST = 1,
+	RGL_RETURN_TYPE_NOT_DIVERGENT = 0,
+	RGL_RETURN_TYPE_FIRST = 1,
+	RGL_RETURN_TYPE_LAST = 2,
 } rgl_return_type_t;
 
 /**

src/gpu/RaytraceRequestContext.hpp

@@ -28,7 +28,7 @@ struct RaytraceRequestContext
 	float nearNonHitDistance;
 	float farNonHitDistance;
 
-	const Mat3x4f* rays;
+	const Mat3x4f* raysWorld;
 	size_t rayCount;
 
 	Mat3x4f rayOriginToWorld;

src/gpu/nodeKernels.cu

@@ -224,7 +224,7 @@ __global__ void kFilterGroundPoints(size_t pointCount, const Vec3f sensor_up_vec
 }
 
 __global__ void kProcessBeamSamplesFirstLast(size_t beamCount, int samplesPerBeam, MultiReturnPointers beamSamples,
-                                             MultiReturnPointers first, MultiReturnPointers last)
+                                             MultiReturnPointers first, MultiReturnPointers last, const Mat3x4f* beamsWorld)
 {
 	LIMIT(beamCount);
 
@@ -244,14 +244,16 @@ __global__ void kProcessBeamSamplesFirstLast(size_t beamCount, int samplesPerBea
 			lastIdx = sampleIdx;
 		}
 	}
+	Vec3f beamOrigin = beamsWorld[beamIdx] * Vec3f{0, 0, 0};
+	Vec3f beamDir = ((beamsWorld[beamIdx] * Vec3f{0, 0, 1}) - beamOrigin).normalized();
 	bool isHit = firstIdx >= 0; // Note that firstHit >= 0 implies lastHit >= 0
 	first.isHit[beamIdx] = isHit;
 	last.isHit[beamIdx] = isHit;
 	if (isHit) {
-		first.xyz[beamIdx] = beamSamples.xyz[beamIdx * samplesPerBeam + firstIdx];
 		first.distance[beamIdx] = beamSamples.distance[beamIdx * samplesPerBeam + firstIdx];
-		last.xyz[beamIdx] = beamSamples.xyz[beamIdx * samplesPerBeam + lastIdx];
 		last.distance[beamIdx] = beamSamples.distance[beamIdx * samplesPerBeam + lastIdx];
+		first.xyz[beamIdx] = beamOrigin + beamDir * first.distance[beamIdx];
+		last.xyz[beamIdx] = beamOrigin + beamDir * last.distance[beamIdx];
 	}
 }
 
@@ -329,7 +331,7 @@ void gpuRadarComputeEnergy(cudaStream_t stream, size_t count, float rayAzimuthSt
 }
 
 void gpuProcessBeamSamplesFirstLast(cudaStream_t stream, size_t beamCount, int samplesPerBeam, MultiReturnPointers beamSamples,
-                                    MultiReturnPointers first, MultiReturnPointers last)
+                                    MultiReturnPointers first, MultiReturnPointers last, const Mat3x4f* beamsWorld)
 {
-	run(kProcessBeamSamplesFirstLast, stream, beamCount, samplesPerBeam, beamSamples, first, last);
+	run(kProcessBeamSamplesFirstLast, stream, beamCount, samplesPerBeam, beamSamples, first, last, beamsWorld);
 }

src/gpu/nodeKernels.hpp

@@ -53,4 +53,4 @@ void gpuRadarComputeEnergy(cudaStream_t stream, size_t count, float rayAzimuthSt
                            const Field<DISTANCE_F32>::type* hitDist, const Field<NORMAL_VEC3_F32>::type* hitNorm,
                            const Field<XYZ_VEC3_F32>::type* hitPos, Vector<3, thrust::complex<float>>* outBUBRFactor);
 void gpuProcessBeamSamplesFirstLast(cudaStream_t stream, size_t beamCount, int samplesPerBeam, MultiReturnPointers beamSamples,
-                                    MultiReturnPointers first, MultiReturnPointers last);
+                                    MultiReturnPointers first, MultiReturnPointers last, const Mat3x4f* beamWorld);

src/gpu/optixPrograms.cu

@@ -55,7 +55,7 @@ extern "C" __global__ void __raygen__()
 		return;
 	}
 
-	Mat3x4f ray = ctx.rays[rayIdx];
+	Mat3x4f ray = ctx.raysWorld[rayIdx];
 	const Mat3x4f rayLocal =
 	    ctx.rayOriginToWorld.inverse() *
 	    ray; // TODO(prybicki): instead of computing inverse, we should pass rays in local CF and then transform them to world CF.
@@ -137,12 +137,16 @@ extern "C" __global__ void __closesthit__()
 	// Hitpoint
 	Vec3f hitObject = Vec3f((1 - u - v) * A + u * B + v * C);
 	const Vec3f hitWorldRaytraced = optixTransformPointFromObjectToWorldSpace(hitObject);
-	const float distance = (hitWorldRaytraced - beamSampleOrigin).length();
+	const Vector<3, double> hwrd = hitWorldRaytraced;
+	const Vector<3, double> hso = beamSampleOrigin;
+	const double distance = (hwrd - hso).length();
 	const Vec3f hitWorldSeenBySensor = [&]() {
 		if (!ctx.doApplyDistortion) {
 			return hitWorldRaytraced;
 		}
-		Mat3x4f undistortedRay = ctx.rays[beamIdx] * makeBeamSampleRayTransform(ctx.beamHalfDivergence, circleIdx, vertexIdx);
+		Mat3x4f sampleRayTf = beamSampleRayIdx == 0 ? Mat3x4f::identity() :
+		                                              makeBeamSampleRayTransform(ctx.beamHalfDivergence, circleIdx, vertexIdx);
+		Mat3x4f undistortedRay = ctx.raysWorld[beamIdx] * sampleRayTf;
 		Vec3f undistortedOrigin = undistortedRay * Vec3f{0, 0, 0};
 		Vec3f undistortedDir = undistortedRay * Vec3f{0, 0, 1} - undistortedOrigin;
 		return undistortedOrigin + undistortedDir * distance;
@@ -186,7 +190,6 @@ extern "C" __global__ void __closesthit__()
 
 	// Save sub-sampling results
 	ctx.mrSamples.isHit[mrSamplesIdx] = true;
-	ctx.mrSamples.xyz[mrSamplesIdx] = hitWorldSeenBySensor;
 	ctx.mrSamples.distance[mrSamplesIdx] = distance;
 	if (beamSampleRayIdx != 0) {
 		return;
@@ -267,7 +270,7 @@ __device__ Mat3x4f makeBeamSampleRayTransform(float halfDivergenceAngleRad, unsi
 
 __device__ void saveNonHitRayResult(float nonHitDistance)
 {
-	Mat3x4f ray = ctx.rays[optixGetLaunchIndex().x];
+	Mat3x4f ray = ctx.raysWorld[optixGetLaunchIndex().x];
 	Vec3f origin = ray * Vec3f{0, 0, 0};
 	Vec3f dir = ray * Vec3f{0, 0, 1} - origin;
 	Vec3f displacement = dir.normalized() * nonHitDistance;

src/graph/NodesCore.hpp

@@ -212,6 +212,9 @@ struct MultiReturnSwitchNode : IPointsNodeSingleInput
 	// Data getters
 	IAnyArray::ConstPtr getFieldData(rgl_field_t field) override
 	{
+		if (returnType == RGL_RETURN_TYPE_NOT_DIVERGENT) {
+			return rtxInput->getFieldData(field);
+		}
 		return rtxInput->getFieldDataMultiReturn(field, returnType);
 	}
 

src/graph/RaytraceNode.cpp

@@ -94,7 +94,7 @@ void RaytraceNode::enqueueExecImpl()
 	    .doApplyDistortion = doApplyDistortion,
 	    .nearNonHitDistance = nearNonHitDistance,
 	    .farNonHitDistance = farNonHitDistance,
-	    .rays = raysPtr,
+	    .raysWorld = raysPtr,
 	    .rayCount = raysNode->getRayCount(),
 	    .rayOriginToWorld = raysNode->getCumulativeRayTransfrom(),
 	    .rayRanges = rayRanges.has_value() ? (*rayRanges)->asSubclass<DeviceAsyncArray>()->getReadPtr() :
@@ -133,8 +133,10 @@ void RaytraceNode::enqueueExecImpl()
 	CHECK_OPTIX(optixLaunch(Optix::getOrCreate().pipeline, getStreamHandle(), pipelineArgsPtr, pipelineArgsSize, &sceneSBT,
 	                        launchDims.x, launchDims.y, launchDims.y));
 
-	gpuProcessBeamSamplesFirstLast(getStreamHandle(), raysNode->getRayCount(), MULTI_RETURN_BEAM_SAMPLES,
-	                               mrSamples.getPointers(), mrFirst.getPointers(), mrLast.getPointers());
+	if (beamHalfDivergence > 0.0f) {
+		gpuProcessBeamSamplesFirstLast(getStreamHandle(), raysNode->getRayCount(), MULTI_RETURN_BEAM_SAMPLES,
+		                               mrSamples.getPointers(), mrFirst.getPointers(), mrLast.getPointers(), raysPtr);
+	}
 }
 
 IAnyArray::ConstPtr RaytraceNode::getFieldDataMultiReturn(rgl_field_t field, rgl_return_type_t type)
@@ -144,15 +146,15 @@ IAnyArray::ConstPtr RaytraceNode::getFieldDataMultiReturn(rgl_field_t field, rgl
 			case XYZ_VEC3_F32: return mrFirst.xyz;
 			case DISTANCE_F32: return mrFirst.distance;
 			case IS_HIT_I32: return mrFirst.isHit;
-			default: throw InvalidPipeline(fmt::format("Multi-Return not supported for this field ({})", toString(field)));
+			default: return getFieldData(field);
 		}
 	}
 	if (type == RGL_RETURN_TYPE_LAST) {
 		switch (field) {
 			case XYZ_VEC3_F32: return mrLast.xyz;
 			case DISTANCE_F32: return mrLast.distance;
 			case IS_HIT_I32: return mrLast.isHit;
-			default: throw InvalidPipeline(fmt::format("Multi-Return not supported for this field ({})", toString(field)));
+			default: return getFieldData(field);
 		}
 	}
 	throw InvalidPipeline(fmt::format("Unknown multi-return type ({})", type));

test/include/helpers/geometryData.hpp

@@ -3,6 +3,7 @@
 #include <rgl/api/core.h>
 
 static constexpr float CUBE_HALF_EDGE = 1.0;
+static constexpr float CUBE_EDGE = 2.0;
 
 static rgl_vec3f cubeVertices[] = {
     {-1, -1, -1}, // 0

test/include/helpers/testPointCloud.hpp

@@ -241,7 +241,7 @@ class TestPointCloud
 	}
 
 	template<rgl_field_t T>
-	std::vector<typename Field<T>::type> getFieldValues()
+	std::vector<typename Field<T>::type> getFieldValues() const
 	{
 		int fieldIndex = std::find(fields.begin(), fields.end(), T) - fields.begin();
 
@@ -252,7 +252,7 @@ class TestPointCloud
 
 		for (int i = 0; i < getPointCount(); i++) {
 			fieldValues.emplace_back(
-			    *reinterpret_cast<typename Field<T>::type*>(data.data() + i * getPointByteSize() + offset));
+			    *reinterpret_cast<const typename Field<T>::type*>(data.data() + i * getPointByteSize() + offset));
 		}
 
 		return fieldValues;