Cuda 學習 - Parallel Reduction（下）
===
* 本文接續上一篇 [Cuda 學習 - Parallel Reduction（上）](https://hackmd.io/lfQVUDnZSZ-j5hbRLENTOA) 做後續的演練
* 在上一篇已完成了 Parallel reduction 的介紹及基本的程式演練但在 Sequential Addressing 上面雖然已完成加速但造成了一半以上的 threads 呈現 ideal 狀態, 這是非常浪費資源的事情

# First Add During Load
* 由於後半的 thread 基本上都為使用到那和不一開始就先往前加到前半呢
* 可以看到直接將 i 的部份前半的 block offset 乘上兩倍並放到 shared memory 上
* 迴圈的 offset 也改由 blockDim.x 的一半開始, 每次新的迭代就除以 2, 並針對小於 offset 的 tid 進行加總

__global__ void reduce3(int *g_idata, int *g_odata) 
{
    extern __shared__ int sdata[];
    // each thread loads one element from global to shared mem
    
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;
    sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
    __syncthreads();

    for (unsigned int s=blockDim.x/2; s>0; s>>=1) {
        if (tid < s) 
        {
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

* 然而實際計算還離原有 Bandwidth 有較大差距, 原因在於 address 的運算及迴圈運算
# Unroll the Last Warp
* 在運算的最後我們只會剩下一個 Warp 去使用, 而當 thread <= 32 可不用在使用 __synthreads 及其於同步控制
* Warp 跟 Warp 之間才需要做同步, 單一個 Warp 內的 thread 本身會執行同時執行同一個指令
* 在 Unroll the last warp 時需要使用 volatile 來做每次資料正確的讀取

__device__ void warpReduce(volatile int* sdata, int tid) 
{
    sdata[tid] += sdata[tid + 32];
    sdata[tid] += sdata[tid + 16];
    sdata[tid] += sdata[tid + 8];
    sdata[tid] += sdata[tid + 4];
    sdata[tid] += sdata[tid + 2];
    sdata[tid] += sdata[tid + 1];
}

__global__ void reduce4(int *g_idata, int *g_odata) 
{
    extern __shared__ int sdata[];
    // each thread loads one element from global to shared mem
    
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;
    sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
    __syncthreads();

    for (unsigned int s=blockDim.x/2; s>32; s>>=1) {
        if (tid < s) 
        {
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();
    }

    // write result for this block to global mem
    if (tid < 32)
    {
        warpReduce(sdata, tid);
    } 
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

# Completely Unrolled
* 剛剛我們只展開了最後一個 Warp , 現在來將 for 迴圈也做展開

__global__ void reduce5(int *g_idata, int *g_odata) 
{
    extern __shared__ int sdata[];
    // each thread loads one element from global to shared mem
    
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;
    sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
    __syncthreads();

    if(blockDim.x>=512 && tid<256)
    {
        sdata[tid] += sdata[tid+256];
    }
    __syncthreads();
    if(blockDim.x>=256 && tid<128)
    {
        sdata[tid] += sdata[tid+128];
    }
    __syncthreads();
    if(blockDim.x>=128 && tid<64)
    {
        sdata[tid] += sdata[tid+64];
    }
    __syncthreads();
    // write result for this block to global mem
    if (tid < 32)
    {
        warpReduce(sdata, tid);
    } 
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

# 計算結果
* 硬體環境
* 筆電 : Dell G15
* GPU : RTX3050 Laptop 4G
* CPU : 12th Gen Intel(R) Core(TM) i5-12500H, 16 cores
* Memory : 16GB
* 計算任務
* 2^22 的 Array 內的元素和
* block : 128
* grid : (2^22/128)/2
* Grid 的部份從 First Add During Load 開始因為將後半的資料往前加, 所以 grid 本身也只需要一半即可完成

| Method | Time (ms)| Bandwidth (GB/s) | Grid | block
| ------------------------------------ | -------- | -----------------|------------|------
| Interleaved Addressing | 0.5519 | 56.6171 |2^22/128 |128
| Interleaved Addressing (Improved) | 0.493769 | 63.2887 |2^22/128 |128
| Sequential Addressing | 0.473715 | 65.9679 |2^22/128 |128
| First Add During Load | 0.341927 | 91.3938 |(2^22/128)/2|128
| Unroll the Last Warp | 0.307297 | 101.693 |(2^22/128)/2|128
| Completely Unrolled | 0.263343 | 118.667 |(2^22/128)/2|128

* 
* 從 Interleaved Addressing 到 Completely Unrolled 已加速 2.09X
# 結論
* 我們完成了 Parallel Reduction 的許多小 trick, 例如善用 Warp 的機制以及盡可能的使用每個 thread 進而不造成資源的浪費, 而在 IA proved 及 SA 須再確認 差異不大的原因
source : https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf
https://blog.csdn.net/langb2014/article/details/51348523
* github : https://github.com/GU-Lin/Cuda_Practice/tree/main/5_Parallel_Reduction
* ###### tags : `Cuda`