Skip to content

Commit 6e50610

Browse files
Copilotchrxh
andcommitted
Add GeometryTests for signal restriction modes in cudaExtractSelectedObjectData and cudaExtractSelectedConnectionData
Co-authored-by: chrxh <[email protected]>
1 parent 0f4831c commit 6e50610

File tree

10 files changed

+262
-0
lines changed

10 files changed

+262
-0
lines changed

source/EngineGpuKernels/GeometryKernelsService.cu

Lines changed: 62 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -207,3 +207,65 @@ void GeometryKernelsService::extractObjectData(
207207
KERNEL_CALL(cudaExtractDetonationEventData, data, renderingData.deviceDetonationEventBuffer, _numDetonationEventVertices);
208208
}
209209
}
210+
211+
std::vector<SelectedObjectVertexData> GeometryKernelsService::testOnly_getSelectedObjectData(SettingsForSimulation const& settings, SimulationData data)
212+
{
213+
auto const& gpuSettings = settings.cudaSettings;
214+
215+
// First count how many selected objects
216+
setValueToDevice(_numSelectedObjects, static_cast<uint64_t>(0));
217+
KERNEL_CALL(cudaExtractSelectedObjectData, data, nullptr, _numSelectedObjects);
218+
cudaDeviceSynchronize();
219+
auto numObjects = copyToHost(_numSelectedObjects);
220+
221+
if (numObjects == 0) {
222+
return {};
223+
}
224+
225+
// Allocate device memory and extract data
226+
SelectedObjectVertexData* deviceBuffer;
227+
CHECK_FOR_CUDA_ERROR(cudaMalloc(&deviceBuffer, numObjects * sizeof(SelectedObjectVertexData)));
228+
229+
setValueToDevice(_numSelectedObjects, static_cast<uint64_t>(0));
230+
KERNEL_CALL(cudaExtractSelectedObjectData, data, deviceBuffer, _numSelectedObjects);
231+
cudaDeviceSynchronize();
232+
233+
// Copy to host
234+
std::vector<SelectedObjectVertexData> result(numObjects);
235+
CHECK_FOR_CUDA_ERROR(cudaMemcpy(result.data(), deviceBuffer, numObjects * sizeof(SelectedObjectVertexData), cudaMemcpyDeviceToHost));
236+
237+
CHECK_FOR_CUDA_ERROR(cudaFree(deviceBuffer));
238+
239+
return result;
240+
}
241+
242+
std::vector<ConnectionArrowVertexData> GeometryKernelsService::testOnly_getConnectionArrowData(SettingsForSimulation const& settings, SimulationData data)
243+
{
244+
auto const& gpuSettings = settings.cudaSettings;
245+
246+
// First count how many connection arrow vertices
247+
setValueToDevice(_numSelectedConnectionVertices, static_cast<uint64_t>(0));
248+
KERNEL_CALL(cudaExtractSelectedConnectionData, data, nullptr, _numSelectedConnectionVertices);
249+
cudaDeviceSynchronize();
250+
auto numVertices = copyToHost(_numSelectedConnectionVertices);
251+
252+
if (numVertices == 0) {
253+
return {};
254+
}
255+
256+
// Allocate device memory and extract data
257+
ConnectionArrowVertexData* deviceBuffer;
258+
CHECK_FOR_CUDA_ERROR(cudaMalloc(&deviceBuffer, numVertices * sizeof(ConnectionArrowVertexData)));
259+
260+
setValueToDevice(_numSelectedConnectionVertices, static_cast<uint64_t>(0));
261+
KERNEL_CALL(cudaExtractSelectedConnectionData, data, deviceBuffer, _numSelectedConnectionVertices);
262+
cudaDeviceSynchronize();
263+
264+
// Copy to host
265+
std::vector<ConnectionArrowVertexData> result(numVertices);
266+
CHECK_FOR_CUDA_ERROR(cudaMemcpy(result.data(), deviceBuffer, numVertices * sizeof(ConnectionArrowVertexData), cudaMemcpyDeviceToHost));
267+
268+
CHECK_FOR_CUDA_ERROR(cudaFree(deviceBuffer));
269+
270+
return result;
271+
}

source/EngineGpuKernels/GeometryKernelsService.cuh

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,9 @@ public:
2222
NumRenderObjects getNumRenderObjects(SettingsForSimulation const& settings, SimulationData data, RealRect const& visibleWorldRect);
2323
void extractObjectData(SettingsForSimulation const& settings, SimulationData data, CudaGeometryBuffers& renderingData, RealRect const& visibleWorldRect);
2424

