CUDA编程学习 (2)——CUDA并行性模型

news/2025/3/17 4:22:26/文章来源:https://www.cnblogs.com/Astron-fjh/p/18503843

CUDA 编程学习(2)—— CUDA并行性模型

1. 基于 kernel 的 SPMD 并行编程

1.1 向量加法 kernel(device 代码)

// Device Code
// Compute vector sum C = A + B
// 每个 thread 执行一次成对加法
__global__ void vecAddKernel(float* A, float* B, float* C, int n)
{int i = threadIdx.x + blockDim.x * blockIdx.x;if(i < n) C[i] = A[i] + B[i];
}

1.2 向量加法 kernel lauch(host 代码)

// Host Code
void vecAdd(float* h_A, float* h_B, float* h_C, int n)
{// 省略了 d_A、d_B、d_C 的 allocate 和 copy// 运行 ceil(n/256.0) blocks,每个 block 包含 256 个 thread// ceil 函数可确保有足够的 thread 覆盖所有元素vecAddKernel<<<ceil(n/256.0), 256>>>(d_A, d_B, d_C, n);
}

表达 ceil 函数的等效方法:

// Host Code
void vecAdd(float* h_A, float* h_B, float* h_C, int n)
{dim3 DimGrid((n-1)/256 + 1, 1, 1);dim3 DimBlock(256, 1, 1);vecAddKernel<<<DimGrid, DimBlock>>>(d_A, d_B, d_C, n);
}

1.3 CUDA 函数声明

image-20240922184519192
  • __device__

    • 在 device 执行,仅从 device 调用
    • 可以和 __host__ 同时调用,不可以和 __global__ 同时调用
  • __global__

    • 在 device 上执行,从 host 中调用(一些特定的 GPU 也可以从 device 上调用)
    • 内核函数必须返回 void,不支持可变参数,不能成为类成员函数
    • 注意用 __global__ 定义的 kernel 是异步的,这意味着 host 不会等待 kernel 执行完就执行下一步
  • __host__

    • 在 host 上执行,仅从 host 上调用,一般省略不写
    • 可以和 __device__,不可以和 __global__ 同时调用,此时函数会在 device 和 host 都编译。

2. 多维 kernel 配置

2.1 多维 grid

image-20240922191607767

2.2 使用二维网格处理图片

image-20240922191928670

在 GPU 计算中,图像被划分为多个小块,每个小块由一个 block 进行处理。每个 block 内部包含多个 thread,这些 thread 协同工作以并行处理图像数据。

2.3 C/C++ 中的行优先布局

image-20240922192347065

在这种布局中,矩阵中的元素按行顺序连续存储在内存中。具体来说,所有行的元素依次排列在一维的内存空间中。

例如,对于一个 4x4 的矩阵,元素 $M_{0,0}, M_{0,1}, M_{0,2}, M_{0,3}$ 按顺序首先存储,然后是 $M_{1,0}, M_{1,1}, M_{1,2}, M_{1,3}$,依此类推。公式 Row * Width + Col 用于计算矩阵中每个元素在一维内存数组中的线性位置。例如图中标注的元素 $M_{2,1}$ 的位置通过公式 $2∗4+1=9$ 计算得到。

2.4 PictureKernel 源代码

__global__ void PictureKernel(float* d_Pin, float* d_Pout, int height, int width)
{// 计算 d_Pin 和 d_Pout 元素的 rowint Row = blockIdx.y * blockDim.y + threadIdx.y;// 计算 d_Pin 和 d_Pout 元素的 columnint Col = blockIdx.x * blockDim.x + threadIdx.x;// 如果在范围内,每个 thread 计算 d_Pout 的一个元素if((Row < height) && (Col < width)){d_Pout[Row * width + Col] = 2.0 * d_Pin[Row * width + Col]; // 将每个 pixel 值缩放 2.0}
}

2.5 Launch PictureKernel 的 host 代码

// 假设图片大小为 m × n
// y 维度上有 m 个 pixel,x 维度上有 n 个 pixel
// 输入 d_Pin 已被 allocate 并 copy 到 device
// 输出 d_Pout 已被 allocate 到 device
dim3 DimGrid((n-1)/16 + 1, (m-1)/16 + 1, 1);
dim3 DimBlock(16, 16, 1);
PictureKernel<<<DimGrid, DimBlock>>>(d_Pin, d_Pout, m, n);

2.6 用 16×16 block 覆盖 62×76 大小的图片

image-20240922194530622

block 中并非所有 thread 都遵循相同的控制流路径。

3. 彩色-灰度图像处理示例

3.1 RGB 转换为灰度

image-20240922195336357

灰度数字图像是指每个 pixel 的值只包含强度信息的图像。

  • 对 $(I,J)$ 处的每个 pixel $(r,g,b)$ 进行处理:

