大模型实现架构 transform refusion dacum模型

Chapter 02 CUDA Progeamming Model

2.0 WHAT’S IN THIS CHAPTER?

• Writing a CUDA program
• Executing a kernel function
• Organizing threads with grids and blocks
• Measuring GPU performance

你可以在很多平系统上写CUDA程序，包括嵌入式设备、平板电脑、笔记本电脑、台式机、工作站以及HPC clustered 系统；
C语言就可以写CUDA程序；
本章将以向量相加和矩阵相加为例子介绍如何写CUDA程序。

2.1 INTRODUCING THE CUDA PROGRAMMING MODEL

编程模型呈现的是计算机架构的抽象，它是程序在硬件上实现的桥梁。
所以CUDA C是编译型语言，不是解释型语言。

通信抽象(Communication Abstraction)是program和programming model implementation之间的边界。

• program：program是写给programming model的，它指示program的各个组件如何共享信息并协调其活动；
• programming model：它是特定计算机架构逻辑层面的体现。它通过汇编语言(compiler/libarary)实现。 一般以体现在编程语言或编程环境中。
• ps：

• 图中看到，编译器(汇编语言)向下通过操作系统控制硬件，向上实现编程模型(programming model)；
• 编程模型向上是applications，也就是the program。
• 所以类比，C语言可以认为是一种编程模型的体现，它也是一种计算机架构(CPU架构)的抽象。我们用C语言开发的applications就是the program；CUDA C语言也是一种编程模型的体现；
• 因此，所谓编程模型狭义地可以理解为一种编程语言(编译型)，广义地可以理解为我们要用到的语法，内存结构，线程结构等这些我们写程序时我们能控制的部分；我们所控制的是硬件，所以编程模型是计算机架构的抽象； CUDA 编程模型包含了CPU和GPU。

CUDA programming model有以下特征：

• 通过一种层次结构在GPU上组织线程；
• 通过一种层次结构在GPU上访问内存。

后面将会有三章内容讲这两个方面。
作为开发人员，可以从三个层面理解并行计算：

• Domain level

• 在算法设计时，主要关心的就是这个层面，如何分解数据和函数，是其能够更好地并行地解决问题。

• Logic level

• 当进入到开发阶段时，你所关心的变成如何组织并发的threads，在此时，你所思考的是逻辑层面，以保证线程和计算正确地解决问题。

• Hardware level

• 在硬件层面，理解threads如何映射到cores，能帮助提升性能。

2.1.1 CUDA Programming Structure

• Host: CPU和它的memory
• Device： GPU和它的memory
  从CUDA 6开始，NVIDIA提供了Unified Memory，架起了host和device memory之间的桥梁，这让你用单一个指针就可以访问CPU和GPU内存。在chapter04中将详细讲，因为理解如何控制host和device memory很重要，暂时先忽略。
• kernel是CUDA programming model的重要组成部分。
  开发者可以用sequential code写kernel，不用管底层创建和管理上千个threads的细节。
  一个典型的CUDA program包含serial code和parallel code。
  一个典型的CUDA program的process flow服从以下形式：

1. 从CPU复制数据到GPU；
2. 调用kernel对GPU memory中的数据进行运算；
3. 从GPU内存复制数据回CPU内存。
   [外链图片转存失败,源站可能有防盗链机制,建议将图片保存下来直接上传(img-eBg7u4E4-1616139777946)(2021-03-04-10-11-46.png)]

2.1.2 Managing Memory

CUDA runtime 提供了functions以分配和释放device memory，下表是standard C 和CUDA C相对应的functions：

一个简单的device memory分级示意图，包含了两个主要的组成部分：global memory和shared memory.chapter04、05会详细讲。

以下通过一个向量相加的例子，认识如何将数据在host和device之间传输：

• 完整代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>


__global__ void sum_array_device(float* a, float* b, float* c, const int n) {
	int i = threadIdx.x;
	if (i < n) c[i] = a[i] + b[i];
}

void initial_data2(float* ip, int size) {
	// generate different seed for random number
	time_t t;
	srand((unsigned int)time(&t));
	for (int i = 0; i < size; i++) {
		ip[i] = (float)(rand() & 0xFF) / 10.0f;
	}
}

void printf_output2(float* a, float* b, float* c, const int n) {
	for (int idx = 0; idx < n; idx++)
		printf("\n %f + %f = %f", a[idx], b[idx], c[idx]);
}

