【博客翻译】让前缀和变得更快

正文

请到「今天看啥」查看全文

numBlocks = (numElements + threadsPerBlock - 1 ) / threadsPerBlock;
kernel_0
<< >>(input, output);
}

该算法实现了上图所示的功能。在每个阶段,我们将偏移量增加两倍。我们会一直这样做,直到到达线程块的中间位置,并且累积元素之间的距离达到线程块大小的一半。注意,我们需要使用 __syncthreads 来避免竞争条件。如果没有这两个同步屏障,可能会出现两个数组元素同时被读写的情况。

使用一个简单的CPU实现来检查程序的正确性是一个很好的做法。

emplate <int threadsPerBlock, int numElements>
void cpu_scan(int *input, int *output) {
  output[0] = input[0];
  for (int i = 1; i < numElements; i++) {
    if (!((i & (threadsPerBlock - 1)) == 0)) {
      output[i] = input[i] + output[i - 1];
    } else {
      output[i] = input[i];
    }
  }
}

该算法给出了正确的结果。不幸的是,它的性能并不理想。这是由于频繁访问全局内存。我们可以通过计算带宽来衡量性能。我们进行了N次读写操作,其中 N = 1 << 30 = 2**30 。上述kernel的测量性能如下:

Bandwidth: 823.944 GB/s
Efficiency: 0.24968

使用共享内存

template <int threadsPerBlock, int numElements>
__global__ void kernel_1(int *input, int *output) {
extern __shared__ int buffer[threadsPerBlock];

constint tid = threadIdx.x;
constint gtid = blockIdx.x * threadsPerBlock + tid;

  buffer[tid] = input[gtid];
  __syncthreads();

#pragma unroll
for (unsignedint offset = 1; offset <= threadsPerBlock / 2; offset <<= 1) {
    int tmp;
    if (tid >= offset) {
      tmp = buffer[tid - offset];
    }
    __syncthreads();

    if (tid >= offset && gtid < numElements) {
      buffer[tid] += tmp;
    }
    __syncthreads();
  }

if (gtid < numElements) {
    output[gtid] = buffer[tid];
  }
}

template <int threadsPerBlock, int numElements>
void launch_kernel_1(int *input, int

请到「今天看啥」查看全文