CUDA程序基本优化

最新推荐文章于 2023-10-21 14:43:53 发布
最新推荐文章于 2023-10-21 14:43:53 发布
阅读量475 收藏 6
点赞数 2
分类专栏： cuda并行程序设计 文章标签： 多线程 cuda GPU编程
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/sinat_38368658/article/details/105381590
版权

写在前面：

  • NVIDIA CUDA初级教程笔记 p11~p12
  • 连接：https://www.bilibili.com/video/BV1kx411m7Fk?p=12

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

文章目录

1.Parallel Reduction 并行规约

在这里插入图片描述

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

//累加存在shared memory内的元素
__shared__ float partialSum[element_num];
	
unsigned int t = threadIdx.x;

//步长：1，2，4...
for (unsigned int stride = 1; stride < blockDim.x; stride*=2)
{
	__syncthreads();//保证每一步做完之后再进行下一步
	//在同一块shared memory里面进行累加
	//当步长增加时，多余的线程在干什么？没事干
	if (t%(2*stride)==0)
	{
		partialSum[t] += partialSum[t + stride];
	}
}

在这里插入图片描述
在这里插入图片描述

在这里插入图片描述

在这里插入图片描述
在这里插入图片描述

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

在这里插入图片描述

在这里插入图片描述
在这里插入图片描述

在这里插入图片描述

每一轮所需要的线程数依然是减半的，但是线程所处的位置不同

在这里插入图片描述

在这里插入图片描述

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

元素求和源代码

#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdlib.h>

/*
 *处理有N个元素的并行规约
 */
# define N 2048
# define threadsPerBlock 512
# define blocksPerGrid (N+threadsPerBlock-1)/threadsPerBlock//4

__global__ void ReductioinSum_0(float* d_a, float* d_partial_sum)
{
	//申请共享内存, 存在于每个block中 
	__shared__ float partialSum[threadsPerBlock];
	
	//确定索引
	unsigned int i = threadIdx.x+blockIdx.x*blockDim.x;
	unsigned int tid = threadIdx.x;

	//传global mem数据到shared memory
	partialSum[tid] = d_a[i];

	//在共享存储器中进行规约
	//步长：1，2，4...
	for (unsigned int stride = 1; stride < blockDim.x; stride*=2)
	{
		__syncthreads();//保证每一步做完之后再进行下一步
		//在同一块shared memory里面进行累加
		//当步长增加时，多余的线程在干什么？没事干
		if (tid%(2*stride)==0)
		{
			partialSum[tid] += partialSum[tid + stride];
		}
	}
	//将当前block的计算结果写回输出数组
	if (tid == 0)
		d_partial_sum[blockIdx.x] = partialSum[0];
}


__global__ void ReductioinSum_01(float* d_a, float* d_partial_sum)
{
	//申请共享内存, 存在于每个block中 
	__shared__ float partialSum[threadsPerBlock];

	//确定索引
	unsigned int i = threadIdx.x + blockIdx.x*blockDim.x;
	unsigned int tid = threadIdx.x;

	//传global mem数据到shared memory
	partialSum[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];
		}
	}
	//将当前block的计算结果写回输出数组
	if (tid == 0)
		d_partial_sum[blockIdx.x] = partialSum[0];
}


int main(void)
{
	float a[N],partial_sum[blocksPerGrid];
	//初始化示例数据
	for (int i = 0; i<N; i++)
	{
		a[i] = i;
	}
	int size = N*sizeof(float);

	//分配显存空间
	float *d_a,*d_partial_sum;
	cudaMalloc((void**)&d_a, size);
	cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);

	cudaMalloc((void**)&d_partial_sum, blocksPerGrid*sizeof(float));
	
	ReductioinSum_01 <<<blocksPerGrid, threadsPerBlock >> >(d_a, d_partial_sum);

	cudaMemcpy(partial_sum, d_partial_sum,blocksPerGrid*sizeof(float), cudaMemcpyDeviceToHost);

	float sum = 0;
	for (int i = 0; i < blocksPerGrid; i++) {
		sum += partial_sum[i];
	}
	printf("all sum = %.2f", sum);
	
	cudaFree(d_a);
	cudaFree(d_partial_sum);

	return 0;
}

2. Warp分割

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

通晓warp分割有助于：减少分支发散，让warp尽早完工

