CUDA 12.4文档2 内核线程架构

本博客参考官方文档进行介绍,全网仅此一家进行中文翻译,走过路过不要错过。

官方网址:https://docs.nvidia.com/cuda/cuda-c-programming-guide/

本文档分成多个博客进行介绍,在本人专栏中含有所有内容:

https://blog.csdn.net/qq_33345365/category_12610860.html

CUDA 12.4为2024年3月2日发表,本专栏开始书写日期2024/4/8,当时最新版本4.1

本人会维护一个总版本,一个小章节的版本,总版本会持续更新,小版本会及时的调整错误和不合理的翻译,内容大部分使用chatGPT 4翻译,部分内容人工调整


开始编辑时间:2024/4/8;最后编辑时间:2024/4/10

5.1 内核Kernels

CUDA C++通过允许程序员定义C++函数,称为内核,当被调用时,这些函数由N个不同的CUDA线程并行执行N次,而不是像常规的C++函数那样只执行一次。

内核使用 global 声明说明符定义,并且给定内核调用的执行该内核的CUDA线程数使用新的 <<<...>>>执行配置语法指定。执行内核的每个线程都被赋予一个在内核内通过内置变量可以访问的唯一线程ID。

作为示例,以下样本代码使用内置变量 threadIdx,将两个大小为N的向量A和B相加,并将结果存储到向量C中:

__global__ void VecAdd(float* A, float* B, float* C)
{int i = threadIdx.x;C[i] = A[i] + B[i];
}
int main()
{...∕∕ Kernel invocation with N threadsVecAdd<<<1, N>>>(A, B, C);...
}

在这里,执行VecAdd()的N个线程中的每一个都执行一次成对的加法。

5.2 线程架构

为了方便,threadIdx是一个三组件向量,因此可以使用一维、二维或三维的线程索引来标识线程,形成一个一维、二维或三维的线程块,称为线程块。这为在像向量、矩阵或体积等域中的元素之间调用计算提供了一种自然的方式。

线程的索引和线程ID之间的关系很直接:对于一维的块,它们是相同的;对于大小为(Dx, Dy)的二维块,索引为(x, y)的线程的线程ID为(x + y Dx);对于大小为(Dx, Dy, Dz)的三维块,索引为(x, y, z)的线程的线程ID为(x + y Dx + z Dx Dy)。

下面的代码就是一个例子,它将两个NxN大小的矩阵A和B加在一起,并将结果存储到矩阵C中:

__global__ void MatAdd(float A[N][N], float B[N][N],float C[N][N])
{int i = threadIdx.x;int j = threadIdx.y;C[i][j] = A[i][j] + B[i][j];
}
int main()
{...∕∕ Kernel invocation with one block of N * N * 1 threadsint numBlocks = 1;dim3 threadsPerBlock(N, N);MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);...
}

每个块的线程数量是有限的,因为所有的线程都应该存在于同一个流式多处理器核心上,并且必须共享这个核心的有限的内存资源。在当前的GPU上,一个线程块可能包含多达1024个线程。

然而,一个内核可以被多个形状相同的线程块执行,这样总的线程数量等于每个块的线程数量乘以块的数量。

块被组织成一个一维、二维或三维的线程块网格,如图4所示。网格中线程块的数量通常由被处理的数据的大小决定,这通常超过了系统中的处理器数量。

在这里插入图片描述

图4:线程块的网格

<<<...>>>语法中指定的每个块的线程数量和每个网格的块数量可以是int类型或dim3类型。可以像上面的例子那样指定二维的块或网格。

网格内的每个块都可以通过一个一维、二维或三维的唯一索引来标识,这个索引在内核中可以通过内置的blockIdx变量访问。线程块的维度在内核中可以通过内置的blockDim变量访问。

下面的代码是将前面的MatAdd()示例扩展为处理多个块的情况:

// 内核定义
__global__ void MatAdd(float A[N][N], float B[N][N],float C[N][N])
{int i = blockIdx.x * blockDim.x + threadIdx.x;int j = blockIdx.y * blockDim.y + threadIdx.y;if (i < N && j < N)C[i][j] = A[i][j] + B[i][j];
}
int main()
{...∕∕ Kernel invocationdim3 threadsPerBlock(16, 16);dim3 numBlocks(N ∕ threadsPerBlock.x, N ∕ threadsPerBlock.y);MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);...
}

