## 3 Parallel ARIA Implementation
### Parallel ARIA Implementation on GPU  
The implementation using the T-table extended Sbox was used to implement ARIA on GPU. The original size of one Sbox is 256 bytes. In the case of the T-table, the size of 1 KB is used. Since ARIA block cipher uses four different Sboxes, the size of the entire table increases from 1 KB to 4 KB. By using T-table, memory usage increases, but it has the advantage of reducing operations in the diffusion layer. When implementing ARIA block ciphers in parallel on GPU, we compare and analyze them using various memory types. The substitution layer introduces a lot of memory access for substitution. Memory accesses cause long latency, and as a result, many memory accesses have a significant impact on performance. The implementation performance depends on the memory access. The Sbox is loaded into different memories showing different latency [24]. In this section, we present four implementations in CUDA codes. Only one Sbox is covered in detail (in code level), but the remaining three tables operate with the same code. Figure 5: Simplified operation of the variable \(T\). First, the basic implementation is to load the Sbox in the _global memory_ (CUDA code is Listing 1). Global memory has the slowest memory access speed, but it has the largest capacity. In CUDA programming, memory is allocated to the GPU (device) through \(cudaMalloc()\). The memory allocated in this way uses global memory. The value of the PC (i.e., host) memory is copied through \(cudaMemcpy()\). ```
1//GlobalMemoryCode
2 constantuint32_ts1[256]={...};
3...
4__global_Aria_device(in,out,s1,...){...}
5__host_Aria_host(in,out,key){
6 roudnkey_gen(key,rk);
7...
8 cudaMalloc(&dev_s1,sizeof(uint32_t)*256);
9...
10 cudaMemcpy(s1,dev_s1,sizeof(uint32_t)*256);
11...
12 Aria_device<<block,thread>>>(in,out,dev_s1,...);
13...
14}
```  
Listing 1: Parallel implementation of ARIA using global memory. The second is _shared memory_ implementation (CUDA code is Listing 2). Shared memory is a memory space provided per block that can be shared and used by multiple threads. In this case, it is possible to access the shared memory in unit of Warp (i.e., 32 threads) at once. Shared memory can be accessed with faster access speed than global memory. However, the memory space provided is small at 48 KB. When using shared memory, it should be implemented with caution against bank conflicts. A bank conflict is a problem in which multiple threads access the same bank. This leads to sequential processing in parallel machine. For this reason, it may show slow performance even with fast access speeds. Shared memory can be initialized through _shared_ in the device function. Afterwards, an additional operation is performed to copy the values from global memory to shared memory. A copy process must be performed for each block. In this case, \(threadIdx.x\), which indicates the thread index of each block, is used [25]. ```
1//SharedMemoryCode
2 constantuint32_ts1[256]={...};
3...
4__global_Aria_device(in,out,dev_s1,...){
5__shared_uint32_tshared_s1[256];
6...
7 if(threadIdx.x < 256){
8 shared_s1[threadIdx.x] =dev_s1[threadIdx.x];
9...
10}
11...
12}
13__host_Aria_host(in,out,key){
14 roudnkey_gen(key,rk);
15...
16 cudaMalloc(&dev_s1,sizeof(uint32_t)*256);
17...
18 cudaMemcpy(s1,dev_s1,sizeof(uint32_t)*256);
19...
20 Aria_device<<block,thread>>>(in,out,dev_s1,...);
21...

We also utilized the technique of copying the Sbox to minimize bank conflicts (CUDA code is Listing 3). Since the Sbox of ARIA is 4 KB, copying the Sbox by the number of banks will exceed the available shared memory size. Therefore, the maximum number of copies of the Sbox to shared memory is 12. However, the number of banks is 32. If we copy it to 12, the table is mixed when sorted in the bank. The bank conflict cannot be controlled if they are sorted as shuffled state. Thus, the best copy is the divisor of 32. In this paper, the implementation was carried out when Sbox was simply used as a shared memory without copying, when it was copied to 8, and, finally, when it was copied to 12.