《CUDA编程》2.CUDA中的线程组织

news2024/11/23 12:29:25

0 来自GPU的hello world

在visua studio 中新建一个CUDA runtime项目,然后把kernel.cu中的代码删掉,输入以下代码

#include"cuda_runtime.h"
#include"device_launch_parameters.h"

#include<stdio.h>

__global__ void hello_from_gpu(void) {
	printf("Hello from GPU!\n");
}

int main(void) {
	hello_from_gpu<<<1, 1>>>();
	cudaDeviceSynchronize();
	return 0;
}

hello_from_gpu<<<1, 1>>>();该条代码可能会飘红标错,可以不用管

输出结果如下
在这里插入图片描述
如果创建了CUDA runtime项目依旧无法直接运行,请确保是否正确安装了CUDA,或者在项目设置中添加以下三个cuda文件路径
在这里插入图片描述
在这里插入图片描述

1 只有主机函数的CUDA程序

CUDA程序的编译器驱动(compilerdriver)nvcc支持编译纯粹的C++代码,所以可以写一个C++代码,然后修改后缀为.cu,然后使用nvcc去编译(nvcc会把C++的部分交给C++编译器去编译,然后其余部分则自己编译

请打开一个记事本,输入以下代码,然后保存为后缀为.cu的文件
在这里插入图片描述

#include<stdio.h>

int main(void)
{
	printf("Hello world!");
	return 0;
}

在该文件目录下打开终端,输入nvcc hello.cu -o hello,会把hello.cu文件,使用nvcc将其编译为一个hello的.exe文件,输入.\hello,输出结果如下
在这里插入图片描述

2 使用核函数的CUDA程序

上一个例子只是使用nvcc编译器进行编译,但是并没有调用GPU,下面将介绍如何在GPU中的输出Hello world!,在《CUDA编程》1.GPU硬件与CUDA环境搭建中已经介绍了,GPU必须在CPU的控制之下,才能进行运算,所以一个典型的、简单的CUDA程序的结构具有下面的形式:
在这里插入图片描述
hsot对device的调用是通过核函数(kernel function)来实现的。

CUDA中的核函数与C++中的函数必须被限定词__global__ 修饰,前后是双下划线,核函数的返回类型必须是void,void和__global__次序随意,下面是一个核函数。

__global__ void hello_from_gpu(void) {
	printf("Hello from GPU!\n");
}

2.1 分析hello world代码

请关注main()函数部分

#include"cuda_runtime.h"
#include"device_launch_parameters.h"

#include<stdio.h>

__global__ void hello_from_gpu(void) {
	printf("Hello from GPU!\n");
}

int main(void) {
	hello_from_gpu<<<1, 1>>>();
	cudaDeviceSynchronize();
	return 0;
}

2.1.1 核函数的调用格式

hello_from_gpu<<<1, 1>>>();与C++格式函数hello_from_gpu();相比,多一个一对三尖括号,里面还有用逗号隔开的两个数字。因为一个GPU中有多个核心,即可以支持多线程(thread),所以host在调用一个核函数时,必须指明需要在设备中指派多少个线程

核函数中的线程常组织为若干线程块(thread_block):三括号中的第一个数字可以看作线程块的个数,第二个数字可以看作每个线程块中的线程数

一个核函数的全部线程块构成一个网格(grid),而线程块的个数就记为网格大小(grid_size)。每个线程块中含有同样数目的线程,该数目称为线程块大小(block_size),则

                   核函数中总的线程数 = Grid_size x Block_size

所以三尖括号意义是<<<Grid_size , Block_size>>>,所以,在上述程序中,主机只指派了设备的一个线程,网格大小和线程块大小都是1,即1 x 1 = 1

核函数中也支持printf(),但只支持<stdio.h>,不支持C++的

cudaDeviceSynchronize();是调用了一个CUDA的runtime的函数API,去掉这个函数就打印不出字符串,作用是同步主机与设备,所以能够促使缓冲区刷新,后面会详解该函数

3 CUDA中的线程组织

3.1 核函数中使用多个线程

一个GPU往往有几千个计算核心,而总的线程数必须至少等于计算核心数时才有可能充分利用GPU中的全部计算资源。

实际上,总的线程数大于计算核心数时才能更充分地利用GPU中的计算资源,因为这会让计算和内存访问之间及不同的计算之间合理地重叠,从而减小计算核心空闲的时间。

假设线程我们要使用8个线程进行运算,代码如下

#include"cuda_runtime.h"
#include"device_launch_parameters.h"

#include<stdio.h>

__global__ void hello_from_gpu(void) {
	printf("Hello from GPU!\n");
}

int main(void) {
	hello_from_gpu<<<2, 4>>>();
	cudaDeviceSynchronize();
	return 0;
}

输出结果如下,会输出8行 Hello from GPU!:
在这里插入图片描述
核函数中代码的执行方式是“单指令-多线程”,即每一个线程都执行同一串指令。

3.2 使用线程索引

每个线程在核函数中都有一个唯一的身份标识,代码中使用了grid_size 、block_size,所以可以使用这两个参数来表明每个线程。

#include"cuda_runtime.h"
#include"device_launch_parameters.h"

#include<stdio.h>

__global__ void hello_from_gpu(void) {
	const int tid = threadIdx.x;//线程id
	const int bid = blockIdx.x;//线程块id
	printf("Hello from GPU! tid:%d bid:%d\n",tid,bid);
}

int main(void) {
	hello_from_gpu<<<2, 4>>>();
	cudaDeviceSynchronize();
	return 0;
}

输出结果如下,便可以看出每一句输出来自于拿哪一个线程块的哪一个线程:
在这里插入图片描述
注意:每个线程块的计算是相互独立的,所以有时候会是bid:1先执行完,有时会是bid:0先执行完

3.3 推广至多维网络

blockIdx 和 threadIdx 是类型为 uint3 的变量,即结构如下:
在这里插入图片描述
即可以用结构体dim3定义“多维”的网格和线程块,例如我们要定义一个2 x 2 x 1的grid和2 x 3 x 1的block,结果如下:在这里插入图片描述
而在核函数中的线程组织图如下所示:
在这里插入图片描述
grid的形状为2 x 2 x 1,所以有4个block;
block的形状是3 x 2 x 1,所以每个block中有6个线程

一个线程块中的线程还可以细分为不同的线程束(threadwarp),目前所有的GPU架构都是32,所以,一个线程束就是连续的32个线程。即一个线程块中第0到第31个线程属于第0个线程束,第32到第63个线程属于第1个线程束,依此类推。
在这里插入图片描述
修改代码如下:

#include"cuda_runtime.h"
#include"device_launch_parameters.h"

#include<stdio.h>

__global__ void hello_from_gpu(void) {
	const int tidx = threadIdx.x;
	const int tidy = threadIdx.y;
	const int bid = blockIdx.x;

	printf("Hello from tid:(%d,%d) bid:%d\n",tidx,tidy,bid);
}

int main(void) {
	hello_from_gpu<<<2, 4>>>();
	cudaDeviceSynchronize();
	return 0;
}

输出结果是:
在这里插入图片描述

3.4 网格与线程块大小的限制

CUDA中对能够定义的网格大小和线程块大小做了限制。

  • 网格大小在x、y和z这3个方向的最大允许值分别为2^31-1、65535和65535。
  • 线程块大小在x、y和z这3个方向的最大允许值分别为1024、1024和64。
  • 线程块总的大小,即blockDim.x、blockDim.y和blockDim.z的乘积不能大于1024,即不管如何定义,一个线程块最多只能有1024个线程。

4 CUDA中的头文件

在使用nvcc编译器驱动编译.cu文件时,将自动包含必要的CUDA头文件,如<cuda.h><cuda_runtime.h>,且为<cuda.h>包含了<stdlib.h>,但在visual studio里编写时,需要手动添加

5 用nvcc编译CUDA程序时,计算能力问题

CUDA的编译器驱动(compilerdriver)nvcc先将全部源代码分离为host代码和device代码。主机代码完整地支持C++语法,但设备代码只部分地支持C++。

nvcc先将设备代码编译为PTX伪汇编代码,再将PTX代码编译为二进制的cubin目标代码。

在将源代码编译为PTX代码时,需要用选项-arch=compute_XY指定一个虚拟架构的计算能力;在将PTX代码编译为cubin代码时,需要用选项-code=sm_ZW指定一个真实架构的计算能力

真实架构的计算能力必须等于或者大于虚拟架构的计算能力

所以
在这里插入图片描述
可以编译,但下面的设置则不能编译,因为虚拟架构计算能力小于真实架构
在这里插入图片描述
本文的程序在编译时并没有通过编译选项指定计算能力。这是因为编译器有一个默认的计算能力。
在这里插入图片描述
设置计算能力的必要性:

  • 兼容性:不同的 GPU 架构有不同的指令集和硬件特性。通过设置计算能力,你可以确保生成的代码能够在特定的 GPU 上运行。如果不设置计算能力,nvcc 默认可能会生成针对最新架构的代码,这可能导致代码在旧的 GPU 上无法运行。
  • 性能优化:不同的 GPU 架构有不同的优化策略。设置计算能力可以确保 nvcc 生成的代码针对特定的 GPU 进行优化,从而提高性能。

以下是一些常见的计算能力及其对应的 GPU 架构:
在这里插入图片描述

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

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

相关文章

VM ware的安装——个人使用

VM ware的安装 Workstation 和 Fusion 对个人使用完全免费&#xff0c;企业许可转向订阅 如果没有官方账号需要注册一个 选择个人下载&#xff0c;会跳转到下一个页面 要勾选同意&#xff0c;才能下载 点击下载之后还会跳转到填写地址的页面&#xff0c;填写完同意后&#x…

合宙LuatOS应用,与时间相关那些事

合宙嵌入式操作系统LuatOS——在蜂窝物联网模组上推出开源二次开发框架&#xff0c;功能齐全性能稳定&#xff0c;可大幅度降低用户的研发成本和研发周期。 在LuatOS中&#xff0c;获取时间函数用得最多的就是os.time()函数了。接下来&#xff0c;我会讲一些与这个函数以及其他…

美团外卖霸王餐系统有哪些优势?如何选择?

外卖霸王餐系统是一种流行的营销策略&#xff0c;旨在通过提供免费或优惠的餐食来吸引新顾客和提升品牌知名度。美团外卖平台本身拥有巨大的流量&#xff0c;美团霸王餐作为美团外卖平台推出的营销服务&#xff0c;能够充分利用这一流量优势。商家在平台上发布霸王餐活动信息&a…

美业门店怎么提升业绩?连锁美业门店管理系统收银系统拓客系统源码

美业门店想要提升业绩&#xff0c;需要考虑多方面的因素&#xff0c;并采取综合性的方法。以下是一些可以考虑的因素和建议&#xff1a; 产品与服务优化&#xff1a; 提供高质量的美容产品和服务&#xff0c;确保顾客满意度。不断更新产品线&#xff0c;引入新的时尚趋势&#…

Ansible流程控制-条件_循环_错误处理_包含导入_块异常处理

文章目录 Ansible流程控制介绍1. 条件判断2. 循环3. 循环控制4. 错误处理5. 包含和导入6. 块和异常处理7. 角色的流程控制*include_tasks、import_tasks_include之间的区别 条件语句再细说且、或、非、是模糊条件when指令的详细使用方法 循环语句再细说如何使用使用item变量结合…

应用targetSdkVersion升级指导

应用targetSdkVersion升级指导 应电信终端产业协会&#xff08;TAF&#xff09;发布的《移动应用软件高 API 等级预置与分发自律公约》&#xff08;以下简称《公约》&#xff09;要求&#xff1a;截止到2019年5月1日所有新发布的应用 API 必须为26或更高&#xff0c;2019年8月…

What is the OpenAI Chat Completion API tools/functions property format?

题意&#xff1a;OpenAI 聊天完成 API 的工具/函数属性格式是什么 问题背景&#xff1a; Is there any clear documentation on the format of OpenAIs Chat Completion API tools/functions object format? I understand its JSON, but there appear to be underlying requi…

tauri开发软件中,使用tauri自带的api用浏览器打开指定的url链接

有能力的可以看官方文档&#xff1a;shell | Tauri Apps 就是使用这个api来打开指定的url链接&#xff0c;要在tauri.config.json中配置打开这个api&#xff1a; 然后在前端页面中导入使用&#xff1a; import { open } from tauri-apps/api/shell; // opens the given URL o…

Cpp类和对象(下)(6)

文章目录 前言一、初始化列表概念使用注意实际运用explicit关键字初始化列表的总结 二、static成员static成员的概念static成员的特性static的一个实用场景 三、友元友元函数友元类 四、内部类概念特性 五、匿名对象六、再次理解封装和面向对象总结 前言 Hello&#xff0c;本篇…

redis学习(013 实战:黑马点评:优惠券秒杀——超卖问题解决方案)

黑马程序员Redis入门到实战教程&#xff0c;深度透析redis底层原理redis分布式锁企业解决方案黑马点评实战项目 总时长 42:48:00 共175P 此文章包含第52p-第p53的内容 文章目录 问题演示使用jmeter测试两百个并发请求 超卖的原因分析解决方案 加锁悲观锁介绍乐观锁介绍乐观锁…

XXL-Job 监控消息队列消息数量预警

1、什么是Basic Authentication认证 Basic Authentication 是一种常用的 HTTP 认证机制&#xff0c;用于保护 Web 资源免受未授权访问。在这种认证方式中&#xff0c;客户端&#xff08;通常是浏览器&#xff09;需要在 HTTP 请求头中提供用户凭据&#xff08;通常是用户名和密…

Leetcode 最小覆盖子串

解题思路&#xff1a; 哈希表存储字符频率&#xff1a;首先统计字符串 t 中每个字符出现的次数。滑动窗口&#xff1a;用两个指针 left 和 right 来标记当前窗口的左右边界&#xff0c;不断右移 right&#xff0c;直到包含了所有 t 中的字符。然后尝试右移 left&#xff0c;缩…

python爬虫/引用requests/基本使用

1.安装requests 进入控制台使用该命令安装requests pip3 install requests 2.对网站使用get请求 这里用对网站进行get请求&#xff0c;然后打印。 import requests //引用requestsresponse requests.get(urlhttps://www.bilibili.com/)print(response.text) 3.对网站使用…

2024全国研究生数学建模竞赛(数学建模研赛)ABCDEF题深度建模+全解全析+完整文章

全国研究生数学建模竞赛&#xff08;数学建模研赛&#xff09;于9月21日8时正式开赛&#xff0c;赛程4天半&#xff0c;咱这边会在开赛后第一时间给出对今年的6道赛题的评价、分析和解答。包括ABCDEF题深度建模全解全析完整文章&#xff0c;详情可以点击底部的卡片来获取哦。 …

座椅空置状态检测系统源码分享

座椅空置状态检测检测系统源码分享 [一条龙教学YOLOV8标注好的数据集一键训练_70全套改进创新点发刊_Web前端展示] 1.研究背景与意义 项目参考AAAI Association for the Advancement of Artificial Intelligence 项目来源AACV Association for the Advancement of Computer…

大模型之基准测试集(Benchmark)-给通义千问2.0做测评的10个权威测基准测评集

引言 在去年(2023)云栖大会上&#xff0c;阿里云正式发布千亿级参数大模型通义千问2.0。据现场介绍&#xff0c;在10个权威测评中&#xff0c;通义千问2.0综合性能超过GPT-3.5&#xff0c;正在加速追赶GPT-4。以下是通义千问在MMLU、C-Eval、GSM8K、HumanEval、MATH等10个主流…

基于Springboot共享充电宝管理系统JAVA|VUE|SSM计算机毕业设计源代码+数据库+LW文档+开题报告+答辩稿+部署教+代码讲解

源代码数据库LW文档&#xff08;1万字以上&#xff09;开题报告答辩稿 部署教程代码讲解代码时间修改教程 一、开发工具、运行环境、开发技术 开发工具 1、操作系统&#xff1a;Window操作系统 2、开发工具&#xff1a;IntelliJ IDEA或者Eclipse 3、数据库存储&#xff1a…

openEuler普通用户su root时Permission denied

openEuler普通用户su root时Permission denied 背景&#xff1a; openEuler默认普通用户是不能通过su切换到root用户的 如果想通过su切换到root&#xff0c;有以下两个解决办法 1、修改/etc/pam.d/su 文件 [rootlocalhost ~]# vim /etc/pam.d/su #修改21行&#xff0c;将“…

视频怎么制作成二维码?视频轻松生成二维码的3步操作

现在很多人为了能够更快捷的实现视频内容的分享&#xff0c;会通过将视频生成二维码的方式&#xff0c;让其他人可以通过扫描二维码来查看视频内容。这种方式不需要用户存储视频&#xff0c;扫码就能够在设备上查看视频&#xff0c;有利于提升查看视频的便捷性&#xff0c;可以…

图片压缩工具免费怎么找?归纳了这几个压缩工具

有哪些图片压缩工具免费&#xff1f;在数字化时代&#xff0c;图像已成为我们生活中不可或缺的一部分。无论是网站设计、社交媒体分享还是文件传输&#xff0c;高质量的图片都扮演着重要的角色。但高质量往往意味着大文件体积&#xff0c;这可能会导致加载速度变慢或存储空间不…