【CUDA 】第4章 全局内存——4.4 核函数可达到的带宽(4对角转置)【补图】

CUDA C编程笔记

  • 第四章 全局内存
    • 4.4 核函数可达到的带宽
        • 4.4.2.4 对角转置【为每个线程分配更独立的任务】
  • 【图】
  • 【补图+说明】

待解决的问题:

第四章 全局内存

4.4 核函数可达到的带宽

4.4.2.4 对角转置【为每个线程分配更独立的任务】

前置条件场景:启用线程块的网格时,线程块会被分配给SM。每个块有唯一的标识符bid,可以按行优先的顺序标注:

int bid = blockIdx.y * gridDim.x + blockIdx.x;//块的标识bid

当启用核函数时,线程块的ID决定分配给SM的顺序,如果所有SM都被占用,剩余的线程块等待有SM空余再分配。但由于线程块完成的速度和顺序不确定,因此可能最初相连的bid也会变得不连续。

下图是笛卡尔坐标系(直角)和对角块坐标系下的块标识顺序。

【图】

对角块坐标系用于确定一维线程块的ID,但访问数据时仍用笛卡尔坐标系。

对角坐标————笛卡尔坐标(直角)的转换

(直角坐标)block_x = (blockIdx.x + blockIdx.y) % gridDim.x;【blockIdx.x对角坐标】
(直角坐标)block_y = blockIdx.x;【blockIdx.y对角坐标】

核函数起始部分:对角坐标到直角坐标的映射计算+直角坐标计算线程索引ix、iy
借助合并读取+交叉写入——>实现转置

//3.对角转置————基于行
//①对角坐标系转直角坐标系②直角坐标系算线程索引ix、iy③转置
__global__ void transposeDiagonalRow(float *out, float *in, const int nx, const int ny){unsigned int blk_y = blockIdx.x;//blk_y直角坐标系,blockIdx.x对角坐标系unsigned int blk_x = (blockIdx.x+blockIdx.y) % gridDim.x;//blk_x直角坐标系,blockIdx.y对角坐标系unsigned int ix = blockDim.x * blk_x + threadIdx.x;//用直角坐标算线程索引unsigned int iy = blockDim.y * blk_y + threadIdx.y;if(ix < nx && iy < ny){out[ix*ny + iy] = in[iy*nx + ix];}
}//3.对角转置————基于列
//在基于行的基础上对换in和out的下标
__global__ void transposeDiagonalRow(float *out, float *in, const int nx, const int ny){unsigned int blk_y = blockIdx.x;//blk_y直角坐标系,blockIdx.x对角坐标系unsigned int blk_x = (blockIdx.x+blockIdx.y) % gridDim.x;//blk_x直角坐标系,blockIdx.y对角坐标系unsigned int ix = blockDim.x * blk_x + threadIdx.x;//用直角坐标算线程索引unsigned int iy = blockDim.y * blk_y + threadIdx.y;if(ix < nx && iy < ny){out[iy*nx + ix] = in[ix*ny + iy];}
}case 6:3.对角转置----基于行kernel = &transposeDiagonalRow;kernelName = "DiagonalRow       ";break;case 7:3.对角转置----基于列kernel = &transposeDiagonalCol;kernelName = "DiagonalCol       ";break;

输出结果如下:

~/cudaC/unit4$ ./4-6.1transposeNsys 6 对角转置-基于行
./4-6.1transposeNsys starting transpose at device 0: NVIDIA GeForce RTX 3090 with matrix nx 2048 ny 2048 with kernel 6
warmup         elapsed 0.000563 sec
DiagonalRow        elapsed 0.000075 sec <<< grid (128,128) block (16,16)>>> effective bandwidth 448.208557 GB~/cudaC/unit4$ ./4-6.1transposeNsys 7 对角转置-基于列
./4-6.1transposeNsys starting transpose at device 0: NVIDIA GeForce RTX 3090 with matrix nx 2048 ny 2048 with kernel 7
warmup         elapsed 0.000561 sec
DiagonalCol        elapsed 0.000064 sec <<< grid (128,128) block (16,16)>>> effective bandwidth 525.139893 GB

查询可得,理论峰值带宽为936 GB/s
基于行的对角是理论峰值的48%↑
基于列的对角是理论峰值的56%↓

结果:对角使得基于行性能提高↑,但使得基于列下降,基于列还是直角坐标性能好

