CUDA编程(八)树状加法

本文介绍了一种利用树状加法在CUDA编程中优化加和操作的方法,通过并行计算显著提高了运算效率。

摘要生成于 C知道 ,由 DeepSeek-R1 满血版支持, 前往体验 >

CUDA编程(八)

树状加法

上一篇博客我们介绍了ShareMemory和Thread同步,最后利用这些知识完成了block内部线程结果的加和,减轻了CPU的负担,结果还是比较令人满意的,但是block的加和工作是使用一个thread0单线程完成的,这点还是有待改进的。

那么这个单线程的加法部分如何解决呢?我们知道GPU上的程序只有并行才能发挥其优势,所以我们自然想到这个加法能不能并行呢?答案当然是可行的,我们可以利用树状加法的方式将加法并行,这也体现了我们之前提到的,一个优秀的CUDA程序是需要一个优秀的算法为基础的。

树状加法

我们传统的加法 a + b + c + d ,只能在一个线程上进行,但是我们也很容易想到,如果把加法分成多步执行,比如先算 a+b,c+d,再把他们的结果相加,通过这样的方式我们就可以把任务分开,也就是可以并行了,这就是树状加法:

这里写图片描述

通过这种方式我们就可以把256个数的加法进行并行了。

树状加法的实现

这里写图片描述

上图是树状加法的一个示意图,示意图中第一排每一个格子就是一个线程的结果,保存在shared[],暂且把shared[0]简写为 sh0,我们可以清楚的看到计算的过程:

sh0=sh0+sh1, sh2=sh2+sh3, sh4=sh4+sh5...

同步

sh0=sh0+sh2;sh4=sh4+sh6...

同步

...

最后结果在sh0里

其实树状加法可以写成一个很简单的while循环:

int offset = 1, mask = 1;

while(offset < THREAD_NUM)
{ 
    if((tid & mask) == 0)  
    {  
         shared[tid] += shared[tid + offset];  
    }  

    offset += offset; 
    mask = offset + mask; 
    __syncthreads(); 

}

下面我们就来看看这个while循环:

注意& 按位“与”,只有1&1 = 1

tid=0时,mask = 1,0&1=0,所以shared[0] = sh0 + sh1,完成第一步的前两个相加。

tid=1时,mask = 1,1&1=1,不作运算。

tid=2时,mask = 1,10&01 = 00,所以shared[2] = sh2 + sh3

tid=3时,mask = 1,11&01 = 01,不作运算。

可以看出来这是第一层的计算

同步之后第二层:

offset=1+1=2,mask=2+1=3;

tid=0时,mask = 3,0&11=0,所以shared[0] = sh0 + sh2,完成第二步的前两个相加。

tid=1时,mask = 3,1&11=1,不作运算。

tid=2时,mask = 3,10&11 = 10,不作运算。

tid=3时,mask = 3,11&11 = 01,不作运算。

tid=4时,mask = 3,100&011 = 000,所以shared[4] = sh4 + sh6

后面都以此类推,直到offset 大于等于线程数就跳出了

最终的结果就在shared[0]内,所以下一步用线程0把结果保存就OK了:

if(tid == 0) { result[bid] = shared[0]; }

所以比起上一版的程序,我们只用改动核函数里面的加和部分就OK了,下面是改好的核函数:

核函数:

// __global__ 函数 (GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{

    //声明一块共享内存
    extern __shared__ int shared[];

    //表示目前的 thread 是第几个 thread(由 0 开始计算)
    const int tid = threadIdx.x;

    //表示目前的 thread 属于第几个 block(由 0 开始计算)
    const int bid = blockIdx.x;

    shared[tid] = 0;

    int i;

    //记录运算开始的时间
    clock_t start;

    //只在 thread 0(即 threadIdx.x = 0 的时候)进行记录,每个 block 都会记录开始时间及结束时间
    if (tid == 0) time[bid] = clock();

    //thread需要同时通过tid和bid来确定,同时不要忘记保证内存连续性
    for (i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {

        shared[tid] += num[i] * num[i] * num[i];

    }

    //同步 保证每个 thread 都已经把结果写到 shared[tid] 里面
    __syncthreads();

    //树状加法
    int offset = 1, mask = 1;

    while (offset < THREAD_NUM)
    {
        if ((tid & mask) == 0)
        {
            shared[tid] += shared[tid + offset];
        }

        offset += offset;
        mask = offset + mask;
        __syncthreads();

    }

    //计算时间,记录结果,只在 thread 0(即 threadIdx.x = 0 的时候)进行,每个 block 都会记录开始时间及结束时间
    if (tid == 0)
    { 
        result[bid] = shared[0];
        time[bid + BLOCK_NUM] = clock(); 
    }

}

