# Prime Time Tactics--Sieve Tweaks and Boosters
## 3 Parallel Sieving Strategies
### Generic Considerations
```
1 kernelvoidSo_simple_local(globaluint8t*outBuff,uint32_tworkBuffSize,localuint8_t*workBuff,uint32_tworkBuffSize
3#ifdef_t_NG_t/8#
4,uint32_tgroup_id_0,uint32_tnum_groups_0,uint64_tglobal_offset_0,uint32_tlocal_id_0,uint32_tlocal_size_0
5_endif
6 constuint32_tlocalID=get_local_id(0),localDim=get_local_size(0);
7 constuint32_tthread_range=workBuffSize/localDim;
8 constuint32_tthread_offset=localID*thread_range;
9 for(uint32_t=0;1<thread_range;1++)
10 workBuff[thread_offset+i] =0;
11 barrier(CLK_LOCAL_MEM_FEENCE);
12
13 constuint32_tgroupID=get_group_id(0),groupNo=get_num_groups(0);
14
15 constuint64_tglobal_offset=get_global_offset(0);
16 constuint32_tlocalDim2=localDim+localDim;
17 constuint32_tgroup_offset=groupID*workBuffSize;
18 constuint32_tgroup_span=workBuffSize=1;
19 constuint64_val_min=global_offset+groupID*group_span;
20 constuint64_val_max=val_min+group_span;
21 constuint32_tstart=(unsigned)sqrt(val_max/2.0f);
22 uint32_t1=istart-localID,delta=2*i+1;
23 for(;;delta-localDim2);
24{
25 uint64_tj=(val_min>i)? (uint64_t)((val_min-i)/(1.0+2*i)):i;
26 uint64_tn=i+j+2*(uint64_t)i*j;
27 while(n<val_min)
28n+=delta;
29 for(unsigned==(unsigned)(n-val_min);nn<group_span;nn+=delta)
30 workBuff[fm]=1;
31 if(i>localDim)
32i-=localDim;
33 else
34 break;
35}
36}barrier(CLK_LOCAL_MEM_FENCE);
37 for(uint32_t i=0;i<thread_range;i+)
38 outBuff[group_offset+thread_offset+i]=workBuff[thread_offset+i];
39}
```
**Algorithm 3** Sieve of Sundaram--segmented kernel

Once we have a decent kernel, the last step in optimizing the sieve is to parallelize the final phase of the sieving process, the actual counting/generation of primes, so that the counting is not longer that the sieving, which should not be difficult; final timings of such an implementation will be driven exclusively by the duration of sieving itself, as executed on the GPU, plus the transfer of data buffers between host and device, which is not negligible. Sometimes, it may be faster to generate initial data directly on the GPU and avoid any data transfer from the host to the device, although modern technologies like PCIe 4/5 and Base Address Register (BAR) Resizing (or Resizable BAR) have improved this issue a lot. The performance may be further improved by flattening the curve inside each segment using the cutoff-transpose technique, as described above, or avoiding unnecessary loops because, for larger is, only a small number of \(j\)s will actually impact a certain segment, and a bucket-like algorithm [26] may benefit the implementation. But perhaps the hardest problem is to avoid collisions between local threads, especially when using on the GPU buffer a 1 bit approach for data compression, as we use on the normal CPU; on the GPU, this is particularly difficult, as one has to be sure that in the same workgroup cycle there are not two work items (local threads) that will try to update the same byte, as this will result in memory races that will keep only the last data value and lose the others. When working with uncompressed buffers (1 byte per value), this issue is not important anymore for Eratosthenes and Sundaram families of sieves, because here, if the values are updated, they are always updated in the same direction: changing from 0/false to 1/true (or vice versa); memory races are not an issue here, as the final result will be correct irrespective of which thread managed to make the update. For the Atkin algorithm, the outcome is contingent on the starting value, and concurrent operations by two or more threads can lead to incorrect results. For instance, in Atkin, if the initial value is 0 and two threads modify it one after the other, the first thread will read 0 and change it to 1, while the second thread will read 1 and switch it back to 0. However, if both threads operate at the same time, each will read the initial value of 0 and both will write back 1, resulting in an erroneous final outcome. And this problem is valid for all sieve families when dealing with bit compression.