线程块的大小为16x16(256个线程),尽管在这种情况下是任意的,但这是一个常见的选择。如前所述,创建足够的块,每个矩阵元素有一个线程。为了简单起见,这个例子假设每个维度的网格线程数可以被该维度的块线程数整除,尽管并非一定是这种情况。

线程块需要独立执行:它们必须能够按任何顺序执行,可以并行也可以串行。这种独立性要求允许线程块按任何顺序在任何数量的核心上调度,如图3所示,使程序员能写出随核心数量扩展的代码。

块内的线程可以通过共享一些共享内存和同步它们的执行来协调内存访问来合作。更精确地说,可以在内核中通过调用__syncthreads()内在函数来指定同步点;__syncthreads()充当一个屏障,所有的块内线程都必须在这里等待,直到允许任何线程继续。共享内存章节给出了一个使用共享内存的例子。除了__syncthreads(),协作组API章节还提供了一整套丰富的线程同步原语。

为了有效的合作,预计共享内存是接近每个处理器核心的低延迟内存(很像L1缓存),并且__syncthreads()预计是轻量级的。

5.2.1 线程块集群 Thread Block Clusters

设计线程数:共享内存上同步 < 线程块集群内同步 < 全局同步

随着NVIDIA计算能力9.0的引入,CUDA编程模型引入了一个叫做线程块群集的可选等级层次,这些都是由线程块构成的。与线程块中的线程保证在流式多处理器上被并行调度类似,线程块群集中的线程块也保证在GPU处理集群(GPC)上进行并行调度。

与线程块类似,群集也以一维、二维或三维的方式组织,如图5所示。群集中的线程块数量可以由用户定义,CUDA中支持以8个线程块为单位的群集大小作为最大限制。注意,在GPU硬件或MIG配置中,如果太小以致不能支持8个多处理器,那么最大群集大小将相应减小。识别这些较小的配置,以及支持线程块群集大小超过8的较大配置,是架构特定的,并可以使用cudaOccupancyMaxPotentialClusterSize API进行查询。

在这里插入图片描述

图5:线程块集群的网格

在使用群集支持启动的内核中,为了兼容性,gridDim变量仍然表示线程块数量的大小。可以通过使用Cluster Group API找到群集中块的等级。

线程块群集可以通过使用__cluster_dims__(X,Y,Z)的编译器时间内核属性或使用CUDA内核启动API cudaLaunchKernelEx在内核中启用。下面的例子展示了如何使用编译器时间内核属性启动一个群集。使用内核属性的群集大小在编译时固定,然后可以使用经典的<<<,>>>来启动内核。如果内核使用编译时群集大小,那么在启动内核时,群集大小不能被修改。

∕∕ 内核定义, Compile time cluster size 2 in X-dimension and 1 in Y and Z dimension
__global__ void __cluster_dims__(2, 1, 1) cluster_kernel(float *input, float* output)
{ }
int main()
{float *input, *output;∕∕ Kernel invocation with compile time cluster sizedim3 threadsPerBlock(16, 16);dim3 numBlocks(N ∕ threadsPerBlock.x, N ∕ threadsPerBlock.y);∕∕ The grid dimension is not affected by cluster launch, and is still enumerated∕∕ using number of blocks.∕∕ The grid dimension must be a multiple of cluster size.cluster_kernel<<<numBlocks, threadsPerBlock>>>(input, output);
}

线程块群集大小也可以在运行时设置,并通过CUDA内核启动API cudaLaunchKernelEx启动内核。下面的代码示例展示了如何使用可扩展API启动一个群集内核。

∕∕ 内核定义
∕∕ No compile time attribute attached to the kernel
__global__ void cluster_kernel(float *input, float* output)
{ }
int main()
{float *input, *output;dim3 threadsPerBlock(16, 16);dim3 numBlocks(N ∕ threadsPerBlock.x, N ∕ threadsPerBlock.y);∕∕ Kernel invocation with runtime cluster size{cudaLaunchConfig_t config = {0};∕∕ The grid dimension is not affected by cluster launch, and is still enumerated∕∕ using number of blocks.∕∕ The grid dimension should be a multiple of cluster size.config.gridDim = numBlocks;config.blockDim = threadsPerBlock;cudaLaunchAttribute attribute[1];attribute[0].id = cudaLaunchAttributeClusterDimension;attribute[0].val.clusterDim.x = 2; ∕∕ Cluster size in X-dimensionattribute[0].val.clusterDim.y = 1;attribute[0].val.clusterDim.z = 1;config.attrs = attribute;config.numAttrs = 1;cudaLaunchKernelEx(&config, cluster_kernel, input, output);}
}

