# CUDA-Optimized GPU Acceleration of 3 GPP 3D Channel Model Simulations for 5G Network Planning
## 5 Channel Emulator Acceleration on GPU  
The channel model output computation is a set of FIR filters, one per path. Consequently, the sampled signal at the receiver can be expressed as the sum over paths of a convolution between the taps of this FIR filter and the channel model input signal. In this study, we used a two-kernel acceleration:  
1. A less computation-intensive kernel computes the FIR coefficients, i.e., the CIR, according to (4). Its pseudocode is shown in Listing 1. 2. A more computation-intensive FIR kernel that applies the coefficients to each input symbol, as in Equation (6). Its pseudocode is shown in Listing 2. In addition to the two kernels, our accelerator also includes a host code that is written in C++ and executed on the host CPU. It is responsible for interacting with the simulation clients via sockets, performing preliminary model configurations and data transfers with the GPU. The architecture of the proposed accelerated channel model is shown in Figure 4.  
It uses CUDA cooperative groups to eliminate the need for block-level synchronization, because each cluster is computed independently. For efficient use of GPU resources, the long chain of computations is split into parts as shown in Listing 1 where the \(SpeedVect\) and \(ClusterVect\) are computed in shared memory. This allows threads to remain active because there is no penalty for context switching. The register-based warp-wise parallel reduction in FIR taps helps improve latency and resource utilization. ```
1__global__void
2 calcTR(\(H_{u,}\), \(F_{Rx},F_{Tx},\)RxLocation, TxLocation)
3{
4//CUDA grid with X, Y, Z blocks
5i=Idx.x; tx=Idx.y; rx=Idx.z; l=threadIdx.x;
6 cta = this_thread_block(); // createtiled blocks
7 thread_tile<32>tile32=tiled_partition<32>(cta);
8//Considering downlink
9//foruplink,arrival,anddepartureparameterswillbeswapped
10\(\mathit{dopplerSpeed}_{n,m}=\mathit{dopplerSpeed}+f_{rx,m}^{T}\times\mathit{ speed}\times\lambda^{-1}//\)asperEquation(3)
11 tile32.sync();
12\(\mathit{Speed}\mathit{Vect}_{n,m}=2\times\pi\times\mathit{dopplerSpeed}_{n,m} \times\mathit{RxLocation}\times\mathit{TxLocation}\)//computefirstpartof CIR
13 tile32.sync();
14\(\mathit{Cluster}_{tx,n}=\mathit{sgrt}(P_{u}/\mathit{NRAY})\times F_{Rx}\times F _{Tx}\times\mathit{XPR}\)
15\(\_\_\mathit{syncthreads}()\); //blocksync
16\(\mathit{H}_{u,k,m}=\mathit{Speed}\mathit{Vect}_{x,m}\times\mathit{Cluster}_{tx,n}\)
17 tile32.sync();
18\(\mathit{H}_{u,k,m}(t)=\mathit{nurpReduceSum}(H_{u,k,m}(t))\)//
19 foreach
20 if(LOS){
21//loadLOSparametersinsharedmem
22\(\mathit{dopplerSpeed}_{n,m}^{\mathit{LOS}}=\mathit{nLoS}\times\mathit{ speed}\times\lambda^{-1}//\)usingLOSparameters
23\(\mathit{Speed}\mathit{Vect}_{n,m}^{\mathit{LOS}}=2\times\mathit{dopplerSpeed}_{n,m} ^{\mathit{LOS}}\times\mathit{RxLocation}\times\mathit{TxLocation}\)
24\(\mathit{Cluster}_{tx,n}^{\mathit{LOS}}=\mathit{sgrt}(P_{u}^{\mathit{LOS}}/ \mathit{NRAY})\times F_{Rx}^{\mathit{LOS}}\times F_{Tx}^{\mathit{LOS}}//\)computeforLOSasinin Equation(2)
25\(\mathit{H}_{u,k,m}^{\mathit{LOS}}=\mathit{sgrt}(P_{u}^{\mathit{LOS}}/ \mathit{NRAY})\times F_{Rx}^{\mathit{LOS}}\times F_{Tx}^{\mathit{LOS}}//\)computeforLOSasinin Equation(2)
26\(\mathit{H}_{u,k,m}^{\mathit{LOS}}=\mathit{sgrt}(P_{u}^{\mathit{LOS}}/ \mathit{NRAY})\times\mathit{Cluster}_{tx,n}^{\mathit{LOS}}//\)computeforLOSasinin
27\(\mathit{H}_{u,k,m}^{\mathit{LOS}}=\mathit{sgrt}(P_{u}^{\mathit{LOS}}/ \mathit{NRAY})\times\mathit{Cluster}_{tx,n}^{\mathit{LOS}}//\)reducinininreg
28\(\mathit{H}_{u,k}=\mathit{sgrt}(P_{u}^{\mathit{LOS}}/ \mathit{NRAY})\times\mathit{Cluster}_{tx,n}^{\mathit{LOS}}//\)reducininreg
29\(\mathit{H}_{u,k}=\mathit{sgrt}(P_{u}^{\mathit{LOS}}/ \mathit{NRAY})\times\mathit{Cluster}_{tx,n}^{\mathit{LOS}}//\)reducininreg
30\(\_\_\mathit{syncthreads}();\)
```  
Listing 1: CUDA calcCIR kernel.