CUDA编程 - 用向量化访存优化 - Cuda elementwise - Add(逐点相加)- 学习记录

Cuda elementwise - Add

  • 一、简介
    • 1.1、ElementWise Add
    • 1.2、 float4 - 向量化访存
  • 二、实践
    • 2.1、如何使用向量化访存
    • 2.1、简单的逐点相加核函数
    • 2.2、ElementWise Add + float4(向量化访存)
    • 2.3、完整代码

一、简介

1.1、ElementWise Add

Element-wise 操作是最基础,最简单的一种核函数的类型,它的计算特点很符合GPU的工作方式:对于每个元素单独做一个算术操作,然后直接输出。

Add 函数 :逐点相加

  • 传入 数组 a,b,c
  • 传入 数据数量 N
  • 传出结果 数组c

1.2、 float4 - 向量化访存

所谓向量化访存,就是一次性读 4 个 float,而不是单单 1 个

要点:

  • 小数据规模情况下,可以不考虑向量化访存的优化方式
  • 大规模数据情况下,考虑使用向量化访存,且最好是缩小grid的维度为原来的1/4,避免影响Occupancy
  • float4 向量化访存只对数据规模大的时候有加速效果,数据规模小的时候没有加速效果

float4的性能提升主要在于访存指令减少了(同样的数据规模,以前需要4条指令,现在只需1/4的指令),指令cache里就能存下更多指令,提高指令cache的命中率。

判断是否用上了向量化访存,是在 nsight compute 看生成的SASS代码里会有没有LDG.E.128 Rx, [Rx.64]或STG.E.128 [R6.64], Rx这些指令的存在。有则向量化成功,没有则向量化失败。

在这里插入图片描述

官方参考链接1
官方参考链接2

二、实践

2.1、如何使用向量化访存

c :

#define FLOAT4(value)  *(float4*)(&(value))

宏解释:

对于一个值,先对他取地址,然后再把这个地址解释成 float4
对于这个 float4的指针,对它再取一个值
这样编译器就可以一次读四个 float

c++ :

#define FLOAT4(value) (reinterpret_cast<float4*>(&(value))[0])

2.1、简单的逐点相加核函数

__global__ void elementwise_add(float* a, float* b, float* c, int N) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < N) c[idx] = a[idx] + b[idx];
}

2.2、ElementWise Add + float4(向量化访存)

__global__ void elementwise_add_float4(float* a, float* b, float *c, int N)
{int idx = (blockDim.x * blockIdx.x + threadIdx.x) * 4;if(idx < N ){float4 tmp_a = FLOAT4(a[idx]);float4 tmp_b = FLOAT4(b[idx]);float4 tmp_c;tmp_c.x = tmp_a.x + tmp_b.x;tmp_c.y = tmp_a.y + tmp_b.y;tmp_c.z = tmp_a.z + tmp_b.z;tmp_c.w = tmp_a.w + tmp_b.w;FLOAT4(c[idx]) = tmp_c;}
}

将核函数写成 float4 的形式的时候,首先要先使用宏定义(参考1.3),其次要注意线程数的变化。

线程数变化原因:因为一个线程可以处理4个float了,所以要减少 四倍的线程。

2.3、完整代码

elementwise_add.cu

