NVIDIA CUDA初级教程（P11）CUDA程序基本优化

文章目录

1.Parallel Reduction并行规约
2.Warp分割

1.Parallel Reduction并行规约

最优性能=有效的数据并行算法+针对GPU架构特性的优化

eg：Parallel Reduction(sum)
这个过程类似篮球锦标赛的淘汰过程：n个元素进行log(n)个回合，如何在CUDA上实现？

//累加存在shared memory内的元素，目的是提升访存性能
__shared__ float partialSum[element_num];unsigned int t = threadIdx.x;//步长：1，2，4...,步长为4，eg：第三行6->22
for (unsigned int stride = 1; stride < blockDim.x; stride*=2)
{__syncthreads();//保证每一层做完再到下一层//在同一块shared memory里面进行累加//当步长增加时，多余的线程在干什么？没事干if (t%(2*stride)==0){partialSum[t] += partialSum[t + stride];}
}

求解过程如下：

如果我们改进这个过程会怎么样？

步长修改为4（eg：第一行0+4），2（eg：第二行4+8），1

//申请共享内存, 存在于每个block中 __shared__ float partialSum[threadsPerBlock];//确定索引unsigned int i = threadIdx.x + blockIdx.x*blockDim.x;unsigned int tid = threadIdx.x;//传global mem数据到shared memorypartialSum[tid] = d_a[i];//在共享存储器中进行规约//步长：4.2.1.for (unsigned int stride = blockDim.x / 2; stride > 0; stride /= 2){__syncthreads();if (tid<stride){partialSum[tid] += partialSum[tid + stride];}}

求解过程如下：

上述两种方法的区别：

第二种可以将提前完成的线程的硬件资源释放，用来做其他的事情。

2.Warp分割

Warp分割：块内线程如何划分wrap

通晓warp分割有助于：减少分支发散，让warp尽早完工
Block被划分为以32为单位的线程组，叫做warp
warp是最基本的调度单元；
Warp里的线程一直执行相同的指令（SIMT）；
每个线程只能执行自己的代码路径；
Fermi SM有2个warp调度器（Tesla has 1）；
warp里设备切换没有时间代价；
许多warps在一起可以隐藏访存延时；
Warp分割的原则是：threadIdx连续增加的一组
（1）一维Block：threadIdx的Warp位置：第n个Warp起始线程ID：32n，结尾线程ID：32(n+1)-1，若块大小不是32的倍数，则最后一个warp将被填充空

（2）二维block：增长threadIdx意味着增长threadIdx.x，始于行threadIdx.y == 0；
（3）三维block：始于threadIdx.z==0，分割为二维block，重复增长threadIdx.z
Warp分支分散会降低性能

例如：给定warpSize=32，以下代码是否有哪个warp存在分支发散
if(threadIdx.x>15)
{// 存在 ×
}所以threadIdx要以warpSize来进行分割
if(threadIdx.x>warpSize-1)
{// 不存在 √
}

Memory Coalescing访存合并

Bank冲突

SM资源动态分割

数据预读

指令混合

循环展开

参考：链接

NVIDIA CUDA初级教程（P11）CUDA程序基本优化相关推荐

Nvidia CUDA初级教程6 CUDA编程一
Nvidia CUDA初级教程6 CUDA编程一视频:https://www.bilibili.com/video/BV1kx411m7Fk?p=7 讲师:周斌 GPU架构概览 GPU特别使用于: ...
Nvidia CUDA初级教程4 GPU体系架构概述
Nvidia CUDA初级教程4 GPU体系架构概述视频:https://www.bilibili.com/video/BV1kx411m7Fk?p=5 讲师:周斌本节内容: 为什么需要GPU 三 ...
Nvidia CUDA初级教程2 并行程序设计概述
Nvidia CUDA初级教程2 并行程序设计概述视频:https://www.bilibili.com/video/BV1kx411m7Fk?p=3 讲师:周斌本节内容: 为什么需要? 怎么做? ...
NVIDIA CUDA初级教程（P2-P3）CPU体系架构概述、并行程序设计概述
文章目录 1.CPU体系架构概述 2.并行程序设计概述 1.CPU体系架构概述现代CPU架构 CPU的定义 (1)执行指令.处理数据的器件:完成基本的逻辑和算术指令 (2)内存接口.外部设备接口 ( ...
Cuda 学习教程：Cuda 程序初始化
Cuda程序初始化目前,cuda里面没有对设备的初始化函数InitDevice(),只能每次调用的api函数的时候,加载设备的上下文,自动进行初始化,这将带来问题: First函数调用的时候,需要自 ...
linux运行cobol语言,COBOL语言初级教程(2)--COBOL程序
在COBOL-85标准前,COBOL程序要求全部使用大写.因此许多旧程序都是全部用大写字母编写的,但COBOL-85允许采用小写字母,因为利用大小写混合用提高程序的可读性. 1.COBOL程序的部(D ...
php 用户认证,用户认证-php初级教程手册,php程序员教程网
在专门?? Web 网站上,常常会需要用户的帐号及密码,也就是身份确认的步骤.早期的 NCSA httpd 服务器并没有提供这项用户确认的功能,Webmaster 只能用手工打造一个身份确认的 CGI ...
最简单、实用的cuda安装教程！！！（nvidia官方渠道下载）
网上教程一大推,讲了一大堆,也没解释原理,实用的没几个,自己总结的,比较简单 note:无需卸载原机器驱动,无需禁用nouveau驱动 attention:以下内容为有sudo权限安装教程,没有sud ...
Ubuntu + nvidia驱动+ cuda安装教程以及重装问题
操作系统:ubuntu 18 nvidia驱动版本: 515 cuda: 11.7 在此方面掉过坑,写此篇文章mark一下.如果遇到同样问题的小伙伴可以参考一下. 第一次在ubuntu上尝试安装nvi ...

NVIDIA CUDA初级教程（P11）CUDA程序基本优化

文章目录

1.Parallel Reduction并行规约

2.Warp分割

NVIDIA CUDA初级教程（P11）CUDA程序基本优化相关推荐

最新文章

热门文章