Thread的层次结构

要弄清楚Thread的层次结构,必须先弄清楚它的参照系。看看“The Futures of Ruby Threading”这段话:“Current stable releases of Ruby use user space threads (also called "green threads"), which means that the Ruby interpreter takes care of everything to do with threads. This is in contrast to kernel threads, where the creation, scheduling and synchronization is done with OS syscalls, which makes these operations costly, at least compared to their equivalents in user space threads.”这里的用户空间线程和内核线程是相对于Ruby解释器来说的。如果把参照系换成操作系统,此处的内核线程只是一个用户空间线程。

常说Java线程是Native thread,是说Java的一个线程映射到操作系统上的一个用户线程。往下随后怎么映射就要看操作系统的实现了。可以看看 Solaris的线程模型这里也有些说明资料。

不过Java线程也并不一定是一一映射的,比如Jikes RVM虚拟机,采用了M:N模型,而不是大家经常看到的1:1,具体资料可以看 这里。BEA的JRockit是两个都有,既支持1:1模型,也支持M:N模型,叫Thin Thread。

在JVM上实现了M:N模型,当然在操作系统上也可以实现M:N模型,只是层次不同,抽象级别不同。象Solaris就实现了M:N模型,不过这个模型在Solaris 9中已经放弃,为什么?实现太难。“It is not to say that a good implementation of the M:N model is impossible,but simply that a good 1:1 implementation is probably sufficient. ”想更进一步了解Solaris多线程的发展历程,看这个pdf吧( www.sun.com/software/whitepapers/)。

了解清楚Thread的层次结构,碰到Green Thread,Native Thread,User space Thread,Kernel Thread时才不会糊涂。
### GPU 中线程的次化组织结构 #### 线程块(Thread Block) 在GPU编程模型中,线程按照一定的层次结构进行组织。最基本的单位是线程(Thread),多个线程可以组合成一个线程块(Block)。每个线程块内的线程能够协作完成特定的任务,并且这些线程之间可以通过共享内存来交换数据[^1]。 #### 多维线性化 当涉及到多维线程时,在划分成Warp之前会先将这些维度映射到一维数组中。具体来说,假设有一个二维或三维网格中的线程,则较大坐标值所在的行会被放置于较小坐标的后面,从而实现一种基于行优先级(linearization by row-major order)的一维表示形式。例如,在一个具有`threadIdx.y=1`的所有线程将会紧接在`threadIdx.y=0`之后排列;而相同y轴索引下的线程则依据x方向上的索引依次排序[^3]。 #### Warp 的概念 为了提高硬件利用率以及简化调度逻辑,CUDA架构进一步把线程划分为更小组件——Warps。通常情况下,一个Warp包含固定数量(比如NVIDIA设备上通常是32个)连续编号的线程。同一Warp内部的所有成员在同一时刻执行同一条指令,即采用单指令流(SIMD),这有助于减少分支预测错误带来的性能损失并增强吞吐量。 #### 流处理器簇(Streaming Multiprocessor, SM) 最终,上述提到的各种粒度级别的线程单元都会被分配给不同的SM去处理。每一个SM都拥有自己的资源集,包括但不限于寄存器文件、缓存以及其他必要的组件用于支持其负责范围内所有活动着的线程的工作需求。值得注意的是,实际运行过程中可能会有来自不同block甚至kernel的不同warps共存在同一个sm当中等待被执行[^2]。 ```cpp // CUDA Kernel Example showing thread organization within a block. __global__ void exampleKernel(float* data){ int idx = blockIdx.x * blockDim.x + threadIdx.x; // Each thread performs an operation on its corresponding element of the array 'data'. if(idx < N){ data[idx] *= 2.0f; } } ```
评论
成就一亿技术人!
拼手气红包6.0元
还能输入1000个字符
 
红包 添加红包
表情包 插入表情
 条评论被折叠 查看
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值