CUDA循环展开

debugcn 发表于 Dev

Zigiggurats

我在CUDA中循环展开时遇到一些问题。

在普通序列号中：

//serial basic:
for(int i = 0; i < n; i++){
    c[i] = a[i] + b[i];}

//serial loop unroll:
for(int i = 0; i < n/4; i++){
    c[i] = a[i] + b[i];
    c[i+1] = a[i+1] + b[i+1];
    c[i+2] = a[i+2] + b[i+2];
    c[i+3] = a[i+3] + b[i+3];}

因此，我认为CUDA循环展开如下所示：

int i = 2*(threadIdx.x + blockIdx.x * gridDim.x);
a[i+0] = b[i+0] + c[i+0];
a[i+1] = b[i+1] + c[i+1];

但是在CUDA手册中，我无法理解展开的示例

这是一个普通的GlobalWrite内核：

__global__ void GlobalWrites( T *out, T value, size_t N )
{
for(size_t i = blockIdx.x*blockDim.x+threadIdx.x;
    i < N;
    i += blockDim.x*gridDim.x ) {
    out[i] = value;
    } 
 }

展开内核：

template<class T, const int n> __global__ void Global_write(T* out, T value, size_t N){
size_t i;
for(i = n*blockDim.x*blockIdx.x + threadIdx.x;
    i < N - n*blockDim.x*blockIdx.x;
    i += n*gridDim.x*blockDim.x;)
    for(int j = 0; j < n; i++){
        size_t index = i + j * blockDim.x;
        outp[index] = value;
    }
for ( int j = 0; j < n; j++ ) {
    size_t index = i+j*blockDim.x;
    if ( index<N ) out[index] = value;
}}

我知道此内核使用较少的块，但也许有人可以解释为什么它会更好（n = 4,10％加速）。

看守人

如果不是很明显，因为n它是模板参数，所以在编译时是恒定的。这意味着编译器可以通过展开自由地优化恒定行程计数循环。因此，对于您提到的n = 4情况，删除模板魔术并手动展开循环是有启发性的：

template<class T> 
__global__ void Global_write(T* out, T value, size_t N)
{
    size_t i;
    for(i = 4*blockDim.x*blockIdx.x + threadIdx.x;
        i < N - 4*blockDim.x*blockIdx.x;
        i += 4*gridDim.x*blockDim.x;) {
            out[i + 0 * blockDim.x] = value;
            out[i + 1 * blockDim.x] = value;
            out[i + 2 * blockDim.x] = value;
            out[i + 3 * blockDim.x] = value;
    }
    if ( i+0*blockDim.x < N ) out[i+0*blockDim.x] = value;
    if ( i+1*blockDim.x < N ) out[i+1*blockDim.x] = value;
    if ( i+2*blockDim.x < N ) out[i+2*blockDim.x] = value;
    if ( i+3*blockDim.x < N ) out[i+3*blockDim.x] = value;
}

展开的内部循环产生四个完全独立的写入，这些写入被合并。正是这种指令级并行性使代码具有更高的指令吞吐量和更高的性能。如果您还没看过的话，我强烈建议您在几年前的GTC会议上推荐瓦西里·沃尔科夫（Vasily Volkov）的“展开平行循环”。他的演讲为为什么这种类型的循环展开是CUDA中的优化提供了理论背景。

本文收集自互联网，转载请注明来源。

如有侵权，请联系[email protected] 删除。