CUDA学习［二］：CUDA编程模型

 CUDA 编程结构

        CUDA编程模型使用由 C 语言拓展生成的注释代码在异构计算系统中执行应用程序。在一个异构环境中包含多个CPU 和 GPU，每个 CPU 和 GPU 的内存都由一条 PCI-Express 总线分隔开。因此我们需要区分以下内容：

        *主机（host）：CPU及其内存（主机内存）

        *设备（device）：GPU及其内存（设备内存）

         为了清楚地指明不同的内存空间，在示例代码中 主机内存中的变量名采用 h_ 开头， 设备内存中的变量名采用 d_ 开头。

         另外，从CUDA 6.0 开始， NVIDIA 提出名为“统一寻址” 的编程模型的改进，他连接了主机内存和设备内存空间，可使用单个指针变量访问 CPU 和 GPU 内存，无须批次之间拷贝数据。

2.1.1 内核 （kernel） 

         内核（kernel）是CUDA 编程模型的一个重要组成部分，其代码在GPU上运行。

         多数情况下，主机可以独立地对设备进行操作，内核一旦被启动，管理权立刻返回主机，释放 CPU 来执行由设备上运行的并行代码实现额外的任务。CUDA编程模型主要是异步的。

        一个典型的CUDA程序包括由并行代码互补的串行代码。串行代码（及任务并行代码）在主机 CPU 上执行，而并行代码在 GPU 上执行。

        主机代码按照 ANSI C 标准进行编写，而设备代码使用 CUDA C 进行编写。 我们可以将所有的代码统一放在一个源文件中，也可以用多个源文件来构建应用程序和库。NVIDA 的 C 编译器（nvcc） 为主机和设备生成可执行代码。

        一个典型的CUDA程序实现流程遵循以下模式：

        1、把数据从CPU内存拷贝到GPU内存

        2、调用核函数对存储在GPU内存中的数据进行操作

        3、将数据从GPU内存传送回到CPU内存

2.1.2 内存管理

        CUDA 编程模型假设系统是由一个主机和一个设备组成的，而且各自拥有独立的内存。核函数（kernel）是在设备端运行的，为了使我们拥有充分的控制权并使系统达到最佳性能，CUDA运行时负责分配和释放设备内存，并且在主机内存和设备内存之间传输数据。

下表是标准的 C 函数和相应的针对内存操作的 CUDA C 函数

标准的C函数	CUDA C 函数	标准的C函数	CUDA C 函数
malloc	cudaMalloc	memset	cudaMemset
memcpy	cudaMemcpy	free	cudaFree

用于执行 GPU 内存分配的是 cudaMalloc 函数，其函数原型为： 

cudaError_t cudaMalloc (void** devPtr, size_t size)

该函数负责向设备分配一定字节的线性内存，并以 devPtr 的形式返回指向分配内存的指针。

 

cudaMemcpy 函数负责主机和设备之间的数据传输，其函数原型为：

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)

此函数从 src 指向的原存储区复制一定数量的字节到dst的目标存储区。复制方向由 kind 指定，其中 kind 有以下几种。

cudaMemcpyHostToHost                        主机到主机

cudaMemcpyHostToDevice                    主机到设备

cudaMemcpyDeviceToHost                    设备到主机

cudaMemcpyDeviceToDevice                 设别到设备

这个函数以同步方式执行，因为在 cudaMemcpy 函数返回以及传输完成之前主机应用程序是阻塞的。除了内核启动之外的CUDA函数都会返回一个错误的枚举类型 cudaError_t。

如果GPU内存分配成功，函数返回 cudaSuccess

否则返回 cudaErrorMemoryAllocation

可以使用以下CUDA运行时函数将错误代码转化为可读的错误信息：

char* cudaGetErrorString(cudaError_t error)    （这个函数和C语言中的strerror函数类似）

 

GPU内存结构主要包含两部分： 全局内存(global memory) 和 共享内存(shared memory)

每个线程(Thread)都有私有的本地内存(local memory)。每个线程块(block)都具有对该块的所有线程可见的共享内存(shared memory)，并且具有与该块相同的生存期。所有线程都可以访问相同的全局内存(Glabal memory)。

所有线程还可以访问另外两个只读内存空间:常量内存空间(constant memory)和纹理内存空间(texture memory)。全局、常量和纹理内存空间针对不同的内存使用进行了优化(参见设备内存访问)。纹理内存还为一些特定的数据格式提供了不同的寻址模式和数据过滤(参见纹理和表面内存)。

 

2.1.3 线程管理

当核函数在主机端启动时，它的执行会移动到设备上，，此时设备会产生大量的线程并且每个线程都执行由核函数指定的语句。如何组织线程呢？CUDA为我们明确了由 线程块（block） 和 线程块网格（Grid） 构成的线程层次结构。

网格（grid）：由一个内核启动所产生的所有线程统称一个网格，同一个中所有的线程共享相同的全局内存空间。一个网格内由多个线程块 (block) 构成。

线程块（block）：一个线程块包含一组线程，同一个线程块内的线程协助可以通过以下的方式来实现：