基于行性能提升的原因:DRAM的并行访问
DRAM分区完成发送给全局内存的请求,设备内存中连续的256字节区域分配到连续的分区。使用直角坐标把线程块映射到——>数据块时,全局内存访问无法均匀分配到整个DRAM分区,发生“分区冲突”:内存请求在部分分区内排队等待,但另一部分分区一直空闲未被调用。
对角坐标映射造成了线程块——>数据块的非线性映射,交叉访问不太可能落到一个独立的分区,这导致速度提升。

最佳性能,一般是(所有活跃warp并发访问的)全局内存被均匀地划分。

【补图+说明】

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

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

相关文章

【玩转 Postman 接口测试与开发2_020】(完结篇)DIY 实战:随书示例 API 项目本地部署保姆级搭建教程(含完整调试过程)

《API Testing and Development with Postman》最新第二版封面 文章目录 最新版《Postman 接口测试与开发实战》示例 API 项目本地部署保姆级搭建教程1 前言2 准备工作3 具体部署3.1 将项目 Fork 到自己名下3.2 创建虚拟环境并安装依赖3.3 初始运行与项目调试 4 示例项目的用法…

2025年02月19日Github流行趋势

项目名称&#xff1a;OmniParser 项目地址url&#xff1a;https://github.com/microsoft/OmniParser 项目语言&#xff1a;Jupyter Notebook 历史star数&#xff1a;12878 今日star数&#xff1a;2153 项目维护者&#xff1a;yadong-lu, ThomasDh-C, aliencaocao, nmstoker, kr…

侯捷 C++ 课程学习笔记:设计模式在面向对象开发中的应用

在侯捷老师的《C 面向对象开发》课程中&#xff0c;除了对面向对象编程的基础特性&#xff08;封装、继承和多态&#xff09;的深入讲解外&#xff0c;还引入了设计模式这一高级主题。设计模式是面向对象编程中的一种最佳实践&#xff0c;能够帮助开发者解决常见的设计问题&…

前七章综合练习

一&#xff0c;拓扑图 二&#xff0c;实验要求 不限 三&#xff0c;实验步骤 第一步&#xff0c;搭建拓扑图 如上 注意&#xff1a; 第二步&#xff0c;配置IP trust&#xff1a; client1 client2 fw untrusrt-1&#xff1a; fw r3 电信DNS 百度web-1 untrust-2&#xf…

个人shell脚本分享

在周一到周五做增量备份&#xff0c;在周六周日做完全备份 #!/bin/bash定义变量 SRC“/path/to/source” # 源目录 BKUP“/backup” # 备份主目录 FUL“KaTeX parse error: Expected EOF, got # at position 22: …ull" #̲ 完全备份目录 INC"BKUP/inc” # 增量备份…

C语言之函数封装技巧

目录 前言 一、函数在源代码中的三种状态 二、函数封装的运用 案例1&#xff1a;实现打印20以内的素数 案例2&#xff1a;存放因子数并返回长度 三、return返回与形参返回 四、<>与“” 五、解耦 总结 前言 在C语言中&#xff0c;函数封装是一种重要的技巧&#…

深度神经网络终极指南:从数学本质到工业级实现(附Keras版本代码)

深度神经网络终极指南&#xff1a;从数学本质到工业级实现&#xff08;附Keras版本代码&#xff09; 为什么深度学习需要重新理解&#xff1f;&#xff08;与浅层模型的本质差异&#xff09; 模型类型参数容量特征学习方式适合问题类型浅层模型102-104手动特征工程低维结构化数…

vue3 + thinkphp 接入 七牛云 DeepSeek-R1/V3 流式调用和非流式调用

示例 如何获取七牛云 Token API 密钥 https://eastern-squash-d44.notion.site/Token-API-1932c3f43aee80fa8bfafeb25f1163d8 后端 // 七牛云 DeepSeek API 地址private $deepseekUrl https://api.qnaigc.com/v1/chat/completions;private $deepseekKey 秘钥;// 流式调用pub…

IIS asp.net权限不足

检查应用程序池的权限 IIS 应用程序池默认使用一个低权限账户&#xff08;如 IIS_IUSRS&#xff09;&#xff0c;这可能导致无法删除某些文件或目录。可以通过以下方式提升权限&#xff1a; 方法 1&#xff1a;修改应用程序池的标识 打开 IIS 管理器。 在左侧导航树中&#x…

代码解读:如何将HunYuan T2V模型训练成I2V模型?

