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服从以下形式：
- 从CPU复制数据到GPU；
- 调用kernel对GPU memory中的数据进行运算；
- 从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.xblockIdx.yblockIdx.zthreadIdx.xthreadIdx.ythreadIdx.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.xblockDim.yblockDim.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
4205
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