运行结果:

这里写图片描述

我们看到比起上一次没用树状加法的144185个周期,这次只用了133738个周期,总的来说这个结果还是非常不错的,甚至和完全不在GPU上加和的程序速度差不多,这是因为,在完全不在 GPU 上进行加总的版本,写入到 global memory 的数据数量很大(8192 个数字),这对效率也会有影响。所以,这一版程序不但在 CPU 上的运算需求降低,在 GPU 上也能跑的更快~

总结:

这篇博客我们主要介绍了怎么去把加法进行并行,利用树状加法,最终并行了之前效率比较差的加和部分,到这里为止,这个程序的一般性优化也做完了,因为程序也很简单,所以很多的方面都无法体现,比如之前提过的GPU运算的一大问题在于精度,还有尽可能减少访存这些方面的优化都没有体现出来,所以下一篇博客我们准备真正向应用CUDA靠拢,去进行矩阵计算~

希望我的博客能帮助到大家~

参考资料:《深入浅出谈CUDA》

<think>好的,我现在需要解释CUDA中的规约(Reduce)操作。首先,我得回忆一下规约的基本概念。规约是一种将数据集通过某种操作(比如加法、最大值、最小值等)合并成单个结果的过程。在CUDA中,高效实现规约对性能至关重要,尤其是在处理大规模数据时。 用户可能想了解CUDA规约的实现原理和优化方法。我需要从基础开始,逐步深入。首先,应该介绍并行规约的基本思想:将数据分成块,每个线程块处理一部分数据,通过逐步合并部分结果来得到最终结果。 然后,应该提到线程块内的规约。使用共享内存来加速数据访问,因为共享内存的访问速度比全局内存快得多。这里可能会涉及线程同步的问题,需要用__syncthreads()来确保所有线程完成数据加载后再进行计算。 接下来,需要解释不同的优化策略。比如顺序访问与交错访问的区别,交错访问可以减少共享内存的bank冲突。Bank冲突会导致多个线程同时访问同一个bank,从而降低性能,因此优化访问模式很重要。 可能还需要提到循环展开,减少分支预测的开销。循环展开可以通过减少循环次数来提高指令级并行度,从而加速计算。 然后,要讨论多级规约。当数据量超过一个线程块能处理的范围时,需要在全局内存中进行多级规约。第一级每个线程块生成一个部分结果,第二级再对这些部分结果进行规约,直到得到最终结果。 另外,需要考虑线程块的数量和线程数的选择。线程块数量通常与SM(流多处理器)的数量相关,以充分利用GPU并行计算能力。每个线程块中的线程数一般是256或512,这取决于GPU架构和共享内存的大小。 可能还需要提到原子操作,但要注意原子操作的效率较低,通常只在最后一步不得已时使用。例如,在全局内存中对部分结果进行原子加操作,但应尽量避免频繁使用。 最后,可以给出一个简单的代码示例,展示如何实现一个加法规约。示例中应包括内核函数的实现,共享内存的使用,以及逐步规约的过程。同时,指出代码中的优化点,比如避免bank冲突和循环展开。 需要注意的是,确保LaTeX公式的正确使用。比如,提到时间复杂度时,使用$$O(\log n)$$来表示对数复杂度。行内公式如每个线程处理$n$个元素,需要用$...$包裹。 可能用户还会关心性能优化的具体数据,比如使用优化后的规约比简单实现快多少倍。这时候可以提到一些实际测试结果,比如优化后的速度提升可达10倍以上,但具体数值取决于GPU架构和数据规模。 此外,需要提醒用户在实际编程中要注意共享内存的大小限制,以及线程块和网格的配置参数。不同的GPU架构可能有不同的最优配置,需要根据具体情况调整。 总结一下,解释CUDA规约的步骤应该是:基本概念、线程块内规约、共享内存使用、优化策略(访问模式、循环展开)、多级规约、线程配置和注意事项。最后通过代码示例和性能考量来巩固理解。</think>在CUDA中,**规约(Reduce)** 是一种将大量数据通过特定操作(如求和、求最大值等)合并为单个结果的并行算法。以下是逐步解析其核心原理和优化方法: --- ### **1. 基本思想** - **目标**:将数组元素通过二元操作(如加法)合并为一个值。 - **并行策略**:将数据划分为块,每个线程块处理局部规约,再对局部结果进行全局规约。 - **时间复杂度**:理想情况下为 $$O(\log n)$$,优于串行的 $$O(n)$$。 --- ### **2. 线程块内规约** #### **(1) 共享内存加速** - **步骤**: 1. 每个线程块将全局数据加载到共享内存。 2. 使用树状规约逐步合并数据(如二分法相加)。 3. 最终结果由线程块的第一个线程写入全局内存。 - **代码片段**: ```cpp __global__ void reduce_kernel(float *d_in, float *d_out) { extern __shared__ float s_data[]; int tid = threadIdx.x; int i = blockIdx.x * blockDim.x + threadIdx.x; s_data[tid] = d_in[i]; __syncthreads(); for (int s = blockDim.x/2; s > 0; s >>= 1) { if (tid < s) { s_data[tid] += s_data[tid + s]; } __syncthreads(); } if (tid == 0) d_out[blockIdx.x] = s_data[0]; } ``` #### **(2) 避免共享内存Bank冲突** - **问题**:连续线程访问同一共享内存bank会导致串行化。 - **优化**:使用交错访问(Strided Access)或调整步长。 --- ### **3. 多级规约** - **场景**:当数据规模超过单线程块容量时。 - **步骤**: 1. 第一级规约:每个线程块生成一个局部结果。 2. 第二级规约:对局部结果再次调用规约内核。 3. 重复直到得到最终结果。 --- ### **4. 优化策略** #### **(1) 循环展开(Unrolling)** - **作用**:减少循环开销,提高指令级并行度。 - **示例**:手动展开循环以减少条件判断。 #### **(2) 向量化访存** - **方法**:使用`float4`或`int4`类型一次性读取多个数据。 #### **(3) 原子操作(慎用)** - **适用场景**:最终合并少量局部结果时。 - **注意**:全局原子操作(如`atomicAdd`)性能较差,尽量避免高频使用。 --- ### **5. 性能影响因素** - **线程块大小**:通常选择256或512线程,需适配共享内存容量。 - **网格大小**:与GPU流多处理器(SM)数量匹配,确保满载。 - **数据对齐**:确保全局内存访问对齐,提高吞吐量。 --- ### **6. 代码示例(优化版)** ```cpp template <unsigned blockSize> __device__ void warpReduce(volatile float* s_data, int tid) { if (blockSize >= 64) s_data[tid] += s_data[tid + 32]; if (blockSize >= 32) s_data[tid] += s_data[tid + 16]; if (blockSize >= 16) s_data[tid] += s_data[tid + 8]; if (blockSize >= 8) s_data[tid] += s_data[tid + 4]; if (blockSize >= 4) s_data[tid] += s_data[tid + 2]; if (blockSize >= 2) s_data[tid] += s_data[tid + 1]; } __global__ void optimized_reduce(float *d_in, float *d_out) { extern __shared__ float s_data[]; int tid = threadIdx.x; int i = blockIdx.x * (blockDim.x * 2) + threadIdx.x; s_data[tid] = d_in[i] + d_in[i + blockDim.x]; __syncthreads(); if (blockSize >= 512) { if (tid < 256) s_data[tid] += s_data[tid + 256]; __syncthreads(); } if (blockSize >= 256) { if (tid < 128) s_data[tid] += s_data[tid + 128]; __syncthreads(); } // ... 类似处理64,32,16等 if (tid < 32) warpReduce<blockSize>(s_data, tid); if (tid == 0) d_out[blockIdx.x] = s_data[0]; } ``` --- ### **7. 性能对比** - **朴素规约**:约5-10 GB/s。 - **优化后规约**:可达50-100 GB/s(取决于GPU架构)。 --- ### **总结** CUDA规约的关键在于: 1. 利用共享内存减少全局访问。 2. 通过树状规约降低计算复杂度。 3. 优化内存访问模式(如避免bank冲突)。 4. 合理配置线程块和网格大小。 实际开发中需结合`nvprof`工具分析性能瓶颈,并针对特定硬件调整参数。
评论 11
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值