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

news/2024/6/29 11:51:43 标签: 性能优化, 编译器, pragma unroll, CUDA, 循环展开

引言：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 unroll
    for(int i = 0;i < count;i++){  
        *sum += i;
    }
}

性能分析与测试接口实现

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

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

编译如下，因为把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 unroll
    for(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 microseconds
                                                   sum = 704982704
cpu pragma unroll cost time: 0 microseconds
                                                   sum = 704982704
cpu pragma unroll_1 cost time: 0 microseconds
                                                   sum = 704982704
cuda origin cost time: 18 microseconds
                                                   sum = 704982704
cuda pragma unroll cost time: 6 microseconds
                                                   sum = 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.niftyadmin.cn/n/5232543.html

相关文章

C# Windows开机后立刻设置地区过一会地区自动变回原来的值

C# Windows开机后立刻设置地区过一会地区自动变回原来的值

一开始使用PowerShell命令执行&#xff0c;发现过一会地区自动变回原来的值。 后来尝试了使用PowerShell脚本执行&#xff0c;并且增加了一些脚本策略&#xff0c;测试通过&#xff0c;不再变回原来的值。 /// <summary> /// 设置区域 /// </summary> /// <para…
阅读更多...
linux docker 使用详解

linux docker 使用详解

docker 相对一vm 来说是更轻量级的 隔离&#xff0c;他和host 共享linux 和硬件 不需要虚拟hw。 在嵌入式中比较常用的是lxc 在linux 下是docker 1. 查看容器的root用户密码 docker logs <容器名or ID> 2>&1 | grep ^User: | tail -n1 因为docker容器启动时…
阅读更多...
vue form表单的封装--使用的是elementUI

vue form表单的封装--使用的是elementUI

该组件是一个动态生成表单的Vue组件&#xff0c;通过传入formList参数和插槽的方式实现表单项的定制化渲染 组件的封装 HCommonFormItem 根据传入的"formList"参数生成一组表单项&#xff0c;并通过插槽&#xff08;slot&#xff09;将表单项渲染出来。 在component…
阅读更多...
ReactDomServer 将react组件转化成html静态标签（SSR服务器渲染）

ReactDomServer 将react组件转化成html静态标签（SSR服务器渲染）

前言&#xff1a; 因为使用图表里面的formatter函数需要原生的html标签&#xff0c;但是本身技术栈是react&#xff0c;所以为了方便&#xff0c;便使用了ReactDomServer api将react组件转化成html原生标签引入&#xff1a; import ReactDomServer from react-dom/server; 使…
阅读更多...
C#中GDI+图形图像绘制(直线、矩形、圆、椭圆、圆弧、扇形、多边形）

C#中GDI+图形图像绘制(直线、矩形、圆、椭圆、圆弧、扇形、多边形）

目录 一、直线 二、矩形 三、椭圆 四、圆 五、圆弧 六、扇形 七、多边形 八、示例源码 一、直线 调用Graphics类中的DrawLine()方法&#xff0c;结合Pen对象可以绘制直线。DrawLine()方法有以下两种构造函数。 第一种用于绘制一条连接两个Point结构的线。当参数pt1的值…
阅读更多...
java FTP客户端获取文件流假死问题

java FTP客户端获取文件流假死问题

依赖 hutool FTP配置 inspection.data.ftp.host172.26.1.41 inspection.data.ftp.port21 inspection.data.ftp.user6c inspection.data.ftp.password6cqq123 inspection.data.ftp.charsetNameGBK FTP配置类 import lombok.Data; import org.springframework.boot.context.pr…
阅读更多...
101.套接字-Socket网络编程3

101.套接字-Socket网络编程3

目录 1.字节序 主机字节序&#xff08;小端&#xff09; 网络字节序&#xff08;大端&#xff09; 字节序转换函数 2.IP地址转换函数 3.套接字地址结构 通用 socket 地址结构 专用 socket 地址结构 Socket套接字的目的是将TCP/IP协议相关软件移植到UNIX类系统中。设计…
阅读更多...
第九节HarmonyOS 常用基础组件1-Text

第九节HarmonyOS 常用基础组件1-Text

一、组件介绍 组件&#xff08;Component&#xff09;是界面搭建与显示的最小单位&#xff0c;HarmonyOS ArkUI声名式为开发者提供了丰富多样的UI组件&#xff0c;我们可以使用这些组件轻松的编写出更加丰富、漂亮的界面。 组件根据功能可以分为以下五大类&#xff1a;基础组件…
阅读更多...
最新文章

Copyright @ 2022~2023