4

NVIDIA CUDA2023春训营（三）CUDA 线程层次结构与线程索引

 1 year ago
source link: https://alex-mcavoy.github.io/nvidia/cuda-spring-bootcamp/acc1b378.html
Go to the source link to view the article. You can view the picture content, updated content and better typesetting reading experience. If the link is broken, please click the button below to view the snapshot at that time.
neoserver,ios ssh client

Reference

线程层次结构

核函数在 device 端执行时，会启动若干线程，一个核函数所启动的所有线程被称为一个线程网格（Thread Grid），同一个线程网格上的线程共享相同的全局内存空间，每个线程网格又可分为若干线程块（Thread Block），每个线程块中又包含若干线程

通过如下语句，可以定义一个 3*2 的 grid，每个 block 中含有 5*3 的 thread 的二维线程组织，其中 gird 和 block 是定义为 dim3 类型的变量，在定义时，缺省值初始化为 1，可以灵活的定义 1-dim、2-dim、3-dim 结构

dim3 grid(3, 2, 1);
dim3 block(5, 3, 1);
kernel_function<<<grid, block>>>(prams...);

需要注意的是，每个 grid 中最多含有 65535 个 block，而每个 block 中最多含有 1024 个线程

相应的线程层次如下图所示

03-1.png

在 CUDA 中，一个线程需要两个内置的坐标变量 (blockIdx，threadIdx) 来标识，它们都是 dim3 类型变量，其中 blockIdx 指明线程所在的 block 在 grid 中的位置，而 threaIdx 指明线程在 block 中的位置

blockIdx 与 threadIdx 可通过内置变量来获得：

  • threadIdx.[x y z]：执行当前核函数的线程在 block 中的索引值
  • blockIdx.[x y z]：执行当前核函数的线程所在的 block 在 grid 中的索引值

例如，threadIdx.x 是执行当前核函数的线程在 block 中的 x 方向的序号，blockIdx.x 是执行当前核函数的线程所在 block 在 grid 中的 x 方向的序号

#include <stdio.h>

__global__ void hello_from_gpu() {
    const int bx = blockIdx.x;
    const int tx = threadIdx.x;
    printf("block: %d, thread: %d\n", bx, tx);
}

int main() {
    hello_from_gpu<<<4, 4>>>();
    cudaDeviceSynchronize();
    return 0;
}

线程唯一索引

唯一索引计算公式

一个线程块上的线程是放在同一个流式多处理器（SM）上的，但是单个 SM 的资源有限，这就使得 block 中的线程数是有限制的，现代 GPUs 的 block 可支持的线程数可达 1024 个，这就使得有时若想要知道一个线程在所有线程中的全局 ID，就必须要知道相应的组织结构，通过以下两个内置变量，可获得相应的组织结构信息

  • blockDim.[x y z]：一个 block 中包含多少个线程
  • gridDim.[x y z]：一个 grid 中包含多少个 block

对于线程唯一索引，需要知道以下三个信息：

  • blockIndex：block 在整个 grid 中的索引（1 维到 3 维）
  • blockSize ：block 的大小，描述其中含有多少个线程
  • threadId ：线程在 block 中的索引（1 维到 3 维）

进而有唯一索引计算公式：index = blockIndex * blockSize + threadIndex

1D Grid, 1D Block

对于 1-dim 的 grid 与 1-dim 的 block，有：

// block在整个一维grid中的索引
int blockIndex = blockIdx.x;
// 一维block的大小
int blockSize = blockDim.x;
// 线程在一维block中的索引
int threadIndex = threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

1D Grid, 2D Block

对于 1-dim 的 grid 与 2-dim 的 block，有：

// block在整个一维grid中的索引
int blockIndex = blockIdx.x;
// 二维block的大小
int blockSize = blockDim.x * blockDim.y;
// 线程在二维block中的索引
int threadIndex = blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

1D Grid, 3D Block

对于 1-dim 的 grid 与 3-dim 的 block，有：

// block在整个一维grid中的索引
int blockIndex = blockIdx.x;
// 三维block的大小
int blockSize = blockDim.x * blockDim.y * blockDim.z;
// 线程在三维block中的索引
int threadIndex = blockDim.x * blockDim.y * threadIdx.z + blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

2D Grid, 1D Block

对于 2-dim 的 grid 与 1-dim 的 block，有：

// block在整个二维grid中的索引
int blockIndex = gridDim.x * blockIdx.y + blockIdx.x
// 一维block的大小
int blockSize = blockDim.x;
// 线程在一维block中的索引
int threadIndex = threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

2D Grid, 2D Block

对于 2-dim 的 grid 与 2-dim 的 block，有：

// block在整个二维grid中的索引
int blockIndex = gridDim.x * blockIdx.y + blockIdx.x
// 二维block的大小
int blockSize = blockDim.x * blockDim.y;
// 线程在二维block中的索引
int threadIndex = blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

2D Grid, 3D Block

对于 2-dim 的 grid 与 3-dim 的 block，有：

// block在整个二维grid中的索引
int blockIndex = gridDim.x * blockIdx.y + blockIdx.x
// 三维block的大小
int blockSize = blockDim.x * blockDim.y * blockDim.z;
// 线程在三维block中的索引
int threadIndex = blockDim.x * blockDim.y * threadIdx.z + blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

3D Grid, 1D Block

对于 3-dim 的 grid 与 1-dim 的 block，有：

// block在整个三维grid中的索引
int blockIndex = gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
// 一维block的大小
int blockSize = blockDim.x;
// 线程在一维block中的索引
int threadIndex = threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

3D Grid, 2D Block

对于 3-dim 的 grid 与 2-dim 的 block，有：

// block在整个三维grid中的索引
int blockIndex = gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
// 二维block的大小
int blockSize = blockDim.x * blockDim.y;
// 线程在二维block中的索引
int threadIndex = blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

3D Grid, 3D Block

对于 3-dim 的 grid 与 3-dim 的 block，有：

// block在整个三维grid中的索引
int blockIndex = gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
// 三维block的大小
int blockSize = blockDim.x * blockDim.y * blockDim.z;
// 线程在三维block中的索引
int threadIndex = blockDim.x * blockDim.y * threadIdx.z + blockDim.x * threadIdx.y + threadIdx.x;
// 线程在整个grid中的索引
int index = blockIndex * blockSize + threadIndex;

report

Recommend

About Joyk


Aggregate valuable and interesting links.
Joyk means Joy of geeK