cuda6

SCAN

前缀和

串行

for(j=1;j<n;j++)
    out[j] = out[j-1] + f(j);

并行

forall(j) { temp[j] = f(j) };
    scan(out, temp);

---

工作效率低下的扫描内核

优化方案

复杂度为 nlog(n) 迭代 log(n)次，每次运算数量级为 n

每次迭代都需要同步已确保输入为最新值

__global__ void work_inefficient_scan_kernel(float *X, float *Y, int InputSize) {
    __shared__ float XY[SECTION_SIZE];
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < InputSize) {XY[threadIdx.x] = X[i];}
        // the code below performs iterative scan on XY
        for (unsigned int stride = 1; stride <= threadIdx.x; stride *= 2) {
        __syncthreads();
            float in1 = XY[threadIdx.x - stride];
            __syncthreads();
            XY[threadIdx.x] += in1;
        }
      __ syncthreads();
      If (i < InputSize) {Y[i] = XY[threadIdx.x];}
}

但仍劣于串行算法

此扫描执行 log(n) 次并行迭代
迭代 do (n-1), (n-2), (n-4),..(n- n/2) 添加每个
总和：n * log(n) - (n-1)  O(n*log(n)) 工作

---

更优的算法

将输入转化为平衡树，从下至上扫描树至根，根即为所有叶子综合，部分和记录在树中。

// XY[2*BLOCK_SIZE] is in shared memory
//上部分
for (unsigned int stride = 1;stride <= BLOCK_SIZE; stride *= 2) {
    int index = (threadIdx.x+1)*stride*2 - 1;
    if(index < 2*BLOCK_SIZE)
        XY[index] += XY[index-stride];
    __syncthreads();
}
  for (unsigned int stride = BLOCK_SIZE/2; stride > 0; stride /= 2) {
        __syncthreads();
        int index = (threadIdx.x+1)*stride*2 - 1;
        if(index+stride < 2*BLOCK_SIZE) {
          XY[index + stride] += XY[index];
        }
  }
    __syncthreads();
  if (i < InputSize) Y[i] = XY[threadIdx.x];

复杂度

规约时执行 log(n)次迭代，每次数量从 n/2 以 1/2 递减至 1，随后反规约迭代 log(n)-1 次。

等比数列，2(n-1)

工作效率高的内核在缩减步骤中执行 log(n) 次并行迭代
迭代做 n/2, n/4,..1 添加
总加：(n-1)  O(n) 工作

它在减少后的反向步骤中执行 log(n)-1 并行迭代
迭代做 2-1, 4-1, …. n/2-1 添加
总加：(n-2) – (log(n)-1)  O(n) 工作

一些讨论

对比

第一种优化方式绝对性能更好，因为只需要一个方向的规约。

第二种优化方式更理想，效率更高，占用资源更少

大量输入

对于两种方案都是将输入分块，得到块总和后再以相同算法扫描

Exclusive Scan

第一个元素赋 0，其他元素后移即可

---

OpenACC

open accelerators

为简化在异构 CPU/GPU 开发而出现

OpenACC 设备结构

运行层次

向量：包括多个线程以锁步形式工作(SIMD/SIMT)

Worker：共同计算一个向量

Gang：一个或多个 Workers 共享资源

占有率计算

GPU 占用率衡量 GPU 计算资源的利用率。

How much parallelism is running / How much parallelism the hardware could run

基本指令(DIRECTIVE)

OpenACC loop directive: gang, worker, vector, seq

The loop directive gives the compiler additional information about the next loop.

• gang – Apply gang-level parallelism to this loop
• worker – Apply worker-level parallelism to this loop
• vector – Apply vector-level parallelism to this loop
• seq – Do not apply parallelism to this loop, run it sequentially

Multiple levels can be applied to the same loop, but the levels must be applied in a top-down order.

指令格式

#pragma acc directive-name [clause-list] new-line
#pragma acc 编译指示名 [子句[ [,] 子句]…] 换行

OpenACC kernels Directive

OpenACC Data movement directive

OpenACC Loop