【CUDA】线程配置

一、 线程层次结构

1.1 认识

GPU 可并行执行工作

Thread:所有线程执行相同的核函数,并行执行

Thread Block:执行在一个Streaming Multiprocessor (SM),同一个Block中的线程可以协作

线程的集合称为块,块的数量很多。每个 block 的线程数是有限制的,因为 block 的所有线程都驻留在同一个流式多处理器内核上,并且必须共享该内核的有限内存资源。在当前 GPU 上,一个线程块最多可以包含 1024 个线程

Thread Grid:一个Grid中的Block可以在多个SM中执行

与给定核函数启动相关联的块的集合被称为网格

GPU 函数称为核函数,核函数通过执行配置启动,执行配置定义了网格中的块数以及每个块中的线程数,网格中的每个块均包含相同数量的线程

启动并行运行的核函数

可通过执行配置指定有关如何启动核函数以在多个 GPU 线程中并行运行的详细信息。即可通过执行配置指定线程组(称为线程块或简称为块)数量以及其希望每个线程块所包含的线程数量

执行配置的语法如下:

<<< NUMBER_OF_BLOCKS, NUMBER_OF_THREADS_PER_BLOCK>>>

启动核函数时,核函数代码由每个已配置的线程块中的每个线程执行

若假设已定义一个名为 someKernel 的核函数:

  • someKernel<<<1, 1>>() 配置为在具有单线程的单个线程块中运行后,将只运行一次
  • someKernel<<<1, 10>>() 配置为在具有10线程的单个线程块中运行后,将运行10次
  • someKernel<<<10, 1>>() 配置为在10个线程块(均具有单线程)中运行后,将运行10次
  • someKernel<<<10, 10>>() 配置为在10个线程块(均具有10线程)中运行后,将运行100次

1.2 线程层次结构变量

  • 网格(Grid):一个网格由多个线程块组成,这些线程块可以在一维、二维或三维空间中排列。网格的大小由 dim3 gridDim 变量指定,其中 gridDim.x、gridDim.y和gridDim.z 分别表示网格在x、y和z轴上的大小

  • 线程块(Block):一个线程块包含多个线程,这些线程在同一个SM(Streaming Multiprocessor)上并发执行。线程块的大小由 dim3 blockDim 变量指定,其中 blockDim.x、blockDim.y和blockDim.z 分别表示线程块在x、y和z轴上的大小

  • 线程(Thread):每个线程块中的线程都有一个唯一的线程ID,由 threadIdx 变量表示。同样,每个线程块在网格中也有一个唯一的块ID,由 blockIdx 变量表示

blockIdx.x 就是当前线程块在网格x轴上的索引。若网格是一维的,blockIdx.x 就足够用来唯一标识每个线程块了。若网格是二维或三维的,还需要使用 blockIdx.y和blockIdx.z 来分别表示线程块在y轴和z轴上的索引

dim3 grid(3,2,1), block(5,3,1)的线程分布示意图:

cuda线程在cuda core上执行,block在sm上执行,grid在整个device上执行

二、 协调并行线程

元素数量与线程数匹配

假设数据位于索引为 0 的向量中,由于某种未知原因,必须映射每个线程以处理向量中的元素

公式 threadIdx.x + blockIdx.x * blockDim.x 可将每个线程映射到向量的元素中

threadIdx.x的取值为0到3,blockIdx.x的取值为0到1,blockDim.x的取值为4

元素数量小于线程数

上述这种场景中,网络中的线程数与元素数量完全匹配,若线程数超过要完成的工作量,该怎么办?尝试访问不存在的元素会导致运行时错误

鉴于 GPU 的硬件特性,所含线程的数量为 32 的倍数的线程块是最理想的选择,此时具备性能上的优势。假设要启动一些线程块且每个线程块中均包含 256 个线程(32 的倍数),并需运行 1000 个并行任务(此处使用极小的数量以便于说明),则任何数量的线程块均无法在网格中精确生成 1000 个总线程,因为没有任何整数值在乘以 32 后可以恰好等于1000

  • 编写执行配置,使其创建的线程数超过执行分配工作所需的线程数
  • 将一个值作为参数传递到核函数 (N) 中,该值表示要处理的数据集总大小或完成工作所需的总线程数
  • 计算网格内的线程索引后(使用 threadIdx + blockIdx * blockDim),请检查该索引是否超过 N,并且只在不超过的情况下执行与核函数相关的工作

以下是编写执行配置的惯用方法示例,适用于 N 和线程块中的线程数已知,但无法保证网格中的线程数和 N 之间完全匹配的情况。可确保网格中至少始终拥有 N 所需的线程数,且超出的线程数至多不会超过 1 个线程块的线程数量:

