diff --git a/docs/livox-gen1-lidar-dap-spec.md b/docs/livox-gen1-lidar-dap-spec.md index 7331c9b..0c24d7a 100644 --- a/docs/livox-gen1-lidar-dap-spec.md +++ b/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) diff --git a/stimBuffApis/livoxGen1/collateDgrams.cl b/stimBuffApis/livoxGen1/collateDgrams.cl index 64ea801..0379a47 100644 --- a/stimBuffApis/livoxGen1/collateDgrams.cl +++ b/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; } } diff --git a/stimBuffApis/livoxGen1/openClCollatingAndMeshingEngine.cpp b/stimBuffApis/livoxGen1/openClCollatingAndMeshingEngine.cpp index f32242d..19b9699 100644 --- a/stimBuffApis/livoxGen1/openClCollatingAndMeshingEngine.cpp +++ b/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>( @@ -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> intensityStimFrame, std::optional> 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(frameAssemblyDesc->numSlots); @@ -623,12 +641,12 @@ bool OpenClCollatingAndMeshingEngine::setupSlotCompactorsArgs( } bool OpenClCollatingAndMeshingEngine::setupCollateDgramsArgs( - StagingBuffer& assemblyBuff, std::optional> intensityStimFrame, std::optional> ambienceStimFrame) { // Extract parameters for collateDgrams kernel - uint32_t slotStride = static_cast(assemblyBuff.slotStrideNBytes); + uint32_t slotStride = static_cast( + 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 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> 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> ambienceStimFrame, + uint32_t nSucceeded) +{ + if (!ambienceStimFrame.has_value()) { return; } + + std::shared_ptr 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( + 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(lowVal)) + { + ++ambienceCount; + } + } + + // Write the ambience count to the ambienceStimFrame + StimulusFrame& ambienceFrame = ambienceStimFrame->get(); + uint32_t* ambienceValue = reinterpret_cast( + ambienceFrame.slotDesc.vaddr); + ambienceValue[0] = ambienceCount; +} + class OpenClCollatingAndMeshingEngine::CompactCollateAndMeshFrameReq : public PostedAsynchronousContinuation { @@ -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(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(ambienceFrame.slotDesc.vaddr); + std::shared_ptr ambienceBuff = + engine.parent.ambienceStimulusBuffer.load(std::memory_order_acquire); + uint32_t lowVal = ambienceBuff->ambienceIntensityLowVal; + uint32_t postrinThreshold = ambienceBuff->postrinInterestThreshold; + + float* averageIntensityAverages = reinterpret_cast( + 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(lowVal)) + { + ++framesMetThreshold; + } } + + // Read the stimFrame value (ambience count) + StimulusFrame& ambienceFrame = context->ambienceStimFrame->get(); + uint32_t* ambienceValue = reinterpret_cast( + 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); } diff --git a/stimBuffApis/livoxGen1/openClCollatingAndMeshingEngine.h b/stimBuffApis/livoxGen1/openClCollatingAndMeshingEngine.h index 410f92d..677809f 100644 --- a/stimBuffApis/livoxGen1/openClCollatingAndMeshingEngine.h +++ b/stimBuffApis/livoxGen1/openClCollatingAndMeshingEngine.h @@ -64,7 +64,6 @@ private: StagingBuffer& assemblyBuff, uint32_t nSucceeded, compactKernelCbFn callback); bool startCollateKernel( - StagingBuffer& assemblyBuff, StagingBuffer& collationBuff, std::optional> intensityStimFrame, std::optional> 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> ambienceStimFrame, + uint32_t nSucceeded); + private: PcloudStimulusProducer& parent; @@ -94,9 +98,11 @@ private: // OpenCL buffers (managed by ComputeManager) std::shared_ptr clAssemblyBufferClBuffer; std::shared_ptr clCollationBufferClBuffer; + std::shared_ptr 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; @@ -142,7 +151,6 @@ private: bool setupSlotCompactorsArgs( StagingBuffer& assemblyBuff, uint32_t nSucceeded); bool setupCollateDgramsArgs( - StagingBuffer& assemblyBuff, std::optional> intensityStimFrame, std::optional> 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; diff --git a/stimBuffApis/livoxGen1/pcloudAmbienceStimulusBuffer.h b/stimBuffApis/livoxGen1/pcloudAmbienceStimulusBuffer.h index c198a06..ecbc3bb 100644 --- a/stimBuffApis/livoxGen1/pcloudAmbienceStimulusBuffer.h +++ b/stimBuffApis/livoxGen1/pcloudAmbienceStimulusBuffer.h @@ -5,8 +5,11 @@ #include #include #include +#include +#include #include #include +#include #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 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(nDgramsPerFrame_) + : static_cast(nDgramsPerFrame_ * 9 / 10); + + postrinInterestThreshold = static_cast( + device::DeviceAttachmentSpec::parseOptionalParamAsIntWithSynonyms( + deviceAttachmentSpec->qualeIfaceApiParams, + postrinInterestThresholdParamNames, + defaultPostrinThreshold)); + // Parse ambienceIntensityLowVal from qualeIfaceApiParams + const std::vector ambienceIntensityLowValParamNames = { + "ambience-intensity-low-val" + }; + ambienceIntensityLowVal = static_cast( + 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 stencils; }; diff --git a/stimBuffApis/livoxGen1/pcloudStimulusProducer.cpp b/stimBuffApis/livoxGen1/pcloudStimulusProducer.cpp index f2cd8b1..f5a6d1a 100644 --- a/stimBuffApis/livoxGen1/pcloudStimulusProducer.cpp +++ b/stimBuffApis/livoxGen1/pcloudStimulusProducer.cpp @@ -55,13 +55,23 @@ static StagingBuffer::IOEngineConstraints openClIntensityInputConstraints( static_cast(sysconf(_SC_PAGE_SIZE))); static StagingBuffer::IOEngineConstraints openClAmbienceInputConstraints( - // slotStartAlignmentByteVal (page alignment) - static_cast(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(sysconf(_SC_PAGE_SIZE)), - // framePadToNBytes (pointer size) + // framePadToNBytes (page alignment) + static_cast(sysconf(_SC_PAGE_SIZE))); + +static StagingBuffer::IOEngineConstraints openClAverageIntensityConstraints( + // slotStartAlignmentByteVal (sizeof(float)) + sizeof(float), + // slotPadToNBytes (sizeof(float)) + sizeof(float), + // frameStartAlignmentByteVal (page alignment) + static_cast(sysconf(_SC_PAGE_SIZE)), + // framePadToNBytes (page alignment) static_cast(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 ambienceHighValParamNames = { - "negtrin-intolerable-threshold", - "negtrin-intolerable", - "intolerable-threshold", - "intolerable" - }; - int ambienceHighValInt = device::DeviceAttachmentSpec - ::parseOptionalParamAsIntWithSynonyms( - deviceAttachmentSpec->qualeIfaceApiParams, - ambienceHighValParamNames, 116); - uint32_t ambienceHighVal = static_cast(ambienceHighValInt); - // Parse n-stencils from qualeIfaceApiParams const std::vector nStencilsParamNames = { "n-stencils" @@ -333,18 +334,10 @@ std::cout << __func__ << ": $$$$$$$ Created PcloudIntensityStimulusBuffer" << st nStencilsParamNames, 1); size_t nStencils = static_cast(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( *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; diff --git a/stimBuffApis/livoxGen1/pcloudStimulusProducer.h b/stimBuffApis/livoxGen1/pcloudStimulusProducer.h index 103a2a1..7581fee 100644 --- a/stimBuffApis/livoxGen1/pcloudStimulusProducer.h +++ b/stimBuffApis/livoxGen1/pcloudStimulusProducer.h @@ -89,6 +89,7 @@ public: StagingBuffer assemblyBuffer; IoUringAssemblyEngine ioUringAssemblyEngine; StagingBuffer collationBuffer; + StagingBuffer averageIntensityBuffer; size_t tempStimulusFrameMem; StimulusFrame tempStimulusFrame; std::atomic> meshStimulusBuffer;