核心判断只有一句：网格（Grid）组织线程块（Block），Block 组织线程（Thread）；全局索引把线程坐标映射到数据位置。① 从“线程能启动”到“线程知道自己做什么”上一篇已经让一个内核（kernel）真正运行起来，但只看见"有线程输出"还不够。现在的问题变成：当一个 kernel 启动很多线程时，每个线程怎样知道自己在整个任务中的位置？CUDA 用Grid → Block → Thread的层次结构组织线程。一次内核启动（kernel launch）对应一个 Grid，Grid 由多个 Block 组成，每个 Block 又包含多个 Thread；Grid 和 Block 都可以是一维、二维或三维的。下面这张图展示了这个层次的嵌套关系：图中从上到下依次是 Grid、Block、Thread 三层：Grid 包含多个 Block，每个 Block 包含多个 Thread，Thread 是其中最小的执行单位——你写的每个线程都在这个层次里拥有确定的位置。一次 kernel launch 会指定一个 Grid，Grid 由多个 Block 组成。Block 是线程的组织和协作单位，Thread 是 CUDA 编程模型中的基本执行单位。这种层次化组织让同一个 kernel 可以用不同规模的 Grid 和 Block 处理不同大小的数据集；线程只需要根据自己的坐标确定要处理的数据位置。如下图所示，多个 Block 组成一个 Grid，每个 Block 又包含多个 Thread。图中的数量用于说明层次，不代表实际程序必须使用同样的规模：这张图重点展示的是组织关系，而不是线程执行先后。线程如何获得坐标，接下来用一个一维例子说明。② 一维坐标：Block 内编号如何变成全局索引上一篇的2, 4已经启动了 2 个 Block、每个 Block 4 个 Thread。本节让每个线程打印自己的三项信息：所在 Block、Block 内编号，以及我们计算出的全局索引。#includecstdio#includecuda_runtime.h// 每个线程计算自己在线性 Grid 中的全局索引__global__voidprint_1d_coordinates(){// blockIdx.x：当前 Block 在 Grid 中的编号// blockDim.x：每个 Block 在 x 方向包含的线程数// threadIdx.x：当前 Thread 在 Block 内的编号intglobal_id=blockIdx.x*blockDim.x+threadIdx.x;printf("block=%d thread=%d global=%d\n",blockIdx.x,threadIdx.x,global_id);}intmain(){// 启动 2 个 Block，每个 Block 包含 4 个 Thread，共 8 个 Threadprint_1d_coordinates2,4