25+
std::vector<SelectedObjectVertexData> testOnly_getSelectedObjectData(SettingsForSimulation const& settings, SimulationData data);
26+
std::vector<ConnectionArrowVertexData> testOnly_getConnectionArrowData(SettingsForSimulation const& settings, SimulationData data);
27+
2528
private:
2629
GeometryKernelsService() = default;
2730

source/EngineGpuKernels/SimulationCudaFacade.cu

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -672,6 +672,24 @@ NumRenderObjects _SimulationCudaFacade::testOnly_getNumRenderObjects()
672672
return result;
673673
}
674674

675+
std::vector<SelectedObjectVertexData> _SimulationCudaFacade::testOnly_getSelectedObjectData()
676+
{
677+
checkAndProcessSimulationParameterChanges();
678+
auto simulationData = getSimulationDataPtrCopy();
679+
auto result = GeometryKernelsService::get().testOnly_getSelectedObjectData(_settings, simulationData);
680+
syncAndCheck();
681+
return result;
682+
}
683+
684+
std::vector<ConnectionArrowVertexData> _SimulationCudaFacade::testOnly_getConnectionArrowData()
685+
{
686+
checkAndProcessSimulationParameterChanges();
687+
auto simulationData = getSimulationDataPtrCopy();
688+
auto result = GeometryKernelsService::get().testOnly_getConnectionArrowData(_settings, simulationData);
689+
syncAndCheck();
690+
return result;
691+
}
692+
675693
void _SimulationCudaFacade::initCuda()
676694
{
677695
log(Priority::Important, "initialize CUDA");

source/EngineGpuKernels/SimulationCudaFacade.cuh

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -113,6 +113,8 @@ public:
113113
void testOnly_resizeArrays(ArraySizesForGpu const& sizeDelta);
114114
bool testOnly_arePointersValid();
115115
NumRenderObjects testOnly_getNumRenderObjects();
116+
std::vector<SelectedObjectVertexData> testOnly_getSelectedObjectData();
117+
std::vector<ConnectionArrowVertexData> testOnly_getConnectionArrowData();
116118

117119
private:
118120
void initCuda();

source/EngineImpl/EngineWorker.cpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -508,6 +508,18 @@ NumRenderObjects EngineWorker::testOnly_getNumRenderObjects()
508508
return _simulationCudaFacade->testOnly_getNumRenderObjects();
509509
}
510510

