CUDA 中的线程组织

目录

1 CUDA 中的 Hello World 程序

2 CUDA 中的线程组织

2.1 线程模型结构

2.2 一维线程模型

 2.3 推广至多维网格

2.4 线程全局索引计算方式

2.4.1 一维网格 一维线程块

2.4.2 二维网格二维线程块

2.4.3 三维网格三维线程块

2.4.4 其余组合方式

2.5 网格和线程块大小的限制

3 CUDA 中的头文件

4 用 nvcc 编译 CUDA 程序

4.1 nvcc 编译流程

4.2 PTX

4.3 GPU 架构与计算能力

4.4 CUDA 程序兼容性问题

4.4.1 指定虚拟架构计算能力

4.4.2 指定真实架构计算能力

4.4.3 指定多个版本编译

4.4.4 nvcc 即时编译

 4.4.5 nvcc 编译默认计算能力


1 CUDA 中的 Hello World 程序

        CUDA 程序的编译器驱动 nvcc 支持编译纯粹的 C++ 代码。一般来说,一个标准的CUDA 程序中既有纯粹的 C++ 代码,又有不属于 C++ 的真正的 CUDA 代码。CUDA 程序的编译器驱动 nvcc 在编译一个CUDA 程序时,会将纯粹的 C++ 代码交给 C++的编译器去处理,它自己则负责编译剩下的部分。CUDA 程序源文件的扩展名是 .cu。

        GPU 只是一个设备,要他工作的话,需要由主机给他下达命令。这个主机就是 CPU。所以,一个正真正利用了 GPU 的 CUDA 程序既有主机代码,又有设备代码。主机对设备的调用是通过核函数来实现的。核函数在 GPU 上进行并行运行。

        CUDA 中的核函数与 C++ 中的函数是类似的,但一个显著的差别是它必须被限定词 __global__ 修饰。另外,核函数的返回值必须是空类型 void。【void 和限定词的次序可随意】

核函数注意事项:

  • 核函数只能访问 GPU 内存
  • 核函数不能使用变长参数
  • 核函数不能使用静态变量
  • 核函数不能使用函数指针
  • 核函数具有异步性
  • 核函数不支持 C++ 的 iostream
#include<stdio.h>

__global__ void hello_from_gpu(){
    printf("Hello World from the GPU!\n");
}

int main(){
    hello_from_gpu<<<1, 1>>>();
    // 同步函数,等待GPU执行完毕
    cudaDeviceSynchronize(); 
    return 0;
}

2 CUDA 中的线程组织

2.1 线程模型结构

核函数的调用格式为

hello_from_gpu<<<1, 1>>>

         主机在调用一个核函数时,必须指明需要在设备中指派多少个线程,否则设备不知道该如何工作。三括号就是用来指明核函数中的线程数目及排列情况的。
        核函数中的线程常组织为若干线程块:三括号中的第一个数字可以看做是线程块的个数,第二个数字可以看做每个线程块中的线程数。一个核函数的全部线程块构成一个网格,而线程块的个数就记为网格大小。每个线程块中含有同样数目的线程,该数目称为线程块大小。
        所以,核函数中总的线程数就等于网格大小乘以线程块大小,而三括号中的两个数字分别为网格大小和线程块大小,即<<<网格大小, 线程块大小>>>。核函数中代码的执行方式为“单指令-多线程”,即每一个线程都执行同一串指令。

2.2 一维线程模型

        每个线程在核函数中都有一个唯一的身份标识。每个线程的唯一标识由这两个<<<网格大小, 线程块大小>>>确定。这两个值(网格大小:grid_size,线程块大小:block_size)保存在内建变量中,即:

gridDim.x:该变量的数值等于执行配置中变量 grid_size 的值

blockDim.x:该变量的数值等于执行配置中变量 block_size 的值

 类似的,在核函数中预定义了如下标识线程的内建变量:

blockIdx.x:该变量指定一个线程在一个网格中的线程块索引值,范围为0~ gridDim.x-1

threadIdx.x:该变量指定一个线程在一个线程块中的线程索引值,范围为0~ blockDim.x-1

#include<stdio.h>

__global__ void hello_from_gpu(){
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;

    const int id = threadIdx.x + blockDim.x * blockIdx.x;
    printf("Hello World from block %d and thread %d, global id %d!\n", bid, tid, id);
}

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

 一维线程模型的唯一ID可以由下面的公式计算得到:

                        ​​​​​​​        id = threadIdx.x+blockIdx.x\times blockDim.x

 2.3 推广至多维网格

        前述所说的四个内建变量,都使用的结构体的成员变量的语法,其中:

  • blockIdx 和 threadIdx 是类型为 uint3 的变量。该类型是一个结构体,具有x、y、z 这3个成员。所以,blockIdx.x 只是3 个成员中的一个,另外两个分别是 blockId.y 和 blockIdx.z。类似的,threadIdx.x 只是3 个成员中的一个,另外两个分别是 threadId.y 和 threadIdx.z。
  • gridDim 和 blockDim是类型为 dim3 的变量。该类型是一个结构体,具有x、y、z 这3个成员。所以,gridDim .x 只是3 个成员中的一个,另外两个分别是 gridDim .y 和 gridDim .z。类似的,blockDim.x 只是3 个成员中的一个,另外两个分别是 blockDim.y 和 blockDim.z。除了和结构体 uint3 有同样的3个成员外,还在使用C++程序的情况下定义了一些成员函数。

这些内建变量都只在核函数中有效,且满足如下的关系:

2.4 线程全局索引计算方式

