CUDA Programming Model--CUDA编程模型
CUDA Programming Model
1 Kernels
CUDA c++ 扩展了c++,允许编程者定义C++ 函数,被称为kernel。每次执行,由N个不同的CUDA线程执行N次。
每个执行内核的线程拥有一额独一无二的线程ID,可以通过内置的threadIdx变量在内核中访问(在块内是唯一的,并不一定是全局唯一的)
// Kernel definition
__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);...
}
这里只用的一个块,因此threadIdx是唯一的
2 thread hierarchy 线程层次
threadIdx 是一个3维向量。所以线程可以使用一维,二维,三维索引标识,形成一维,二维,三维的线程块。
线程索引和线程ID直接相关:
- 一维 他俩相同
- 二维 对于(Dx,Dy)(D_x,D_y)(Dx,Dy)的块, 线程索引为 (x,y)的线程ID是 x+yDxx+yD_xx+yDx
- 三维 对于(Dx,Dy,Dz)(D_x,D_y,D_z)(Dx,Dy,Dz)的块, 索引为(x,y,z)的线程ID为(x+yDx+zDxDy)(x+yD_x+zD_xD_y)(x+yDx+zDxDy)
// Kernel definition
__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);...
}
dim3可以看成是包含三个无符号整数(x,y,z)成员的结构体变量(主要为了确定线程位置)。在定义时,缺省值初始化为1,dim2就是(x,y,1)。一个线程需要两个内置的坐标变量(blockIdx,threadIdx)来唯一标识,它们都是dim3类型变量,其中blockIdx指明线程所在grid中的位置,而threaIdx指明线程所在block中的位置
- 块内的线程必须在同一个处理器核心中并宫向该核心有限的存储器资源
- 一个线程块可以包含多达1024个线程
- 一个内核可被多个同样大小的线程块执行,所以总的线程数等于每个块内的线程数乘以线程块数
- 线程块被组织成一维、二维、三维的 grid
- 线程块内线程数和网格内线程块数由<<< … >>>语法确定,参数可以是整形或者dim3类型
- grid 内的每个块可以通过一维二维三维索引唯一确定,索引通过blockIdx查到,块的维度由blockDim访问
// Kernel definition
__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);...
}
线程块必须独立执行:而且能够以任意顺序,串行或者并行执行。这种独立性要求使得线程块可以以任何顺序在任意数目核心上调度
- 块内线程可通过共享存储器和同步执行协作,共享存储器可以共享数据,同步执行可以协调存储器访问。
2.1 Thread Block Clusters 线程块簇
CUDA 编程引入了一个可选的层次,由线程块组成的 线程块簇
- 同一个簇的线程块,可以保证被一个 GPU Processing Cluster(GPC) co-scheduled
- 簇 也可以有三个维度
- 一个簇中的线程块数量可以自定义,但最多只能有8个线程块
- 使用 cluster_dims(X,Y,Z) 和 API cudaLaunchKernelEx 启用cluster,cluster的size是固定的在编译前就确定了
// Kernel definition
// 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);
}
- cudaLaunchKernelEx 可以运行时改变参数
// Kernel definition
// 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);cluster_kernel<<<numBlocks, threadsPerBlock>>>();// 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);}
}
- 属于同一个cluster的线程块可以方位分布式共享内存
3 Memory Hierarchy
- CDUA 线程可以从多个存储空间访问数据,每个线程有私有的local memory
- 每个线程块有对所有块内县城可见的共享内存,共享内存的生命周期和线程块相同
- 每个cluster 中的线程块可以对其他块的宫向内存进行读写和原子操作
- 每个线程都可以访问global memory
- 对于所有线程,有两个额外的 read-only 内存空间 常量和纹理存储器空间。全局,常量和纹理存储器空间为不同的存储器用途作了优化
4 Heterogeneous Programming 异构编程
- CUDA编程模型假设CUDA线程在物理上独立的设备上执行,设备作为主机的协处理器运行C++程序
- kernel 在GPU上执行,剩余的C++程序在CPU上执行
- CUDA编程模型也假设host和device 分别维护各自的变成空间
- 程序通过调用CUDA 运行时,来管理对内核可见的全局、常量和纹理存储器空间
- 统一内存通过managed memory提供 host和device的内存空间的桥
- CPU GPU都可以访问menaged memory ,作为一个单一的连贯的拥有同一地址空间的内存
5 Asynchronous SIMT Programming Model 异步SIMT 编程模型
- CUDA 编程中线程是最低的计算或者内存抽象层次
- 异步编程模型定义了 为了线程间同步的 Asynchronous Barrier的行为。
- 模型也定义了 cuda::memcpy_aync 在GPU计算时,可以异步地在global memory 中移动数据
Asychronous Barrier 异步屏障
- 引入了GPU std::barrier 的实现,允许用户自定义barrier对象范围
简单的同步模式
- 使用syncthreads() 或者group.sync()
#include <cooperative_groups.h>__global__ void simple_sync(int iteration_count) {auto block = cooperative_groups::this_thread_block();for (int i = 0; i < iteration_count; ++i) {/* code before arrive */block.sync(); /* wait for all threads to arrive here *//* code after wait */}
}
所有线程都在同步点 block.sync() 处被阻塞,知道所有线程都到达这个点
5.1 Asynchronous Operations 异步操作
- 一个异步操作定义为,由CUDA线程启动并由另一个线程异步执行地操作
- 在格式良好的程序中,一个或多个 CUDA 线程同步由异步操作
这里没看明白
6 Compute Capability 计算能力
- 设备的计算能力由版本号表示,有时也称为“SM version”
CUDA Programming Model--CUDA编程模型相关推荐
- CUDA刷新器:CUDA编程模型
CUDA刷新器:CUDA编程模型 CUDA Refresher: The CUDA Programming Model CUDA,CUDA刷新器,并行编程 这是CUDA更新系列的第四篇文章,它的目标是 ...
- cuda编程_CUDA刷新器:CUDA编程模型
CUDA刷新器:CUDA编程模型 CUDA Refresher: The CUDA Programming Model CUDA,CUDA刷新器,并行编程 这是CUDA更新系列的第四篇文章,它的目标是 ...
- 【GPGPU编程模型与架构原理】第二章 2.1 计算模型
本章介绍以CUDA和OpenCL 并行编程中的一些核心架构概念来展示GPGPU的计算.编程和存储模型.本章还介绍虚拟指令集和机器指令集,逐步揭开GPGPU体系结构的面纱. 2.1 计算模型 计算模 ...
- CUDA C++ Programming Guide——编程模型
本章通过概述CUDA编程模型在C ++中的使用方式,介绍了其主要概念. 编程接口中给出了CUDA C ++的广泛描述.本章和下章使用的向量加法示例的完整代码可以在vectorAdd CUDA示例中找到 ...
- GPU硬件结构和编程模型(源于nvidia的CUDA文档)
GPU的硬件结构 GPU通过一个可扩展的多线程流式多处理器(SMs)构建.一个multiprocessor可以在同一时间处理上百个线程.为了管理这些线程,使用一个特殊的结构SIMT.利用单线程中指令级 ...
- CUDA C 编程指导(二):CUDA编程模型详解
CUDA编程模型详解 本文以vectorAdd为例,通过描述C在CUDA中的使用(vectorAdd这个例子可以在CUDA sample中找到.)来介绍CUDA编程模型的主要概念.CUDA C的进一步 ...
- 一种基于CUDA标准的异构并行编程模型开发简介
一种基于CUDA标准的异构并行编程模型开发简介 目录 一.绪论 1.1研究背景及意义 1.2目标平台体系结构简介 二.HPPA基本组成结构 三.编译工具链开发 3.1 拆分工具HPCufe开发 3.2 ...
- CUDA是Nvidia开发的一种并行计算平台和编程模型,用于在其自己的GPU(图形处理单元)上进行常规计算
https://baike.baidu.com/item/CUDA/1186262?fr=aladdin CUDA是Nvidia开发的一种并行计算平台和编程模型,用于在其自己的GPU(图形处理单元)上 ...
- GPU — CUDA 编程模型
目录 文章目录 目录 GPGPU CUDA 编程模型 CUDA 的架构 CUDA 的工作原理 Grid.Block.Thread Warp GPGPU GPU 起初是用来处理图像的,但是后来人们发现其 ...
最新文章
- Ubuntu 彻底卸载 OpenCV
- hihocoder offer收割编程练习赛11 B 物品价值
- Java8(jdk1.8)中文档注释处理工具javadoc的环境参量配置及使用方法
- 一套基础自动化部署搭建过程
- VIM 多行注释与取消
- 微信从原版到现在所有界面图片_微信突然宣布:现在能改微信号了,所有人都能改...
- C#中Attribute的继承
- 关于SQL Server 数据库归档的一些思考和改进
- python使用redis在实际场景使用_Python使用Redis实现作业调度系统(超简单)
- java netty聊天室_netty实现消息中心(二)基于netty搭建一个聊天室
- CSS中的未定义行为,浏览器的差异(一)
- 云栖专辑 | 阿里开发者们的第6个感悟:享受折磨 1
- 计算机视觉中 RNN 应用于目标检测
- python小波分解与重构_python - 使用pyWavelets进行多级局部小波重构 - 堆栈内存溢出...
- 基于机器学习的GitHub敏感信息泄露监控
- 计算机拆装与维修技能综述,综述虚拟机在计算机硬件组装与维护教学中的应用...
- react调试工具Reactdevelopertools
- 基于java的田径运动会报名系统
- Android 开源项目和文章集合(更新:2022.03.21)
- 军事指挥系统时间同步解决方案
热门文章
- jquery表格插件jqgrid
- 2020年起重机司机(限桥式起重机)考试及起重机司机(限桥式起重机)答案解析
- 微信支付失败中关于“签名错误”的解决方案
- reddit_如何减少Reddit的吸吮
- Logo就是商标吗?别再被扰乱视线了
- easyui图片放大功能(取巧)
- 计算机前端总线频率,前端总线频率
- 健身与不健身五年后的差别?你可不能轻易忽视!
- [TYVJ 1927] 『Citric II』一道防AK好题 · 模拟
- 关于将EXCEL文件导入到MYSQL数据库的一些方法