# Boundary-Aware Concurrent Queue: A Fast and Scalable Concurrent FIFO Queue on GPU Environments
## 3 BACQ Proposal
### Data Structure  
The key ingredient behind BACQ is a special circular buffer queue structure, which can be reused for further insertions instead of reallocating a new buffer. The source code of this structure is provided in Listing 1 from Lines 1-5. The key components of this data structure are outlined below:  
1. Headand Tail: These counters are key elements to establish the priority of operations. Every enqueue and dequeue requests atomically increment tail and head, respectively. These counters represent the positions within the circular buffer queue after a modulo operation and have the capability to increase infinitely. Figure 1 shows a circular buffer structure that bases on BACQ array with size \(N\). Head and Tail point to the next dequeue and enqueue elements after a modulo operation with the queue size.
2. Items: This is the actual centralized ring buffer to store queue elements.
3. Tickets: This array is designed to assign a ticket for each request and aligns them with the respective buffer index position.  
Listings 2 and 3 provide the pseudocode for the BACQ algorithm. Our algorithm addresses the queue access latency, fair publication order and scalability problems in GPU concurrent queues.  
Figure 1: Circular buffer queue structure with a size of N. “Head % queue_size” points to the next dequeue element, and “Tail % queue_size” points to the next enqueue element.  
```
1 Enqueue(BACQ *q, item){
2 if((q->head == 0 && q->tail+1 == N) || (q->tail - q->head) == N)
3{ returnQUEE_FULL;
4}
5 active +- coalesced_threads();
6 guaranteesSpace +- ((q->tail-q->head) + active.size() > N)? (N - (q-> tail-q->head)):active.size();
7 if(active.thread_rank()==0)
8{
9 warpBaseIndex +- atomicAdd(q->tail,guaranteeSpace);
10}
11 else if(active.thread_rank() > garanteeSpace)
12 returnQUEE_FULL;
13}
14
15 localTail + (active.shelf(warpBaseIndex,0)+active.thread_rank());
16 tailBufferIndex - (active.shelf(warpBaseIndex,0) + active.thread_rank())% N;
17
18 if(localTail > q->head)
19{
20 ticket - 2 *(active.shelf(warpBaseIndex,0) + active.thread_rank()) +1)/N);
21}
22
23 else
24 ticket - 2*(active.shelf(warpBaseIndex,0)+active.thread_rank()) +1)/N)+1;
25}
26
27
28 else
29 ticket - 2*(active.shelf(warpBaseIndex,0)+active.thread_rank()+1) /N)+1;
30}
31
32
33 closToFullQueue +- (tailBufferIndex > N - 32)? true : false;
34 if(closeToFullQueue)
35{
36 ticket +- (2 *(active.shelf(warpBaseIndex,0)+active.thread_rank() +1)/N)+1;
37}
38 while(threadIdx.x < 32) {
39 if(q->tickets[tailIndex]!= ticket)
40{
41 continue;
42}
43
44
45
46
47 ull Old +- atomicExch(q->items[tailIndex],item);
48
49 atomicExch(q->tickets[tailIndex], ticket+1);
50
51 break;
52
53
54 return0;
55}
```  
Listing 3: Dequeue algorithm of BACQ.  
```
1 Depuense(BACQ *q){
2 if((q->head > q->tail) || (q->tail - q->head) == 0)
4{
5 returnEWPT_QUEE;
6}
7
8 active +- coalesced_threads();
9
10
11
12
13 if(active.thread_rank()==0)
14{
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
99
100
97
111
122
133
145
156
167
178
189
190
191
192
193
194
195
196
197
198
201
2197
2198
202
210
223
231
241
252
263
274
285
297
298
299
300
3101



21 localHead +- (active.shfl(warpBaseIndex,0) + active.thread_rank(+1));
22 headBufferIndex +- (active.shfl(warpBaseIndex,0) + active.thread_rank() +1)N;
25 if(localHead < q->tail) {  ticket - 2*((active.shfl(warpBaseIndex,0)*active.thread_rank()+1) /N)+1;
26} closeToEmptyQueue - (headBufferIndex < 2)? true : false;
37 ticket += cloneToEmptyQueue? ((2*active.shfl(warpBaseIndex,0) + active.thread_rank()+1) /N) + 1 : ((2*active.shfl(warpBaseIndex,0)+active.thread_rank()) /N);
38 while(threadIdx.x < 32) {  if(q->tickets(headBufferIndex)+ticket) {  continue;
39 }  deqItem = atomicExch(q->items[headBufferIndex], SENTINL);
40 atomicExch(q->tickets[headBufferIndex], ticket+1);
41 break;
42 } returnRetVal;
43 }
```