在这里插入图片描述

在这里插入图片描述

  • Block被划分为以32为单位的线程组，叫做warp
  • warp是最基本的调度单元
  • Warp里的线程一直执行相同的指令（SIMT）
  • 每个线程只能执行自己的代码路径
  • Fermi SM有2个warp调度器（Tesla has 1）
  • warp里设备切换没有时间代价
  • 许多warps在一起可以隐藏访存延时

Warp分割的原则是：threadIdx连续增加的一组

在这里插入图片描述
在这里插入图片描述

以行为主元的情况：

在这里插入图片描述
在这里插入图片描述

warp分支分散：会降低性能

在这里插入图片描述

例如：给定warpSize=32，以下代码是否有哪个warp存在分支发散

if(threadIdx.x>15)
{
  // 存在 ×
}

if(threadIdx.x>warpSize-1)
{
  // 不存在 √
}

在这里插入图片描述

第二种更好

在这里插入图片描述
在这里插入图片描述

在第一轮，右边的warp2和warp3可以腾出来做其他的事情
在这里插入图片描述
在第2轮，右边的wap1 warp2和warp3可以腾出来做其他的事情
在这里插入图片描述

3.Memory Coalesing 访存合并

上文了解了整个线程的调度和warp的切分后，我们来关注存储优化问题。

CPU-GPU 数据传输最小化

  • Host <- - > device数据传输带宽远低于global memory
  • 减少传输
    • 中间数据直接在GPU分配，操作，释放
    • 有时更适合在GPU进行重复运算
    • 如果没有减少数据传输的话，将CPU代码一直到GPU可能无法提升性能
  • 组团传输
    • 大块传输好于小块：10微妙延迟，8GB/s => 如果数据小于80KB，性能将受延迟支配
  • 内存传输与计算时间重叠
    • 双缓存

Coalesing合并

Global memory 延时：400 ~ 800 cycles
最重要的影响因子！
在Fermi，global memory默认缓存于 一级缓存L1
通过给nvcc指令设置参数 ”-Xptxas -dlcm=cg“可以绕过一级缓存L1：只缓存于二级缓存L2
如果缓存：warp的读写请求落到L1 cache line，只需一次传输
#transaction = #L1 line accessed
如果没有缓存：有一些合并原则
但是传输大小可以减至32字节块

Memory Coalescing 合并访存

在这里插入图片描述
“相邻的人搬相邻的砖哈哈哈哈”
在这里插入图片描述
在这里插入图片描述在这里插入图片描述
合并举例

  • 小型kernel拷贝数据时的有效带宽
    • 偏移和步长对性能的影响
  • 2款GPUs
    • GTX 280：compute capability 1.3；峰值带宽141GB/2
    • FX 5600：compute capability 1.0；峰值带宽77GB/s

偏移量的影响
在这里插入图片描述
步长的影响
在这里插入图片描述
在这里插入图片描述

Shared memory

  • 比global memory快上百倍
  • 可以通过缓存数据减少global memory访存次数
  • 线程可以同通过shared memory协作
  • 用来避免不满足合并条件的访存
    • 读入shared memory重排顺序，从而支持合并寻址

4.Bank 冲突

### shared memory架构

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
按行读，按列写，读写总有一个是不合并的
在这里插入图片描述
在这里插入图片描述
读和写都实现合并访存
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

Texture 纹理

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

总结

如果遵行一些简单的原则，GPU硬件在数据可并行计算问题上，可以达到很好的性能：

  • 有效利用并行性
  • 尽可能合并内存访问
  • 利用shared memory
  • 开发其他存储空间
    • Texture
    • Constant
  • 减少bank冲突

5.SM资源动态分割

在这里插入图片描述
在片描述
在这里插入图片描述
Performance Cliff: 增加资源用量后导致并行性急剧下降。例如,增加寄存器数量，除非为了隐藏global memory访存延迟
在这里插入图片描述

kernel 启动参数配置介绍

Grid Size试探法

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

Occupancy 占用率

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

6.数据预读

用来隐藏访存延时的一个小手段
在这里插入图片描述

float m = Md[i];//read global memory
float f = a*b+c*d;//执行指令，不依赖于读内存的操作
float f2 = m*f;//在上一行执行足够多的warp隐藏访存延时以后，再使用global memory

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在上图中
load next tile into register //为循环的下一次迭代预读
accumulate dot product //这些指令被足够多的线程执行，从而隐藏了预读内存产生的延时