在具有9.0计算能力的GPU中,群集中的所有线程块都保证在单个GPU处理集群(GPC)上进行协同调度,并允许群集中的线程块使用Cluster Group API cluster.sync()执行硬件支持的同步。群集组还提供成员函数,使用num_threads()num_blocks()API分别查询群集组大小,以线程数量或块数量表示。群集组中的线程或块的排名可以使用dim_threads()dim_blocks()API分别进行查询。

属于群集的线程块可以访问分布式共享内存。群集中的线程块具有读取、写入和执行分布式共享内存中任何地址的原子操作的能力。分布式共享内存章节给出了在分布式共享内存中执行直方图的示例。

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

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

相关文章

C++设计模式:单例模式(十)

1、单例设计模式 单例设计模式&#xff0c;使用的频率比较高&#xff0c;整个项目中某个特殊的类对象只能创建一个 并且该类只对外暴露一个public方法用来获得这个对象。 单例设计模式又分懒汉式和饿汉式&#xff0c;同时对于懒汉式在多线程并发的情况下存在线程安全问题 饿汉…

【优选算法专栏】专题十三:队列+宽搜(一)

本专栏内容为&#xff1a;算法学习专栏&#xff0c;分为优选算法专栏&#xff0c;贪心算法专栏&#xff0c;动态规划专栏以及递归&#xff0c;搜索与回溯算法专栏四部分。 通过本专栏的深入学习&#xff0c;你可以了解并掌握算法。 &#x1f493;博主csdn个人主页&#xff1a;小…

智能边缘自动化:HDMI接口钡铼ARM工业电脑实践案例

一款具备HDMI接口的高性能ARM工业计算机应运而生&#xff0c;为实现在工业4.0时代的关键数据实时处理与可视化管理提供了强有力的硬件支撑。这款计算机依托其独特的边缘计算能力&#xff0c;完美解决了工业环境中大规模数据传输至云端的高延迟问题&#xff0c;成功实现了OT&…

Asp .Net Core 系列:集成 Refit 和 RestEase 声明式 HTTP 客户端库

背景 .NET 中 有没有类似 Java 中 Feign 这样的框架&#xff1f;经过查找和实验&#xff0c;发现 在 .NET 平台上&#xff0c;虽然没有直接的 Feign 框架的端口&#xff0c;但是有一些类似的框架和库&#xff0c;它们提供了类似的功能和设计理念。下面是一些在 .NET 中用于声明…

基于springboot实现宠物咖啡馆平台管理系统项目【项目源码+论文说明】计算机毕业设计

基于springboot实现宠物咖啡馆平台演示 摘要 随着信息技术在管理上越来越深入而广泛的应用&#xff0c;管理信息系统的实施在技术上已逐步成熟。本文介绍了基于Spring Boot的宠物咖啡馆平台的设计与实现的开发全过程。通过分析基于Spring Boot的宠物咖啡馆平台的设计与实现管理…

ExoPlayer停止更新,建议升级到AndroidX Media3

1. 大家常用的ExoPlayer地址&#xff1a;GitHub - google/ExoPlayer: An extensible media player for Android ExoPlayer是谷歌官方提供的媒体播放库&#xff0c;大家在开发项目中经常使用ExoPlayer播放音视频&#xff0c;谷歌官方已经明确表示该库在2024-04-03停止更新&…

【图论】详解链式前向星存图法+遍历法

细说链式前向星存图法 首先要明白&#xff0c;链式前向星的原理是利用存边来进行模拟图。 推荐左神的视频–建图、链式前向星、拓扑排序 比方说有这样一张图&#xff0c;我们用链式前向星来进行模拟时&#xff0c;可以将每一条边都进行编号&#xff0c;其中&#xff0c;红色的…

MAC: 自己制作https的ssl证书(自己签发免费ssl证书)(OPENSSL生成SSL自签证书)

MAC: 自己制作https的ssl证书(自己签发免费ssl证书)(OPENSSL生成SSL自签证书) 前言 现在https大行其道, ssl又是必不可少的环节. 今天就教大家用开源工具openssl自己生成ssl证书的文件和私钥 环境 MAC电脑 openssl工具自行搜索安装 正文 1、终端执行命令 //生成rsa私钥&…

自动化 单元测试Test

XCTest测试框架(单元测试XCTests、性能测试XCPPerformanceTests、用户界面测试XCUItests) 单元测试XCTests&#xff1a;测试应用中事件或逻辑是否预期工作。 用户界面测试XCUItests&#xff1a;测试用户与应用的UI交互(如点击按钮、滑动屏幕)。 性能测试XCPPerformanceTests&am…

