# CUDA-Optimized GPU Acceleration of 3 GPP 3D Channel Model Simulations for 5G Network Planning
## 5 Channel Emulator Acceleration on GPU
```
1__global__void
2 applyFIR(y(t),x(t),H_{u,x}(t),H_{u,y}(t-\tau),\)cirBuf,pos)
3{
4//CUDAgridwithX,Y,Zblocks
5i=Idx.x;tx=Idx.y;rx=Idx.z; l=threadIdx.x;
6 cta = this_thread_block(); //createtiledblocks
7 thread_tile<32>tile32=tiled_partition<32>(cta);
8\(\mathit{AH}(t)=\mathit{H}_{tx}(t)-\mathit{H}_{u,k}(t-\tau)\);//usesharedmem
9\(\mathit{rcgCirBuf}_{x}=\mathit{cirBuf}_{x}/\)loadcircularbuffer
10\(\_\_\mathit{syncthreads}();\)//blocksync

Figure 4: The 3 Rd Generation Partnership Project (3 GPP) three-dimensional (3D) channel model on graphics processing unit (GPU). \begin{tabular}{l l} \(11\) & \(x_{l}=x(t)_{i,tx}\); // load received symbolionsharedmem \\ \(12\) & \(index=pos_{l}\); // read clusterpositionfromconstmem \\ \(13\) & \(tapV=tapV+AH(t)+H_{us}(t)\) // warp-widetapvector \\ \(14\) & \(tile32\_sync()\); // warp-wisesoftsync \\ \(15\) & \(tap=regCirBuf_{index}\times x_{l}\) // interpolationlines \\ \(16\) & \(tile32\_sync()\); \\ \(17\) & \(acc_{l}=uarpReduceSum_{l}(tap)\); // reductioninregs \\ \(18\) & \(tile32\_sync()\); \\ \(19\) & \(y(t)_{rx}=y(t)_{rx}+acc_{l}\); // accumulateoverRx \\ \(20\) \\ \end{tabular} 

## 6 Results and Discussion

The baseline CPU performance was determined using an Intel Core i7-6900K @3.2 GHz CPU. The baseline channel model is implemented in C++ and runs as a MEX C++ function within a MATLAB R2021a environment. The performance of the channel model is evaluated using the benchmark values in Table 1. To evaluate the performance for link-level simulations, we consider two CDL profiles, i.e., CDL-B for NLOS clusters and CDL-D for LOS clusters. Figure 5 illustrates various MIMO antenna element configurations for single-polarized antennas in Figure 4(a),b and dual-polarized arrays in Figure 5(c),d on transmitter and receiver end for CDL-B profile. Similarly, the same is reported for CDL-D in Figure 6 where Figure 5(a),b show antenna patterns for single-polarized arrays and Figure 5(c),d illustrate dual-polarized antennas on transmitting and receiving end. The accelerator discussed in this paper was developed using the CUDA development tools [27], targeting the NVIDIA GeForce GTX 1070 GPU [28] which features 1920 CUDA cores, 120 texture mapping units (TMUs), 1.5 MB of shared memory, 4 MB of local memory, 8 GB of GDDR5 memory, and 15 SMs. Figure 5: MIMO antenna configuration in CDL-B profile for NLOS clusters. We compare the performance of the GPU accelerator with an FPGA implementation [19], which was developed using the _Vitis Unified Software Platform_[29] for the _AMD Alveo U280_[30]. The FPGA used in [19] is based on the same \(16\,\mathrm{nm}\) technology node as the GPU and contains 9024 digital signal processing (DSP) blocks, 41 MB of on-chip static RAM, 1,303,680 look-up tables, and \(8\,\mathrm{GB}\) of high bandwidth memory (HBM2). Thus, its computational power is comparable to that of the GPU used in this work, because (1) a DSP unit can be used to implement a single-precision (SP) multiply and add, and (2) in [19] are used only \(1/3\) of the total FPGA resources so that the kernel can fit on one chiplet to avoid routing problems. The primary goal of this work is to reduce the overall execution time of the channel model under resource constraints. We report the achieved performance for the kernels in Listings 1 and 2 on GPU platforms. To analyze the performance for both LOS and NLOS scenarios, we consider CDL-B and CDL-D profiles and uplink and downlink connection types (Tables 2 and 3). Table 2 reports the execution latency for various combination of \(R_{x}\) and \(T_{x}\) antenna elements considering NLOS clusters in CDL-B profile for the parameters listed in Table 1.

Figure 7 illustrates a comparison of link-level simulation latency on CPU and GPU platforms in the two CDL profiles. The values on the horizontal axis represent the number of receiving and transmitting antenna elements, whereas the vertical axis denotes the total execution time in logarithmic scale.