CUDA之——深入理解threadIdx
来源:互联网 发布:linux解压缩zip文件 编辑:程序博客网 时间:2024/05/22 07:09
摘要
本文主要讲述CUDA的threadIdx。
1. Grid,Block和Thread三者的关系
其中,一个grid包含多个blocks,这些blocks的组织方式可以是一维,二维或者三维。任何一个block包含有多个Threads,这些Threads的组织方式也可以是一维,二维或者三维。举例来讲:比如上图中,任何一个block中有10个Thread,那么,Block(0,0)的第一个Thread的ThreadIdx是0,Block(1,0)的第一个Thread的ThreadIdx是11;Block(2,0)的第一个Thread的ThreadIdx是21,......,依此类推,不难整理出其中的映射公式(表达式已在代码中给出)。
2. GridID,BlockID,ThreadID三者的关系
ThreadID是线性增长的,其目的是用于在硬件和软件上唯一标识每一个线程。CUDA程序中任何一个时刻,每一个线程的ThreadIdx都是特定唯一标识的!grid,block的划分方式不同,比如一维划分,二维划分,或者三维划分。显然,Threads的唯一标识ThreadIdx的表达方式随着grid,block的划分方式(或者说是维度)而不同。下面通过程序给出ThreadIdx的完整的表达式。其中,由于使用的时候会考虑到GPU内存优化等原因,代码可能也会有所不同,但是threadId的计算的表达式是相对固定的。
/**************************************************************/// !!!!!!!!!!!!!!注意!!!!!!!!!!!!!!!!/**************************************************************/// grid划分成a维,block划分成b维,// 等价于// blocks是a维的,Threads是b维的。// 这里,本人用的是第一中说法。/**************************************************************/// 情况1:grid划分成1维,block划分为1维。__device__ int getGlobalIdx_1D_1D() {int threadId = blockIdx.x *blockDim.x + threadIdx.x;return threadId;}// 情况2:grid划分成1维,block划分为2维。__device__ int getGlobalIdx_1D_2D() { int threadId = blockIdx.x * blockDim.x * blockDim.y+ threadIdx.y * blockDim.x + threadIdx.x;return threadId; }// 情况3:grid划分成1维,block划分为3维。__device__ int getGlobalIdx_1D_3D() { int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z + threadIdx.z * blockDim.y * blockDim.x+ threadIdx.y * blockDim.x + threadIdx.x;return threadId;}// 情况4:grid划分成2维,block划分为1维。__device__ int getGlobalIdx_2D_1D() { int blockId = blockIdx.y * gridDim.x + blockIdx.x;int threadId = blockId * blockDim.x + threadIdx.x;return threadId;}// 情况5:grid划分成2维,block划分为2维。__device__ int getGlobalIdx_2D_2D() {int blockId = blockIdx.x + blockIdx.y * gridDim.x;int threadId = blockId * (blockDim.x * blockDim.y)+ (threadIdx.y * blockDim.x) + threadIdx.x;return threadId;}// 情况6:grid划分成2维,block划分为3维。__device__ int getGlobalIdx_2D_3D() {int blockId = blockIdx.x + blockIdx.y * gridDim.x;int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)+ (threadIdx.z * (blockDim.x * blockDim.y))+ (threadIdx.y * blockDim.x) + threadIdx.x;return threadId;}// 情况7:grid划分成3维,block划分为1维。__device__ int getGlobalIdx_3D_1D() {int blockId = blockIdx.x + blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;int threadId = blockId * blockDim.x + threadIdx.x;return threadId;}// 情况8:grid划分成3维,block划分为2维。__device__ int getGlobalIdx_3D_2D() {int blockId = blockIdx.x + blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;int threadId = blockId * (blockDim.x * blockDim.y)+ (threadIdx.y * blockDim.x) + threadIdx.x;return threadId;}// 情况9:grid划分成3维,block划分为3维。__device__ int getGlobalIdx_3D_3D() {int blockId = blockIdx.x + blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)+ (threadIdx.z * (blockDim.x * blockDim.y))+ (threadIdx.y * blockDim.x) + threadIdx.x;return threadId;}
3. GPU Threads与CPU Threads的比较
GPU Threads的生成代价小,是轻量级的线程;CPU Threads的生成代价大,是重量级的线程。CPU Threads虽然生成的代价高于GPU Threads,但其执行效率高于GPU Threads,所以GPU Threads无法在个体的比较上取胜,只有在数量上取胜。在这个意义上来讲,CPU Threads好比是一头强壮的公牛在耕地,GPU Threads好比是1000头弱小的小牛在耕地。因此,为了保证体现GPU并行计算的优点,线程的数目必须足够多,通常至少得用上1000个GPU线程或者更多才够本,才能很好地体现GPU并行计算的优点!
4. GPU Threads的线程同步
线程同步是针对同一个block中的所有线程而言的,因为只有同一个block中的线程才能在有效的机制中共同访问shared memory。要知道,由于每一个Thread的生命周期长度是不相同的,Thread对Shared Memory的操作可能会导致读写的不一致,因此需要线程的同步,从而保证该block中所有线程同时结束。
- CUDA之——深入理解threadIdx
- CUDA学习笔记之CUDA初步理解
- 技术积累 — IIC之深入理解
- CUDA学习笔记09: 深入理解CUDA的Warp
- 深入理解CUDA点积运算
- CUDA 编程 之 基本概念理解
- 《深入理解计算机系统》之链接——读书笔记
- 《深入理解计算机系统》之链接——实例
- 深入理解Activity——Token之旅
- 深入理解Java注解——注解之基本概念
- JAVA基础之——深入理解java多态性
- 深入理解Activity——Token之旅
- 深入理解 RecyclerView 系列之——ItemAnimator
- 深入理解Activity——Token之旅
- 深入理解css之 —— absolute定位
- 深入理解盒子模型——CSS之BFC详解
- 软件设计之——深入理解日志系统
- 张鑫旭慕课之——深入理解absolute(上)
- Form元素application/x-www-form-urlencoded
- Memcached的基础
- 移动设备开发注意事项
- Android中资源文件中id重用的的问题
- 终端创建.html文件, 移到xcode中
- CUDA之——深入理解threadIdx
- 数据库视频总结(三)
- 剑指offer题解 二叉树的镜像
- CI框架之无限极分类
- 机器学习方向的一些专家大神
- JavaWeb之Filter(过滤器)
- CAS操作
- CSS 隐藏页面元素的 5 种方法介绍
- 为什么静态成员、静态方法中不能用this和super关键字