OClCollMeshEngn: Produce ambience count; set postrin threshold

We modify the semantics/meaning of the ambience stim feature.
It now represents the number of frames whose average intensity
is below the ambienceLowVal.

We can now implement the postrin as the event wherein the number
of frames whose intensity <= ambienceLowVal exceeds
postrin-interest-threshold.

docs/livox-gen1-lidar-dap-spec.md

@@ -33,7 +33,7 @@ Each stim-buff-api is designed to work with specific stim-iface libraries that u
 
 ### 2. Point Cloud Ambience Data Device (Interoceptor)
 
-**Purpose**: Provides ambience data from the LiDAR point cloud, counting high-intensity points per slot.
+**Purpose**: Provides ambience data from the LiDAR point cloud, computing average intensity per frame and counting frames that qualify for intrinsic stimuli.
 
 **Syntax**:
 ```
@@ -41,16 +41,18 @@ Each stim-buff-api is designed to work with specific stim-iface libraries that u
 ```
 
 **Stim-Buff-API**: `livoxGen1-pcloud`
-**Quale-Iface-API**: `pcloudAmbience` - Counts points with intensity >= threshold per slot
+**Quale-Iface-API**: `pcloudAmbience` - Computes average intensity per frame and exports postrin/negtrin based on configurable thresholds
 
-**Postrin Support** (for pcloudAmbience quale-iface-api):
-The `pcloudAmbience` quale-iface-api can supply postrins (positive intrinsic
-stimuli). The default threshold values for postrin importance levels are:
-- **Stupefaction threshold**: 0 (postrin-stupefaction-threshold=0)
-- **Postrin distraction threshold**: 10 (postrin-distraction=10)
-- **Postrin interest threshold**: 30 (postrin-interest-threshold=30)
-These thresholds can be configured via `quale-iface-api-params` to override
-the defaults.
+**Intrinsic Stimuli Support** (for pcloudAmbience quale-iface-api):
+The `pcloudAmbience` quale-iface-api exports both a postrin and a negtrin whose
+thresholds are configurable via standard quale-iface-api-params:
+- **Postrin interest threshold**: Configurable via `postrin-interest-threshold` / `postrin-interest` (default: 30)
+- **Negtrin interest threshold**: Configurable via `negtrin-interest-threshold` / `negtrin-interest` (default: 30)
+
+Frames whose average intensity is <= postrin-interest-threshold qualify for
+postrin, and frames whose average intensity is >= negtrin-interest-threshold
+qualify for negtrin. The ambience count written to each stimFrame is the number
+of frames that qualify for either intrinsic.
 
 ### 3. Point Cloud Coordinate Data Device (Extrospector)
 

stimBuffApis/livoxGen1/collateDgrams.cl

@@ -25,8 +25,7 @@ __kernel void collate(
 	__global uchar* assembly,
 	__global float* collation,
 	__global float* intensityBuffer,
-	__global uint* ambienceBuffer,
-	uint ambienceHighVal,
+	__global float* averageIntensityBuffer,
 	uint slotStride,
 	uint nPointsPerSlot,
 	uint nDgramsPerFrame)