7.指令混合

指令吞吐量优化
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

8.循环展开

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
缺点：可扩展性差，blocksize变化会带来影响。

Elylicery
关注
  • 2
    点赞
  • 踩
  • 6
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
CUDA基本优化方法
nn1997729的博客
06-24 3190
一、基于线程的优化方法 1.选取合适的gridDim和blockDim blockDim最好为32的整数倍：因为执行指令的基本单位为线程束，线程束内的所有线程统一执行广播下来的命令，而线程束的线程数量基本为32。当block被分到SM中去，其会被划分为多个线程束，若blockDim!=线程束内线程数整数倍，则会造成线程的浪费。 2.减少存在分支的if 因为线程束中的所有线程执行同一条的指令，若出现存在出现分支的if，即线程束内的部分线程符合if条件，部分符合else条件，因cuda编译不具备分支预测能力，则
Tesla架构下的CUDA程序优化
darkstorm2111203的专栏
07-30 5409
CUDA优化的最终目的是：在最短的时间内，在允许的误差范围内完成给定的计算任务。在这里，“最短的时间”是指整个程序运行的时间，更侧重于计算的吞吐量，而不是单个数据的延迟。在开始考虑使用GPU和CPU协同计算之前，应该先粗略的评估使用CUDA是否能达到预想的效果，包括以下几个方面：精度：目前GPU的单精度性能要远远超过双精度性能，整数乘法、求模、求余等运算的指令吞吐量也较为有限。在科学计算中，由
CUDA程序优化
lanyijia的专栏
05-30 1664
<br />CUDA程序优化应该考虑的点：<br />精度：只在关键步骤使用双精度，其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡；<br />           延迟：需要首先缓冲一部分数据，缓冲的大小应该可以保证每个内核程序处理的一批数据能够让GPU慢负荷工作；<br />           计算量：计算量太小的程序使用CUDA很不合算；当需要计算的问题的计算密集度很低的时候，执行计算的时间远远比IO花费的时间短，整个程序的瓶颈出现在PCI-E带宽上。<br />优秀的CUDA程序特征：<b
CUDA学习笔记之程序优化
yujiflying的专栏
12-16 972
CUDA程序优化 CUDA程序优化应该考虑的点： 精度：只在关键步骤使用双精度，其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡；            延迟：需要首先缓冲一部分数据，缓冲的大小应该可以保证每个内核程序处理的一批数据能够让GPU慢负荷工作；            计算量：计算量太小的程序使用CUDA很不合算；当需要计算的问题的计算密集度很低的时候，执行计算的时间远远比
cuda优化策略
u010794523的专栏
08-28 1007
cuda优化策略2 1. memory coalescing，保证内存融合。因为global memory在CC为1.x上是按照half wrap进行访问读写的，而在2.x上是按照wrap进行访问读写的。在显存中，有多个存储器控制器，负责对显存的读写，因此，一定要注意存储器控制器的负载均衡问题。每一个存储器控制器所控制的那片显存中的地址空间称为一个分区。连续的256Byte数据位
CUDA程序的优化
wufenxia的专栏
10-21 398
CUDA程序的优化 CUDA程序员优化的目的，是以最短的时间，在允许的误差范围内完成给定的计算任务。在这里，“最短的时间”是指整个程序的运行时间，更侧重于计算的吞吐量，而不是单个数据的延迟。在开始考虑GUP和CPU协同计算之前，应该先粗略地评估使用CUDA是否能达到预想效果，包括以下几个方面： 1.精度 目前，GPU的单精度计算性能要远远超过双精度计算性能，整数乘法、除...
CUDA——性能优化（一）
01-20
本文将深入探讨CUDA全局内存的合并访问以及与之相关的Warp基础知识，旨在帮助开发者更好地理解和优化CUDA程序。 首先，我们要明白Warp的概念。Warp是CUDA中基本的执行单元，由32个线程组成，它们在同一时间执行相同...
cuda_exp:示例CUDA程序
04-21
CUDA（Compute Unified Device Architecture）是由NVIDIA推出的一种并行计算平台和编程模型...通过分析和运行这个示例，学习者可以直观地了解CUDA程序的基本框架和工作流程，为进一步深入学习和应用CUDA编程打下基础。
mul.zip_cuda程序_矩阵乘法
09-20
9. **性能分析**：CUDA程序的性能优化通常需要借助NVIDIA的Nsight工具或其他性能分析工具，监控和分析CUDA事件，找出瓶颈并进行优化。 10. **CUDA SDK和CUDA Toolkit**：CUDA开发通常需要安装CUDA SDK和CUDA ...
cuda程序示例代码.rar
05-24
这个“cuda程序示例代码.rar”压缩包包含了CUDA编程的基础、进阶以及实用性示例，是学习和理解CUDA编程的宝贵资源。 1. **CUDA入门代码**：这部分通常包含简单的CUDA HelloWorld程序，用于介绍CUDA的基本结构，如...
cuda立体匹配——中值滤波
06-20
综上所述，本课题深入探讨了立体匹配的基本原理和实现技术，结合CUDA的并行计算能力，优化了计算效率，并使用中值滤波来提高匹配结果的质量。这是一个典型的计算机视觉和并行计算的交叉研究，对于理解和应用高性能...
CUDA程序性能调优
yu132563的专栏
12-28 833
介绍了GPU的结构以及资源的控制要素（GPU硬件结构和程序具体参数设置_yu132563的专栏-CSDN博客），下面就可以对CUDA进行程序的调优，从而在不同的GPU上面运行同一个CUDA程序的参数设置方法。 对于一个CUDA kernel function而言，其通常由如下几个部分组成： kernel function paras local variables shared memory with __syncthreads__ call device function call loop/
CUDA: CUDA程序优化步骤
923723914
10-12 253
在《runtime API创建CUDA程序》中，我们做了一个计算一大堆数字的平方和的程序。不过，我们也提到这个程序的执行效率并不理想。当然，实际上来说，如果只是要做计算平方和的动作，用CPU做会比用 GPU 快得多。这是因为平方和的计算并不需要太多运算能力，所以几乎都是被内存带宽所限制。因此，光是把数据复制到显卡内存上的这个动作，所需要的时间，可能已经和直接在CPU上进行计算差不多了。 　　不过...
CUDA第一个程序优化（若干个线程）
yanghangjun的专栏
05-17 997
#include #include #include#define TOTAL_NUM 50000#define Thread_num 500bool CUDA_initial(void){ int i; int device_count; if( cudaGetDeviceCount(&device_count) ) {  printf(" There is zero device beyond
CUDA编程入门系列（十一）CUDA程序优化技巧
qq_45788429的博客
10-21 231
我们在设计CUDA程序的时候，要对线程块的个数进行考虑。因为GPU中流处理器SM的数量是相对固定的，所以我们应该尽量的将多个block放到同一个SM当中（至少保证每个SM中都有一个块），使得SM时刻处于活跃状态。因为SM是以warp为单位的，那么我们就要尽量的保证块的大小是32的倍数，使得所有的warp中所有的线程都处于活跃状态。如果块上的线程多的话，就可以隐藏一些内存的延迟。但是，如果一个块上的线程越多，每个线程拥有的寄存器大小就越小。优化原则：最大化并行执行。共享内存：防止bank conflict。
CUDA程序优化[转]
03-23 72
本文章转载自CSDN，原博客地址http://blog.csdn.net/ouczoe/article/details/5137063 CUDA程序优化 CUDA程序优化应该考虑的点： 精度：只在关键步骤使用双精度，其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡； 延迟：需要首先缓冲一部分数据，缓冲的大小应该可以保证每个内核程序处理的一批数据能够让G...
CUDA学习之CUDA程序优化
weixin_30425949的博客
06-25 101
一 测量程序运行时间 1主机端测时 由于CUDA API函数都是异步的，为了准确测量CUDA调用运行的时间，首先要使用cudaThreadSynchronize(),同步cpu与gpu之后，才能结束测时。 2设备端测时 使用clock()函数，这个函数测的结果是一个block在gpu中上下文保持的时间，而不是实际执行的时间。（一般大于执行时间） 二任务划分 1原则 划分为cp...
5G智慧农业大数据数字乡村乡村振兴整体建设方案.pdf
最新发布
08-24
5G智慧农业大数据数字乡村乡村振兴整体建设方案.pdf
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值