1、同步 2、共享内存

不同块内的线程不能协作。

线程依靠以下两个坐标向量来区分彼此：

blockIdx（线程块在线程网格内的索引）

threadIdx（块内的线程索引）

当执行一个核函数时，CUDA运行时为每个线程分配坐标变量 blockIdx 和 threadIdx。该坐标向量是基于 uint3 定义的CUDA内置的向量类型，是一个包含 3 个无符号整数的结构，可以通过x 、y、z 三个字段来指定。

blockIdx.x	blockIdx.y	blockIdx.z
threadIdx.x	threadIdx.y	threadIdx.z

 

 

CUDA可以组织三维的网格和块。上图展示了一个线程层次结构的示例。其结构是一个包含了二维块的二维网格。

网格和块的维度由两个内置变量指定：

blockDim（线程块的维度，每个线程块中的线程数）

gridDim（线程网格的维度，每个线程网格中的线程块数）

他们是 dim3 类型变量，基于 uint3 定义的整数型向量，用来表示维度。当定义一个dim3类型的变量时，所有未指定的元素都被初始化为 1。dim3 类型变量中的每个组件可以通过它的 x，y，z字段获得。

blockDim.x

blockDim.y

blockDim.z

 

线程组织小结：在CUDA程序中有两组不同的网格和块变量：手动定义的 dim3 数据类型 和 预定义的 uint3 数据类型。

在主机端，作为内核调用的一部分，你可以使用 dim3 数据类型定义一个网格和块的维度。

当执行核函数时，CUDA运行时会生成相应的内置预初始化的网格、块和线程变量，它们在核函数内均可以被访问到且为uint3类型。

手动定义的 dim3 类型的网格和块变量仅在主机端可见，而 uint3 类型的内置预初始化的网格和块变量仅在设备端可见。

 

简单例子：查看两种变量的作用

#include <stdlib.h>
#include <iostream>
#include <time.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
using namespace std;

__global__ void checkIdex(void)
{
	//输出每个线程的 线程索引 块索引 块维度 网格维度
	printf("threadIdx(%d, %d, %d), blockIdx(%d, %d, %d), threadDim(%d, %d, %d), blockDim(%d, %d, %d)\n", threadIdx.x, threadIdx.y, threadIdx.z, blockIdx.x, blockIdx.y, blockIdx.z, blockDim.x, blockDim.y, blockDim.z, gridDim.x, gridDim.y, gridDim.z);
}

int main()
{
	int nElem = 6;			//数据量大小

	dim3 block(3);     //一维的线程块， 一个线程块包含三个线程
	dim3 grid((nElem + block.x - 1) / block.x);			//一维 grid 中 需要分配多少个block

	printf("grid.x %d, grid.y %d, grid.z %d\n", grid.x, grid.y, grid.z);				//自定义的 dim3 grid
	printf("block.x %d, block.y %d, block.z %d\n", block.x, block.y, block.z);			//自定义的 dim3 block

	checkIdex <<<grid, block>>> ();				//调用核函数

	cudaDeviceReset();

	return 0;

}

对于一个给定的数据大小，确定网格和块尺寸的一般步骤为：

①确定块的大小

②在已知数据大小和块大小的基础上计算网格维度

要确定块的尺寸，通常需要考虑：

①内核的性能特性

②GPU资源的限制

 

 2.1.4 启动一个CUDA核函数

CUDA内核调用是对C语言函数调用语句的延伸，<<< >>> 运算符内是核函数的执行配置

kernel_name <<<grid, block>>> (argument list)

利用执行配置可以指定线程在GPU上调度运行的方式。配置的第一个值 （grid）是网格维度，也就是启动块的数目。第二个值（block）是块维度，也就是每个块中线程的数目。

例如：假设有一个32个元素的数据元素，每8个元素一个块，需要启动4个块

kernel_name <<<4, 8>>> (argument list)

由于数据是在全局内存中是线性存储的，因此可以使用变量 blockIdx.x 和 threadIdx.x 来进行以下操作：

①在网格中表示一个唯一的线程

②建立线程和数据元素之间的映射关系

 

 

2.1.5 编写核函数

核函数是在设备端执行的代码。在核函数中，需要为一个线程规定要进行的计算以及进行的数据访问。当核函数被调用时，许多不同的CUDA线程并行执行同一个计算任务。

例如：用__global__ 声明定义核函数

__global__ void kernel_name(argument list)

函数类型限定符指定一个函数在主机上执行还是在设备上执行，以及可被主机调用还是被设备调用。

函数类型限定符

限定符	执行	调用	备注
__global__	在设备端执行	可从主机端调用 也可以从计算能力为3的设备中调用	必须有一个void的返回类型
__device__	在设备端执行	仅能从设备端调用
__host__	在主机端执行	仅能从主机端调用	可以省略

 

__device__ 和 __host__限定符可以一齐使用，这样函数可以同时在主机和设备端进行。

CUDA核函数的限制：

①只能访问设备内存

②必须具有void返回类型

③不支持可变数量的参数

④不支持静态变量

⑤显示异步行为