Diffusion models代码解读:入门与实战 前言:HunYuan T2V模型出来很久了,但是想要训练成I2V的模型还是有点难度。此外,还有很多预训练视频模型都是T2V的,可以借鉴本文的方法加入参考图作为条件,并严格保持视频的第一帧与Image一样。 目录 Patch Image Padding Channel …

windows事件倒计时器与提醒组件

widgets 这是桌面组件前端开源组件&#xff0c;作者称&#xff1a;项目还在持续完善中&#xff0c;目前包含键盘演示、抖音热榜、喝水提醒、生日列表、待办事项、倒计时、灵动通知、打工进度等多个组件 有vue编程能力的可以自己做组件 百度网盘 夸克网盘 桌面组件 | Ca…

汽车零部件工厂如何通过工业一体机实现精准控制

在汽车制造行业中&#xff0c;零部件的精度和质量直接关系到整车的性能与安全。随着汽车工业的快速发展&#xff0c;汽车零部件工厂对生产过程的精准控制提出了更高的要求。传统的生产管理模式往往依赖人工操作和分散的系统&#xff0c;难以满足现代汽车零部件工厂的需求。而工…

BMS保护板测试仪:电池安全与性能的坚实守护者

在新能源汽车、储能系统、电动工具等电池驱动型产品日益普及的今天&#xff0c;电池的安全性和性能成为了人们关注的焦点。而BMS保护板测试仪作为电池管理系统&#xff08;BMS&#xff09;中不可或缺的一部分&#xff0c;为电池的安全运行提供了有力保障。 BMS保护板测试仪的重…

Django的初步使用

1.安装Django pip install django 验证是否安装成功&#xff1a; $ python3 Python 3.8.10 (default, Jan 17 2025, 14:40:23) [GCC 9.4.0] on linux Type "help", "copyright", "credits" or "license" for more information. >…

(前端基础)CSS(一)

了解 Cascading Style Sheet&#xff1a;层叠级联样式表 CSS&#xff1a;表现层&#xff08;美化网页&#xff09;如&#xff1a;字体、颜色、边框、高度、宽度、背景图片、网页定位、网页浮动 css优势&#xff1a; 内容和表现分离网页结构表现统一&#xff0c;可以实现复用…

CASAIM与韩国DOOSAN集团达成合作,开启工业制造自动化检测新篇

近日&#xff0c;CASAIM与韩国知名跨国企业斗山集团&#xff08;DOOSAN&#xff09;达成战略合作&#xff0c;联合打造CASAIM全自动化智能检测系统&#xff0c;助力斗山集团全面提升产品质量检测精度与效率&#xff0c;完成智能化检测升级&#xff0c;保持在全球市场竞争中的领…

矛盾(WEB)

##解题思路 打开靶场就是一段自相矛盾的代码&#xff0c;既要num是数字类型&#xff0c;又要判断为1 这种情况我们会想到弱类型的编程语言&#xff0c;插件查看过后&#xff0c;php就是弱类型的语言&#xff0c;此处并非是严格相等&#xff0c;只是 因此可以根据弱类型编程语言…

[AI]docker封装包含cuda cudnn的paddlepaddle PaddleOCR

封装,启动时需要在GPU服务器上 显卡驱动 cuda等下载需要注册账号 环境Ubuntu 24.04 LTS镜像uvicorn-gunicorn-fastapi:python3.8显卡支持版本CUDA Version: 12.5cuda版本cuda_12.4.0_550.54.14_linux.runcudnn版本PaddlePaddle-GpuPaddleOCR 启动进入容器 docker run --gpus …

idea连接gitee(使用idea远程兼容gitee)

文章目录 先登录你的gitee拿到你的邮箱找到idea的设置选择密码方式登录填写你的邮箱和密码登录成功 先登录你的gitee拿到你的邮箱 具体位置在gitee–>设置–>邮箱管理 找到idea的设置 选择密码方式登录 填写你的邮箱和密码 登录成功

VisionTransformer(ViT)与CNN卷积神经网络的对比

《------往期经典推荐------》 一、AI应用软件开发实战专栏【链接】 项目名称项目名称1.【人脸识别与管理系统开发】2.【车牌识别与自动收费管理系统开发】3.【手势识别系统开发】4.【人脸面部活体检测系统开发】5.【图片风格快速迁移软件开发】6.【人脸表表情识别系统】7.【…