​CUDA学习笔记(三)CUDA简介

news/2024/4/19 20:12:17/

本篇博文转载于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

下图展示了不同配置的性能;

 

关于性能的分析将在之后的博文中总结,现在只是了解下,本文在于掌握线程组织的方法。

 


http://www.ppmy.cn/news/1168692.html

相关文章

Unity之ShaderGraph如何实现触电电流效果

前言 之前使用ASE做过一个电流效果的shader&#xff0c;今天我们通过ShaderGraph来实现一个电流效果。 效果如下&#xff1a; 关键节点 Simple Noise&#xff1a;根据输入UV生成简单噪声或Value噪声。生成的噪声的大小由输入Scale控制。 Power&#xff1a;返回输入A的结果…

STM32不使用 cubeMX实现外部中断

这篇文章将介绍如何不使用 cubeMX完成外部中断的配置和实现。 文章目录 前言一、文件加入工程二、代码解析exti.cexti.hmain.c 注意&#xff1a;总结 前言 实验开发板&#xff1a;STM32F103C8T6。所需软件&#xff1a;keil5 &#xff0c; cubeMX 。实验目的&#xff1a;如何不…

蓝牙音视频远程控制协议(AVRCP)介绍

零.声明 本专栏文章我们会以连载的方式持续更新&#xff0c;本专栏计划更新内容如下&#xff1a; 第一篇:蓝牙综合介绍 &#xff0c;主要介绍蓝牙的一些概念&#xff0c;产生背景&#xff0c;发展轨迹&#xff0c;市面蓝牙介绍&#xff0c;以及蓝牙开发板介绍。 第二篇:Trans…

wireshark 中无线帧的类型和过滤规则对照表

帧类型 过滤器语法 Management frame wlan.fc.type 0 Control frame wlan.fc.type 1 Data frame wlan.fc.type 2 Association request wlan.fc.type_subtype 0x00 Association response wlan.fc.type_subtype 0x01 Reassociation request wlan.fc.type_subty…

【C++】一些C++11特性

C特性 1. 列表初始化1.1 {}初始化1.2 initializer_list 2. 声明2.1 auto2.2 typeid2.3 decltype2.4 nullptr 3. STL3.1 新容器3.2 新接口 4. 右值引用5. 移动构造与移动赋值6. lambda表达式7. 可变参数模板8. 包装器9. bind 1. 列表初始化 1.1 {}初始化 C11支持所有内置类型和…

C语言字符串

字符串的作用 作用&#xff1a; 用于表示一串字符 如&#xff1a; hello world等等 字符串常量及输出 字符串常量&#xff1a; 双引号引起的内容 称为字符串常量&#xff0c;如 ” hello world “ 输出方式&#xff1a; printf&#xff08; ” 字符串内容 “&#xff09;&am…

PHP的基础知识点解析

目录 一、语法 二、变量 三、数据类型 四、运算符 五、条件语句 六、循环结构 七、函数 八、数组 PHP是一种流行的开源服务器端脚本语言&#xff0c;用于Web开发。它易于学习和使用&#xff0c;并具有许多强大的功能和特性。本文将介绍PHP的一些基本知识点&#xff0c;…

Qt文件系统模型

创建文件系统模型&#xff1a;QFileSystemModel* model new QFileSystemModel(this); 设置根目录&#xff1a;model->setRootPath(QDir::currentPath()); 为视图设置模型&#xff1a; ui.treeView->setModel(model);ui.listView->setModel(model);ui.tableView-&g…

【软考】11.2 开发方法/产品线/软件复用/逆向工程

《信息系统开发方法》 结构化方法&#xff08;生命周期法&#xff09; 自顶向下、逐步求精和模块化设计遵循“用户第一”原则 三部分有机组合&#xff1a; a. 结构化分析&#xff08;SA&#xff09; b. 结构化设计&#xff08;SD&#xff09; c. 结构化程序设计&#xff08;SP…

