CUDA C 编程指导（二）：CUDA编程模型详解

最新推荐文章于 2024-08-15 19:57:38 发布

成长Bar

最新推荐文章于 2024-08-15 19:57:38 发布

阅读量3.9k

点赞数

分类专栏： CUDA 文章标签： CUDA 编程模型核函数线程结构异构结构

本文链接：https://blog.csdn.net/FRD2009041510/article/details/43278029

版权

本文深入介绍了CUDA编程模型，包括核函数的概念，线程的层次结构，存储结构的细节，异构编程的工作原理，计算能力的定义，以及通过vectorAdd实例分析。CUDA C通过核函数实现并行计算，线程组织为一维、二维或三维结构，内存层次包括本地、共享、全局、常量和纹理内存，设备和主机之间通过CUDA runtime进行数据交互。

摘要由CSDN通过智能技术生成

CUDA编程模型详解

本文以vectorAdd为例，通过描述C在CUDA中的使用（vectorAdd这个例子可以在CUDA sample中找到。）来介绍CUDA编程模型的主要概念。CUDA C的进一步描述可以参考《Programming Interface》。

主要内容包括：

1、Kernels（核函数）

2、Thread Hierarchy（线程结构）

3、Memory Hierarchy（存储结构）

4、Heterogeneous Programming（异构编程）

5、Compute Capability（计算能力）

6、实例分析

下面分别进行详细地讲解。

1、Kernels（核函数）

CUDA C通过允许编程人员自定义C功能函数（叫做kernels）来扩展C。此功能可以理解为：以并行方式由N个不同的CUDA线程编译N次的程序相当于由常规C编译一次一样。

核函数用_global_来声明定义，对于一个给定的内核调用，执行该内核的CUDA线程数量用一个<<<...>>>执行配置语法（详见C Language Extensions）来指定。每一个执行该内核的线程会分配一个唯一的线程ID，通过内置的threadIdx变量在内核中访问。

下面给出一个例子进行说明：

//功能：把两个大小为N的向量A和向量B相加，并把结果存储在向量C中。
// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
	int i = threadIdx.x;
	C[i] = A[i] + B[i];
}
int main()
{
	...
	// Kernel invocation with N threads
	VecAdd<<<1, N>>>(A, B, C);
	...
}

此处，N个线程中的每一个线程都会执行VecAdd（），是一个双向加法的例子。

2、Thread Hierarchy（线程结构）

为了方便，threadIdx是一个“三组成”向量，因此，线程能够被一维、二维、三维的线程索引识别，组成一维、二维、三维的线程块。这提供了一个很自然的方式，即在一个向量、矩阵或册领域调用整个元素进行计算。

线程的索引及其ID以一种简单的方式互相关：对于一维块，它们等价；对于大小为（Dx,Dy）的二维块，一个线程的索引（x,y）的线程ID是（x+yDx）；对于大小为（Dx,Dy,Dz）的三维线程，一个线程的索引（x,y,z）的线程ID是（x+yDx+zDxDy）。

下面给出一个例子：

//功能：将大小为NxN的两个矩阵A和B相加，把结果存储在矩阵C中。
// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
	int i = threadIdx.x;
	int j = threadIdx.y;
	C[i][j] = A[i][j] + B[i][j];
}
int main()
{
	...
	// Kernel invocation with one block of N * N * 1 threads
	int numBlocks = 1;
	dim3 threadsPerBlock(N, N);
	MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
	...
}

对于每个块，此处对线程的数量有一个限制，因为一个块的所有线程都依赖于同一个处理器的核，它们必须共享此核有限的内存资源。在目前的GPUs，一个线程块最大可包含1024个线程。

然而，一个核函数可以被多个等型的线程块执行，因此线程的总数等于每一个块线程数乘以块数量。

块可以被组织为一维、二维或者三维的线程块网格，如下图所示，在一个网格内线程块的数目经常由将要处理的数据大小或者系统中的处理器核数决定，当然也可以大大超过此数目。