【CUDA】Branch Divergence and Unrolling Loop

目录

一、避免分支发散

1.1 并行规约问题

1.2 并行规约中的发散

二、UNrolling Loops

---

一、避免分支发散

控制流有时依赖于 thread 索引。同一个warp中，一个条件分支可能导致性能很差。通过重新组织数据获取模式可以减少或避免 warp divergence。具体问题查看下面这篇文章：

【CUDA】Warp解析-优快云博客https://blog.youkuaiyun.com/GG_Bruse/article/details/143772619

1.1 并行规约问题

若要计算一个数组N个元素的和，使用CPU编程实现十分容易

int sum = 0;
for (int i = 0; i < N; ++i)
    sum += array[i]

若数组中的元素非常多，应用并行计算可以大大提高效率。鉴于加法交换律等性质，这个求和过程可以以元素的任意顺序来进行：

• 将输入数组切割成很多小的块
• 用 thread 来计算每个块的和
• 对这些块的结果再求和得最终结果

数组的切割主旨是：用 thread 求数组中按一定规律配对的的两个元素和，然后将所有结果组合成一个新的数组，然后再次求配对两元素和，多次迭代，直到数组中只有一个结果

比较直观的两种实现方式是：

1. Neighbored pair：每次迭代都是相邻两个元素求和
2. Interleaved pair：按一定跨度配对两个元素

对于N个元素的数组，该过程需N - 1次求和、步。InterleavedPair的跨度是半个数组长度

上述讲的这类问题术语为 reduction problem。Parallel reduction（并行规约）是指迭代减少操作，是并行算法中非常关键的一种操作

1.2 并行规约中的发散

Neighbored pair

在这个kernel中，有两个global memory array，一个用来存放数组所有数据，另一个用来存放部分和。所有 block 独立的执行求和操作

#include <iostream>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>
using namespace std;

// 1 << 28 == 268435456
__global__ 
void reduceNeighbored(int* inputData, int *outputData, unsigned int N) 
{
    unsigned int threadIndex = threadIdx.x; // 0 - 511
    unsigned int index = blockIdx.x * blockDim.x + threadIdx.x; // 0 - 524287, 512, 0 - 511 即 0 - 268435455
    int* iData = inputData + blockIdx.x * blockDim.x; // inputData + (0 - 268434944) 即指向每个block的起始位置
    if (index >= N) return;
    // stride = (1 - 511) 1, 2, 4, ..., 500
    for (int stride = 1; stride < blockDim.x; stride *= 2) 
    {
        if ((threadIndex % (2 * stride)) == 0) // 每次取threadIndex为偶数的thread
            iData[threadIndex] += iData[threadIndex + stride];
        __syncthreads();
    }
    if (threadIndex == 0) outputData[blockIdx.x] = iData[0]; // 放入计算出的数据 
}

void GPU_ReduceNeighbored(int* dInputData, int* dOutputData, int size, dim3 grid, dim3 block)
{
    reduceNeighbored<<<grid, block>>>(dInputData, dOutputData, size);
}

因为无法让所有的block同步，所以最后将所有block的结果送回 host 进行串行计算

主函数代码：

void GPU_ReduceNeighbored(int* inputData, int *outputData, int N, dim3 grid, dim3 block);

long long Seconds()
{
    // 获取当前时间点
    auto now