Cuda on Arm 冬令营第一天学习内容

January 10, 2022

简述

今天的主要学习内容包括课程说明，Cuda和Gpu简介，简单的Linux操作，简单的Cuda编程。

简单复习几个并行化指标、MakeFile和Cuda编程内容

并行化指标--加速比

加速比的理解

加速比主要是为了对比并行Gpu计算的情况下比串行Cpu计算代码速度有多大的提升，其计算公式在下图中：

按照文字进行理解，假定有80%的内容可以进行并行话，那么

Speedup = 1 / (1 - 0.8) = 5

对于并行的加速比计算公式如下：

加速比的极限

加速比是存在上限的，根据上图我们可以知道，当N越大的时候，P/N就会越小，也就表示分母越小，从而加速比越大。

但是，P/N也受到P的制约，由于S + P = 1，所以S的大小也决定着加速比的上限，按照p=.50、p=.90、p=.99进行加速比计算举例，结果如图：

可以看到加速比是没能随着N的增加而继续增加的。

MakeFile

个人理解

MakeFile是编译C，C++的一个文件，按照MakeFile内的内容对部分的文件进行编译或者文件的生成。

简单指令示例

使用make执行MakeFile文件。如果MakeFile内写有Clean等内容，可以使用make clean执行清理操作

OBJ = main.o add.o sub.o mul.o div.o
CC = gcc

app: $(OBJ)
    $(CC) -o app $(OBJ)

main.o: main.c
    $(CC) -c main.c

add.o: add.c
    $(CC) -c add.c

sub.o:sub.c
    $(CC) -c sub.c

mul.o: mul.c
    $(CC) -c mul.c

div.o: div.c
    $(CC) -c div.c

.PHONY : clean
clean :
    -rm $(OBJ) app

从上面的代码可以看出，临时变量写在最开始，使用的时候使用$()进行包裹指代，而clean应该写在最后，以免影响make指令的执行顺序。

其中的.PHONY : clean中应该关注.PHONY这一项，因为这个关键字主要是为了避免指令和文件名冲突，假如目录下同样存在clean文件，那么make clean在没有.PHONY的情况下会报错（可以自行尝试）。

另外需要注意，换行之后尽量使用TAB进行缩进。

Cuda代码示例

课件中同样提供了Cuda编译使用的一个示例：

TEST_SOURCE = hello_cuda.cu

TARGETBIN := ./hello_cuda

CC = /usr/local/cuda/bin/nvcc

$(TARGETBIN):$(TEST_SOURCE)
    $(CC)  $(TEST_SOURCE) -o $(TARGETBIN)

.PHONY:clean
clean:
    -rm -rf $(TARGETBIN)

可以看到基本的内容和上面是一样的，只是编译的文件从xxx.c变成了xxx.cu，编译指令从gcc变成了nvcc。其他并无不同。

Cuda编程初识

课程内容比较多，就简单挑着自己的理解去写了。

数据传输

在进行代码之前，先对cuda的数据传输进行一定的了解，由于是操作GPU进行计算，所以主要的数据其实都应该导入到GPU的缓存中进行，包括部分指令也会在GPU上操作。

所以可以看到如下图所示：

在运行执行结束之后，运算结果讲重新搬回CPU中，进行下一步的操作。如下图：

所以这几步操作会体现在代码中会有一定的复制、转移的操作，同时内存申请和释放也会比单纯的串行程序使用更多一步GPU的内存申请和释放。

运行位置

在CUDA编程中，程序执行和调用的位置会有两个，一个是CPU（也称Host），另外一个则是GPU（也称Device）。函数模式，则一共有三种，分别是：__device__，__global__，__host__。先看下图：

根据上图我们可以看到__device__很明显应该是在GPU设备中被调用然后执行，而__global__则是从CPU到GPU的过度函数，甚至就是使用GPU的函数，而__host__则是单纯的CPU函数，也就是常规的C函数。

CUDA程序编译

在上面的MakeFile中已经进行了简单的差异描述，但是，为了保证编译后的文件的可用，GPU编译比CPU编译多了一个参数，通常是为了指定架构，例如arch = 60，由于走神，这里没太懂，所以放一下PPT图片：

PPT

NVPROF：程序分析软件

主要是三种模式，每种模式都有不同的显示内容：

nvprof a.exe

nvprof --print-gpu-trace a.exe

nvprof --print-api-trace a.exe

CUDA线程调度

在线程调度这一部分，主要有三个层级Grid，Block，Thread，其中Thread是常见的线程，而Block是为了管理Thread而出现的，在常规的GPU core中，一个Block通常拥有32个Thread。

而Grid则是管理Block的，在线程层次和线程调度入门中，个人认为有两个地方较为重要：一是线程的组织结构，二是线程索引。

组织结构

从Thread开始，首先我们应该注意，Thread会被wrap调度器以32个为一组，组成一个wrap，但是Block才是管理thread的，所以Block_si ze应该是以32作为倍数，比如32或者64，以求最大化使用算力。

而接下来，由Thread组成了Block，在基础的情况中中，<<<BLOCKS,BLOCK_SIZE>>>表示的是申请BLOCKS个Block，每个Block内部有
BLOCK_SIZE个Thread，这时候情况较为简单，Grid的应该是一个1XBLOCKS的长矩阵，每个Block也是一个1xBLOCK_SIZE的长矩阵。

