CUDA学习笔记(三)CUDA简介
本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/,仅用于学习。
前言
线程的组织形式对程序的性能影响是至关重要的,本篇博文主要以下面一种情况来介绍线程组织形式:
- 2D grid 2D block
线程索引
矩阵在memory中是row-major线性存储的:

在kernel里,线程的唯一索引非常有用,为了确定一个线程的索引,我们以2D为例:
- 线程和block索引
- 矩阵中元素坐标
- 线性global memory 的偏移
首先可以将thread和block索引映射到矩阵坐标:
ix = threadIdx.x + blockIdx.x * blockDim.x
iy = threadIdx.y + blockIdx.y * blockDim.y
之后可以利用上述变量计算线性地址:
idx = iy * nx + ix

上图展示了block和thread索引,矩阵坐标以及线性地址之间的关系,谨记,相邻的thread拥有连续的threadIdx.x,也就是索引为(0,0)(1,0)(2,0)(3,0)...的thread连续,而不是(0,0)(0,1)(0,2)(0,3)...连续,跟我们线代里玩矩阵的时候不一样。
现在可以验证出下面的关系:
thread_id(2,1)block_id(1,0) coordinate(6,1) global index 14 ival 14
下图显示了三者之间的关系:

代码
int main(int argc, char **argv) {printf("%s Starting...\n", argv[0]);// set up deviceint dev = 0;cudaDeviceProp deviceProp;CHECK(cudaGetDeviceProperties(&deviceProp, dev));printf("Using Device %d: %s\n", dev, deviceProp.name);CHECK(cudaSetDevice(dev)); // set up date size of matrixint nx = 1<<14;int ny = 1<<14;int nxy = nx*ny;int nBytes = nxy * sizeof(float);printf("Matrix size: nx %d ny %d\n",nx, ny);// malloc host memoryfloat *h_A, *h_B, *hostRef, *gpuRef;h_A = (float *)malloc(nBytes);h_B = (float *)malloc(nBytes);hostRef = (float *)malloc(nBytes);gpuRef = (float *)malloc(nBytes);// initialize data at host sidedouble iStart = cpuSecond();initialData (h_A, nxy);initialData (h_B, nxy);double iElaps = cpuSecond() - iStart;memset(hostRef, 0, nBytes);memset(gpuRef, 0, nBytes);// add matrix at host side for result checksiStart = cpuSecond();sumMatrixOnHost (h_A, h_B, hostRef, nx,ny);iElaps = cpuSecond() - iStart;// malloc device global memoryfloat *d_MatA, *d_MatB, *d_MatC;cudaMalloc((void **)&d_MatA, nBytes);cudaMalloc((void **)&d_MatB, nBytes);cudaMalloc((void **)&d_MatC, nBytes);// transfer data from host to devicecudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice);cudaMemcpy(d_MatB, h_B, nBytes, cudaMemcpyHostToDevice);// invoke kernel at host sideint dimx = 32;int dimy = 32;dim3 block(dimx, dimy);dim3 grid((nx+block.x-1)/block.x, (ny+block.y-1)/block.y);iStart = cpuSecond();sumMatrixOnGPU2D <<< grid, block >>>(d_MatA, d_MatB, d_MatC, nx, ny);cudaDeviceSynchronize();iElaps = cpuSecond() - iStart;printf("sumMatrixOnGPU2D <<<(%d,%d), (%d,%d)>>> elapsed %f sec\n", grid.x,grid.y, block.x, block.y, iElaps);// copy kernel result back to host sidecudaMemcpy(gpuRef, d_MatC, nBytes, cudaMemcpyDeviceToHost);// check device resultscheckResult(hostRef, gpuRef, nxy);// free device global memorycudaFree(d_MatA);cudaFree(d_MatB);cudaFree(d_MatC);// free host memoryfree(h_A);free(h_B);free(hostRef);free(gpuRef);// reset devicecudaDeviceReset();return (0);
}
编译运行:
$ nvcc -arch=sm_20 sumMatrixOnGPU-2D-grid-2D-block.cu -o matrix2D $ ./matrix2D
输出:
./a.out Starting... Using Device 0: Tesla M2070 Matrix size: nx 16384 ny 16384 sumMatrixOnGPU2D <<<(512,512), (32,32)>>> elapsed 0.060323 sec Arrays match.
接下来,我们更改block配置为32x16,重新编译,输出为:
sumMatrixOnGPU2D <<<(512,1024), (32,16)>>> elapsed 0.038041 sec
可以看到,性能提升了一倍,直观的来看,我们会认为第二个配置比第一个多了一倍的block所以性能提升一倍,实际上也确实是因为block增加了。但是,如果你继续增加block的数量,则性能又会降低:
sumMatrixOnGPU2D <<< (1024,1024), (16,16) >>> elapsed 0.045535 sec
下图展示了不同配置的性能;

