# Performance Optimization of Object Tracking Algorithms in OpenCV on GPUs
## 4 Performance Optimization of the OpenCV Object Tracking Algorithm
### Improving the Occupancy
#### 4.1.1 Mapping Data from the VGPR to LDS Memory  
The VGPR usage is the most influential factor in limiting the occupancy. Even though the size of the work-group is large enough to supply a sufficient number of wavefronts, if the VGPR usage of the kernel is high, only a few wavefronts can be active. Hence, when LDS memory usage does not restrict the occupancy, variables can be moved from the VGPR to LDS memory to relieve the VGPR usage. However, mapping data from the VGPR to LDS memory can lead to occupancy degradation due to the LDS memory shortage. In addition, all work-items in a work-group share LDS memory data. Therefore, data selection to map to LDS memory must be carefully performed. We demonstrate the properties of variables that can and cannot be mapped to LDS memory by using the code in Figure 8.  
LDS memory can contain variables that store the same data within the work-group because work-items in a work-group share LDS memory. This is the most effective approach to lowering the VGPR usage and not significantly expanding the LDS memory usage. If  
Figure 8: Part of the kernel code of the runHaarClassifier for variable allocation [41].  
a variable is not the same across work-items, the variable movement from the VGPR to LDS requires a vector array, rather than a scalar variable in the LDS memory. For example, functions get_group_id() and get_global_size() return a value that is static and shared in a work-group. Therefore, variables storing the return values of them can be mapped to LDS memory, such as groupIdx and ngroups in Lines 10 and 11 in Figure 8.  
In addition, LDS memory can include variables storing the same calculation results across all work-items. Two types of calculations correspond to returning the same result. The first is the calculation of a constant and a kernel function argument. The examples are the sumstep and invarea variables in Lines 10 and 17 in the figure. The second is the variable storing the result of calculation only among the arguments of the kernel function. The example is the normarea variable in Line 16. Furthermore, LDS memory can contain variables storing the result obtained by calculating between the above variables and non-private data of a work-item.  
Work-item private data must be excluded from storing in LDS memory. For example, the lx and ly variables correspond to private data returned by the get_local_id() function in Lines 5 and 6 in the figure.  
Table 4 shows the runHaarClassifier kernel's occupancy according to transferring variables from the VGPR to LDS memory. _Baseline_ means the case without any variable movement. When only one variable is moved to LDS memory, the kernel uses 51 to 59 VGPRs (next five rows in the table). As shown in Table 1, the four wavefronts are active in this VGPR usage range. Because a CU provides ten wavefront buffers, the occupancy is 40% in this case. Furthermore, the table demonstrates that the more variables are moved to LDS memory, the less the VGPR usage becomes. The minimum VGPR usage is 44 when six variables are mapped to LDS memory (the bottom row). For 44 VGPR usage, the occupancy is 50%. Therefore, the runHaarClassifier kernel can achieve the maximum occupancy of 50% when moving the six variables from the VGPR to LDS memory. Moving two variables, nofs and nofs0, to LDS memory also acquires 50% of occupancy, as shown in the table (the row with boldface). Therefore, transferring only these two variables to LDS memory can achieve the maximum achievable occupancy while minimizing the increase in LDS memory usage. As described in Section 2.5, the variables are mapped to LDS memory by using the __local indicator front of the variables.