$$
grayPixel[I,J] = 0.21r + 0.71g + 0.07*b
$$

  • 这只是一个点积 $<[r,g,b],[0.21,0.71,0.07]>$,常数与输入的 RGB 空间有关。
image-20240922195717728

3.2 RGB 转灰度代码

# define CHANNELS 3 // 我们有 3 个与 RGB 相对应的通道
// 输入图像以无符号字符 [0, 255] 编码
__global__ void colorConvert(unsigned char* grayImage, unsigned char* rgbImage, int width, int height)
{int x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;if(x < width && y < height){// 获取灰度图像的 1D 坐标int grayOffeset = y * width + x;// 可以认为 RGB 图像的 columns 是灰度图像的 CHANNEL 倍int rgbOffset = grayOffset * CHANNELS;unsigned char r = rgbImage[rgbOffset];		// red   value for pixelunsigned char g = rgbImage[rgbOffset + 1];	// green value for pixelunsigned char b = rgbImage[rgbOffset + 2];	// blue  value for pixel// 进行重新缩放并存储// 我们用浮点常量进行乘法运算grayImage[grayOffset] = 0.21f*r + 0.71f*g + 0.07f*b;}
}

4. 图像模糊示例

4.1 图像模糊

image-20240922213837446
  • 模糊块(Blurring Box)
image-20240922214028289

4.2 作为 2D Kernel 的图像模糊代码

__global__ void blurKernel(unsigned char* in, unsigned char* out, int w, int h)
{int Col = blockIdx.x * blockDim.x + threadIdx.x;int Row = blockIdx.y * blockDim.y + threadIdx.y;if(Col < w && Row < h){int pixVal = 0;int pixels = 0;// 获取周围 2xBLUR_SIZE x 2xBLUR_SIZE 方框的平均值for(int blurRow = -BLUR_SIZE; blurRow < BLUR_SIZE+1; ++blurRow){for(int blurCol = -BLUR_SIZE; blurCol < BLUR_SIZE+1; ++blurCol){int curRow = Row + blurRow;int curCol = Col + blurCol;// 验证图像像素是否有效if(curRow > -1 && curRow < h && curCol > -1 && curCol < w){pixVal += in[curRow * w + curCol];pixels++; // 跟踪累计总数中的像素数}}}   // 写出新的 pixel 值out[Row * w + Col] = (unsigned char)(pixVal / pixels);}
}

5. 线程调度(Thread Scheduling)

5.1 透明扩展性(Transparent Scalability)

image-20240922221151805

  • block 是 GPU 的工作单元,它们可以在任何顺序下执行,彼此之间没有执行依赖关系。
  • 硬件可随时自由地将 block 分配给任何处理器
    • 一个 CUDA kernel 可以在任意数量的并行处理器上扩展,无论处理器数量如何,Block的分配是动态的

5.2 示例:执行 thread block

image-20240922221911870
  • thread 按 block 粒度分配给流式多处理器(SM,Streaming Multiprocessors)

    • 在资源允许的情况下,每个 SM 最多可连接 32 个 block
    • Volta SM 最多可使用 2048 个 thread
      • 可能是 $256\ (threads/block)\ *\ 8\ blocks$
      • 或 $512\ (threads/block)\ *\ 4\ blocks$ 等
  • SM 维护 thread / block idx #s:每个SM都会跟踪它所管理的线程和线程块索引,以便于调度和管理。

  • SM 管理和调度 thread 执行

5.3 作为调度单位的 Warp

  • 每个 block 以 32-thread Warp 执行(每个 warp 包含 32 个线程)
    • 实施决策,而非 CUDA 编程模型的一部分
    • Warp 是 SM 中的调度单位
    • Warp 中的 thread 以 SIMD 方式执行
    • 未来的 GPU 可能会在每个 warp 中使用不同数量的 thread

5.3.1 Warp 示例

  • 如果一个 SM 有 3 个 block,每个区块有 256 个 thread,那么一个 SM 有多少个 Warp?
    • 每个 block 分为 $256/32 = 8$ 个 Warp
    • 共有 $8 * 3 = 24$ 个 Warp
image-20240922225501027

5.4 线程调度

  • SM 实现零开销 Warp 调度
    • 下一条指令的操作数已准备就绪的 Warp 可执行(如果数据未准备好(例如由于内存延迟),该 warp 将会停滞,SM 会调度另一个 warp 执行)
    • 符合条件的 warp 会根据优先调度策略被选择执行
    • 当一个 warp 被选中执行时,warp 内的所有线程执行相同的指令(这是因为 GPU 采用 SIMT(单指令多线程) 模型,虽然每个线程处理的数据可能不同,但它们共享相同的指令流,使得并行执行更加高效。)

5.5 block 粒度考虑因素

  • 对于使用多个 block 进行矩阵乘法运算的 Volta,每个 block 应该有 $44$、$88$ 还是 $30*30$ 个 thread?
    • 对于 $4*4$,每个 block 有 16 个 thread。每个 SM 最多可容纳 2048 个 thread,也就是 128 个 block。然而,每个 SM 最多只能容纳 32 个 block,因此每个 SM 只能容纳 512 个 thread!
    • 对于 $8*8$,每个 block 有 64 个 thread。由于每个 SM 最多可容纳 2048 个 thread,因此它最多可容纳 32 个 block 并实现满负荷运行,除非有其他资源方面的考虑因素。
    • 对于 $30*30$,每个 block 将有 900 个 thread。对于 Volta 来说,一个 SM 只能容纳 2 个 block,因此只能使用 $1800/2048$ 个 SM thread 容量。

参考文献

[1] CUDA编程入门极简教程 - 知乎 (zhihu.com)

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

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

相关文章

SLF4J 中的适配器模式

什么是适配器模式 适配器模式中,适配器包装不兼容指定接口的对象,来实现不同兼容指定接口。 SLF4J 中的适配器模式 SLF4J 是一个日志门面系统,其中提供了统一的 Logger 等接口,许多框架都会面向 SLF4J 打印日志,这样就不会和具体的日志框架耦合在一起,框架使用者也就能够…

Pbootcms留言“提交成功”的提示语怎么修改

要在 PbootCMS 中修改留言“提交成功”的提示语,可以按照以下步骤操作:定位文件:打开 apps/home/controller/MessageController.php 文件。查找代码段:在文件中找到大约第 103 行的代码段,该段代码如下:if ($this->model->addMessage($data)) {session(lastsub, ti…

Protues中51单片机按键无法复位(已解决)

前言 昨晚用 Protues 搭建了 51 的最小系统电路,在实物中好用的复位电路,到仿真里不能正常复位了。 51 单片机是高电平复位,所以在运行时 RST 引脚应该是低电平,但在仿真中 RST 引脚一直保持高电平,导致按下按键也不能复位单片机。解决方法 我在网上搜索的解决方法一共有两…

南昌航空大学-22207316-涂高杰-JAVA第一次blog作业

一.前言 本学期新增JAVA的面向程序设计课程,为增加学生编写能力开始了本学期的PTA作业,以及接下来我将根据我的实际情况总结前三次PTA题目集中最后一题并讲诉自己对Java的学习心得。从这三次PTA作业中学习到的了对ArrayLis、Vector等自动增长的数组的使用方法,学习到了许多Ja…

安装网站出现404 not found如何解决?

遇到404 Not Found错误时,可以尝试以下几个步骤来解决问题:检查URL:确认输入的网址是否正确,包括大小写和拼写。 检查是否有遗漏或多余的字符。清除浏览器缓存:有时候旧的缓存数据会导致页面加载错误,清除缓存后重新尝试访问。刷新页面:使用F5键或点击浏览器的刷新按钮重…

pbootcms模板上线推广百度竞价后打不开网站出现404错误

PbootCMS V3.2.5 版本中为了增强安全性或优化URL结构,加入了对URL参数的严格判断。当URL中包含?但不符合特定条件(如/?tag=、/?page=、/?ext_)时,系统会自动返回404错误页面。这种做法虽然有助于防止一些非法请求,但也可能导致合法的请求被误判为无效,特别是对于那些…

分布式数据库TDSQL搭建

TDSQL介绍TDSQL是腾讯基于MySQL/Mariadb社区版本打造的一款金融级分布式数据库集群方案,目前腾讯主推TDSQL MySQL版。TDSQL MySQL版具备强一致高可用、全球部署架构、分布式水平扩展、高性能、企业级安全等特性,同时提供智能 DBA、自动化运营、监控告警等配套设施,为客户提供…

BUUCTF_BUU SQL COURSE 1

BUUCTF_BUU SQL COURSE 1 打开实例发现“热点”及“登录”两个选项 根据题目提示,sql注入,尝试在登录界面寻找注入点,无果 接着进入热点界面,发现三篇新闻,依次点击发现url变化,burp抓包发现是通过对id值的控制访问不同界面,由此发现注入点 1.判断注入类型,输入1,2-1 ?…

openGussDb企业版5.0.1搭建

openGussDb企业版5.0.1搭建官方文档地址:https://docs-opengauss.osinfra.cn/zh/docs/5.0.0/docs/ReleaseNotes/Releasenotes.html软件包&客户端下载地址:https://opengauss.org/zh/download/环境准备1、软件环境要求软件类型配置描述linux操作系统ARM:openEuler 20.03L…

USB协议详解第20讲(USB包-帧首包SOF)

1.包的四种类型 根据包的组成把包分为四种类型,分别是、帧首包SOF(Start of Frame)、命令包(Token)、数据包(Data)、握手包(Handshake),如下图(大家一定要把PID类型和包类型分开)。4种PID类型和4种包类型(按照组成分类)的区别如下,大家注意区分。2.SOF包组成 我…