关于性能的分析将在之后的博文中总结,现在只是了解下,本文在于掌握线程组织的方法。
相关文章:
CUDA学习笔记(三)CUDA简介
本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/,仅用于学习。 前言 线程的组织形式对程序的性能影响是至关重要的,本篇博文主要以下面一种情况来介绍线程组织形式: 2D grid 2D block 线程索引 矩阵在memory中是row-major线性…...
RK3568笔记三:基于ResNet18的Cifar-10分类识别训练部署
若该文为原创文章,转载请注明原文出处。 本篇文章参考的是野火-lubancat的rk3568教程,本篇记录了在正点原子的ATK-DLK3568部署。 一、介绍 ResNet18 是一种卷积神经网络,它有 18 层深度,其中包括带有权重的卷积层和全连接层。它…...
块状数据结构学习笔记
分块 分块的思想和珂朵莉树很类似,就是把原序列分成若干个块,对块进行操作的奇妙思想。复杂度通常带根号。分块的块长也有讲究,通常对于大小为 n n n 的数组,取距离 n \sqrt n n 最近的 2 2 2 的幂数或直接取 n \sqrt n n…...
DOM4J解析.XML文件
<?xml version"1.0" encoding"utf-8" ?> <books><book id"SN123123413241"><name>java编程思想</name><author>华仔</author><price>9.9</price></book><book id"SN1234…...
黑豹程序员-架构师学习路线图-百科:MVC的演变终点SpringMVC
MVC发展史 在我们开发小型项目时,我们代码是混杂在一起的,术语称为紧耦合。 如最终写ASP、PHP。里面既包括服务器端代码,数据库操作的代码,又包括前端页面代码、HTML展现的代码、CSS美化的代码、JS交互的代码。可以看到早期编程就…...
二、BurpSuite Intruder暴力破解
一、介绍 解释: Burp Suite Intruder是一款功能强大的网络安全测试工具,它用于执行暴力破解攻击。它是Burp Suite套件的一部分,具有高度可定制的功能,能够自动化和批量化执行各种攻击,如密码破解、参数枚举和身份验证…...
solidworks 2024新功能之-让您的工作更加高效
您可以创建杰出的设计,并将这些杰出的设计将融入产品体验中。为了帮您简化和加快由概念到成品的产品开发流程,SOLIDWORKS 2024 涵盖全新的用户驱动型增强功能,致力于帮您实现更智能、更快速地与您的团队和外部合作伙伴协同工作。 SOLIDWORKS…...
华为eNSP配置专题-VRRP的配置
文章目录 华为eNSP配置专题-VRRP的配置0、参考文档1、前置环境1.1、宿主机1.2、eNSP模拟器 2、基本环境搭建2.1、基本终端构成和连接 2.VRRP的配置2.1、PC1的配置2.2、接入交换机acsw的配置2.3、核心交换机coresw1的配置2.4、核心交换机coresw2的配置2.5、配置VRRP2.6、配置出口…...
LuatOS-SOC接口文档(air780E)--lcd - lcd驱动模块
常量 常量 类型 解释 lcd.font_opposansm8 font 8号字体 lcd.font_unifont_t_symbols font 符号字体 lcd.font_open_iconic_weather_6x_t font 天气字体 lcd.font_opposansm10 font 10号字体 lcd.font_opposansm12 font 12号字体 lcd.font_opposansm16 font…...
敏捷是怎么提高工作效率的
敏捷管理是一门极力减少不必要工作量的艺术。 谷歌、亚马逊、苹果、微信、京东等全球 500 强企业都在用的管理方法,适用于各行各业,被盛赞为应获“管理学的诺贝尔奖”。 它专注于让员工不受种种杂事的羁绊,激发个体斗志,释放出巨大…...
【C++】哈希的应用 -- 布隆过滤器
文章目录 一、布隆过滤器提出二、布隆过滤器概念三、布隆过滤器哈希函数个数的选择四、布隆过滤器的实现1.布隆过滤器的插入2.布隆过滤器的查找3.布隆过滤器删除4.完整代码实现 五、布隆过滤器总结1.布隆过滤器优点2.布隆过滤器缺陷3.布隆过滤器的应用4.布隆过滤器相关面试题 一…...
如何在Git中修改远程仓库地址
原文(可不登录复制代码):如何在Git中修改远程仓库地址-北的杂货间 Git是广泛使用的分布式版本控制系统,它允许开发者在本地仓库上工作,并将更改上传到远程仓库。然而,有时候你可能需要修改远程仓库的地址&…...
Go语言的sync.Once()函数
sync.Once 是 Go 语言标准库 sync 包提供的一个类型,它用于确保一个函数只会被执行一次,即使在多个 goroutine 中同时调用。 sync.Once 包含一个 Do 方法,其签名如下: func (o *Once) Do(f func()) Do 方法接受一个函数作为参数…...
修改 Stable Diffusion 使 api 接口增加模型参数
参考:https://zhuanlan.zhihu.com/p/644545784 1、修改 modules/api/models.py 中的 StableDiffusionTxt2ImgProcessingAPI 增加模型名称 StableDiffusionTxt2ImgProcessingAPI PydanticModelGenerator("StableDiffusionProcessingTxt2Img",StableDiff…...
微信小程序自定义组件及会议管理与个人中心界面搭建
一、自定义tabs组件 1.1 创建自定义组件 新建一个components文件夹 --> tabs文件夹 --> tabs文件 创建好之后win7 以上的系统会报个错误:提示代码分析错误,已经被其他模块引用,只需要在 在project.config.json文件里添加两行配置 &…...
UiPath:一家由生成式AI驱动的流程自动化软件公司
来源:猛兽财经 作者:猛兽财经 总结: (1)UiPath(PATH)的股价并没有因为生成式AI的炒作而上涨,但很可能会成为主要受益者。 (2)即使在严峻的宏观环境下,UiPath的收入还在不…...
使用AI编写测试用例——详细教程
随着今年chatGPT的大热,每个行业都试图从这项新技术当中获得一些收益我之前也写过一篇测试领域在AI技术中的探索:软件测试中的AI——运用AI编写测试用例现阶段AI还不能完全替代人工测试用例编写,但是如果把AI当做一个提高效率的工具ÿ…...
又哭又笑,这份面试宝典要是早遇到就好了
01、算法原理 选择排序(Selection sort)是一种简单直观的排序算法。 第一次从待排序的数据元素中选出最小(或最大)的一个元素,存放在序列的起始位置,然后再从剩余的未排序元素中寻找到最小(大)元素&#…...
订单30分钟自动关闭的五种解决方案
1 前言 在开发中,往往会遇到一些关于延时任务的需求。例如 生成订单30分钟未支付,则自动取消生成订单60秒后,给用户发短信 对上述的任务,我们给一个专业的名字来形容,那就是延时任务 。那么这里就会产生一个问题,这…...
【vSphere 8 自签名 VMCA 证书】企业 CA 签名证书替换 vSphere VMCA CA 证书Ⅰ—— 生成 CSR
目录 替换拓扑图证书关系示意图说明 & 关联博文1. 默认证书截图2. 使用 certificate-manager 生成CSR2.1 创建存放CSR的目录2.2 记录PNID和IP2.3 生成CSR2.4 验证CSR 参考资料 替换拓扑图 证书关系示意图 本系列博文要实现的拓扑是 说明 & 关联博文 因为使用企业 …...
论文解读:交大港大上海AI Lab开源论文 | 宇树机器人多姿态起立控制强化学习框架(二)
HoST框架核心实现方法详解 - 论文深度解读(第二部分) 《Learning Humanoid Standing-up Control across Diverse Postures》 系列文章: 论文深度解读 + 算法与代码分析(二) 作者机构: 上海AI Lab, 上海交通大学, 香港大学, 浙江大学, 香港中文大学 论文主题: 人形机器人…...
C++:std::is_convertible
C++标志库中提供is_convertible,可以测试一种类型是否可以转换为另一只类型: template <class From, class To> struct is_convertible; 使用举例: #include <iostream> #include <string>using namespace std;struct A { }; struct B : A { };int main…...
shell脚本--常见案例
1、自动备份文件或目录 2、批量重命名文件 3、查找并删除指定名称的文件: 4、批量删除文件 5、查找并替换文件内容 6、批量创建文件 7、创建文件夹并移动文件 8、在文件夹中查找文件...
sqlserver 根据指定字符 解析拼接字符串
DECLARE LotNo NVARCHAR(50)A,B,C DECLARE xml XML ( SELECT <x> REPLACE(LotNo, ,, </x><x>) </x> ) DECLARE ErrorCode NVARCHAR(50) -- 提取 XML 中的值 SELECT value x.value(., VARCHAR(MAX))…...
Spring Cloud Gateway 中自定义验证码接口返回 404 的排查与解决
Spring Cloud Gateway 中自定义验证码接口返回 404 的排查与解决 问题背景 在一个基于 Spring Cloud Gateway WebFlux 构建的微服务项目中,新增了一个本地验证码接口 /code,使用函数式路由(RouterFunction)和 Hutool 的 Circle…...
力扣-35.搜索插入位置
题目描述 给定一个排序数组和一个目标值,在数组中找到目标值,并返回其索引。如果目标值不存在于数组中,返回它将会被按顺序插入的位置。 请必须使用时间复杂度为 O(log n) 的算法。 class Solution {public int searchInsert(int[] nums, …...
如何在网页里填写 PDF 表格?
有时候,你可能希望用户能在你的网站上填写 PDF 表单。然而,这件事并不简单,因为 PDF 并不是一种原生的网页格式。虽然浏览器可以显示 PDF 文件,但原生并不支持编辑或填写它们。更糟的是,如果你想收集表单数据ÿ…...
安卓基础(aar)
重新设置java21的环境,临时设置 $env:JAVA_HOME "D:\Android Studio\jbr" 查看当前环境变量 JAVA_HOME 的值 echo $env:JAVA_HOME 构建ARR文件 ./gradlew :private-lib:assembleRelease 目录是这样的: MyApp/ ├── app/ …...
用机器学习破解新能源领域的“弃风”难题
音乐发烧友深有体会,玩音乐的本质就是玩电网。火电声音偏暖,水电偏冷,风电偏空旷。至于太阳能发的电,则略显朦胧和单薄。 不知你是否有感觉,近两年家里的音响声音越来越冷,听起来越来越单薄? —…...
纯 Java 项目(非 SpringBoot)集成 Mybatis-Plus 和 Mybatis-Plus-Join
纯 Java 项目(非 SpringBoot)集成 Mybatis-Plus 和 Mybatis-Plus-Join 1、依赖1.1、依赖版本1.2、pom.xml 2、代码2.1、SqlSession 构造器2.2、MybatisPlus代码生成器2.3、获取 config.yml 配置2.3.1、config.yml2.3.2、项目配置类 2.4、ftl 模板2.4.1、…...
