版本记录

前言

很多做视频和图像的，相信对这个框架都不是很陌生，它渲染高级3D图形，并使用GPU执行数据并行计算。接下来的几篇我们就详细的解析这个框架。感兴趣的看下面几篇文章。
1. Metal框架详细解析（一）—— 基本概览
2. Metal框架详细解析（二） —— 器件和命令（一）
3. Metal框架详细解析（三） —— 渲染简单的2D三角形（一）
4. Metal框架详细解析（四） —— 关于GPU Family 4（一）
5. Metal框架详细解析（五） —— 关于GPU Family 4之关于Imageblocks（二）
6. Metal框架详细解析（六） —— 关于GPU Family 4之关于Tile Shading（三）
7. Metal框架详细解析（七） —— 关于GPU Family 4之关于光栅顺序组（四）
8. Metal框架详细解析（八） —— 关于GPU Family 4之关于增强的MSAA和Imageblock采样覆盖控制（五）
9. Metal框架详细解析（九） —— 关于GPU Family 4之关于线程组共享（六）
10. Metal框架详细解析（十） —— 基本组件（一）
11. Metal框架详细解析（十一） —— 基本组件之器件选择 - 图形渲染的器件选择（二）
12. Metal框架详细解析（十二） —— 基本组件之器件选择 - 计算处理的设备选择（三）
13. Metal框架详细解析（十三） —— 计算处理（一）
14. Metal框架详细解析（十四） —— 计算处理之你好，计算（二）

About Threads and Threadgroups - 关于线程和线程组

了解Metal如何组织计算处理工作负载。

Overview - 概览

回想一下Hello Compute，当您调度计算传递时，Metal会在1D，2D或3D网格上执行您的内核函数。 网格中的每个点代表内核函数的单个实例，称为线程。 例如，在图像处理中，网格通常是线程的2D矩阵 - 表示整个图像 - 每个线程对应于正被处理的图像的单个像素。

线程被组织成一起执行的线程组，并且可以共享公共的内存块。 虽然有时内核函数的设计使得线程彼此独立运行，但线程组中的线程在其工作集上进行协作也很常见。

Identification of Threads by Position in Grid - 按网格中的位置识别线程

图1显示了计算内核处理的图像如何划分为线程组以及每个线程组如何由各个线程组成。 每个线程处理一个像素。

Figure 1 A grid divided into threadgroups that are composed of individual threads.

线程可以通过它在网格中的位置来识别；这是一个独特的位置，允许你的内核函数为每个线程做一些不同的事情。 下面的Hello Compute中的示例内核函数显示了线程在网格中的位置如何作为参数传递给函数。 在这种情况下，参数gid是表示2D坐标的向量，用于读取和写入纹理中的特定位置。

kernel void
grayscaleKernel(texture2d<half, access::read>  inTexture  [[texture(AAPLTextureIndexInput)]],
                texture2d<half, access::write> outTexture [[texture(AAPLTextureIndexOutput)]],
                uint2                          gid        [[thread_position_in_grid]])
{
    if((gid.x >= outTexture.get_width()) || (gid.y >= outTexture.get_height()))
    {
        return;
    }
    half4 inColor  = inTexture.read(gid);
    half  gray     = dot(inColor.rgb, kRec709Luma);
    outTexture.write(half4(gray, gray, gray, 1.0), gid);
}

[[thread_position_in_grid]]是属性限定符。 可通过双方括号语法识别的属性限定符允许将内核参数绑定到资源和内置变量，在这种情况下，线程在网格中的位置对应内核函数。

例如，给定一个16 x 16线程的网格划分为2 x 4线程组的8 x 4线程，单个线程（如图2中的红色所示）在网格中的位置为（9,10）：

Figure 2 The position of a single thread in a 16 x 16 grid.

Identification of Threads by Position in Threadgroup - 线程组中按位置识别线程

线程组中线程的位置也可用作属性限定符[[thread_position_in_threadgroup]]，并且线程组在网格中的位置可用作[[threadgroup_position_in_grid]]。

根据网格的形状，这些位置属性可以是标量值，也可以是两元素或三元素矢量。 在2D网格的情况下，位置属性是双元素向量，原点位于左上角。

图2中标识的线程位于线程组中，网格中的位置为（1,2），其在该线程组中的位置为（1,2），如图3所示：

Figure 3 The position of a single thread in a threadgroup

使用以下代码，您还可以根据线程在线程组中的位置以及线程组在网格中的大小和位置来计算线程在网格中的位置：

kernel void 
myKernel(uint2 threadgroup_position_in_grid   [[ threadgroup_position_in_grid ]],
         uint2 thread_position_in_threadgroup [[ thread_position_in_threadgroup ]],
         uint2 threads_per_threadgroup        [[ threads_per_threadgroup ]]) 
{
    
    uint2 thread_position_in_grid = 
        (threadgroup_position_in_grid * threads_per_threadgroup) + 
        thread_position_in_threadgroup;
}

SIMD Groups - SIMD组

线程组中的线程进一步组织成单指令，多数据（SIMD）组，也称为warps或wavefronts，它们并发执行。 SIMD组中的线程执行相同的代码。避免编写可能导致内核功能发散的代码；也就是说，遵循不同的代码路径。分歧的典型示例是使用if语句引起的。即使SIMD组中的单个线程采用与其他线程不同的路径，该组中的所有线程也会执行两个分支，并且该组的执行时间是两个分支的执行时间的总和。

线程组到SIMD组的划分由Metal定义。它在内核执行期间，在具有相同启动参数的给定内核的调度之间以及在调度内从一个线程组到另一个线程组保​​持不变。

SIMD组中的线程数由计算管道状态对象的 threadExecutionWidth返回。属性限定符允许您访问线程组中SIMD组的标量索引，以及SIMD组中线程的标量索引：

• [[simdgroup_index_in_threadgroup]]
  • SIMD组在其线程组中的唯一标量索引，也称为通道ID(lane ID)。
• [[thread_index_in_simdgroup]]
  • SIMD组中线程的唯一标量索引。

虽然线程组可以是多维的，但SIMD组是1D。 因此，SIMD组中线程的位置是所有线程组形状的标量值。 SIMD组大小保持不变，不受线程组大小的影响。

例如，使用图2中所示的相同16 x 16网格，线程执行宽度为16，单个8 x 4线程组由2个SIMD组组成。 因为SIMD组包含16个线程，所以每个SIMD组在线程组中构成2行：

Figure 4 A threadgroup composed of 2 SIMD groups

图5中以红色显示的线程的[[simdgroup_index_in_threadgroup]]值为1，[[thread_index_in_simdgroup]]值为1：

Figure 5 The position of a single thread in a SIMD group

后记

本篇主要讲述了关于线程和线程组，感兴趣的给个赞或者关注~~~