【Cuda Runtime】dim3线程块与网格尺寸与核函数几种调用的方式

文章目录

前言

CUDA（Compute Unified Device Architecture）是NVIDIA开发的一种并行计算架构，使得开发人员可以利用GPU进行通用计算。为了有效地利用GPU的并行处理能力，CUDA引入了网格（Grid）和线程块（Block）的概念。这些概念帮助开发人员将复杂的计算任务分解成多个可以并行执行的小任务，从而提高计算效率。本文将详细介绍什么是网格尺寸（Grid Size）和线程块（Block Size），以及如何使用dim3来定义它们。

网格尺寸（Grid Size）

在CUDA编程中，网格是由多个线程块组成的二维或三维结构。网格尺寸表示网格中线程块的数量，每个线程块包含一组线程。通过调整网格尺寸，可以控制在GPU上同时执行的线程块的数量，从而影响计算任务的并行度。

定义网格尺寸

网格尺寸由dim3结构体定义，它包含三个无符号整数成员：x、y和z，分别表示网格在每个维度上的尺寸。通常，我们使用二维或三维网格，但可以通过将不需要的维度设置为1来简化。

示例代码

// 定义一个2D网格，每个维度上有16个线程块
dim3 gridSize(16,16);

线程块（Block Size）

线程块是CUDA中最小的计算单位，一个线程块包含多个线程。线程块的尺寸表示一个线程块中线程的数量，它同样由dim3结构体定义。通过调整线程块的尺寸，可以控制单个线程块中的线程数量，从而影响计算任务在单个线程块内的并行度。

定义线程块尺寸

线程块尺寸由dim3结构体定义，包含三个无符号整数成员：x、y和z，分别表示线程块在每个维度上的尺寸。通常，我们使用一维、二维或三维线程块。

示例代码

// 定义一个2D线程块，每个维度上有8个线程
dim3 blockSize(8,8);

网格尺寸和线程块

在CUDA编程中，计算任务被分解成许多可以并行执行的小任务。这些小任务被组织成线程块（Block）和网格（Grid）。为了让你更好地理解它们，可以把线程块和网格想象成一个学校里的班级和年级。

线程块（Block）

线程块就像一个班级，一个班级里有很多学生（线程）。每个线程块中的线程可以相互合作，分享数据并同步工作。线程块的尺寸表示一个线程块中有多少个线程。你可以定义线程块的尺寸，使它适合你的计算任务。

例如：

dim3 blockSize(8,8);// 这表示每个线程块有8 x 8 = 64个线程

网格（Grid）

网格是由多个线程块组成的，就像一个年级由多个班级组成。网格的尺寸表示网格中有多少个线程块。通过调整网格的尺寸，你可以控制GPU上同时执行的线程块数量，从而影响计算任务的并行度。

例如：

dim3 gridSize(16,16);// 这表示网格中有16 x 16 = 256个线程块

线程块与网格

同一个线程块中的所有线程是会一起执行的
但是网格中的线程块不一定一起执行，网格中的线程块会被分配到不同的多处理器（Streaming Multiprocessors, SMs）上执行。由于每个GPU有多个SM，并且每个SM可以同时执行多个线程块，所以理论上，一些线程块可能会同时开始执行，但这并不意味着所有线程块会同时执行。这取决于GPU的硬件资源和调度机制。

结合使用

当你定义了网格尺寸和线程块尺寸后，CUDA会根据这些信息将计算任务分配到GPU上执行。下面是一个简单的示例，展示如何结合使用网格尺寸和线程块尺寸来启动CUDA内核：