【LLM】大模型微调,压缩,量化,部署(还在缓慢更新

前段时间很忙一直没时间follow最近的大模型工作&#xff0c;最近几天闲一点了…这个可能会出现整理不全或者是结果没跑完的情况&#xff0c;我尽量快一点&#xff08;如果最近没啥事的话&#xff09;&#xff0c;有啥想法可以在评论区d一下我。 LLM排行榜 &#xff1a; https:/…

Ubuntu - 用户和权限

sudo sudo&#xff08;Super User Do&#xff09;是在Linux和Unix系统中用于执行具有超级用户&#xff08;root&#xff09;权限的命令的命令。它允许普通用户以特权身份运行特定命令&#xff0c;通常需要输入密码以确认其身份。 sudo 是一种安全的方式&#xff0c;用于限制哪…

CSS中 通过自定义属性(变量)动态修改元素样式(以 el-input 为例)

传送门&#xff1a;CSS中 自定义属性&#xff08;变量&#xff09;详解 1. 需求及解决方案 需求&#xff1a;通常我们动态修改 div 元素的样式&#xff0c;使用 :style 和 :class 即可&#xff1b;但想要动态修改 如&#xff1a;Element-ui 中输入框&#xff08;input&#x…

uniapp map地图实现marker聚合点,并点击marker触发事件

1.uniapp官方文档说明 2.关键代码片段 // 仅调用初始化&#xff0c;才会触发 on.("markerClusterCreate", (e) > {})this._mapContext.initMarkerCluster({enableDefaultStyle: false, // 是否使用默认样式zoomOnClick: true, // 点击聚合的点&#xff0c;是否…

React Hooks批量更新问题

React 版本17.0.2 import React, { useState } from react;const Demo () > {const [count, setCount] useState(0);const [count1, setCount1] useState(0);const [count2, setCount2] useState(0);console.log(Demo);const add async () > {await 10;setCount(c…

SQL INSERT INTO 语句(在表中插入)

SQL INSERT INTO 语句 INSERT INTO 语句用于向表中插入新的数据行。 SQL INSERT INTO 语法 INSERT INTO 语句可以用两种形式编写。  第一个表单没有指定要插入数据的列的名称&#xff0c;只提供要插入的值&#xff0c;即可添加一行新的数据&#xff1a; INSERT INTO table_n…

IIC基础知识

文章目录 寻址方式 寻址方式

Linux/Ubuntu 安装 Java运行环境

linux下安装Java运行环境 1、下载安装包 .tar.gz 先在官网下载 JDK 点击这里 在这里要选择对应的 JDK 版本&#xff0c;一般我们目前选择JDK8 点击这里 2、在 /usr/local/ 目录下创建Java文件夹 cd /usr/local/ mkdir java3、将下载的文件通过FTP程序上传到刚刚创建的Java文…

网络协议--RARP:逆地址解析协议

5.1 引言 具有本地磁盘的系统引导时&#xff0c;一般是从磁盘上的配置文件中读取IP地址。但是无盘机&#xff0c;如X终端或无盘工作站&#xff0c;则需要采用其他方法来获得IP地址。 网络上的每个系统都具有唯一的硬件地址&#xff0c;它是由网络接口生产厂家配置的。无盘系统…

java基础--transient关键字减少序列化

前置内容 1.序列化 今天在看ArrayList源码的时候看到了这个&#xff0c;我之前应该是看过的&#xff0c;但是忘记了。现在在总结一下。 transient 用在类的属性上&#xff0c;不能修饰其他的。 作用&#xff1a;在序列化的时候transient修饰的属性不能被序列化 用途 在一些…

蓝桥杯 (猜生日、棋盘放麦子、MP3储存 C++)

思路&#xff1a; 1、用循环。 2、满足条件&#xff0c;能整除2012、3、12且month等于6、day<30 #include<iostream> using namespace std; int main() {for (int i 19000101; i < 20120312; i){int month i / 100 % 100;int day i % 100;if (i % 2012 0 &…