编译代码性能优化实践：理解循环展开(pragma unroll)

引言：CUDA的矩阵乘优化经常见到 pragma unroll 的使用，本文通过简单的示例，展示了CPU和CUDA对循环展开前后的性能表现，来通俗理解循环展开的优化策略。

一、什么是循环展开？

        简单理解：将代码中的for循环展开，减少循环次数；循环展开的本质是，利用CPU指令级并行，来降低循环的开销，当然，同时也有利于指令流水线的高效调度

优点

  • 提高缓存命中（cache hit）率，增加循环体内语句并发执行的可能性（需要循环体内语句不相关）；
  • 减少分支预测失败的可能性，提高性能

缺点

  • 程序代码膨胀、代码可读性降低
  • 消耗较多寄存器缓存（SM里的寄存器大小是有限的，SM会根据一个块需要消耗的寄存器大小和线程的个数去分配该SM上块的个数，当一个SM连一个块都分配不了时，就会导致内核启动不了）

二、循环展开的使用

        循环展开在CPU和CUDA端都可以使用，但在CPU端可以由程序员手动实现，也可以通过成熟的编译器实现优化。# pragma unroll 是常用在CUDA编程的核函数中对for循环展开的使用方法。

        下面通过计算0-100000个数字累加的和为例，展示CPU和CUDA下的对循环展开使用的理解。

CPU端

1）原始不展开

void test_cpu_1(int count, const char* name)
{  int sum = 0;auto start = std::chrono::system_clock::now();for(int i = 0;i < count;i++){  sum += i;}auto end = std::chrono::system_clock::now();auto dura = std::chrono::duration_cast<std::chrono::microseconds> (end - start);std::cout << name <<" cost time: "<< dura.count() << " microseconds" << std::endl;printf("                                                   sum = %d\n",sum);
}

2）循环展间隔4次

void test_cpu_2(int count, const char* name)
{int sum = 0;auto start = std::chrono::system_clock::now();for(int i=0; i<count; i+=4){sum += i;sum += i+1;sum += i+2;sum += i+3;}auto end = std::chrono::system_clock::now();auto dura = std::chrono::duration_cast<std::chrono::microseconds> (end - start);std::cout << name <<" cost time: "<< dura.count() << " microseconds" << std::endl;printf("                                                   sum = %d\n",sum);}

3）循环展开间隔4次，优化循环内的数据依赖关系

        上面虽然实现了循环展开，但是循环体内是的4行代码之间共用sum地址, 所以是有先后依赖的，如果我们把他们之间的依赖关系去掉，则能进一步提升代码性能。

void test_cpu_3(int count, const char* name)
{int sum = 0;int sum1=0,sum2=0,sum3=0, sum4=0;auto start = std::chrono::system_clock::now();for(int i=0;i < count;i+=4){sum1 += i;sum2 += i+1;sum3 += i+2;sum4 += i+3;}sum = sum1+sum2+sum3+sum4;auto end = std::chrono::system_clock::now();auto dura = std::chrono::duration_cast<std::chrono::microseconds> (end - start);std::cout << name <<" cost time: "<< dura.count() << " microseconds" << std::endl;printf("                                                   sum = %d\n",sum);}

CUDA端

CUDA则主要对比使用# pragma unroll前后的区别。

1）原始不展开

__global__ void progam_kernel1(int* sum, int count)
{for(int i = 0;i < count;i++){  *sum += i;}}

2）使用循环展开

__global__ void progam_kernel2(int* sum, int count)
{#pragma unrollfor(int i = 0;i < count;i++){  *sum += i;}
}

性能分析与测试接口实现

        上面各种对比的方法测试时间如下，可以看到CPU端循环展开比原始不展开时间减少接近一半，而优化后的循环展开时间又减少将近一半。CUDA端使用pragma unroll后，时间减少三分之二。