__global__ voidmyKernel(){// 这里是内核代码}intmain(){
    dim3 gridSize(16,16);// 定义网格尺寸
    dim3 blockSize(8,8);// 定义线程块尺寸

    myKernel<<<gridSize, blockSize>>>();// 启动CUDA内核cudaDeviceSynchronize();// 同步设备return0;}

核函数中得到该线程的grid，block和thread 的id

在 CUDA 编程中，获取线程在网格（grid）中的位置、线程块（block）中的位置，以及线程在块内的位置，对于编写高效的 GPU 程序非常重要。CUDA 提供了内建变量来帮助你获取这些信息。

内建变量：

CUDA 提供了一些内建变量来获取线程和线程块的索引信息：

1. **threadIdx**：表示当前线程在其线程块中的索引。
2. **blockIdx**：表示当前线程块在网格中的索引。
3. **blockDim**：表示线程块的维度（即线程块中线程的数量）。
4. **gridDim**：表示网格的维度（即网格中线程块的数量）。

这些内建变量都是内建的

struct

类型，你可以用它们来计算线程和线程块的全局索引。

当然，下面是对

blockDim

和

blockIdx

的通俗易懂的介绍，帮助你理解它们的区别：

blockDim

（线程块维度）：

• **什么是 blockDim**：- blockDim 是一个内建变量，表示当前线程块的大小，即线程块中线程的数量和布局。- 想象你在一个大工厂里工作，每个工厂车间（线程块）有一定数量的工人（线程）。blockDim 就是用来描述这个车间有多少个工人，以及他们在车间里的排列方式。
• 如何使用：- blockDim.x 表示车间在 X 方向上的工人数（线程数）。- blockDim.y 表示车间在 Y 方向上的工人数（线程数）。- blockDim.z 表示车间在 Z 方向上的工人数（线程数）。

blockIdx

（线程块索引）：

• **什么是 blockIdx**：- blockIdx 是一个内建变量，表示当前线程块在整个网格中的位置。- 继续用工厂的比喻，整个工厂（网格）里有很多车间（线程块）。blockIdx 就是用来告诉你当前车间在整个工厂中的位置。例如，它告诉你这个车间是工厂中的第几个车间。
• 如何使用：- blockIdx.x 表示当前车间在工厂中 X 方向的位置（车间的编号）。- blockIdx.y 表示当前车间在工厂中 Y 方向的位置（车间的编号）。- blockIdx.z 表示当前车间在工厂中 Z 方向的位置（车间的编号）。

总结：

• blockDim 关注的是 单个线程块 内的布局和大小（每个车间内部有多少工人，以及他们的排列方式）。
• blockIdx 关注的是 整个网格 中线程块的位置（你在工厂的哪个车间）。

示例：

假设你有一个二维网格的 CUDA 程序，其中每个线程块是一个 2x2 的二维网格。

• blockDim：- blockDim.x = 2（每个线程块的宽度有 2 个线程）- blockDim.y = 2（每个线程块的高度有 2 个线程）
• blockIdx：- blockIdx.x = 1（当前线程块在网格的 X 方向上是第 1 个车间）- blockIdx.y = 0（当前线程块在网格的 Y 方向上是第 0 个车间）

这意味着，如果你有一个 4x4 的线程网格（共 16 个线程块，每个块内 2x2 线程），

blockDim

描述了每个线程块的大小，而

blockIdx

告诉你这个线程块在整个网格中的具体位置。

获取线程和线程块的索引

1. 线程块内的索引：- threadIdx.x：线程在其线程块中的 X 维索引（对于一维线程块）。- threadIdx.y：线程在其线程块中的 Y 维索引（对于二维线程块）。- threadIdx.z：线程在其线程块中的 Z 维索引（对于三维线程块）。
2. 线程块在网格中的索引：- blockIdx.x：线程块在网格中的 X 维索引（对于一维网格）。- blockIdx.y：线程块在网格中的 Y 维索引（对于二维网格）。- blockIdx.z：线程块在网格中的 Z 维索引（对于三维网格）。
3. 线程块的维度：- blockDim.x：线程块在 X 维的大小（线程块的宽度）。- blockDim.y：线程块在 Y 维的大小（线程块的高度）。- blockDim.z：线程块在 Z 维的大小（线程块的深度）。
4. 网格的维度：- gridDim.x：网格在 X 维的大小（网格的宽度）。- gridDim.y：网格在 Y 维的大小（网格的高度）。- gridDim.z：网格在 Z 维的大小（网格的深度）。

示例：获取线程和线程块的索引

下面的示例展示了如何在核函数中获取线程和线程块的索引，并计算每个线程的全局索引：

#include<cuda_runtime.h>#include<iostream>

__global__ voidindexKernel(){// 获取线程在块内的索引int tid_x = threadIdx.x;int tid_y = threadIdx.y;int tid_z = threadIdx.z;// 获取线程块在网格中的索引int block_x = blockIdx.x;int block_y = blockIdx.y;int block_z = blockIdx.z;// 获取线程块的维度int block_dim_x = blockDim.x;int block_dim_y = blockDim.y;int block_dim_z = blockDim.z;// 获取网格的维度int grid_dim_x = gridDim.x;int grid_dim_y = gridDim.y;int grid_dim_z = gridDim.z;// 计算线程的全局索引int global_idx_x = tid_x + block_x * block_dim_x;int global_idx_y = tid_y + block_y * block_dim_y;int global_idx_z = tid_z + block_z * block_dim_z;printf("Thread (%d, %d, %d) in Block (%d, %d, %d) in Grid (%d, %d, %d)\n",
           tid_x, tid_y, tid_z,
           block_x, block_y, block_z,
           grid_dim_x, grid_dim_y, grid_dim_z);printf("Global Index: (%d, %d, %d)\n",
           global_idx_x, global_idx_y, global_idx_z);}intmain(){
    dim3 threadsPerBlock(4,4,1);
    dim3 blocksPerGrid(2,2,1);
    
    indexKernel<<<blocksPerGrid, threadsPerBlock>>>();cudaDeviceSynchronize();return0;}

解释：

