CUDA编程之线程模型
时间:2022-07-22
本文章向大家介绍CUDA编程之线程模型,主要内容包括其使用实例、应用技巧、基本知识点总结和需要注意事项,具有一定的参考价值,需要的朋友可以参考一下。
CUDA编程之线程模型
CUDA线程模型概述
线程模型
CUDA线程层次
线程层次——二维Block
线程层次——三维Block
通过上图线程层次可划分为:
网格(Grid)
- 一Kernel映射一网格
- 网格在设备上执行
- 划分为线程块
线程块(Block)
- 发射到SM上执行
- 利用共享存储器通信
- 划分为线程
线程(Thread)
- 映射到SP上执行
五个内建变量
运行时获得网格和块的尺寸及线程索引等信息。
-
gridDim
:包含三个元素x, y, z的结构体,表示网格在x,y,z方向上的尺寸,对应于执行配置中的第一个参数。 -
blockDim
:包含三个元素x, y, z的结构体,表示块在x,y,z方向上的尺寸,对应于执行配置的第二个参数 -
blockIdx
:包含三个元素x, y, z的结构体,分别表示当前线程所在块在网格中x, y, z方向上的索引 -
threadIdx
:包含三个元素x, y, z的结构体,分别表示当前线程在其所在块中x, y, z方向上的索引 -
warpSize
:表明warp
的尺寸,在计算能力1.0的设备中,这个值是24,在1.0以上的设备中,这个值是32。
Kernel分配线程
一个kernel结构如下:Kernel<<>>(param1, param2, …)
- Dg:
grid
的尺寸,说明一个grid
含有多少个block
,为dim3
类型,一个grid
最多含有65535 * 65535 * 65535个block
,Dg.x
,Dg.y
,Dg.z
最大值为65535
; - Db:
block
的尺寸,说明一个block
含有多少个thread
,为dim3
类型,一个block
最多含有1024
(cuda2.x版本)个threads
,Db.x
和Db.y
最大值为1024,Db.z
最大值64;(举个例子,一个block
的尺寸可以是:1024 * 1 * 1 | 256 * 2 * 2 | 1 * 1024 * 1 | 2 * 8 * 64 | 4 * 4 * 64
等) - Ns:可选参数,如果
kernel
中由动态分配内存的shared memory
,需要在此指定大小,以字节为单位; - S:可选参数,表示该
kernel
处在哪个流当中。
CUDA向量加法深入理解grid、block、thread的关系及thread索引的计算
CUDA编程流程
- CPU在GPU上分配内存:cudaMalloc;
- CPU把数据发送到GPU:cudaMemcpy;
- CPU在GPU上启动内核(kernel),它是自己写的一段程序,在每个线程上运行;
- CPU把数据从GPU取回:cudaMemcpy;
- CPU释放GPU上的内存。
CUDA向量加法源代码
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i] + b[i];
}
int main()
{
// 定义向量长度
const int arraySize = 16;
// 向量a与b初始化填值
int a[arraySize], b[arraySize], c[arraySize];
for (int i = 0; i < arraySize; i++) {
a[i] = b[i] = i;
}
// 向量c初始化为零
int c[arraySize] = { 0 };
// 向量加法并行化计算
cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addWithCuda failed!");
return 1;
}
//打印输出结果
for (int i = 0; i < arraySize; i++) { // 打印出来方便观察
cout << c[i] << " ";
}
cout << endl;
//退出之前调用cudaDeviceReset,以便分析和跟踪工具(如Nsight和Visual Profiler)显示完整的跟踪。
cudaStatus = cudaDeviceReset();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceReset failed!");
return 1;
}
system("pause");
return 0;
}
//CUDA实现向量加法操作.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
// 初始化设备端地址
int *dev_a = 0;
int *dev_b = 0;
int *dev_c = 0;
cudaError_t cudaStatus;
// 选择0号GPU进行计算
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}
// 为两个输入向量和一个输出向量申请显存
cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
// 将主机端的内存拷贝到设备端的显存
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
// 在GPU上启动一个内核,为每个元素启动一个线程
dim3 grid(1, 1, 1), block(size, 1, 1);
addKernel<<<grid, block>>>(dev_c, dev_a, dev_b);
// 检查启动内核的任何错误
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addKernel launch failed: %sn", cudaGetErrorString(cudaStatus));
goto Error;
}
// cudaDeviceSynchronize等待内核完成并返回(同步)
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!n", cudaStatus);
goto Error;
}
// 将GPU显存上的输出向量复制到主机内存。
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
基于源码的线程分配分析
方式一:grid(1, 1, 1), block(size, 1, 1)
首先定义一个线程如下图所示:
线程示意图
然后直观解释程序中的线程设置
dim3 grid(1, 1, 1), block(size, 1, 1); // 设置参数
在这段代码中,我们设置参数为线程格(grid)中只有一个一维的块(block),该block的x维度上有16个线程。图示如下:
因此,按照此线程分配情况执行代码如下:
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i] + b[i];
}
方式二:grid(1, 1, 1), block(8, 2, 1)
dim3 grid(1, 1, 1), block(8, 2, 1); // 设置参数
直观图示:
索引执行代码
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.y * blockDim.x + threadIdx.x; // 使用了threadIdx.x, threadIdx.x, blockDim.x
c[i] = a[i] + b[i];
}
方式三: grid(16, 1, 1), block(1, 1, 1)
dim3 grid(16, 1, 1), block(1, 1, 1); // 设置参数
直观图示:
索引执行代码
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = blockIdx.x;
c[i] = a[i] + b[i];
}
方式四: grid(4, 1, 1), block(4, 1, 1)
dim3 grid(4, 1, 1), block(4, 1, 1); // 设置参数
直观图示:
索引执行代码
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = blockIdx.x * gridDim.x + threadIdx.x;
c[i] = a[i] + b[i];
}
方式五: grid(2, 2, 1), block(2, 2, 1)
dim3 grid(2, 2, 1), block(2, 2, 1); // 设置参数
直观图示:
索引执行代码
__global__ void addKernel(int *c, const int *a, const int *b)
{
// 在第几个块中 * 块的大小 + 块中的x, y维度(几行几列)
int i = (blockIdx.y * gridDim.x + blockIdx.x) * (blockDim.x * blockDim.y) + threadIdx.y * blockDim.y + threadIdx.x;
c[i] = a[i] + b[i];
}
参考
[CUDA基础(1):操作流程与kernel概念]https://www.cnblogs.com/hankeyyh/p/6580427.html
[【CUDA】grid、block、thread的关系及thread索引的计算]https://blog.csdn.net/hujingshuang/article/details/53097222
- JavaScript 教程
- JavaScript 编辑工具
- JavaScript 与HTML
- JavaScript 与Java
- JavaScript 数据结构
- JavaScript 基本数据类型
- JavaScript 特殊数据类型
- JavaScript 运算符
- JavaScript typeof 运算符
- JavaScript 表达式
- JavaScript 类型转换
- JavaScript 基本语法
- JavaScript 注释
- Javascript 基本处理流程
- Javascript 选择结构
- Javascript if 语句
- Javascript if 语句的嵌套
- Javascript switch 语句
- Javascript 循环结构
- Javascript 循环结构实例
- Javascript 跳转语句
- Javascript 控制语句总结
- Javascript 函数介绍
- Javascript 函数的定义
- Javascript 函数调用
- Javascript 几种特殊的函数
- JavaScript 内置函数简介
- Javascript eval() 函数
- Javascript isFinite() 函数
- Javascript isNaN() 函数
- parseInt() 与 parseFloat()
- escape() 与 unescape()
- Javascript 字符串介绍
- Javascript length属性
- javascript 字符串函数
- Javascript 日期对象简介
- Javascript 日期对象用途
- Date 对象属性和方法
- Javascript 数组是什么
- Javascript 创建数组
- Javascript 数组赋值与取值
- Javascript 数组属性和方法
- 用阻塞队列,再系一次鞋带
- I2C总线架构 之 设备驱动
- kail 安装及卸载 docker【亲测可用】
- mac 登录远程服务器(常规ssh+免密快捷方式)
- git常用操作--分支同步master 本地库提交到远程分支
- mac苹果 配置maven settings文件【注意,与win环境不一样!!!】
- mac小程序开发 本地调试 安装 npm modules
- Deeplearning.ai 课程笔记第一部分:神经网络与深度学习
- 如何从Node.js开始-Visual Studio2017
- Rust竟然没有异常处理?
- MAC系统 JDK 卸载及彻底删除
- 轻松学Pytorch –Mask-RCNN图像实例分割
- docker安装awvs13
- 远程连接kail Permission denied或者refused【已解决】
- 设计模式~原始模型模式(二)