版权声明:本文为博主原创文章,未经博主允许不得转载。 https://blog.csdn.net/u011988573/article/details/69787614
CUDA编程模型相关笔记[1]
- 学习书籍:[1] 刘金硕等.基于CUDA的并行程序设计.科学出版社.2014
- CUDA编程模型:
- 如上图所示,一个完整的CUDA程序由主机代码和设备代码两部分组成。主机代码部分在主机端CPU上串行执行,是普通的C代码;设备代码部分在设备端GPU上并行执行,称为内核(kernel)。kernel函数不是一个完整的程序,而是任务中能被分解为并行执行的步骤的集合。CPU执行的串行程序负责kernel启动之前进行数据准备和设备初始化的工作,以及在kernel之间进行一些串行计算。GPU执行的并行部分是在被称为grid和block的两个层次并行中完成的,即每个kernel函数存在两个层次的并行:网格(grid)中的线程块(block)间并行和线程块中的线程(thread)间并行。
内核函数
一个完整的CUDA程序包括了在CPU端执行的串行代码和在GPU端执行的并行代码。CUDA程序的编写流程如下:
CUDA程序主要包括以下几个步骤:
初始化GPU
- CUDA程序首次调用Runtime函数时会初始化设备。初始化时,Runtime函数为系统中的每个设备建立一个上下文,该上下文作为设备的主要上下文,被应用的主机线程共享。
主机端本地数据的准备工作
- 所要准备的数据主要指待放入GPU上执行的计算任务,即函数的输入参数集。
为输入参数和输出参数分配显存空间
在显存中可以分配的空间有两种:线性存储器和CUDA数组。- 线性存储器指的是全局存储器,可以分配一维、二维或三维的线性存储空间。其中,一维数组的分配函数为
cudaMalloc()
,二维或三维的线性存储器分配函数为cudaMallocPitch()
和cudaMalloc3D()
。 - CUDA数组可以通过
cudaMalloc3DArray()
分配一维、二维或三维的CUDA数组,而cudaMallocArray()
一般用于分配二维CUDA数组。
- 线性存储器指的是全局存储器,可以分配一维、二维或三维的线性存储空间。其中,一维数组的分配函数为
将输入参数从主机端复制到显存
- 一维、二维、三维陷性存储器分别使用
cudaMemcpy()
、cudaMemcpy2D()
、cudaMemcpy3D()
进行主机端和设备端的数据传输。 - 由
cudaMalloc3DArray()
分配的CUDA数组使用cudaMemcpy3D()
完成与其他CUDA数组或者线性内存的数据传输。
- 一维、二维、三维陷性存储器分别使用
内核启动设置
- CUDA程序使用类似于
kernelFunc<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C);
的语句来启动kernelFunc
函数。其中,kernelFunc
为函数名,<<<>>>
运算符中的blocksPerGrid
和threadsPerBlock
是主机端告诉设备运行时如何启动kernelFunc
函数。blocksPerGrid
表示一个grid中有多少个并行block,threadsPerBlock
表示一个block中有多少个并行thread。(d_A, d_B, d_C)
为kernelFunc
函数的函数参数,和普通C函数一样。 - 在上一篇博客的代码中,执行内核的每个线程都会被分配一个独特的线程ID,可通过
blockDim.x * blockIdx.x + threadIdx.x
在内核中访问此ID(blockDim
指块的维度,类型为dim3;blockIdx
指网格内的块索引,类型为uint3;threadIdx
指块内的线程索引,类型为uint3)。在该程序中,GPU将自动产生blocksPerGrid * threadsPerBlock
个(即1*256个)ID不同的CUDA线程,相当于有256个执行加法命令的线程同时进行操作,从而总体时间缩短了256倍。
- CUDA程序使用类似于
将输入函数从显存复制到主机端
- 该步骤和将输入参数从主机端复制到显存几乎一样,唯一的区别在于,复制时的方向参数不同。将输入参数从主机端复制到显存的复制方向参数为
cudaMemcpyHostToDevice
,而将输入函数从显存复制到主机端的复制方向参数为cudaMemcpyDeviceToHost
。
- 该步骤和将输入参数从主机端复制到显存几乎一样,唯一的区别在于,复制时的方向参数不同。将输入参数从主机端复制到显存的复制方向参数为
释放在设备端分配的显存空间
- 释放线性存储器使用
cudaFree()
函数 - 释放CUDA数组使用
cudaFreeArray()
函数。
- 释放线性存储器使用
线程层次
- 在内核函数中,通过线程的索引来访问线程ID。在一维数组中,可以用一维线程块中的
threadIdx
直接指向相应的ID的线程,但是在二维数据、三维数据中却不相同。对于大小为(Dx,Dy)
的二维块,索引为(x,y)
的线程的ID是(x+yDx)
;对于大小为(Dx,Dy,Dz)
的三维块,索引为(x,y,z)
的线程的ID是(x+yDx+zDxDy)
。 - 针对不同维度的数组的索引方式,可以定义出二维、三维的线程块,去应对不同情况的数据并行方式。