云手机解决海外社媒运营的诸多挑战

随着海外社交媒体运营的兴起&#xff0c;如何有效管理多个账户成为了一项挑战。云手机作为一种新兴的解决方案&#xff0c;为海外社媒运营带来了前所未有的便利。 云手机的基本原理是基于云计算和虚拟化技术&#xff0c;允许用户在物理手机之外创建和使用多个虚拟手机。这种创新…

校园论坛系统

文章目录 校园论坛系统一、项目演示二、项目介绍三、10000字论文参考四、系统部分功能截图五、部分代码展示六、底部获取项目和10000字论文参考&#xff08;9.9&#xffe5;&#xff09; 校园论坛系统 一、项目演示 校园论坛系统 二、项目介绍 基于springbootvue的前后端分离…

SpringBoot3 + uniapp 对接 阿里云0SS 实现上传图片视频到 0SS 以及 0SS 里删除图片视频的操作(最新)

SpringBoot3 uniapp 对接 阿里云0SS 实现上传图片视频到 0SS 以及 0SS 里删除图片视频的操作 最终效果图uniapp 的源码UpLoadFile.vuedeleteOssFile.jshttp.js SpringBoot3 的源码FileUploadController.javaAliOssUtil.java 最终效果图 uniapp 的源码 UpLoadFile.vue <tem…

AI大模型引领未来智慧科研暨ChatGPT自然科学高级应用

以ChatGPT、LLaMA、Gemini、DALLE、Midjourney、Stable Diffusion、星火大模型、文心一言、千问为代表AI大语言模型带来了新一波人工智能浪潮&#xff0c;可以面向科研选题、思维导图、数据清洗、统计分析、高级编程、代码调试、算法学习、论文检索、写作、翻译、润色、文献辅助…

机器学习实训 Day1

线性回归练习 Day1 手搓线性回归 随机初始数据 import numpy as np x np.array([56, 72, 69, 88, 102, 86, 76, 79, 94, 74]) y np.array([92, 102, 86, 110, 130, 99, 96, 102, 105, 92])from matplotlib import pyplot as plt # 内嵌显示 %matplotlib inlineplt.scatter…

设计模式——责任链模式13

责任链模式 每个流程或事物处理 像一个链表结构处理。场景由 多层部门审批&#xff0c;问题分级处理等。下面体现的是 不同难度的问题由不同人进行解决。 设计模式&#xff0c;一定要敲代码理解 传递问题实体 /*** author ggbond* date 2024年04月10日 07:48*/ public class…

数据结构-----链表

目录 1.顺序表经典算法 &#xff08;1&#xff09;移除元素 &#xff08;2&#xff09;合并数组 2.链表的创建 &#xff08;1&#xff09;准备工作 &#xff08;2&#xff09;建结构体 &#xff08;3&#xff09;链表打印 &#xff08;4&#xff09;尾插数据 &#xff…

【unity】【C#】UGUI组件

文章目录 UI是什么对UI初步认识 UI是什么 UI是用户界面&#xff08;User Interface&#xff09;的缩写&#xff0c;它是用户与软件或系统进行交互的界面。UI设计旨在提供用户友好的界面&#xff0c;使用户能够轻松地使用软件或系统。UI设计包括界面的布局、颜色、字体、图标等…

Github Benefits 学生认证/学生包 新版申请指南

本教程适用于2024年之后的Github学生认证申请&#xff0c;因为现在的认证流程改变了很多&#xff0c;所以重新进行了总结这方面的指南。 目录 验证教育邮箱修改个人资料制作认证文件图片转换Base64提交验证 验证教育邮箱 进入Email settings&#xff0c;找到Add email address…

Java集合List

List特有方法 经典多态写法 // 经典的多态写法 List<String> list new ArrayList<>();常用API&#xff1a;增删改查 // 添加元素 list.add("Java"); // 添加元素到指定位置 list.add(0, "Python");// 获取元素 String s list.get(0);// 修改…

Docker容器嵌入式开发:在Ubuntu上配置Postman和flatpak

在 Ubuntu 上配置 Postman 可以通过 Snap 命令完成&#xff0c;以下是所有命令的总结&#xff1a; sudo snap install postmansudo snap install flatpak在 Ubuntu 上配置 Postman 和 Flatpak 非常简单。以下是一些简单的步骤&#xff1a; 配置 Flatpak 安装 Flatpak&#x…