#include <stdio.h>
#include <stdlib.h>
#include <float.h>
#include <vector>
#include<assert.h>
#include <algorithm>
#include <cublas_v2.h>
#include <cuda_runtime.h>#define FLOAT4(value)  *(float4*)(&(value))#define checkCudaErrors(func)               \
{                                   \cudaError_t e = (func);         \if(e != cudaSuccess)                                        \printf ("%s %d CUDA: %s\n", __FILE__,  __LINE__, cudaGetErrorString(e));        \
}// ElementWise Add  
// elementwise_add<<<CeilDiv(N, block_size), block_size>>>(d_A, d_B, d_C, N);
// a: Nx1, b: Nx1, c: Nx1, c = elementwise_add(a, b)
__global__ void elementwise_add(float* a, float* b, float* c, int N) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < N) c[idx] = a[idx] + b[idx];
}__global__ void elementwise_add_float4(float* a, float* b, float *c, int N)
{int idx = (blockDim.x * blockIdx.x + threadIdx.x) * 4;if(idx < N ){float4 tmp_a = FLOAT4(a[idx]);float4 tmp_b = FLOAT4(b[idx]);float4 tmp_c;tmp_c.x = tmp_a.x + tmp_b.x;tmp_c.y = tmp_a.y + tmp_b.y;tmp_c.z = tmp_a.z + tmp_b.z;tmp_c.w = tmp_a.w + tmp_b.w;FLOAT4(c[idx]) = tmp_c;}
}template <typename T> 
inline T CeilDiv(const T& a, const T& b) {return (a + b - 1) / b;
}int main(){size_t block_size = 128;size_t N =  32 * 1024 * 1024;size_t bytes_A = sizeof(float) * N;size_t bytes_B = sizeof(float) * N;size_t bytes_C = sizeof(float) * N;float* h_A = (float*)malloc(bytes_A);float* h_B = (float*)malloc(bytes_B);float* h_C = (float*)malloc(bytes_C);for( int i = 0; i < N; i++ ){h_A[i] = i / 666;}for( int i = 0; i < N; i++ ) {h_B[i] = i % 666;}float* d_A;float* d_B;float* d_C;checkCudaErrors(cudaMalloc(&d_A, bytes_A));checkCudaErrors(cudaMalloc(&d_B, bytes_B));checkCudaErrors(cudaMalloc(&d_C, bytes_C));checkCudaErrors(cudaMemcpy( d_A, h_A, bytes_A, cudaMemcpyHostToDevice));checkCudaErrors(cudaMemcpy( d_B, h_B, bytes_B, cudaMemcpyHostToDevice));cudaEvent_t start, stop;checkCudaErrors(cudaEventCreate(&start));checkCudaErrors(cudaEventCreate(&stop));float msec = 0;int iteration = 1;checkCudaErrors(cudaEventRecord(start));for(int i = 0; i < iteration; i++){elementwise_add<<<CeilDiv(N, block_size), block_size>>>(d_A, d_B, d_C, N);                   //elementwise_add_float4<<<CeilDiv(N, block_size), block_size/4>>>(d_A, d_B, d_C, N);          //elementwise_add_float4<<<CeilDiv(N/4, block_size), block_size>>>(d_A, d_B, d_C, N);}checkCudaErrors(cudaEventRecord(stop));checkCudaErrors(cudaEventSynchronize(stop));checkCudaErrors(cudaEventElapsedTime(&msec, start, stop));printf("elementwise add takes %.5f msec\n", msec/iteration);checkCudaErrors(cudaMemcpy(h_C, d_C, bytes_C, cudaMemcpyDeviceToHost));for(int i = 0; i < N; i++){double err = fabs(h_C[i] - (h_A[i] + h_B[i]));if(err > 1.e-6) {printf("wrong answer!\n");break;}}cudaFree(d_A);cudaFree(d_B);cudaFree(d_C);free(h_A);free(h_B);free(h_C);return 0;
}

编译和运行:

nvcc -o elementwise_add elementwise_add.cu 
./elementwise_add

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

如若内容造成侵权/违法违规/事实不符,请联系瞎胡闹网进行投诉反馈,一经查实,立即删除!

相关文章

ref()和reactive()

setup(){} ref() 声明响应式状态&#xff0c;ref() 接收参数&#xff0c;并将其包裹在一个带有 .value 属性的 ref 对象中返回&#xff0c;要在组件模板中访问 ref&#xff0c;从组件的 setup() 函数中声明并返回它们。 ref 的好处是&#xff0c;与普通变量不同&#xff0c…

画板CHECK-检测

一, 检查组件是否会压住板上的组件. LBXR20会压住板上的组件. 会压住431管子.431的5V上下电阻要在.用三极管降压到5V时, 记得记得431的输出要放在电容的前面才可以增加反馈的时间.插件的丝印方面或贴片的丝印方面要注意正确,一般都是向外的,今天换那个new-609时,昨天晚上发出的…

【Web】CTFSHOW XXE刷题记录(全)

目录 web373 web374 web375 web376 web377 web378 前置知识先看这篇文章&#xff1a;XXE漏洞学习 用的多的就是外部实体声明和参数实体声明 web373 有回显的xxe <!DOCTYPE test [ <!ENTITY xxe SYSTEM "file:///flag"> ]> <z3r4y> <ct…

C# 学习第二弹

一、变量 存储区&#xff08;内存&#xff09;中的一个存储单元 &#xff08;一&#xff09;变量的声明和初始化 1、声明变量——根据类型分配空间 ①声明变量的方式 —变量类型 变量名 数值&#xff1b; —变量类型 变量名&#xff1b; 变量名 数值&#xff1b; —变…

IP 电话

1 IP 电话概述 IP 电话是在互联网上传送多媒体信息。 多个英文同义词&#xff1a; VoIP (Voice over IP) Internet Telephony VON (Voice On the Net) 1.1 狭义的和广义的 IP 电话 狭义的 IP 电话&#xff1a;指在 IP 网络上打电话。 广义的 IP 电话&#xff1a;不仅仅是…

智慧公厕让社区生活更美好

随着科技的迅猛发展&#xff0c;城市管理、城市服务均使用科技化的手段进行升级改造&#xff0c;社区生活更美好赋予全新的智慧效能&#xff0c;其中智慧公厕也成为了城市环卫设施的新宠。智慧公厕以物联网、互联网、大数据、云计算、5G通信、自动化控制等技术为核心&#xff0…

【mysql】 1819 - Your password does not satisfy the current policy requirements

创建mysql账户密码时候提示&#xff1a; 1819 - Your password does not satisfy the current policy requirements 1819-您的密码不符合当前策略要求 下面是执行的sql DROP DATABASE IF EXISTS company;CREATE DATABASE company CHARACTER SET utf8mb4 ;grant all on com…

大语言模型推理加速技术:计算加速篇

原文&#xff1a;大语言模型推理加速技术&#xff1a;计算加速篇 - 知乎 目录 简介 Transformer和Attention 瓶颈 优化目标 计算加速 计算侧优化 KVCache Kernel优化和算子融合 分布式推理 内存IO优化 Flash Attention Flash Decoding Continuous Batching Page…

thinkphp6定时任务

这里主要是教没有用过定时任务没有头绪的朋友, 定时任务可以处理一些定时备份数据库等一系列操作, 具体根据自己的业务逻辑进行更改 直接上代码 首先, 是先在 tp 中的 command 方法中声明, 如果没有就自己新建一个, 代码如下 然后就是写你的业务逻辑 执行定时任务 方法写好了…

降压型DC电源模块的特点与优势

BOSHIDA 降压型DC电源模块的特点与优势 降压型DC电源模块是一种将输入电压降低到需要的输出电压的电源模块。其特点与优势如下&#xff1a; 1. 输入电压范围广泛&#xff1a;降压型DC电源模块可以适应不同的输入电压范围&#xff0c;可以在不同的电压环境下使用。 2. 输出电…

YOLOv9图像标注和格式转换

一、软件安装 labelimg安装&#xff08;anaconda&#xff09; 方法一、 pip install labelImg 方法二、 pip install PyQt5 -i https://pypi.tuna.tsinghua.edu.cn/simple/ pip install pyqt5-tools -i https://pypi.tuna.tsinghua.edu.cn/simple/ pip install lxml -i ht…

前后端项目宝塔linux部署(springboot,vue,python)

宝塔linux安装就省略了&#xff0c;网上一堆 1.部署后端 1.首先把自己项目里面打包好的的jar包上传到服务器随便一个地方&#xff0c;我这里就上传到www/wwwroot下面了&#xff0c;宝塔的文件页面可以很便携上传 2.然后到下面这个页面 选那个java环境管理装个jdk&#xff…

2024022601-数据库语言SQL

数据库语言SQL SQL的发展 1974年&#xff0c;由Boyce和Chamberlin提出 1975~1979&#xff0c;IBM San Jose Research Lab的关系数据库管理系统原型System R实施了这种语言 SQL-86是第一个SQL标准 SQL-89、SQL-92(SQL2)、SQL-99(SQL3) 非过程化语言 SQL语言进行数据库操作…

Java数据结构---初识集合框架

目录 一、什么是集合框架 二、集合框架的重要性 三、背后涉及的数据结构及算法 1.什么是数据结构 2.容器背后对应的数据结构 3.相关的Java知识 4.什么是算法 一、什么是集合框架 Java 集合框架 Java Collection Framework &#xff0c;又被称为容器 container &#xff0…

详解顺序结构滑动窗口处理算法

&#x1f380;个人主页&#xff1a; https://zhangxiaoshu.blog.csdn.net &#x1f4e2;欢迎大家&#xff1a;关注&#x1f50d;点赞&#x1f44d;评论&#x1f4dd;收藏⭐️&#xff0c;如有错误敬请指正! &#x1f495;未来很长&#xff0c;值得我们全力奔赴更美好的生活&…

python Matplotlib Tkinter-->tab切换2

环境 python:python-3.12.0-amd64 包: matplotlib 3.8.2 pillow 10.1.0 import matplotlib.pyplot as plt from matplotlib.backends.backend_tkagg import FigureCanvasTkAgg, NavigationToolbar2Tk import tkinter as tk import tkinter.ttk as ttk# 创建自定义工具栏类 c…

nginx---------------重写功能 防盗链 代理 (五)

一、重写功能 rewrite Nginx服务器利用 ngx_http_rewrite_module 模块解析和处理rewrite请求&#xff0c;此功能依靠 PCRE(perl compatible regular expression)&#xff0c;因此编译之前要安装PCRE库&#xff0c;rewrite是nginx服务器的重要功能之一&#xff0c;重写功能(…

各类几何像差

1、离焦&#xff1a; 简单理解&#xff0c;离焦就是成像面不在焦点处&#xff1a; 越远&#xff0c;越模糊 2、球差&#xff1a; 球差是由于透镜中心区域和边缘区域对电磁波会聚能力不同而造成的。远轴电磁波通过透射时被折射得比近轴电磁波要厉害得多&#xff0c;因而由同一…

vue3使用elementPlus进行table合并处理

elementPlus中table合并部分列 虚拟数据中公司下有多个客户&#xff0c;公司一样的客户&#xff0c;公司列需要合并&#xff0c;客户如果一样也需要合并进行展示&#xff0c;效果展示 const tableData ref([])自定定义自已想要的数据&#xff0c;一般都是通过接口拿到 //table…

后端:跨端轻量JavaScript引擎的实现与探索

一、JavaScript 1.JavaScript语言 JavaScript是ECMAScript的实现,由ECMA 39(欧洲计算机制造商协会39号技术委员会)负责制定ECMAScript标准。 ECMAScript发展史: 时间版本说明1997年7月ES1.0 发布当年7月&#xff0c;ECMA262 标准出台1998年6月ES2.0 发布该版本修改完全符合…