@@ -55,8 +54,9 @@ __kernel void collate(
 	uint intensityBaseOffset = slotIndex * nPointsPerSlot;
 	DBG_PRINTF("Running kernel: about to process points in slot.\n");
 
-	// Initialize ambience counter for this work item
-	uint ambienceCount = 0;
+	// Initialize running average calculation for this work item
+	float intensitySum = 0.0f;
+	uint validPointCount = 0;
 
 	// Process based on data type using nested ifs (outer) with loops (inner)
 	if (dataType == 0)
@@ -97,9 +97,11 @@ __kernel void collate(
 			if (intensityBuffer != NULL) {
 				intensityBuffer[intensityBaseOffset + i] = intensity;
 			}
-			// Count high intensity values for ambience buffer
-			if (intensity >= (float)ambienceHighVal) {
-				++ambienceCount;
+			// Accumulate intensity for average calculation (exclude points where XYZ=0)
+			if (!(x == 0.0f && y == 0.0f && z == 0.0f))
+			{
+				intensitySum += intensity;
+				++validPointCount;
 			}
 			// Don't write intensity to collation buffer
 		}
@@ -143,9 +145,11 @@ __kernel void collate(
 			if (intensityBuffer != NULL) {
 				intensityBuffer[intensityBaseOffset + i] = intensity;
 			}
-			// Count high intensity values for ambience buffer
-			if (intensity >= (float)ambienceHighVal) {
-				++ambienceCount;
+			// Accumulate intensity for average calculation (exclude points where XYZ=0)
+			if (!(x == 0.0f && y == 0.0f && z == 0.0f))
+			{
+				intensitySum += intensity;
+				++validPointCount;
 			}
 			// Don't write intensity to collation buffer
 		}
@@ -190,9 +194,11 @@ __kernel void collate(
 			if (intensityBuffer != NULL) {
 				intensityBuffer[intensityBaseOffset + pointIndex] = intensity1;
 			}
-			// Count high intensity values for ambience buffer
-			if (intensity1 >= (float)ambienceHighVal) {
-				++ambienceCount;
+			// Accumulate intensity for average calculation (exclude points where XYZ=0)
+			if (!(x1 == 0.0f && y1 == 0.0f && z1 == 0.0f))
+			{
+				intensitySum += intensity1;
+				++validPointCount;
 			}
 			// Don't write intensity to collation buffer
 			++pointIndex;
@@ -225,9 +231,11 @@ __kernel void collate(
 			if (intensityBuffer != NULL) {
 				intensityBuffer[intensityBaseOffset + pointIndex] = intensity2;
 			}
-			// Count high intensity values for ambience buffer
-			if (intensity2 >= (float)ambienceHighVal) {
-				++ambienceCount;
+			// Accumulate intensity for average calculation (exclude points where XYZ=0)
+			if (!(x2 == 0.0f && y2 == 0.0f && z2 == 0.0f))
+			{
+				intensitySum += intensity2;
+				++validPointCount;
 			}
 			// Don't write intensity to collation buffer
 			++pointIndex;
@@ -273,9 +281,11 @@ __kernel void collate(
 			if (intensityBuffer != NULL) {
 				intensityBuffer[intensityBaseOffset + pointIndex] = intensity1;
 			}
-			// Count high intensity values for ambience buffer
-			if (intensity1 >= (float)ambienceHighVal) {
-				++ambienceCount;
+			// Accumulate intensity for average calculation (exclude points where XYZ=0)
+			if (!(x1 == 0.0f && y1 == 0.0f && z1 == 0.0f))
+			{
+				intensitySum += intensity1;
+				++validPointCount;
 			}
 			// Don't write intensity to collation buffer
 			++pointIndex;
@@ -308,9 +318,11 @@ __kernel void collate(
 			if (intensityBuffer != NULL) {
 				intensityBuffer[intensityBaseOffset + pointIndex] = intensity2;
 			}
-			// Count high intensity values for ambience buffer
-			if (intensity2 >= (float)ambienceHighVal) {
-				++ambienceCount;
+			// Accumulate intensity for average calculation (exclude points where XYZ=0)
+			if (!(x2 == 0.0f && y2 == 0.0f && z2 == 0.0f))
+			{
+				intensitySum += intensity2;
+				++validPointCount;
 			}
 			// Don't write intensity to collation buffer
 			++pointIndex;
@@ -343,9 +355,11 @@ __kernel void collate(
 			if (intensityBuffer != NULL) {
 				intensityBuffer[intensityBaseOffset + pointIndex] = intensity3;
 			}
-			// Count high intensity values for ambience buffer
-			if (intensity3 >= (float)ambienceHighVal) {
-				++ambienceCount;
+			// Accumulate intensity for average calculation (exclude points where XYZ=0)
+			if (!(x3 == 0.0f && y3 == 0.0f && z3 == 0.0f))
+			{
+				intensitySum += intensity3;
+				++validPointCount;
 			}
 			// Don't write intensity to collation buffer
 			++pointIndex;
@@ -353,8 +367,10 @@ __kernel void collate(
 	}
 	// Unsupported data types are silently ignored
 
-	// Write ambience count for this work item (once at the end)
-	if (ambienceBuffer != NULL) {
-		ambienceBuffer[slotIndex] = ambienceCount;
+	// Write average intensity for this work item (once at the end)
+	if (averageIntensityBuffer != NULL)
+	{
+		averageIntensityBuffer[slotIndex] = (validPointCount > 0) ?
+			(intensitySum / (float)validPointCount) : 0.0f;
 	}
 }

stimBuffApis/livoxGen1/openClCollatingAndMeshingEngine.cpp

@@ -37,8 +37,10 @@ slotCompactorProgram(nullptr), collateProgram(nullptr),
 slotCompactorKernel(nullptr), collateKernel(nullptr),
 clAssemblyBufferClBuffer(nullptr),
 clCollationBufferClBuffer(nullptr),
+clAverageIntensityBufferClBuffer(nullptr),
 clAssemblyBuffer(nullptr),
 clCollationBuffer(nullptr),
+clAverageIntensityBuffer(nullptr),
 shouldAcceptRequests(false),
 compactIsRunning(false),
 collateIsRunning(false),
@@ -47,8 +49,11 @@ assemblyBufferPtr(nullptr),
 assemblyBufferSize(0),
 collationBufferPtr(nullptr),
 collationBufferSize(0),
+averageIntensityBufferPtr(nullptr),
+averageIntensityBufferSize(0),
 mappedAssemblyBuffer(nullptr),
 mappedCollationBuffer(nullptr),
+mappedAverageIntensityBuffer(nullptr),
 frameAssemblyDesc(nullptr)
 {
 }
@@ -87,11 +92,15 @@ bool OpenClCollatingAndMeshingEngine::setup()
 	// Get StagingBuffer memory pointers from parent
 	struct iovec assemblyIov = parent.assemblyBuffer.getClEngineIovec();
 	struct iovec collationIov = parent.collationBuffer.getClEngineIovec();
+	struct iovec averageIntensityIov = parent.averageIntensityBuffer
+		.getClEngineIovec();
 
 	assemblyBufferPtr = assemblyIov.iov_base;
 	assemblyBufferSize = assemblyIov.iov_len;
 	collationBufferPtr = collationIov.iov_base;
 	collationBufferSize = collationIov.iov_len;
+	averageIntensityBufferPtr = averageIntensityIov.iov_base;
+	averageIntensityBufferSize = averageIntensityIov.iov_len;
 
 	// Get FrameAssemblyDesc from assembly buffer
 	frameAssemblyDesc = static_cast<std::shared_ptr<FrameAssemblyDesc>>(
@@ -133,13 +142,27 @@ bool OpenClCollatingAndMeshingEngine::setup()
 		goto cleanup;
 	}
 
+	clAverageIntensityBufferClBuffer = smoHooksPtr
+		->ComputeManager_createUseHostPtrBuffer(
+			averageIntensityBufferPtr, averageIntensityBufferSize,
+			CL_MEM_WRITE_ONLY);
+
+	if (!clAverageIntensityBufferClBuffer)
+	{
+		std::cerr << __func__ << ": failed to create average intensity buffer"
+			<< std::endl;
+		goto cleanup;
+	}
+
 	// Cache cl_mem handles for the device we're using
 	clAssemblyBuffer = clAssemblyBufferClBuffer
 		->getAssociatedBufferHandleForDevice(computeDevice);
 	clCollationBuffer = clCollationBufferClBuffer
 		->getAssociatedBufferHandleForDevice(computeDevice);
+	clAverageIntensityBuffer = clAverageIntensityBufferClBuffer
+		->getAssociatedBufferHandleForDevice(computeDevice);
 
-	if (!clAssemblyBuffer || !clCollationBuffer)
+	if (!clAssemblyBuffer || !clCollationBuffer || !clAverageIntensityBuffer)
 	{
 		std::cerr << __func__ << ": failed to get buffer handles for device"
 			<< std::endl;
@@ -213,6 +236,12 @@ void OpenClCollatingAndMeshingEngine::finalize()
 	// Release OpenCL buffers via smo hooks
 	if (smoHooksPtr && smoHooksPtr->ComputeManager_releaseUseHostPtrBuffer)
 	{
+		if (clAverageIntensityBufferClBuffer)
+		{
+			smoHooksPtr->ComputeManager_releaseUseHostPtrBuffer(
+				clAverageIntensityBufferClBuffer);
+			clAverageIntensityBufferClBuffer.reset();
+		}
 		if (clCollationBufferClBuffer)
 		{
 			smoHooksPtr->ComputeManager_releaseUseHostPtrBuffer(
@@ -229,6 +258,7 @@ void OpenClCollatingAndMeshingEngine::finalize()
 
 	// Reset cached cl_mem handles
 	clCollationBuffer = nullptr;
+	clAverageIntensityBuffer = nullptr;
 	clAssemblyBuffer = nullptr;
 
 	// Release kernels
@@ -272,6 +302,8 @@ void OpenClCollatingAndMeshingEngine::finalize()
 	assemblyBufferSize = 0;
 	collationBufferPtr = nullptr;
 	collationBufferSize = 0;
+	averageIntensityBufferPtr = nullptr;
+	averageIntensityBufferSize = 0;
 	frameAssemblyDesc = nullptr;
 }
 
@@ -367,7 +399,6 @@ bool OpenClCollatingAndMeshingEngine::startCompactKernel(
 }
 
 bool OpenClCollatingAndMeshingEngine::startCollateKernel(
-	StagingBuffer& assemblyBuff, StagingBuffer& collationBuff,
 	std::optional<std::reference_wrapper<StimulusFrame>> intensityStimFrame,
 	std::optional<std::reference_wrapper<StimulusFrame>> ambienceStimFrame,
 	collateKernelCbFn callback)
@@ -376,13 +407,18 @@ bool OpenClCollatingAndMeshingEngine::startCollateKernel(
 	collateKernelCb = std::move(callback);
 
 	// Validate buffers callable
-	auto validateBuffers = [this, &assemblyBuff, &collationBuff]() {
-		struct iovec assemblyIov = assemblyBuff.getClEngineIovec();
-		struct iovec collationIov = collationBuff.getClEngineIovec();
+	auto validateBuffers = [this]() {
+		struct iovec assemblyIov = parent.assemblyBuffer.getClEngineIovec();
+		struct iovec collationIov = parent.collationBuffer.getClEngineIovec();
+		struct iovec averageIntensityIov = parent.averageIntensityBuffer
+			.getClEngineIovec();
+
 		if (assemblyIov.iov_base != assemblyBufferPtr
 			|| assemblyIov.iov_len != assemblyBufferSize
 			|| collationIov.iov_base != collationBufferPtr
-			|| collationIov.iov_len != collationBufferSize)
+			|| collationIov.iov_len != collationBufferSize
+			|| averageIntensityIov.iov_base != averageIntensityBufferPtr
+			|| averageIntensityIov.iov_len != averageIntensityBufferSize)
 		{
 			throw std::runtime_error(
 				std::string(__func__) + ": buffer mismatch - buffers have changed");
@@ -390,12 +426,9 @@ bool OpenClCollatingAndMeshingEngine::startCollateKernel(
 	};
 
 	// Setup args callable
-	auto setupArgs = [
-		this, &assemblyBuff, intensityStimFrame, ambienceStimFrame
-		]()
+	auto setupArgs = [this, intensityStimFrame, ambienceStimFrame]()
 	{
-		return setupCollateDgramsArgs(
-			assemblyBuff, intensityStimFrame, ambienceStimFrame);
+		return setupCollateDgramsArgs(intensityStimFrame, ambienceStimFrame);
 	};
 
 	/**	EXPLANATION:
@@ -422,6 +455,14 @@ bool OpenClCollatingAndMeshingEngine::startCollateKernel(
 	}
 
 	unmapCollationBuffer();
+	if (!mapAverageIntensityBuffer(CL_MAP_WRITE))
+	{
+		std::cerr << __func__ << ": failed to map average intensity buffer"
+			<< std::endl;
+		return false;
+	}
+
+	unmapAverageIntensityBuffer();
 
 	// Map/unmap intensity buffer if it exists
 	if (intensityStimFrame.has_value())
@@ -446,29 +487,6 @@ bool OpenClCollatingAndMeshingEngine::startCollateKernel(
 		}
 	}
 
-	// Map/unmap ambience buffer if it exists
-	if (ambienceStimFrame.has_value())
-	{
-		StimulusFrame& ambienceFrame = ambienceStimFrame->get();
-		cl_mem ambienceClBuffer = ambienceFrame.clBuffer
-			->getAssociatedBufferHandleForDevice(computeDevice);
-
-		if (ambienceClBuffer)
-		{
-			void* mappedAmbienceBuffer = nullptr;
-			if (!mapBuffer(
-				ambienceClBuffer, ambienceFrame.slotDesc.nBytes,
-				CL_MAP_WRITE_INVALIDATE_REGION, mappedAmbienceBuffer))
-			{
-				std::cerr << __func__ << ": failed to map ambience buffer"
-					<< std::endl;
-				return false;
-			}
-
-			unmapBuffer(ambienceClBuffer, mappedAmbienceBuffer);
-		}
-	}
-
 	// Calculate global work size (just num slots in the frame)
 	size_t globalWorkSize = static_cast<uint32_t>(frameAssemblyDesc->numSlots);
 
@@ -623,12 +641,12 @@ bool OpenClCollatingAndMeshingEngine::setupSlotCompactorsArgs(
 }
 
 bool OpenClCollatingAndMeshingEngine::setupCollateDgramsArgs(
-	StagingBuffer& assemblyBuff,
 	std::optional<std::reference_wrapper<StimulusFrame>> intensityStimFrame,
 	std::optional<std::reference_wrapper<StimulusFrame>> ambienceStimFrame)
 {
 	// Extract parameters for collateDgrams kernel
-	uint32_t slotStride = static_cast<uint32_t>(assemblyBuff.slotStrideNBytes);
+	uint32_t slotStride = static_cast<uint32_t>(
+		parent.assemblyBuffer.slotStrideNBytes);
 
 	// Calculate nPointsPerSlot from device return mode
 	if (!parent.device)
@@ -675,14 +693,22 @@ bool OpenClCollatingAndMeshingEngine::setupCollateDgramsArgs(
 		return false;
 	}
 
-	// Set ambience buffer argument (arg 3)
-	cl_mem ambienceClBuffer = nullptr;
-	if (ambienceStimFrame.has_value())
-	{
-		ambienceClBuffer = ambienceStimFrame->get().clBuffer
-			->getAssociatedBufferHandleForDevice(computeDevice);
+	// Set average intensity buffer argument (arg 3)
+	/**	EXPLANATION:
+	 * We only pass the average intensity buffer argument to the collate kernel
+	 * when ambienceStimFrame is present. This is because the collate kernel
+	 * only needs the average intensity buffer if ambience processing is
+	 * requested (i.e., the ambience stimulus buffer is attached). If no
+	 * ambienceStimFrame is supplied, we skip passing the buffer to avoid
+	 * unnecessary work.
+	 */
+	cl_mem averageIntensityClBuffer = nullptr;
+	if (ambienceStimFrame.has_value()) {
+		averageIntensityClBuffer = clAverageIntensityBuffer;
 	}
-	err = clSetKernelArg(collateKernel, 3, sizeof(cl_mem), &ambienceClBuffer);
+	err = clSetKernelArg(
+		collateKernel, 3, sizeof(cl_mem), &averageIntensityClBuffer);
+
 	if (err != CL_SUCCESS)
 	{
 		std::cerr << __func__ << ": failed to set kernel arg 3: " << err
@@ -690,17 +716,7 @@ bool OpenClCollatingAndMeshingEngine::setupCollateDgramsArgs(
 		return false;
 	}
 
-	// Set ambienceHighVal argument (arg 4)
-	uint32_t ambienceHighVal = 0;
-	std::shared_ptr<PcloudAmbienceStimulusBuffer> ambienceBuff = nullptr;
-	if (ambienceStimFrame.has_value()
-		&& (ambienceBuff = parent.ambienceStimulusBuffer.load(
-			std::memory_order_acquire)))
-	{
-		ambienceHighVal = ambienceBuff->ambienceHighVal;
-	}
-
-	err = clSetKernelArg(collateKernel, 4, sizeof(uint32_t), &ambienceHighVal);
+	err = clSetKernelArg(collateKernel, 4, sizeof(uint32_t), &slotStride);
 	if (err != CL_SUCCESS)
 	{
 		std::cerr << __func__ << ": failed to set kernel arg 4: " << err
@@ -708,15 +724,15 @@ bool OpenClCollatingAndMeshingEngine::setupCollateDgramsArgs(
 		return false;
 	}
 
-	err = clSetKernelArg(collateKernel, 5, sizeof(uint32_t), &slotStride);
+	err = clSetKernelArg(collateKernel, 5, sizeof(uint32_t), &nPointsPerSlot);
 	if (err != CL_SUCCESS)
 	{
-		std::cerr << __func__ << ": failed to set kernel arg 3: " << err
+		std::cerr << __func__ << ": failed to set kernel arg 5: " << err
 			<< std::endl;
 		return false;
 	}
 
-	err = clSetKernelArg(collateKernel, 6, sizeof(uint32_t), &nPointsPerSlot);
+	err = clSetKernelArg(collateKernel, 6, sizeof(uint32_t), &nDgramsPerFrame);
 	if (err != CL_SUCCESS)
 	{
 		std::cerr << __func__ << ": failed to set kernel arg 6: " << err
@@ -724,14 +740,6 @@ bool OpenClCollatingAndMeshingEngine::setupCollateDgramsArgs(
 		return false;
 	}
 
-	err = clSetKernelArg(collateKernel, 7, sizeof(uint32_t), &nDgramsPerFrame);
-	if (err != CL_SUCCESS)
-	{
-		std::cerr << __func__ << ": failed to set kernel arg 7: " << err
-			<< std::endl;
-		return false;
-	}
-
 	return true;
 }
 
@@ -779,6 +787,8 @@ void OpenClCollatingAndMeshingEngine::collateKernelComplete(
 	std::optional<std::reference_wrapper<StimulusFrame>> ambienceStimFrame,
 	bool isFinalizing)
 {
+	(void)ambienceStimFrame;
+
 	cl_map_flags mapFlags;
 	/**	EXPLANATION:
 	 * Technically we should only need to do this if we plan to read the
@@ -791,6 +801,10 @@ void OpenClCollatingAndMeshingEngine::collateKernelComplete(
 		unmapCollationBuffer();
 	}
 
+	if (mapAverageIntensityBuffer(mapFlags)) {
+		unmapAverageIntensityBuffer();
+	}
+
 	// Map/unmap intensity buffer if it exists
 	if (intensityStimFrame.has_value())
 	{
@@ -810,25 +824,6 @@ void OpenClCollatingAndMeshingEngine::collateKernelComplete(
 		}
 	}
 
-	// Map/unmap ambience buffer if it exists
-	if (ambienceStimFrame.has_value())
-	{
-		StimulusFrame& ambienceFrame = ambienceStimFrame->get();
-		cl_mem ambienceClBuffer = ambienceFrame.clBuffer
-			->getAssociatedBufferHandleForDevice(computeDevice);
-
-		if (ambienceClBuffer)
-		{
-			void* mappedAmbienceBuffer = nullptr;
-			if (mapBuffer(
-				ambienceClBuffer, ambienceFrame.slotDesc.nBytes,
-				CL_MAP_READ, mappedAmbienceBuffer))
-			{
-				unmapBuffer(ambienceClBuffer, mappedAmbienceBuffer);
-			}
-		}
-	}
-
 	clFlush(computeDevice->commandQueue);
 
 	// Stop only collate kernel
@@ -952,6 +947,56 @@ bool OpenClCollatingAndMeshingEngine::unmapCollationBuffer()
 	return true;
 }
 
+bool OpenClCollatingAndMeshingEngine::mapAverageIntensityBuffer(
+	cl_map_flags mapFlags
+	)
+{
+	return mapBuffer(
+		clAverageIntensityBuffer, averageIntensityBufferSize, mapFlags,
+		mappedAverageIntensityBuffer);
+}
+
+bool OpenClCollatingAndMeshingEngine::unmapAverageIntensityBuffer()
+{
+	unmapBuffer(clAverageIntensityBuffer, mappedAverageIntensityBuffer);
+	mappedAverageIntensityBuffer = nullptr;
+	return true;
+}
+
+void OpenClCollatingAndMeshingEngine::produceAmbienceStimulusFrame(
+	std::optional<std::reference_wrapper<StimulusFrame>> ambienceStimFrame,
+	uint32_t nSucceeded)
+{
+	if (!ambienceStimFrame.has_value()) { return; }
+
+	std::shared_ptr<PcloudAmbienceStimulusBuffer> ambienceBuff =
+		parent.ambienceStimulusBuffer.load(std::memory_order_acquire);
+	if (!ambienceBuff) { return; }
+
+	uint32_t lowVal = ambienceBuff->ambienceIntensityLowVal;
+
+	// Read average intensity values from averageIntensityBuffer
+	float* averageIntensityAverages = reinterpret_cast<float*>(
+		averageIntensityBufferPtr);
+
+	// Count frames whose average intensity is <= lowVal (postrin only)
+	uint32_t ambienceCount = 0;
+	for (uint32_t i = 0; i < nSucceeded; ++i)
+	{
+		float avg = averageIntensityAverages[i];
+		if (avg <= static_cast<float>(lowVal))
+		{
+			++ambienceCount;
+		}
+	}
+
+	// Write the ambience count to the ambienceStimFrame
+	StimulusFrame& ambienceFrame = ambienceStimFrame->get();
+	uint32_t* ambienceValue = reinterpret_cast<uint32_t*>(
+		ambienceFrame.slotDesc.vaddr);
+	ambienceValue[0] = ambienceCount;
+}
+
 class OpenClCollatingAndMeshingEngine::CompactCollateAndMeshFrameReq
 :	public PostedAsynchronousContinuation<compactCollateAndMeshFrameReqCbFn>
 {
@@ -1070,7 +1115,6 @@ public:
 		engine.collateKernelStartTime = std::chrono::high_resolution_clock::now();
 
 		bool success = engine.startCollateKernel(
-			engine.parent.assemblyBuffer, engine.parent.collationBuffer,
 			context->intensityStimFrame, context->ambienceStimFrame,
 			std::bind(
 				&CompactCollateAndMeshFrameReq
@@ -1126,11 +1170,18 @@ public:
 			return;
 		}
 
+		uint32_t nSucceeded = context->frameAssemblyResult.nSucceeded.load();
+
+		// Produce ambience frame if ambience buffer is attached
+		if (context->ambienceStimFrame.has_value())
+		{
+			engine.produceAmbienceStimulusFrame(
+				context->ambienceStimFrame, nSucceeded);
+		}
+
 		int returnMode = static_cast<int>(engine.parent.device->currentReturnMode);
 		size_t pointsPerDgram = livoxProto1::Device::getNPointsPerDgram(
 			returnMode);
-
-		uint32_t nSucceeded = context->frameAssemblyResult.nSucceeded.load();
 		size_t totalPoints = nSucceeded * pointsPerDgram;
 
 		// Count points with intensity greater than 116
@@ -1149,16 +1200,40 @@ public:
 			}
 		}
 
-		// Sum up ambience counts from ambience buffer
-		uint32_t ambienceCountSum = 0;
+		// Print all averages above thresholds from average intensity buffer
 		if (context->ambienceStimFrame.has_value())
 		{
-			StimulusFrame& ambienceFrame = context->ambienceStimFrame->get();
-			uint32_t* ambienceCounts = reinterpret_cast<uint32_t*>(ambienceFrame.slotDesc.vaddr);
+			std::shared_ptr<PcloudAmbienceStimulusBuffer> ambienceBuff =
+				engine.parent.ambienceStimulusBuffer.load(std::memory_order_acquire);
+			uint32_t lowVal = ambienceBuff->ambienceIntensityLowVal;
+			uint32_t postrinThreshold = ambienceBuff->postrinInterestThreshold;
+
+			float* averageIntensityAverages = reinterpret_cast<float*>(
+				engine.averageIntensityBufferPtr);
+			// Count frames that met the postrin threshold
+			uint32_t framesMetThreshold = 0;
 			for (uint32_t i = 0; i < nSucceeded; ++i)
 			{
-				ambienceCountSum += ambienceCounts[i];
+				float avg = averageIntensityAverages[i];
+				if (avg <= static_cast<float>(lowVal))
+				{
+					++framesMetThreshold;
+				}
 			}
+
+			// Read the stimFrame value (ambience count)
+			StimulusFrame& ambienceFrame = context->ambienceStimFrame->get();
+			uint32_t* ambienceValue = reinterpret_cast<uint32_t*>(
+				ambienceFrame.slotDesc.vaddr);
+			uint32_t stimFrameValue = ambienceValue[0];
+
+			bool meetsPostrinThreshold = (framesMetThreshold >= postrinThreshold);
+
+			std::cout << __func__ << ": frames met threshold=" << framesMetThreshold
+				<< ", stimFrame value=" << stimFrameValue
+				<< ", postrin threshold=" << postrinThreshold
+				<< ", meets postrin=" << (meetsPostrinThreshold ? "yes" : "no")
+				<< std::endl;
 		}
 
 		std::cout << __func__ << ": intensityRingBufferIndex="
@@ -1170,8 +1245,7 @@ public:
 			<< ", pointsPerDgram=" << pointsPerDgram
 			<< ", nSucceeded=" << nSucceeded
 			<< ", totalPoints=" << totalPoints
-			<< ", highIntensityCount=" << highIntensityCount
-			<< ", ambienceCountSum=" << ambienceCountSum << std::endl;
+			<< ", highIntensityCount=" << highIntensityCount << std::endl;
 
 		callOriginalCallback(success);
 	}

stimBuffApis/livoxGen1/openClCollatingAndMeshingEngine.h

@@ -64,7 +64,6 @@ private:
 		StagingBuffer& assemblyBuff, uint32_t nSucceeded,
 		compactKernelCbFn callback);
 	bool startCollateKernel(
-		StagingBuffer& assemblyBuff, StagingBuffer& collationBuff,
 		std::optional<std::reference_wrapper<StimulusFrame>> intensityStimFrame,
 		std::optional<std::reference_wrapper<StimulusFrame>> ambienceStimFrame,
 		collateKernelCbFn callback);
@@ -81,6 +80,11 @@ public:
 	std::chrono::milliseconds getCompactKernelDuration() const;
 	std::chrono::milliseconds getCollateKernelDuration() const;
 
+	// Produce ambience frame from average intensity data
+	void produceAmbienceStimulusFrame(
+		std::optional<std::reference_wrapper<StimulusFrame>> ambienceStimFrame,
+		uint32_t nSucceeded);
+
 private:
 	PcloudStimulusProducer& parent;
 
@@ -94,9 +98,11 @@ private:
 	// OpenCL buffers (managed by ComputeManager)
 	std::shared_ptr<smo::compute::ClBuffer> clAssemblyBufferClBuffer;
 	std::shared_ptr<smo::compute::ClBuffer> clCollationBufferClBuffer;
+	std::shared_ptr<smo::compute::ClBuffer> clAverageIntensityBufferClBuffer;
 	// Cached cl_mem handles for the device we're using
 	cl_mem clAssemblyBuffer;
 	cl_mem clCollationBuffer;
+	cl_mem clAverageIntensityBuffer;
 
 	// State tracking
 	SpinLock shouldAcceptRequestsLock;
@@ -111,9 +117,12 @@ private:
 	size_t assemblyBufferSize;
 	void* collationBufferPtr;
 	size_t collationBufferSize;
+	void* averageIntensityBufferPtr;
+	size_t averageIntensityBufferSize;
 	// Mapped buffer pointers (for zero-copy synchronization)
 	void* mappedAssemblyBuffer;
 	void* mappedCollationBuffer;
+	void* mappedAverageIntensityBuffer;
 
 	// Frame descriptor (cached from setup)
 	std::shared_ptr<FrameAssemblyDesc> frameAssemblyDesc;
@@ -142,7 +151,6 @@ private:
 	bool setupSlotCompactorsArgs(
 		StagingBuffer& assemblyBuff, uint32_t nSucceeded);
 	bool setupCollateDgramsArgs(
-		StagingBuffer& assemblyBuff,
 		std::optional<std::reference_wrapper<StimulusFrame>> intensityStimFrame,
 		std::optional<std::reference_wrapper<StimulusFrame>> ambienceStimFrame);
 
@@ -156,6 +164,8 @@ private:
 	bool unmapAssemblyBuffer();
 	bool mapCollationBuffer(cl_map_flags mapFlags = CL_MAP_READ);
 	bool unmapCollationBuffer();
+	bool mapAverageIntensityBuffer(cl_map_flags mapFlags = CL_MAP_READ);
+	bool unmapAverageIntensityBuffer();
 
 	// Forward declaration for continuation class
 	class CompactCollateAndMeshFrameReq;

stimBuffApis/livoxGen1/pcloudAmbienceStimulusBuffer.h

@@ -5,8 +5,11 @@
 #include <cstdint>
 #include <list>
 #include <cstddef>
+#include <vector>
+#include <string>
 #include <user/stimulusBuffer.h>
 #include <user/stagingBuffer.h>
+#include <user/deviceAttachmentSpec.h>
 #include "lg1PcloudAmbienceStencil.h"
 
 namespace smo {
@@ -30,15 +33,52 @@ public:
 		const StagingBuffer::IOEngineConstraints& outputEngineConstraints,
 		const SmoCallbacks& callbacks,
 		cl_mem_flags flags,
-		uint32_t ambienceHighVal_,
 		size_t nStencils_, size_t nDgramsPerFrame_)
 	:	StimulusBuffer(
 			parent, deviceAttachmentSpec, histbuffMs,
 			inputEngineConstraints, outputEngineConstraints,
 			callbacks, flags),
-	ambienceHighVal(ambienceHighVal_),
 	nStencils(nStencils_)
 	{
+		// Parse postrinInterestThreshold from qualeIfaceApiParams
+		const std::vector<std::string> postrinInterestThresholdParamNames = {
+			"postrin-interest-threshold",
+			"postrin-interest"
+		};
+
+		/**	EXPLANATION:
+		 * The default postrin threshold is determined as follows:
+		 * We want 90% of the scanned points to have intensity under the
+		 * lowVal to consider the ambience criterion met.
+		 *
+		 * If nDgramsPerFrame_ (the number of datagrams per frame) is less
+		 * than 10, we require that all frames (100%) be under the lowVal.
+		 * This is because, for such small sample sizes, calculating 90%
+		 * does not yield a meaningful integer; for example, 90% of 7 is
+		 * 6.3, but we must count whole frames that meet the threshold. By
+		 * using nDgramsPerFrame_ as the threshold in this case, we ensure
+		 * logical, all-or-nothing evaluation at low sample counts while
+		 * maintaining an approximate 90% requirement for larger frame
+		 * sizes.
+		 */
+		uint32_t defaultPostrinThreshold = (nDgramsPerFrame_ < 10)
+			? static_cast<uint32_t>(nDgramsPerFrame_)
+			: static_cast<uint32_t>(nDgramsPerFrame_ * 9 / 10);
+
+		postrinInterestThreshold = static_cast<uint32_t>(
+			device::DeviceAttachmentSpec::parseOptionalParamAsIntWithSynonyms(
+				deviceAttachmentSpec->qualeIfaceApiParams,
+				postrinInterestThresholdParamNames,
+				defaultPostrinThreshold));
+		// Parse ambienceIntensityLowVal from qualeIfaceApiParams
+		const std::vector<std::string> ambienceIntensityLowValParamNames = {
+			"ambience-intensity-low-val"
+		};
+		ambienceIntensityLowVal = static_cast<uint32_t>(
+			device::DeviceAttachmentSpec::parseOptionalParamAsIntWithSynonyms(
+				deviceAttachmentSpec->qualeIfaceApiParams,
+				ambienceIntensityLowValParamNames, 8));
+
 		// Construct stencils and add to list (FIFO behavior)
 		for (size_t i = 0; i < nStencils; ++i) {
 			stencils.emplace_back(nDgramsPerFrame_);
@@ -54,7 +94,8 @@ public:
 	PcloudAmbienceStimulusBuffer& operator=(PcloudAmbienceStimulusBuffer&&) = default;
 
 public:
-	uint32_t ambienceHighVal;
+	uint32_t postrinInterestThreshold;
+	uint32_t ambienceIntensityLowVal;
 	size_t nStencils;
 	std::list<LG1PcloudAmbienceStencil> stencils;
 };

stimBuffApis/livoxGen1/pcloudStimulusProducer.cpp

@@ -55,13 +55,23 @@ static StagingBuffer::IOEngineConstraints openClIntensityInputConstraints(
 	static_cast<size_t>(sysconf(_SC_PAGE_SIZE)));
 
 static StagingBuffer::IOEngineConstraints openClAmbienceInputConstraints(
-	// slotStartAlignmentByteVal (page alignment)
-	static_cast<size_t>(sysconf(_SC_PAGE_SIZE)),
-	// slotPadToNBytes: This is dynamically calculated based on the return mode.
+	// slotStartAlignmentByteVal (sizeof(void*))
+	sizeof(void*),
+	// slotPadToNBytes (sizeof(PcloudAmbienceStimulusValue))
 	sizeof(PcloudAmbienceStencil::PcloudAmbienceStimulusValue),
 	// frameStartAlignmentByteVal (page alignment)
 	static_cast<size_t>(sysconf(_SC_PAGE_SIZE)),
-	// framePadToNBytes (pointer size)
+	// framePadToNBytes (page alignment)
+	static_cast<size_t>(sysconf(_SC_PAGE_SIZE)));
+
+static StagingBuffer::IOEngineConstraints openClAverageIntensityConstraints(
+	// slotStartAlignmentByteVal (sizeof(float))
+	sizeof(float),
+	// slotPadToNBytes (sizeof(float))
+	sizeof(float),
+	// frameStartAlignmentByteVal (page alignment)
+	static_cast<size_t>(sysconf(_SC_PAGE_SIZE)),
+	// framePadToNBytes (page alignment)
 	static_cast<size_t>(sysconf(_SC_PAGE_SIZE)));
 
 PcloudStimulusProducer::PcloudStimulusProducer(
@@ -85,6 +95,10 @@ collationBuffer(
 	StagingBuffer::IOEngineConstraints::openClInputConstraints,
 	StagingBuffer::IOEngineConstraints::openClInputConstraints,
 	nDgramsPerStagingBufferFrame),
+averageIntensityBuffer(
+	openClAverageIntensityConstraints,
+	openClAverageIntensityConstraints,
+	nDgramsPerStagingBufferFrame),
 tempStimulusFrameMem(0),
 tempStimulusFrame(
 	FrameAssemblyDesc::SlotDesc{
@@ -310,19 +324,6 @@ std::cout << __func__ << ": $$$$$$$ Created PcloudIntensityStimulusBuffer" << st
 	}
 	else if (qualeIfaceApi == "pcloudAmbience")
 	{
-		// Parse ambienceHighVal from qualeIfaceApiParams (temporary, undocumented)
-		const std::vector<std::string> ambienceHighValParamNames = {
-			"negtrin-intolerable-threshold",
-			"negtrin-intolerable",
-			"intolerable-threshold",
-			"intolerable"
-		};
-		int ambienceHighValInt = device::DeviceAttachmentSpec
-			::parseOptionalParamAsIntWithSynonyms(
-				deviceAttachmentSpec->qualeIfaceApiParams,
-				ambienceHighValParamNames, 116);
-		uint32_t ambienceHighVal = static_cast<uint32_t>(ambienceHighValInt);
-
 		// Parse n-stencils from qualeIfaceApiParams
 		const std::vector<std::string> nStencilsParamNames = {
 			"n-stencils"
@@ -333,18 +334,10 @@ std::cout << __func__ << ": $$$$$$$ Created PcloudIntensityStimulusBuffer" << st
 				nStencilsParamNames, 1);
 		size_t nStencils = static_cast<size_t>(nStencilsInt);
 
-		/* Calculate slotStrideNBytes:
-		 * nDgramsPerStagingBufferFrame * sizeof(uint32_t)
-		 */
-		size_t slotStrideNBytes = this->nDgramsPerStagingBufferFrame
-			* sizeof(uint32_t);
-		// Reuse openClAmbienceInputConstraints, only modify slotPadToNBytes
-		openClAmbienceInputConstraints.slotPadToNBytes = slotStrideNBytes;
-
 		auto ambienceBuffer = std::make_shared<PcloudAmbienceStimulusBuffer>(
 			*this, deviceAttachmentSpec, histbuffMs,
 			openClAmbienceInputConstraints, openClAmbienceInputConstraints,
-			*smoHooksPtr, CL_MEM_READ_WRITE, ambienceHighVal,
+			*smoHooksPtr, CL_MEM_READ_WRITE,
 			nStencils, this->nDgramsPerStagingBufferFrame);
 
 std::cout << __func__ << ": $$$$$$$ Created PcloudAmbienceStimulusBuffer" << std::endl;

stimBuffApis/livoxGen1/pcloudStimulusProducer.h

@@ -89,6 +89,7 @@ public:
 	StagingBuffer assemblyBuffer;
 	IoUringAssemblyEngine ioUringAssemblyEngine;
 	StagingBuffer collationBuffer;
+	StagingBuffer averageIntensityBuffer;
 	size_t tempStimulusFrameMem;
 	StimulusFrame tempStimulusFrame;
 	std::atomic<std::shared_ptr<MeshStimulusBuffer>> meshStimulusBuffer;