cpu origin cost time: 1079 microsecondssum = 704982704
cpu pragma unroll cost time: 678 microsecondssum = 704982704
cpu pragma unroll_1 cost time: 374 microsecondssum = 704982704
cuda origin cost time: 18 microsecondssum = 704982704
cuda pragma unroll cost time: 6 microsecondssum = 704982704

编译如下，因为把kernel函数写在一起了，所以用.cu为后缀命名。

nvcc -o test test_performance.cu

下面是总体实现的代码 

// file name: test_performance.cu
#include <stdio.h>
#include <cuda_runtime.h>
#include "cublas_v2.h"
#include <unistd.h>
#include <chrono>
#include <iostream>
#include <string>
using namespace std;void test_cpu_1(int count, const char* name)
{  int sum = 0;auto start = std::chrono::system_clock::now();for(int i = 0;i < count;i++){  sum += i;}auto end = std::chrono::system_clock::now();auto dura = std::chrono::duration_cast<std::chrono::microseconds> (end - start);std::cout << name <<" cost time: "<< dura.count() << " microseconds" << std::endl;printf("                                                   sum = %d\n",sum);
}void test_cpu_2(int count, const char* name)
{int sum = 0;auto start = std::chrono::system_clock::now();for(int i=0; i<count; i+=4){sum += i;sum += i+1;sum += i+2;sum += i+3;}auto end = std::chrono::system_clock::now();auto dura = std::chrono::duration_cast<std::chrono::microseconds> (end - start);std::cout << name <<" cost time: "<< dura.count() << " microseconds" << std::endl;printf("                                                   sum = %d\n",sum);}void test_cpu_3(int count, const char* name)
{int sum = 0;int sum1=0,sum2=0,sum3=0, sum4=0;auto start = std::chrono::system_clock::now();for(int i=0;i < count;i+=4){sum1 += i;sum2 += i+1;sum3 += i+2;sum4 += i+3;}sum = sum1+sum2+sum3+sum4;auto end = std::chrono::system_clock::now();auto dura = std::chrono::duration_cast<std::chrono::microseconds> (end - start);std::cout << name <<" cost time: "<< dura.count() << " microseconds" << std::endl;printf("                                                   sum = %d\n",sum);}__global__ void progam_kernel1(int* sum, int count)
{for(int i = 0;i < count;i++){  *sum += i;}}__global__ void progam_kernel2(int* sum, int count)
{#pragma unrollfor(int i = 0;i < count;i++){  *sum += i;}
}void test_cuda_1(int count, const char* name)
{int sum =0;int* g_sum;cudaMalloc((void **)&g_sum, sizeof(int) * 1);cudaMemcpy(g_sum, &sum, 1 * sizeof(int),cudaMemcpyHostToDevice);auto start = std::chrono::system_clock::now();progam_kernel1<<<1,1>>>(g_sum, count); //调用核函数auto end = std::chrono::system_clock::now();auto dura = std::chrono::duration_cast<std::chrono::microseconds> (end - start);std::cout << name <<" cost time: "<< dura.count() << " microseconds" << std::endl;cudaMemcpy(&sum, g_sum, sizeof(int) * 1, cudaMemcpyDeviceToHost);printf("                                                   sum = %d\n",sum);cudaFree(g_sum); }void test_cuda_2(int count, const char* name)
{int sum =0;int* g_sum;cudaMalloc((void **)&g_sum, sizeof(int) * 1);cudaMemcpy(g_sum, &sum, 1 * sizeof(int),cudaMemcpyHostToDevice);auto start = std::chrono::system_clock::now();progam_kernel2<<<1,1>>>(g_sum, count); //调用核函数auto end = std::chrono::system_clock::now();auto dura = std::chrono::duration_cast<std::chrono::microseconds> (end - start);std::cout << name <<" cost time: "<< dura.count() << " microseconds" << std::endl;cudaMemcpy(&sum, g_sum, sizeof(int) * 1, cudaMemcpyDeviceToHost);printf("                                                   sum = %d\n", sum);cudaFree(g_sum);  }void test_performance()
{int count =100000;std::string s1 ="cpu origin";std::string s2 = "cpu pragma unroll";std::string s21 = "cpu pragma unroll_1";std::string s3 = "cuda origin";std::string s4 = "cuda pragma unroll";test_cpu_1(count, s1.c_str());test_cpu_2(count, s2.c_str());test_cpu_3(count, s21.c_str());test_cuda_1(count, s3.c_str());test_cuda_2(count, s4.c_str());}int main(int argc, char *argv[]) 
{test_performance();return 0;}

借助编译器的性能优化

        程序员针对CPU端编写代码时候，可以使用上面的循环展开实现，实际上在c/c++的编译器已经非常成熟，针对这种代码都有对应的优化策略。在实际项目部署时候，可以开启编译器自动优化选项，帮助我们进一步提升代码性能。

        比如，本次测试我写了CMakeLists.txt脚本，添加编译器优化的参数后执行结果如下。CPU端和未开启编译器优化相比，时间性能有了很大的提升。手动增加的循环展开的代码时间也大大降低了。

cpu origin cost time: 31 microsecondssum = 704982704
cpu pragma unroll cost time: 0 microsecondssum = 704982704
cpu pragma unroll_1 cost time: 0 microsecondssum = 704982704
cuda origin cost time: 18 microsecondssum = 704982704
cuda pragma unroll cost time: 6 microsecondssum = 704982704

上面未开启编译器优化的输出：

cpu origin cost time: 1079 microseconds
                                                   sum = 704982704
cpu pragma unroll cost time: 678 microseconds
                                                   sum = 704982704
cpu pragma unroll_1 cost time: 374 microseconds
                                                   sum = 704982704
cuda origin cost time: 18 microseconds
                                                   sum = 704982704
cuda pragma unroll cost time: 6 microseconds
                                                   sum = 704982704

在CMakeLists.txt添加了如下一行：

set(CMAKE_CXX_FLAGS "${CMAKE_C_FLAGS} -O1 -Wall")

参考：

C++性能榨汁机之循环展开 - 知乎

【CMAKE】c++代码编译加速以及优化项_cmake 编译优化-CSDN博客

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mzph.cn/news/188856.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

MySQL索引有哪些优缺点

MySQL索引有哪些优缺点

文章目录 什么是索引&#xff1f;索引有哪些优缺点&#xff1f; 索引有哪几种类型&#xff1f; 什么是索引&#xff1f; 索引是一种特殊的文件(InnoDB数据表上的索引是表空间的一个组成部分)&#xff0c;它们包含着对数据表里所有记录的引用指针。 索引是一种数据结构。数据库…
阅读更多...
shell编程系列(8)-使用sed处理文本

shell编程系列(8)-使用sed处理文本

文章目录 引言sed用法详解在文本中定位打印文本替换文本删除文本新增文本 结语 引言 在日常工作学习中我们都会遇到要编辑文本的场景&#xff0c;例如我们要用vim或者nano等命令去编辑代码&#xff0c;处理文本文件等&#xff0c;这些命令的特点都是需要我们进行交互式的实时处…
阅读更多...
使用visual Studio MFC 平台实现对灰度图添加椒盐噪声，并进行均值滤波与中值滤波

使用visual Studio MFC 平台实现对灰度图添加椒盐噪声，并进行均值滤波与中值滤波

平滑处理–滤波 本文使用visual Studio MFC 平台实现对灰度图添加椒盐噪声&#xff0c;并进行均值滤波与中值滤波 关于其他MFC单文档工程可参考 01-Visual Studio 使用MFC 单文档工程绘制单一颜色直线和绘制渐变颜色的直线 02-visual Studio MFC 绘制单一颜色三角形、渐变颜色边…
阅读更多...
Visual Studio通过ClaudiaIDE插件设置背景图片

Visual Studio通过ClaudiaIDE插件设置背景图片

首先&#xff0c;在VS菜单栏上选择扩展-管理扩展&#xff0c;搜索插件为 ClaudiaIDE&#xff0c; 下载完成之后&#xff0c;关闭VS&#xff0c;点击Modify按钮安装&#xff1a; 等待安装完成&#xff0c;进入 VS , 打开 工具----选项---- ClauDiaIDE 界面 这个是背景色调 我选的…
阅读更多...
Mybatis-Plus实体类注解怎么用

Mybatis-Plus实体类注解怎么用

TableName 用在实体类上&#xff0c;指定实体类对应的表名称。 TableName(value "表名") TableId 用在属性上&#xff0c;指定主键字段的名称和类型。主键字段的名称一般是id&#xff0c;类型为自增。 TableId(value "id", type IdType.AUTO) TableFi…
阅读更多...
2023-12-01 AndroidR 系统在root目录下新建文件夹和创建链接，编译的时候需要修改sepolicy权限

2023-12-01 AndroidR 系统在root目录下新建文件夹和创建链接，编译的时候需要修改sepolicy权限

一、想在android 系统的根目录下新建一个tmp 文件夹&#xff0c;建立一个链接usr链接到data目录。 二、在system/core/rootdir/Android.mk里面的LOCAL_POST_INSTALL_CMD 增加 dev proc sys system data data_mirror odm oem acct config storage mnt apex debug_ramdisk tmp …
阅读更多...
共享单车停放（简单的struct结构运用）

共享单车停放（简单的struct结构运用）

本来不想写这题的&#xff0c;但是想想最近沉迷玩雨世界&#xff0c;班长又问我这题&#xff0c;就草草写了一下 代码如下&#xff1a; #include<stdio.h> #include<math.h> struct parking{int distance;int remain;int speed;int time;int jud; }parking[50]; …
阅读更多...
2021年11月10日 Go生态洞察：Twelve Years of Go

2021年11月10日 Go生态洞察：Twelve Years of Go

&#x1f337;&#x1f341; 博主猫头虎&#xff08;&#x1f405;&#x1f43e;&#xff09;带您 Go to New World✨&#x1f341; &#x1f984; 博客首页——&#x1f405;&#x1f43e;猫头虎的博客&#x1f390; &#x1f433; 《面试题大全专栏》 &#x1f995; 文章图文…
阅读更多...
爬虫-响应状态码篇

爬虫-响应状态码篇

常见的状态码及其原因: 状态码说 明详 情100继续请求者应当继续提出请求。服务器已收到请求的一部分&#xff0c;正在等待其余部分101切换协议请求者已要求服务器切换协议&#xff0c;服务器已确认并准备切换200成功服务器已成功处理了请求201已创建请求成功并且服务器创建了新…
阅读更多...
前端：实现div的隐藏与显示

前端：实现div的隐藏与显示

效果 完整代码 <!DOCTYPE html> <html lang"en"><head><meta charset"UTF-8"><meta http-equiv"X-UA-Compatible" content"IEedge"><meta name"viewport" content"widthdevice-widt…
阅读更多...
零基础怎么样才能学好平面设计？优漫教育

零基础怎么样才能学好平面设计？优漫教育

对于很多同学来讲&#xff0c;设计师是一个很让人羡慕的职业。近年来&#xff0c;不少人都加入到了平面设计的学习浪潮里来。那么对于零基础的学员可不可以学呢&#xff1f;零基础学员又应该怎么做才能学好平面设计呢&#xff1f;下面就和小编一起来看一下吧&#xff01; 零基…
阅读更多...
2021年12月14日 Go生态洞察：Go 1.18 Beta 1 发布与泛型的引入

2021年12月14日 Go生态洞察：Go 1.18 Beta 1 发布与泛型的引入

&#x1f337;&#x1f341; 博主猫头虎&#xff08;&#x1f405;&#x1f43e;&#xff09;带您 Go to New World✨&#x1f341; &#x1f984; 博客首页——&#x1f405;&#x1f43e;猫头虎的博客&#x1f390; &#x1f433; 《面试题大全专栏》 &#x1f995; 文章图文…
阅读更多...
C语言:求十个数中的平均数

C语言:求十个数中的平均数

分析&#xff1a; 程序中定义了一个average函数&#xff0c;用于计算分数的平均值。该函数接受一个包含10个分数的数组作为参数&#xff0c;并返回平均值。在主函数main中&#xff0c;首先提示输入10个分数&#xff0c;然后使用循环读取输入的分数&#xff0c;并将它们存储在名…
阅读更多...
【BEV感知 LSS方案】Lift-Splat-Shoot（LSS）

【BEV感知 LSS方案】Lift-Splat-Shoot（LSS）

前言 LSS全称是Lift-Splat-Shoot&#xff0c;它先从车辆周围的多个摄像头拍摄到的图像进行特征提取&#xff0c;在特征图中估计出每个点的深度&#xff0c;然后把这些点“提升”到3D空间中。 接着&#xff0c;这些3D信息被放置到一个网格上&#xff0c;最后将这些信息“拍扁”…
阅读更多...
07_Collection集合1

07_Collection集合1

集合体系结构 Collection 代表单列集合&#xff0c;每个元素&#xff08;数据&#xff09;只包含一个值 Map 代表双列集合&#xff0c;每个元素包含两个值&#xff08;键值对&#xff09; Collection 框架集合 Collection 概述 List 系列集合&#xff1a;添加的元素是有序…
阅读更多...
Lock还是Synchronized怎么选

Lock还是Synchronized怎么选

需要加锁的时候是使用Lock还是synchronized关键字 开头先说结论&#xff0c;需要看加的锁是不是需要超时时间&#xff0c;超时没获取到的取消&#xff0c;这种需要使用Lock&#xff0c;Synchronized不支持超时时间的设置&#xff0c;那么其他的呢&#xff0c; 并发量大的时候使…
阅读更多...
本地Lambda（SAM LI）+ MySQL（Docker）环境构筑注意点

本地Lambda（SAM LI）+ MySQL（Docker）环境构筑注意点

目录构成 mysql8 ├─data ├─logs └─docker├─docker-compose.yml├─.env├─config└─my.cnf .env DB_NAMEtest_db ROOT_DB_PASSroot_password DB_USERtest_user DB_PASStest_password DB_PORT3306 TZAsia/Tokyo docker-compose.yml version: "3.6" ser…
阅读更多...
OSI七层模型与TCP/IP四层模型

OSI七层模型与TCP/IP四层模型

一、OSI七层模型简述 OSI 模型的七层是什么&#xff1f;在 OSI 模型中如何进行通信&#xff1f;OSI 模型有哪些替代方案&#xff1f; TCP/IP 模型关于专有协议和模型的说明 二、七层模型详解&#xff08;DNS、CDN、OSI&#xff09; 状态码DNS nslookup命令 CDN whois命令 …
阅读更多...
Java中的异常你了解多少？

Java中的异常你了解多少？

目录 一.认识异常二.异常分类三.异常的分类1.编译时异常2.运行时异常 四.异常的处理1.LYBL&#xff1a;事前防御型2.EAFP&#xff1a;事后认错型 五.异常的抛出Throw注意事项 六.异常的捕获1.异常的捕获2.异常声明throws3.try-catch捕获并处理 七.自定义异常 一.认识异常 在Jav…
阅读更多...
Python实现FA萤火虫优化算法优化BP神经网络分类模型(BP神经网络分类算法)项目实战

Python实现FA萤火虫优化算法优化BP神经网络分类模型(BP神经网络分类算法)项目实战

说明&#xff1a;这是一个机器学习实战项目&#xff08;附带数据代码文档视频讲解&#xff09;&#xff0c;如需数据代码文档视频讲解可以直接到文章最后获取。 1.项目背景 萤火虫算法&#xff08;Fire-fly algorithm&#xff0c;FA&#xff09;由剑桥大学Yang于2009年提出 , …
阅读更多...
最新文章

Copyright @ 2022~2023