int main_sum_array_device() {
	int nElem = 1024;
	size_t nBytes = nElem * sizeof(float);
	//分配主机内存
	float* h_A, * h_B, * h_C;
	h_A = (float*)malloc(nBytes);
	h_B = (float*)malloc(nBytes);
	h_C = (float*)malloc(nBytes);

	//在主机上初始化数组h_A和h_B
	initial_data2(h_A, nElem);
	initial_data2(h_B, nElem);

	//在GPU上分配内存
	float* d_A, * d_B, * d_C;
	cudaMalloc((float**)&d_A, nBytes);
	cudaMalloc((float**)&d_B, nBytes);
	cudaMalloc((float**)&d_C, nBytes);

	//将主机上的数组复制到GPU分配好的内存上
	cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
	cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);

	//在CPU上调用kernel,完成数组相加
	dim3 block(nElem);
	dim3 grid(1);
	sum_array_device << <grid, block >> > (d_A, d_B, d_C, nElem);

	//将GPU上的计算结果复制到主机上
	cudaMemcpy(h_C, d_C, nBytes, cudaMemcpyDeviceToHost);

	//打印计算结果
	printf_output2(h_A, h_B, h_C, nElem);

	//释放主机内存
	free(h_A);
	free(h_B);
	free(h_C);

	//释放GPU内存
	cudaFree(d_C);
	cudaFree(d_A);
	cudaFree(d_B);

	return(0);
}

2.1.3 Organizing Threads

主机上调用了kernel函数之后，程序执行会移步到GPU上，生成大量的threads，每个threads执行kernel函数指定它们做的事情，因此知道如何组织threads在CUDA编程中很重要。
CUDA暴露了一个两个级别的thread分级：grids of blocks和blocks of threads,如下图：

由一次kernel函数调用所产生的所有threads的集合称为grid；在一个grid中的所有threads共享一块global memory；一个grid由很多个thread block组成；一个thread block包含了一组可以相互协作的threads，它们通过：

• Block-local synchronization
• Block-local shared memory

进行协作。在不同thread block中间的thread是不能相互协作的。
Threads通过以下坐标变量来区分彼此：

• blockIdx (block index within a grid)
• threadIdx (thread index within a block)

这些变量是内置的，预初始化的变量，可以通过kernel函数直接访问；当执行一个kernel函数的时候，坐标变量blockIdx、threadIdx会被CUDA runtime分配给每一个thread；你可以根据这个坐标将数据分给不同的threads。

坐标变量的类型是uint3，是一个CUDA内置的向量类型，继承了原始的int类型，它们是两个structure(结构体)分别包含了3个uint(无符号整型)类型的变量：

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

CUDA在三维上组织girds和blocks。前面的图就展示了一个包含2维block的2维grid。block和gird的维度通过两个内置的变量指定：

• blockDim (block dimension, measured in threads)
• gridDim (grid dimension, measured in blocks)

这两个变量是dim3类型，一个以uint3类型为基础的整数向量类型，用以指定维度。每一个dim3类型的组件都可以通过x、y、z字段访问，例如：

• blockDim.x
• blockDim.y
• blockDim.z

任何没有指定的字段都初始化为1。

一般情况下，一个grid由2维的blocks组成，block由3维的threads组成，如下图例子：

在一个CUDA项目中，有两组独立的grid和block变量：

• 手动定义的dim3类型数据
• 预定义的uint3类型数据

在主机端，开发者用dim3类型数据定义grid和block的维度，作为驱动kernel的一部分；当开始执行kernel的时候，CUDA runtime会生成对应的在kernel内部可以访问的内置的，预初始化uint3类型的grid、block和thread变量。dim3类型的变量只能在主机端可见，uint3类型的变量只能在device端可见。

比如一个1D的block，包含3个threads，一个1D的grid，包含若干个blocks：

dim3 block(3);
dim3 grid((nElem+block.x-1)/block.x);

CUDA 区别于其他编程模型的一个特征就是，它暴露了一个两个级别的thread分级，因为grid和block的维度会影响性能，暴露这个简单的abstraction提供了开发者一个额外的性能优化途径。

有一些grid和block的维度的限制， 其中一个主要的受限因素是硬件的计算资源，比如rigisters，shared memory等等。

grid和block是thread分级逻辑层面的体现。

2.1.4 Launching a CUDA Kernel

C语言的函数调用：function_name (argument list);

CUDA kernel的调用仅需在括号前加三个三角括号：kernel_name<<<grid, block>>>(argument list);

括号中第一个参数是gird，指定了grid的维度和启动的block的个数；第二个参数是block的维度，指定了每一个block里面包含的thread的个数。