• threadIdx：表示线程在当前线程块中的位置。通过 threadIdx.x, threadIdx.y, 和 threadIdx.z 获取。
• blockIdx：表示线程块在网格中的位置。通过 blockIdx.x, blockIdx.y, 和 blockIdx.z 获取。
• blockDim：表示线程块的维度。通过 blockDim.x, blockDim.y, 和 blockDim.z 获取。
• gridDim：表示网格的维度。通过 gridDim.x, gridDim.y, 和 gridDim.z 获取。

核函数的不同调用方式

使用"网格+线程块的调用方式"

使用"网格+线程块的调用方式"就是这种：

myKernel<<<gridSize, blockSize>>>();

使用共享内存

使用共享内存的方式调用核函数

kernel<<<numBlocks, numThreads, sharedMemSize>>>(args...);

numBlocks：线程块的数量。
numThreads：每个线程块中的线程数量。
sharedMemSize：每个线程块需要的动态共享内存的大小（以字节为单位）。这个参数对于共享内存的使用是必要的，但在静态共享内存（如使用

__shared__

关键字声明的共享内存）中，这个参数可以为 0

核函数中定义共享内存

在 CUDA 编程中，共享内存是一种特殊类型的内存，它在一个线程块（block）内的所有线程之间是共享的。共享内存比全局内存（global memory）要快得多，因为它位于 L1 或 L2 缓存中，具有较低的访问延迟。共享内存的主要目的是在同一线程块内的线程之间提供快速的数据交换和通信。

共享内存的基本特性

1. 线程块级别共享：- 共享内存是线程块级别的，即所有线程块中的线程可以访问同一块共享内存，但不同线程块的共享内存是相互独立的。
2. 访问速度：- 共享内存的访问速度比全局内存快得多，因为它在 GPU 的缓存中，有助于减少延迟和带宽限制。
3. 大小限制：- 共享内存的大小是有限的，通常每个线程块的共享内存大小由设备的硬件限制。使用时需要考虑这些限制，以避免超出设备的共享内存容量。

在核函数中使用共享内存

共享内存的声明通常使用

__shared__

关键字，例如：

__global__ voidexampleKernel(){
    __shared__ float sharedMem[256];// 核函数中的代码}

多线程访问共享内存

在此之前，我们需要知道，GPU的多线程其实只做一个非常非常小的事情，而且一般一个线程只处理一个简单的任务，所以其实我们可以不需要锁内存，因为他们根本不会写它范围外的内存，和其他线程一起干就会导致运行效率的变慢

当多个线程在同一时刻访问共享内存时，必须注意以下几个方面，以避免数据竞争和不一致：

1. 同步：- 数据竞争：共享内存中的数据可以被同一线程块内的多个线程访问和修改。如果多个线程同时写入同一位置，可能会导致数据竞争。- 线程同步：使用 __syncthreads() 可以确保线程块中的所有线程在进行数据操作时，所有的线程都到达了同步点，确保共享内存的数据一致性。__syncthreads() 是一个线程块级别的同步原语，它会阻塞线程块中的所有线程，直到所有线程都执行到 __syncthreads()。__global__ voidexampleKernel(float*data){ __shared__ float sharedMem[256];int tid = threadIdx.x;// 加载数据到共享内存 sharedMem[tid]= data[tid];__syncthreads();// 等待所有线程完成共享内存的加载// 在共享内存中进行计算 sharedMem[tid]*=2.0f;__syncthreads();// 等待所有线程完成计算// 将结果写回全局内存 data[tid]= sharedMem[tid];}
2. 内存银行冲突：- 共享内存的访问在硬件上是分银行（banks）的，每个银行能够独立访问。如果多个线程同时访问同一银行，可能会导致内存银行冲突，从而降低性能。设计时需要注意线程对共享内存的访问模式，以减少银行冲突。
3. 内存地址冲突：- 避免线程同时读写共享内存的同一位置，尤其是写操作。读操作通常不会产生冲突，但写操作需要同步以避免不一致的数据。

示例

下面是一个简单的例子，演示了如何在核函数中使用共享内存：

__global__ voidsharedMemoryExample(float*input,float*output){
    __shared__ float sharedData[256];int tid = threadIdx.x;// 读取全局内存到共享内存
    sharedData[tid]= input[tid];__syncthreads();// 确保所有线程都完成了数据加载// 对共享内存中的数据进行一些计算
    sharedData[tid]*=2.0f;__syncthreads();// 确保所有线程都完成了计算// 将结果写回全局内存
    output[tid]= sharedData[tid];}

总结

网格尺寸和线程块尺寸是CUDA编程中的关键概念，它们定义了计算任务在GPU上的并行执行方式。通过合理地选择网格尺寸和线程块尺寸，可以充分利用GPU的计算能力，提高计算效率。在实际编程中，需要根据具体的计算任务和硬件条件，灵活地调整这些参数，以达到最佳性能。理解和掌握dim3结构的使用，对于进行高效的CUDA编程至关重要。