511+
std::vector<SelectedObjectVertexData> EngineWorker::testOnly_getSelectedObjectData()
512+
{
513+
EngineWorkerGuard access(this);
514+
return _simulationCudaFacade->testOnly_getSelectedObjectData();
515+
}
516+
517+
std::vector<ConnectionArrowVertexData> EngineWorker::testOnly_getConnectionArrowData()
518+
{
519+
EngineWorkerGuard access(this);
520+
return _simulationCudaFacade->testOnly_getConnectionArrowData();
521+
}
522+
511523
void EngineWorker::resetTimeIntervalStatistics()
512524
{
513525
_simulationCudaFacade->resetTimeIntervalStatistics();

source/EngineImpl/EngineWorker.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -127,6 +127,8 @@ class EngineWorker
127127
void testOnly_resizeArrays(ArraySizesForGpu const& sizeDelta);
128128
bool testOnly_arePointersValid();
129129
NumRenderObjects testOnly_getNumRenderObjects();
130+
std::vector<SelectedObjectVertexData> testOnly_getSelectedObjectData();
131+
std::vector<ConnectionArrowVertexData> testOnly_getConnectionArrowData();
130132

131133
private:
132134
void resetTimeIntervalStatistics();

source/EngineImpl/SimulationFacadeImpl.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -416,3 +416,13 @@ NumRenderObjects _SimulationFacadeImpl::testOnly_getNumRenderObjects()
416416
{
417417
return _worker.testOnly_getNumRenderObjects();
418418
}
419+
420+
std::vector<SelectedObjectVertexData> _SimulationFacadeImpl::testOnly_getSelectedObjectData()
421+
{
422+
return _worker.testOnly_getSelectedObjectData();
423+
}
424+
425+
std::vector<ConnectionArrowVertexData> _SimulationFacadeImpl::testOnly_getConnectionArrowData()
426+
{
427+
return _worker.testOnly_getConnectionArrowData();
428+
}

source/EngineImpl/SimulationFacadeImpl.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -114,6 +114,8 @@ class _SimulationFacadeImpl : public _SimulationFacade
114114
void testOnly_resizeArrays(ArraySizesForGpu const& sizeDelta) override;
115115
bool testOnly_arePointersValid() override;
116116
NumRenderObjects testOnly_getNumRenderObjects() override;
117+
std::vector<SelectedObjectVertexData> testOnly_getSelectedObjectData() override;
118+
std::vector<ConnectionArrowVertexData> testOnly_getConnectionArrowData() override;
117119

118120
private:
119121
bool _selectionNeedsUpdate = false;

source/EngineInterface/SimulationFacade.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -130,6 +130,8 @@ class _SimulationFacade
130130
virtual void testOnly_resizeArrays(ArraySizesForGpu const& sizeDelta) = 0;
131131
virtual bool testOnly_arePointersValid() = 0;
132132
virtual NumRenderObjects testOnly_getNumRenderObjects() = 0;
133+
virtual std::vector<SelectedObjectVertexData> testOnly_getSelectedObjectData() = 0;
134+
virtual std::vector<ConnectionArrowVertexData> testOnly_getConnectionArrowData() = 0;
133135

134136
protected:
135137
static SimulationFacade _instance;

source/EngineTests/GeometryTests.cpp

Lines changed: 149 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -268,3 +268,152 @@ TEST_F(GeometryTests, copyBuffers_mixedCellsAndParticles)
268268
EXPECT_EQ(2u, numObjects.cells);
269269
EXPECT_EQ(3u, numObjects.energyParticles);
270270
}
271+
272+
// Signal restriction tests for cudaExtractSelectedObjectData
273+
274+
TEST_F(GeometryTests, selectedObjectData_noRestriction_inactive)
275+
{
276+
auto cell = CellDescription().id(1).pos({100.0f, 100.0f});
277+
cell._signalRestriction._mode = SignalRestrictionMode_Inactive;
278+
cell._signalRestriction._baseAngle = 45.0f;
279+
cell._signalRestriction._openingAngle = 90.0f;
280+
281+
auto data = Description().cells({
282+
cell,
283+
CellDescription().id(2).pos({101.0f, 100.0f}),
284+
});
285+
data.addConnection(1, 2);
286+
_simulationFacade->setSimulationData(data);
287+
288+
// Select cell 1 using position-based selection
289+
_simulationFacade->setSelection({99.0f, 99.0f}, {100.5f, 101.0f});
290+
291+
auto selectedData = _simulationFacade->testOnly_getSelectedObjectData();
292+
ASSERT_EQ(1u, selectedData.size());
293+
EXPECT_EQ(0, selectedData[0].hasSignalRestriction); // Inactive mode = no restriction
294+
}
295+
296+
TEST_F(GeometryTests, selectedObjectData_hasRestriction_active)
297+
{
298+
auto cell = CellDescription().id(1).pos({100.0f, 100.0f});
299+
cell._signalRestriction._mode = SignalRestrictionMode_Active;
300+
cell._signalRestriction._baseAngle = 45.0f;
301+
cell._signalRestriction._openingAngle = 90.0f;
302+
303+
auto data = Description().cells({
304+
cell,
305+
CellDescription().id(2).pos({101.0f, 100.0f}),
306+
});
307+
data.addConnection(1, 2);
308+
_simulationFacade->setSimulationData(data);
309+
310+
// Select cell 1 using position-based selection
311+
_simulationFacade->setSelection({99.0f, 99.0f}, {100.5f, 101.0f});
312+
313+
auto selectedData = _simulationFacade->testOnly_getSelectedObjectData();
314+
ASSERT_EQ(1u, selectedData.size());
315+
EXPECT_EQ(1, selectedData[0].hasSignalRestriction); // Active mode = has restriction
316+
}
317+
318+
TEST_F(GeometryTests, selectedObjectData_hasRestriction_conditional)
319+
{
320+
auto cell = CellDescription().id(1).pos({100.0f, 100.0f});
321+
cell._signalRestriction._mode = SignalRestrictionMode_Conditional;
322+
cell._signalRestriction._baseAngle = 45.0f;
323+
cell._signalRestriction._openingAngle = 90.0f;
324+
325+
auto data = Description().cells({
326+
cell,
327+
CellDescription().id(2).pos({101.0f, 100.0f}),
328+
});
329+
data.addConnection(1, 2);
330+
_simulationFacade->setSimulationData(data);
331+
332+
// Select cell 1 using position-based selection
333+
_simulationFacade->setSelection({99.0f, 99.0f}, {100.5f, 101.0f});
334+
335+
auto selectedData = _simulationFacade->testOnly_getSelectedObjectData();
336+
ASSERT_EQ(1u, selectedData.size());
337+
EXPECT_EQ(1, selectedData[0].hasSignalRestriction); // Conditional mode = has restriction
338+
}
339+
340+
// Signal restriction tests for cudaExtractSelectedConnectionData
341+
342+
TEST_F(GeometryTests, connectionData_noRestriction_inactive_bothDirections)
343+
{
344+
auto cell1 = CellDescription().id(1).pos({100.0f, 100.0f});
345+
cell1._signalRestriction._mode = SignalRestrictionMode_Inactive;
346+
347+
auto cell2 = CellDescription().id(2).pos({101.0f, 100.0f});
348+
cell2._signalRestriction._mode = SignalRestrictionMode_Inactive;
349+
350+
auto data = Description().cells({cell1, cell2});
351+
data.addConnection(1, 2);
352+
_simulationFacade->setSimulationData(data);
353+
354+
// Select both cells using position-based selection
355+
_simulationFacade->setSelection({99.0f, 99.0f}, {102.0f, 101.0f});
356+
357+
auto connectionData = _simulationFacade->testOnly_getConnectionArrowData();
358+
ASSERT_EQ(2u, connectionData.size()); // 2 vertices per connection line
359+
// arrowFlags: bit 0 = arrow to cell1, bit 1 = arrow to cell2
360+
// Both cells have no restriction, so signals can flow both ways (flags = 3)
361+
EXPECT_EQ(3, connectionData[0].arrowFlags);
362+
EXPECT_EQ(3, connectionData[1].arrowFlags);
363+
}
364+
365+
TEST_F(GeometryTests, connectionData_withRestriction_active_restrictedDirection)
366+
{
367+
auto cell1 = CellDescription().id(1).pos({100.0f, 100.0f});
368+
cell1._signalRestriction._mode = SignalRestrictionMode_Active;
369+
// Use baseAngle = 90 and openingAngle = 90 to point away from connection
370+
// Connection angle is 0 (first connection), so range [45+180, 135+180] = [225, 315] doesn't include 0
371+
cell1._signalRestriction._baseAngle = 90.0f;
372+
cell1._signalRestriction._openingAngle = 90.0f;
373+
374+
auto cell2 = CellDescription().id(2).pos({101.0f, 100.0f});
375+
cell2._signalRestriction._mode = SignalRestrictionMode_Inactive;
376+
377+
auto data = Description().cells({cell1, cell2});
378+
data.addConnection(1, 2);
379+
_simulationFacade->setSimulationData(data);
380+
381+
// Select both cells using position-based selection
382+
_simulationFacade->setSelection({99.0f, 99.0f}, {102.0f, 101.0f});
383+
384+
auto connectionData = _simulationFacade->testOnly_getConnectionArrowData();
385+
ASSERT_EQ(2u, connectionData.size());
386+
// Cell1 has restriction that blocks signal to cell2 (connection angle 0 is outside range [225,315])
387+
// Cell2 has no restriction, so signal can flow to cell1
388+
// Expected: arrow to cell1 (bit 0 = 1), no arrow to cell2 (bit 1 = 0) => flags = 1
389+
EXPECT_EQ(1, connectionData[0].arrowFlags);
390+
EXPECT_EQ(1, connectionData[1].arrowFlags);
391+
}
392+
393+
TEST_F(GeometryTests, connectionData_withRestriction_conditional_restrictedDirection)
394+
{
395+
auto cell1 = CellDescription().id(1).pos({100.0f, 100.0f});
396+
cell1._signalRestriction._mode = SignalRestrictionMode_Conditional;
397+
// Use baseAngle = 90 and openingAngle = 90 to point away from connection
398+
cell1._signalRestriction._baseAngle = 90.0f;
399+
cell1._signalRestriction._openingAngle = 90.0f;
400+
401+
auto cell2 = CellDescription().id(2).pos({101.0f, 100.0f});
402+
cell2._signalRestriction._mode = SignalRestrictionMode_Inactive;
403+
404+
auto data = Description().cells({cell1, cell2});
405+
data.addConnection(1, 2);
406+
_simulationFacade->setSimulationData(data);
407+
408+
// Select both cells using position-based selection
409+
_simulationFacade->setSelection({99.0f, 99.0f}, {102.0f, 101.0f});
410+
411+
auto connectionData = _simulationFacade->testOnly_getConnectionArrowData();
412+
ASSERT_EQ(2u, connectionData.size());
413+
// Conditional mode should render the same as Active mode for arrow directions
414+
// Cell1 has restriction that blocks signal to cell2
415+
// Cell2 has no restriction, so signal can flow to cell1
416+
// Expected: arrow to cell1 (bit 0 = 1), no arrow to cell2 (bit 1 = 0) => flags = 1
417+
EXPECT_EQ(1, connectionData[0].arrowFlags);
418+
EXPECT_EQ(1, connectionData[1].arrowFlags);
419+
}

0 commit comments

Comments
 (0)