指定了这两个参数就意味着指定了这个kernel总共包含的threads，同时也指定了kernel所使用的thread布局。

同一个block中的threads可以相互协作，不同的block不行。对于同一个问题，你可以用不同的grid和block布局在组织threads。

• 比如一个32个元素的数据，你可以让一个block包含8个threads，这样就需要4个blocks：
  kernel_name<<<4,8>>>(argument list);



• 
  

因为数据是线性地存储在global memory中的，所以可以用内置的blockIdx.x和threadIdx.x来确定grid中的一个thread，从而建立一个theads和data元素之间的一一对应的映射。

• 如果你让一个block包含了32个threads，这时只需要一个block：
  kernel_name<<<1,32>>>(...);
• 如果一个block只有一个thread，就需要32个block：
  kernel_name<<<32,1>>>(...);

一次kernel的调用时和主机的thread是异步的(asynchronous)，kernel调用后，控制立马还给了主机；当然你可以用以下代码强制主机等待kernel执行完：cudaError_t cudaDeviceSynchronize(void);；有一些CUDA runtime的API是默认同步的(synchronization)，比如cudaMeccpy(...)，就默认主机代码必须等待kernel将数据复制到主机才能执行。

2.1.5 Writing Your Kernel

在一个kernel函数中，你定义一个thread的计算以及这个thread对数据的访问；当执行这个kernel函数的时候，许多不同的threads会并行地执行相同的计算；kernel定义会用__global__声明，且必须有一个void的返回类型：
__global__ vioid kernel_name(...);

函数类型限定符：

kernel函数限制：

以一个例子作为说明：

C code on host:

void sumArraysOnHost(float *A, float *B, float *C, const int N) {
 for (int i = 0; i < N; i++)
 C[i] = A[i] + B[i];
}

CUDA C code on host:

__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
 int i = threadIdx.x;
 C[i] = A[i] + B[i];
}

和C语言代码相比，CUDA C的kernel里面，循环不见了，内置的thread坐标变量代替了数组的索引，而且也没有引用N，因为kernel已经隐式地定义了调用N个threads。

2.1.6 Verifying Your Kernel

你需要一个HOST函数来验证kernel函数有没有算对；
或者将kernel函数设置为<<<1,1>>>，这就和sequential code一样。

2.1.7 Handling Errors

2.1.8 Compiling and Executing

2.2 TIMING YOUR KERNEL

2.3 ORGANIZING PARALLEL THREADS

用一个矩阵相加的例子说明。

2.3.1 Indexing Matrices with Blocks and Threads

一般地，一个矩阵是以行为主的方法线性地存储在global memory里面的。比如一个8x6的矩阵：

在一个矩阵相加的kernel里面，通常每个thread会获得一个数据元素进行计算，第一步需要做的就是通过thread index 访问global memory里面的数据；一般有三种索引需要你将它们相互联系：

• Thread and block index
• Coordinate of a given point in the matrix
• Offset in linear global memory

对于给定的一个thread，你可以先通过将Thread and block index和Coordinate of a given point in the matrix一一对应以获得相对于Thread and block index的global memory中的偏移:

ix = threadIdx.x + blockIdx.x * blockDim.x;
iy = threadIdx.y + blockIdx.y * blockDim.y;

然后将Coordinate of a given point in the matrix和global memory的地址对应：

• idx = iy * nx + ix

三者之间的关系：(注意：ix是矩阵的列)

2.3.2 Summing Matrices with a 2D Grid and 2D Blocks

__global__ void sumMatrixOnGPU2D(float *MatA, float *MatB, float *MatC, int nx, int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy * nx + ix;

    if (ix < nx && iy < ny)
        MatC[idx] = MatA[idx] + MatB[idx];
}

2.3.3 Summing Matrices with a 1D Grid and 1D Blocks

__global__ void sumMatrixOnGPU1D(float* MatA, float* MatB, float* MatC,
	int nx, int ny) {
	unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
	if (ix < nx) {
		for (int iy = 0; iy < ny; iy++) {
			int idx = iy * nx + ix;
			MatC[idx] = MatA[idx] + MatB[idx];
		}
	}
}

2.3.4 Summing Matrices with a 2D Grid and 1D Blocks

__global__ void sumMatrixOnGPUMix(float *MatA, float *MatB, float *MatC, int nx, int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = blockIdx.y;
    unsigned int idx = iy * nx + ix;

    if (ix < nx && iy < ny)
        MatC[idx] = MatA[idx] + MatB[idx];
}

2.4 MANAGING DEVICES

pass