@@ -2767,7 +2767,7 @@ __device__ void advancePSAPVerticesDevice(int vertexIdx, int maxEventsPerEpoch,
27672767 int &queueLengthHistoryQueueEnd = allVerticesDevice->queueLengthHistoryBufferEnd_ [vertexIdx];
27682768 allVerticesDevice->queueLengthHistory_ [vertexIdx][queueLengthHistoryQueueEnd]
27692769 = stepsPerEpoch * (1 - (allVerticesDevice->vertexQueuesFront_ [vertexIdx] >= queueEndIndex))
2770- + allVerticesDevice->vertexQueuesFront_ [vertexIdx] - queueEndIndex;// queueSize
2770+ + allVerticesDevice->vertexQueuesFront_ [vertexIdx] - queueEndIndex; // queueSize
27712771 queueLengthHistoryQueueEnd = (queueLengthHistoryQueueEnd + 1 ) % stepsPerEpoch;
27722772 allVerticesDevice->queueLengthHistoryNumElementsInEpoch_ [vertexIdx]++;
27732773 // EventBuffer::insertEvent
@@ -3059,26 +3059,45 @@ __global__ void maybeTakeCallFromEdge(int totalVertices, uint64_t stepsPerEpoch,
30593059
30603060 // Compute the size of the destination queue
30613061 uint64_t queueFrontIndex = allVerticesDevice->vertexQueuesFront_ [dstIndex];
3062- int queueFull = (int )((1 - (queueFrontIndex >= allVerticesDevice->vertexQueuesEnd_ [dstIndex])) * (allVerticesDevice->numTrunks_ [dstIndex] + 1 ) + queueFrontIndex - allVerticesDevice->vertexQueuesEnd_ [dstIndex]) >= (allVerticesDevice->numTrunks_ [dstIndex] - allVerticesDevice->busyServers_ [dstIndex]);
3063- allVerticesDevice->droppedCalls_ [dstIndex] += queueFull && (!allEdgesDevice->isRedial_ [edgeIdx]);
3064- allVerticesDevice->receivedCalls_ [dstIndex] += queueFull && (!allEdgesDevice->isRedial_ [edgeIdx]);
3062+ int queueFull
3063+ = (int )((1 - (queueFrontIndex >= allVerticesDevice->vertexQueuesEnd_ [dstIndex]))
3064+ * (allVerticesDevice->numTrunks_ [dstIndex] + 1 )
3065+ + queueFrontIndex - allVerticesDevice->vertexQueuesEnd_ [dstIndex])
3066+ >= (allVerticesDevice->numTrunks_ [dstIndex] - allVerticesDevice->busyServers_ [dstIndex]);
3067+ allVerticesDevice->droppedCalls_ [dstIndex]
3068+ += queueFull && (!allEdgesDevice->isRedial_ [edgeIdx]);
3069+ allVerticesDevice->receivedCalls_ [dstIndex]
3070+ += queueFull && (!allEdgesDevice->isRedial_ [edgeIdx]);
30653071 if (!queueFull) {
30663072 // Transfer call to destination
30673073 // We throw an error if the buffer is full
3068- if (((queueFrontIndex + 1 ) % (allVerticesDevice->numTrunks_ [dstIndex] + 1 )) == allVerticesDevice->vertexQueuesEnd_ [dstIndex]) {
3069- printf (" ERROR: Vertex queue is full. Vertex ID [%d] Front Index [%" PRIu64 " ] End Index [%" PRIu64 " ] Buffer size [%" PRIu64 " ]\n " , dstIndex, queueFrontIndex, allVerticesDevice->vertexQueuesEnd_ [dstIndex], (allVerticesDevice->numTrunks_ [dstIndex] + 1 ));
3074+ if (((queueFrontIndex + 1 ) % (allVerticesDevice->numTrunks_ [dstIndex] + 1 ))
3075+ == allVerticesDevice->vertexQueuesEnd_ [dstIndex]) {
3076+ printf (" ERROR: Vertex queue is full. Vertex ID [%d] Front Index [%" PRIu64
3077+ " ] End Index [%" PRIu64 " ] Buffer size [%" PRIu64 " ]\n " ,
3078+ dstIndex, queueFrontIndex, allVerticesDevice->vertexQueuesEnd_ [dstIndex],
3079+ (allVerticesDevice->numTrunks_ [dstIndex] + 1 ));
30703080 return ;
30713081 }
30723082 // Insert the new element and increment the front index
3073- allVerticesDevice->vertexQueuesBufferVertexId_ [dstIndex][queueFrontIndex] = allEdgesDevice->vertexId_ [edgeIdx];
3074- allVerticesDevice->vertexQueuesBufferTime_ [dstIndex][queueFrontIndex] = allEdgesDevice->time_ [edgeIdx];
3075- allVerticesDevice->vertexQueuesBufferDuration_ [dstIndex][queueFrontIndex] = allEdgesDevice->duration_ [edgeIdx];
3076- allVerticesDevice->vertexQueuesBufferX_ [dstIndex][queueFrontIndex] = allEdgesDevice->x_ [edgeIdx];
3077- allVerticesDevice->vertexQueuesBufferY_ [dstIndex][queueFrontIndex] = allEdgesDevice->y_ [edgeIdx];
3078- allVerticesDevice->vertexQueuesBufferPatience_ [dstIndex][queueFrontIndex] = allEdgesDevice->patience_ [edgeIdx];
3079- allVerticesDevice->vertexQueuesBufferOnSiteTime_ [dstIndex][queueFrontIndex] = allEdgesDevice->onSiteTime_ [edgeIdx];
3080- allVerticesDevice->vertexQueuesBufferResponderType_ [dstIndex][queueFrontIndex] = allEdgesDevice->responderType_ [edgeIdx];
3081- allVerticesDevice->vertexQueuesFront_ [dstIndex] = (queueFrontIndex + 1 ) % (allVerticesDevice->numTrunks_ [dstIndex] + 1 );
3083+ allVerticesDevice->vertexQueuesBufferVertexId_ [dstIndex][queueFrontIndex]
3084+ = allEdgesDevice->vertexId_ [edgeIdx];
3085+ allVerticesDevice->vertexQueuesBufferTime_ [dstIndex][queueFrontIndex]
3086+ = allEdgesDevice->time_ [edgeIdx];
3087+ allVerticesDevice->vertexQueuesBufferDuration_ [dstIndex][queueFrontIndex]
3088+ = allEdgesDevice->duration_ [edgeIdx];
3089+ allVerticesDevice->vertexQueuesBufferX_ [dstIndex][queueFrontIndex]
3090+ = allEdgesDevice->x_ [edgeIdx];
3091+ allVerticesDevice->vertexQueuesBufferY_ [dstIndex][queueFrontIndex]
3092+ = allEdgesDevice->y_ [edgeIdx];
3093+ allVerticesDevice->vertexQueuesBufferPatience_ [dstIndex][queueFrontIndex]
3094+ = allEdgesDevice->patience_ [edgeIdx];
3095+ allVerticesDevice->vertexQueuesBufferOnSiteTime_ [dstIndex][queueFrontIndex]
3096+ = allEdgesDevice->onSiteTime_ [edgeIdx];
3097+ allVerticesDevice->vertexQueuesBufferResponderType_ [dstIndex][queueFrontIndex]
3098+ = allEdgesDevice->responderType_ [edgeIdx];
3099+ allVerticesDevice->vertexQueuesFront_ [dstIndex]
3100+ = (queueFrontIndex + 1 ) % (allVerticesDevice->numTrunks_ [dstIndex] + 1 );
30823101 // Record that we received a call
30833102 allVerticesDevice->receivedCalls_ [dstIndex]++;
30843103 allEdgesDevice->isAvailable_ [edgeIdx] = true ;
0 commit comments