// Assume `N` is known
int N = 100000;
// Assume we have a desire to set `threads_per_block` exactly to `256`
size_t threads_per_block = 256;
// Ensure there are at least `N` threads in the grid, but only 1 block's worth extra
size_t number_of_blocks = (N + threads_per_block - 1) / threads_per_block;some_kernel<<<number_of_blocks, threads_per_block>>>(N);
__global__ some_kernel(int N)
{int idx = threadIdx.x + blockIdx.x * blockDim.x;if (idx < N) // Check to make sure `idx` maps to some value within `N`{// Only do work if it does}
}

元素数量大于线程数

数据元素数量往往会大于网格中的线程数。在此情况下,线程无法只处理一个元素

以编程方式解决此问题的其中一种方法是使用网格跨度循环,在网格跨度循环中,线程的第一个元素依旧使用 threadIdx.x + blockIdx.x * blockDim.x 计算得出。然后,线程会按网格中的线程数 (blockDim.x * gridDim.x) 向前迈进,直至其数据索引超出数据元素的数量,所有线程均按此种方式运作,如此便会涵盖所有元素

__global void kernel(int *a, int N)
{int indexWithinTheGrid = threadIdx.x + blockIdx.x * blockDim.x;int gridStride = gridDim.x * blockDim.x;for (int i = indexWithinTheGrid; i < N; i += gridStride){// do work on a[i];}
}

二维情况获取线程index

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

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

相关文章

爬虫-------字体反爬

目录 一、了解什么是字体加密 二. 定位字体位置 三. python处理字体 1. 工具库 2. 字体读取 3. 处理字体 案例1&#xff1a;起点 案例2&#xff1a;字符偏移&#xff1a; 5请求数据 - 发现偏移量 5.4 多套字体替换 套用模板 版本1 版本2 四.项目实战 1. 采集目…

transformer模型写诗词

项目源码获取方式见文章末尾&#xff01; 600多个深度学习项目资料&#xff0c;快来加入社群一起学习吧。 《------往期经典推荐------》 项目名称 1.【基于CNN-RNN的影像报告生成】 2.【卫星图像道路检测DeepLabV3Plus模型】 3.【GAN模型实现二次元头像生成】 4.【CNN模型实现…

【计算机网络】章节 知识点总结

一、计算机网络概述 1. 计算机网络向用户提供的两个最重要的功能&#xff1a;连通性、共享 2. 因特网发展的三个阶段&#xff1a; 第一阶段&#xff1a;从单个网络 ARPANET 向互联网发展的过程。1983 年 TCP/IP 协议成为 ARPANET 上的标准协议。第二阶段&#xff1a;建成三级…

【微服务】不同微服务之间用户信息的获取和传递方案

如何才能在每个微服务中都拿到用户信息&#xff1f;如何在微服务之间传递用户信息&#xff1f; 文章目录 概述利用微服务网关做登录校验网关转微服务获取用户信息openFeign传递微服务之间的用户信息 概述 要在每个微服务中获取用户信息&#xff0c;可以采用以下几种方法&#…

【p2p、分布式,区块链笔记 Torrent】WebTorrent 的lt_donthave插件

扩展实现 https://github.com/webtorrent/lt_donthave/blob/master/index.js /*! lt_donthave. MIT License. WebTorrent LLC <https://webtorrent.io/opensource> */// 导入所需模块 import arrayRemove from unordered-array-remove // 用于从数组中删除元素的函数 i…

兰空图床配置域名访问

图床已经创建完毕并且可以访问了&#xff0c;但是使用IP地址多少还是差点意思&#xff0c;而且不方便记忆&#xff0c;而NAT模式又没法直接像普通服务器一样DNS解析完就可以访问。 尝试了很多办法&#xff0c;nginx配置了半天也没配好&#xff0c;索性直接重定向&#xff0c;反…

Sophos | 网络安全

在 SophosLabs 和 SophosAI 的威胁情报、人工智能和机器学习的支持下&#xff0c;Sophos 提供广泛的高级产品和服务组合&#xff0c;以保护用户、网络和端点免受勒索软件、恶意软件、漏洞利用、网络钓鱼和各种其他网络攻击。Sophos 提供单一的集成式基于云的管理控制台 Sophos …

STM32外设之SPI的介绍

### STM32外设之SPI的介绍 SPI&#xff08;Serial Peripheral Interface&#xff09;是一种高速的&#xff0c;全双工&#xff0c;同步的通信总线&#xff0c;主要用于EEPROM、FLASH、实时时钟、AD转换器等外设的通信。SPI通信只需要四根线&#xff0c;节约了芯片的管脚&#x…

基于 Transformer 的语言模型

基于 Transformer 的语言模型 Transformer 是一类基于注意力机制&#xff08;Attention&#xff09;的模块化构建的神经网络结构。给定一个序列&#xff0c;Transformer 将一定数量的历史状态和当前状态同时输入&#xff0c;然后进行加权相加。对历史状态和当前状态进行“通盘…

图数据库| 2 、大数据的演进和数据库的进阶——从数据到大数据、快数据,再到深数据

时至今日&#xff0c;大数据已无处不在&#xff0c;所有行业都在经受大数据的洗礼。但同时我们也发现&#xff0c;不同于传统关系型数据库的表模型&#xff0c;现实世界是非常丰富、高维且相互关联的。此外&#xff0c;我们一旦理解了大数据的演进历程以及对数据库进阶的强需求…

深度学习笔记10-多分类

多分类和softmax回归 在多分类问题中&#xff0c;一个样本会被划分到三个或更多的类别中&#xff0c;可以使用多个二分类模型或一个多分类模型&#xff0c;这两种方式解决多分类问题。 1.基于二分类模型的多分类 直接基于二分类模型解决多分类任务&#xff0c;对于多分类中的每…

一篇文章入门docker!

文章目录 DockerUbuntu 下 docker 安装安装docker运行docker Docker的常用命令帮助命令镜像命令容器命令其他常用命令小结 分层理解一、Docker镜像的分层结构二、UnionFS与镜像分层三、镜像层的具体内容四、镜像分层的好处五、容器层与镜像层的关系 如何提交一个自己的镜像 Doc…

鸿蒙(Harmony)实现滑块验证码

在Android和ios两端已经使用的滑块验证码框架还未适配鸿蒙版&#xff0c;于是需要自己去实现类似如下的滑块验证码&#xff1a; 那么实现这样的验证码主要涉及到几个内容&#xff1a; 1、自定义弹窗 2、base64图片转换 3、滑动组件与滑块的联动&#xff0c;以及横移距离转换…

什么是嵌入式操作系统?

什么是嵌入式操作系统? 想象一下&#xff0c;如果一个智能设备&#xff0c;比如你口袋里的智能手机&#xff0c;是一个有头脑的机器人&#xff0c;那么嵌入式操作系统&#xff08;Embedded Operating System&#xff0c;简称EOS&#xff09;就相当于这个机器人的大脑。它告诉机…

后台管理系统窗体程序:评论管理

目录 评论管理的功能介绍&#xff1a; 1、进入页面 2、页面内的各种功能设计 &#xff08;1&#xff09;网页内的表格 &#xff08;2&#xff09;拒绝按钮&#xff0c;批准按钮 &#xff08;3&#xff09;删除按钮 &#xff08;4&#xff09;页面翻页跳转按钮 一、网页设计​…

nginx代理 proxy_pass

一、location 包含 location /api/ {proxy_pass http://127.0.0.1:85;} 二、location 不包含 location /api/ {proxy_pass http://127.0.0.1:85/;} 三、locaion 包含 location /api {proxy_pass http://127.0.0.1:85;}四、location 包含 location /api {proxy_pass http://127.…

InnoDB 存储引擎<七>通用表空间+临时表空间

目录 通⽤表空间 - General Tablespace 临时表空间 - Temporary Tablespaces 通⽤表空间 - General Tablespace 对应磁盘上的文件需要用户手动创建 1.通⽤表空间的作⽤和特性&#xff1f; 解答问题&#xff1a; 1.作用&#xff1a;可以把数据量比较小且强相关的表&#xff…

乐维网管平台(五):如何精准定位网络终端设备

在当今数字化高度发展的时代&#xff0c;网络已经成为企业和组织运营的关键基础设施。而在网络管理领域&#xff0c;终端定位技术正发挥着越来越重要的作用。 一、什么是终端定位 终端定位是网络管理中的关键环节&#xff0c;从本质上讲&#xff0c;它是一种精确确定网络终端…

企业邮箱后缀设置指南,轻松融入公司品牌

邮箱后缀指""后域名&#xff0c;本文介绍如何添加公司名作为后缀&#xff0c;以Zoho邮箱为例&#xff0c;需注册账号、购买域名、配置DNS、添加自定义域名、创建账号。Zoho邮箱安全可靠、个性化定制、易于管理&#xff0c;提供不同定价方案&#xff0c;并给出客户端配…

【D3.js in Action 3 精译_039】4.3 D3 面积图的绘制方法及其边界标签的添加

当前内容所在位置&#xff1a; 第四章 直线、曲线与弧线的绘制 ✔️ 4.1 坐标轴的创建&#xff08;上篇&#xff09; 4.1.1 D3 中的边距约定&#xff08;中篇&#xff09;4.1.2 坐标轴的生成&#xff08;中篇&#xff09; 4.1.2.1 比例尺的声明&#xff08;中篇&#xff09;4.1…