接下来是由dim3构成的高纬度线程模型，举例来说，我们可以由dim3 grid_size(3,3,3)定义出一个3x3x3的三维矩阵，在这个矩阵中，一共有27个Block，
此时，这个Grid的维度，也就是gridDim.x、gridDim.y、gridDim.z三个值，均为3，也就是blockIdx的x、y、z最大值均为2（下标从0开始）。
在这样时候，你可以把整个Grid视为一个常见的3x3的魔方，魔方上面的每一个方块都是一个Block。

而我们再定义一个dim3 block_size(3,3,3)之后，每个Block也变成了一个3x3的小魔方，此时每个Thread是组成这个魔方的一份子。

然后使用<<<grid_size,blocksize>>，我们得到的线程总数应该是729个(是 3x3x3 * 3x3x3的结果)。

线程索引

先看一句求线程索引的代码(注意，此时cuda的调度维度是<<<BLOCKS,BLOCK_SIZE>>>):

    int n = blockDim.x * blockIdx.x + threadIdx.x;

其中，Dim 表示的是维度，例如Block是一个10x10的矩阵，那么BlockDim.x和BlockDim.y都应该是10。

而Idx则表示索引，也就是所在位置。例如Grid组成了一个5x4的矩阵，而某个Block位于（3，2）位置，那么BlockIdx.x是2，BlockIdx.y是3。

所以我们再来看n的求得过程，其中blockDim.x * blockIdx.x作用是求得前面的Block一共占了多少个位置，然后再加上在Block内部ThreadIdx.x
来求得自己所在位置。

接下来重新读Idx的那段话，一个很奇怪的地方在于，X是2，而Y是3，这里应该这样思考，X表示的是横坐标，但是放到矩阵上，表示的其实是第几列，而Y则相反，
表示的是纵坐标，在矩阵上表达的是第几行，所以在求一个线程在一个巨型矩阵中的位置的时候，列坐标应该是用

row = threadIdx.x + BlockIdx.x* BlockDim.x

而行坐标则是应该用

col = threadIdx.y + BlockIdx.y * BLockDim.y

CUDA操作代码

接下来是一个Cuda的向量相加的代码示范：

#include <math.h>
#include <stdio.h>

/***
 * 
 */
void __global__ add(const double *x, const double *y, double *z, int count)
{
    //根据所处的block和block内的线程号来确定自己处理的是数组中的哪一个值
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    if( n < count)
    {
        z[n] = x[n] + y[n];
    }

}

/***
 * 
 */
void check(const double *z, const int N)
{
    bool error = false;
    for (int n = 0; n < N; ++n)
    {
        if (fabs(z[n] - 3) > (1.0e-10))
        {
            error = true;
        }
    }
    printf("%s\n", error ? "Errors" : "Pass");
}

/***
 * 
 */
int main(void)
{
    const int N = 100000000;
    const int M = sizeof(double) * N;
    double *h_x = (double*) malloc(M);
    double *h_y = (double*) malloc(M);
    double *h_z = (double*) malloc(M);

    for (int n = 0; n < N; ++n)
    {
        h_x[n] = 1;
        h_y[n] = 2;
    }
    //申请GPU内存并且进行数据拷贝
    double *d_x, *d_y, *d_z;
    cudaMalloc((void **)&d_x, M);
    cudaMalloc((void **)&d_y, M);
    cudaMalloc((void **)&d_z, M);
    // cudaMemcpuHostToDevice 表示数据从CPU转向GPU
    cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
    // 
    const int block_size = 128;
    const int grid_size = (N + block_size - 1) / block_size;
    // 调用GPU进行并行计算，这里是一个一维的grid和一维的block，分别都只有x
    add<<<grid_size, block_size>>>(d_x, d_y, d_z,N);
    // cudaMemDeviceToHost 表示数据从GPU转向CPU
    cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
    check(h_z, N);

    free(h_x);
    free(h_y);
    free(h_z);
    //进行CUDA内存释放
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    return 0;
}

为了更清楚的对比，在下面放上了单纯CPU的向量相加的写法：

#include <math.h>
#include <stdlib.h>
#include <stdio.h>

void add(const double *x, const double *y, double *z, const int N)
{
    for (int n = 0; n < N; ++n)
    {
        z[n] = x[n] + y[n];
    }
}

void check(const double *z, const int N)
{
    bool has_error = false;
    for (int n = 0; n < N; ++n)
    {
        if (fabs(z[n] - 3) > (1.0e-10))
        {
            has_error = true;
        }
    }
    printf("%s\n", has_error ? "Errors" : "Pass");
}


int main(void)
{
    const int N = 100000000;
    const int M = sizeof(double) * N;
    double *x = (double*) malloc(M);
    double *y = (double*) malloc(M);
    double *z = (double*) malloc(M);

    for (int n = 0; n < N; ++n)
    {
        x[n] = 1;
        y[n] = 2;
    }

    add(x, y, z, N);
    check(z, N);

    free(x);
    free(y);
    free(z);
    return 0;
}

结尾

困了困了，后面的重点写的有点敷衍了。明天有精力再修改一下。