2.4.1 一维网格 一维线程块

 定义网格(grid)和线程块(block)的尺寸:

// 其他维度,不写则默认为1
dim3 grid_size(4);
dim3 block_size(8);

 全局唯一标识为

int id = blockDim.x * blockIdx.x + threadIdx.x;

2.4.2 二维网格二维线程块

  定义网格(grid)和线程块(block)的尺寸:

// 其他维度,不写则默认为1
dim3 grid_size(2,2);
dim3 block_size(4,4);

全局唯一标识为

int blockId = blockIdx.x + blockId.y * gridDim.x;

int threadId = threadIdx.y * blockDim.x + threadIdx.x;

int id = blockId * (blockDim.x * blockDim.y) + threadId;

2.4.3 三维网格三维线程块

   定义网格(grid)和线程块(block)的尺寸:

// 其他维度,不写则默认为1
dim3 grid_size(2,2,2);
dim3 block_size(4,4,2);

 全局唯一标识为

int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;

int threadId = (threadIdx.z * (blockDim.x * blockDim.y)) + (threadIdx.y * blockDim.x) + threadIdx.x;

int id = blockId * (blockDim.x *blockDim.y * blockDim.z) + threadId;

2.4.4 其余组合方式

 

2.5 网格和线程块大小的限制

网格大小限制:

线程块大小限制:

不管如何定义,一个线程块最多只能有1024个线程。

3 CUDA 中的头文件

        在使用 nvcc 编译器驱动编译 .cu 文件时,将自动包含必要的 CUDA 头文件,如<cuda.h>和<cuda_runtime.h>。因为 <cuda.h> 包含 <stdlib.h>,顾用 nvcc 编译 CUDA 程序时甚至不用在 .cu文件中包含<stdlib.h>。

4 用 nvcc 编译 CUDA 程序

4.1 nvcc 编译流程

  • nvcc 先将全部源码分离为主机代码和设备代码;
  • nvcc 先将设备代码编译为PTX (Parallel Thread Execution)伪汇编代码,再将PTX代码编译为二进制的 cubin 目标代码;
  • 在将源代码编译为 PTX 代码时,需要用选项 -arch=compute_XY 指定一个虚拟架构的计算能力,用以确定代码中能够使用的 CUDA 功能;
  • 在将 PTX 代码编译为 cubin 代码时,需要用选项 -code=sm_ZW 指定一个真实架构的计算能力,用以确定可执行文件能够使用的 GPU。
  • 真实架构的计算能力一定要大于或等于虚拟架构的计算能力

4.2 PTX

        PTX (Parallel Thread Execution)是 CUDA 平台为基于 GPU 的通用计算而定义的虚拟机和指令集。

        nvcc 编译命令总是使用两个体系结构:一个是虚拟的中间体系结构,另一个是实际的 GPU 体系结构。

        虚拟架构更像是对应用所需的GPU功能的声明。虚拟架构应该尽可能选择低——适配更多实际 GPU,真实架构应该尽可能选择高——充分发挥 GPU 性能。

4.3 GPU 架构与计算能力

         每款 GPU 都有用于标识“计算能力”的版本号。形式为 X.Y,X标识主版本号,代表 GPU 的架构, Y标识次版本号。【GPU 的型号与 GPU的架构是两种不同的概念,型号为Tesla,其架构可以为伏特架构】。

        并非 GPU 的计算能力越高,性能就越强。

4.4 CUDA 程序兼容性问题

        C/C++源码编译为PTX时,可以指定虚拟架构的计算能力,用来确定代码中能够使用的CUDA功能。C/C++源码转化为PTX这一步骤与GPU硬件无关。

4.4.1 指定虚拟架构计算能力

        编译指令:

-arch=compute_XY

XY:第一个数字X代表计算能力的主版本号,第二个数字Y代表计算能力的次版本号。主版本用于确定 GPU 的架构,次版本号用于确定对应版本的计算能力。

4.4.2 指定真实架构计算能力

-code=sm_XY

注意事项:

  • 二进制cubin代码,大版本之间不兼容!!!
  • 指定真实架构计算能力的时候必须指定虚拟架构计算能力!!!
  • 指定的真实架构计算能力必须大于或等于虚拟架构计算能力!!!

4.4.3 指定多个版本编译

        可以同时指定多种计算能力,是的编译出来的可执行文件可以在多 GPU 中执行。

同时指定多组计算能力:

-gencode=arch=compute_35, code=sm_35

-gencode=arch=compute_50, code=sm_50

-gencode=arch=compute_60, code=sm_60

-gencode=arch=compute_70, code=sm_70

编译出来的可执行文件将包括 4 个二进制版本,这样的二进制文件称为胖二进制文件。 

4.4.4 nvcc 即时编译

在运行可执行文件时,从保留的PTX代码临时编译出cubin文件

在可执行文件中保留PTX代码,nvcc编译指令指定所保留的PTX代码虚拟架构,指令为

-gencode arch=compute_60, code=compute_60

注意:

(1)两个计算能力都是虚拟架构计算能力

(2)两个虚拟架构计算能力必须一致

简化编译选项

-arch=sm_XY                         

---->

-gencode arch=compute_XY, code=sm_XY

-gencode arch=compute_XY, code=compute_XY

 4.4.5 nvcc 编译默认计算能力

        不同版本CUDA编译器在编译CUDA代码时,都有一个默认计算能力。

  • 0
    点赞
  • 3
    收藏
    觉得还不错? 一键收藏
  • 2
    评论
评论 2
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值