【参加CUDA线上训练营】--CUDA并行计算基础

最新推荐文章于 2023-02-12 21:35:00 发布

yyc强军梦

最新推荐文章于 2023-02-12 21:35:00 发布

阅读量383

点赞数 1

文章标签： linux 运维服务器 Powered by 金山文档

本文链接：https://blog.csdn.net/m0_50768402/article/details/128971548

版权

CUDA调动硬件

CPU遇到的瓶颈：

时钟频率墙：硬件不能越来越小，可以做，但是不能用，因为会产生散热问题

存储器：单位面积的存储东西越来越多，存储器的速度和大小并不会又很快的提升，不能满足处理器的发展。

GPU应运而生：

只能向多核及并行系统发展，顺势而生的GPU---Graphics Processing Unit

流多处理器(stream multiprocessor,sm)

GPU是一个协处理器，他要和CPU进行合作处理事情。CPU是线性执行的，GPU是并行处理的。

并行计算是同时应用多个计算资源解决一个计算问题，涉及多个计算资源或处理器

问题被分解为多个离散的部分，可以同时处理（并行）

每个部分可以由一系列指令完成

最好是计算密集的任务：通信和计算开销比例合适；不要受制于访存带宽；这些都能在基于ARM平台的Jetson NANO上完成！

3.1 异构计算

Host CPU和内存（host memory） ---做逻辑控制

Device GPU和显存（device memory）---做并行化计算

3.2 CUDA安装

适用设备：所有包含NVIDIA GPU的服务器，工作站，个人电脑，嵌入式设备等电子设备

软件安装：

Windows：https://docs.nvidia.com/cuda/cuda-installation-guide-microsoft-windows/index.html只需要安装一个.exe的可执行程序

Linux：https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html按照上面的教程，需要6/7个步骤即可

Jetson：https://developer.nvidia.com/embedded/jetpack直接利用NVIDIA SDK Manager或者SD image进行刷机即可

软件安装：

查看当前设备中GPU状态：服务器、工作站、个人电脑---nvidia-smi；Jetson等设备：Jtop

查看当前设备参数

在CUDA sample中1_Utilities/deviceQuery文件夹下的deviceQuery程序。以Ubuntu为例，deviceQuery程序在:/user/local/cuda/sample/1_Utilities/deviceQuery

3.3 CUDA程序的编写

名词解释：

执行：就是说程序在哪里执行

调用：就是用的时候在哪里调的

__global__可以在GPU上调用，也可以在CPU上调用；__device__可以在GPU上调用；__device__和_host_执行空间说明符一起使用，说明该函数可以被CPU和GPU一起编译。

__global__ 可以在GPU上执行，可以在GPU上编译也可以在CPU上编译

__device__ 只能在GPU上执行

__host__ 只能在CPU上执行

（图片里的意思：设备指的就是GPU，主机指的就是CPU）

//Kernel function
__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);//调用核函数
    //调用核函数需要一个执行设置：<<<1,N>>> 定义调用了多少个线程
}

3.4 CUDA程序编译

NVCC会把CPU代码和GPU代码分别编译，然后可以把这两部分的编译结果合成一个可执行程序。

//nvcc向后编译，现实版本的要比虚拟版本的架构更高一点
//gpu-architecture要比gpu-code低才可以
nvcc x.cu --gpu-architecture=compute_50 --gpu-code=sm_50,sm_52 
nvcc x.cu --gpu-architecture=compute_50
nvcc x.cu\
    --generate-code arch=compute_50,code=sm_50\
    --generate-code arch=compute_50.code-sm_52\
    --generate-code arch=computer_53,code=sm_53
nvcc x.cu\
    --generate-code
arch=computer_50,code=[sm_50,sm_52]\
    --generate-code arch=computer_53,code=sm_53

跟gcc编译差不多，hello_from_gpu.cu编译成hello_from_gpu.o，然后再把hello_from_gpu.o和hello_cuda_main.cu编译成hello_from_gpu，这就是编译结束了。

3.5 利用NVPEOF查看程序执行情况

Kernel Timeline输出的是以gpu kernel为单位的一段时间的运行时间线，我们可以通过它观察GPU在什么时候有闲置或者利用不够充分的行为，更准确地定位优化问题。nvprof是nvidia提供地用于生成gpu timline的工具，其为cuda toolkit的自带工具。

nvprof -o out.nvvp a.exe

模式一：nvprof a.exe

模式二：nvprof --print-gpu-trace a.exe（可以查看更细节的东西）

模式三：nvprof --print-api-trace a.exe （查看从头到尾的运行细节）

3.6 CUDA线程层次

Thread:sequential execution unit

所有线程执行相同的和函数

并行执行

Thread Block:a group of threads（多个线程组合在一起为一个Block）

执行在一个Streaming Multiprocessor（SM）

同一个Block中的线程可以协作

Thread Grid: a collection of thread blocks（多个块组合在一起为一个Grid）

一个Grid当中的Block可以在多个SM中执行

【HelloFromGPU<<<?,?>>>();第一个问号是Grid里有多少个Block，第二个问号是Block里有多少个线程】

如何定位线程的具体位置呢？

答：

-threadIdx.[x y z]是执行当前kernel函数的线程在block中的索引值（threadIdx.x是1，threadIdx.y是0）

-blockIdx.[x y z]是指执行当前kernel函数的线程所在block，在grid中的索引值（blockIdx.x是1，blockIdx.y是1）

-blockDim.[x y z]表示一个block中包含多少个线程（blockDim.x是5，blockDim.y是3）

-gridDim.[x y z]表示一个grid中包含多少个block（gridDim.x是3，gridDim.y是2）

上述图二：add<<<1,4>>>(a,b,c)申请的所有线程数都会执行上面那一个程序。然后a,b,c都是数组，如果是在CPU里，那就是一个循环，第一步：c[0]=a[0]+b[0];第二步：c[1]=a[1]+b[1]；但是GPU里面就是同时申请多个线程，然后同时运算。

Blocks好像也不是必须的

增加了一个层次

3.7 执行流程

加载核函数

将Grid分配到一个Device

根据<<<..>>>内的执行设置的第一个参数，Giga threads engine将block分配到SM中。一个Block内的线程一定会在同一个SM内，一个SM可以有很多Block

根据<<<..>>>内的执行设置的第二个参数，Wrap调度器会调用线程。

Wrap调度器为了提高运行效率，会将每32个线程分为一组，称作一个warp

每个Wrap会被分配到32个core上运行

3.8 CUDA线程索引

Parallelizable problem: c=a+b; a,b,c are vectors of length N

void vecAdd(int n,int*a,int*b,int*c)
{
    for(int i=0;i<n;i++)
    {
        c[i]=a[i]+b[i];
    }
}

//CUDA implementation的main函数：
void main()
{
    int size=N*sizeof(int);    //空间大小
    int *a,*b,*c;
    a=(int*)malloc(size);      //给向量a申请空间
    b=(int*)malloc(size);      
    c=(int*)malloc(size);
    memset(c,0,size);          //初始化内存
    init_rand_f(a,N);          //赋值
    init_rand_f(b,N);
    
    vecAdd(N,a,b,c);
}

int main(void)
{
    size_t size=N*sizeof(int);                        //N是数据个数，乘sizeof(int)可以得到数据大小
    int *h_a,*h_b;int *d_a,*d_b,*d_c;
    h_a=(int*)malloc(size);
    h_b=(int*)malloc(size);
    ...
    cudaMalloc((void**)&d_a,size);
    cudaMalloc((void**)&d_b,size);
    cudaMalloc((void**)&d_c,size);
    //以下，第一个是目的地，第二个是源，第三个是数据大小，第四个是传输方向，从哪里传到哪里
    cudaMemcpy(d_a,h_a,size,cudaMemcpyHostToDevice);//把a里的数据从CPU传到GPU
    cudaMemcpy(d_b,h_b,size,cudaMemcpyHostToDevice);//把b里的数据从CPU传到GPU
    vectorAdd<<<grid,block>>>(d_a,d_b,d_c,N);       //d_a等，是Device Pointers
    cudaMemcpy(h_c,d_c,size,cudaMemcpyDeviceToHost);//执行完核函数之后，把数据从GPU传到CPU
    
    //以下是释放数据存储
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    return 0;
}

【如何设置Gridsize & Blocksize？】
//这些只是推荐申请，没有什么标准答案
block_size=128;
gride_size=(N+block_size-1)/block_size;

那么我们每个Block可以申请多少个线程？

第五行：每个block最大有1024个线程；第六行：不是说三个维度可以同时到最大值，而是说这几个的乘积不能超过1024。

可以同时执行的线程数是一个Warps（一个Warps里有32个线程），一个Block里面可以有多个Warps，

【如果数据过多，线程不够用怎么办？】

//在每个线程里都有一个for循环
__global__add(const double*x,const double *y,double  *z,int n)
{
    int index=blockDim.x*blockIdx.x+threadIdx.x;
    int stride=blockDim.x*gridDim.x;
    for(;index<n;index+=stride)
    {
        z[index]=x[index]+y[index];
    }
}

3.9 实验--CUDA编程模型

3.9.1 编写第一个CUDA程序

关键词："__global__" , <<<...>>> , .cu

在当前的目录下创建一个名为hello_cuda.cu的文件，编写第一个Cuda程序：

当我们编写一个hello_word程序的时候，我们通常会这样写：

#include <stdio.h>

void hello_from_cpu()
{
    printf("Hello World from the CPU!\n");
}

int main(void)
{
    hello_from_cpu();
    return 0;
}

如果我们要把它改成调用GPU的时候，我们需要在void hello_from_cpu()之前加入 __global__标识符，并且在调用这个函数的时候添加<<<...>>>来设定执行设置

在当前的目录下创建一个名为hello_cuda.cu的文件，更改上述程序，将它改为在GPU上执行的程序，

#include <stdio.h>
__global__ void hello_from_gpu()
{
    printf("Hello World from the GPU!\n");
}
int main(void)
{
    hello_from_gpu<<<1,1>>>();
    cudaDeviceSynchronize();
    return 0;
}

3.9.2、编写完成后，我们开始编译并执行程序，在这里我们可以利用nvcc进行编译

/usr/local/cuda/bin/nvcc -arch=compute_72 -code=sm_72 hello_cuda.cu -o hello_cuda -run

3.9.3、这里我们也可以利用编写Makefile的方式来进行编译

# 先写一个Makefile脚本
TEST_SOURCE = hellocuda.cu       # 源文件
TARGETBIN := ./hello_cuda        # 目标文件
CC = /usr/local/cuda/bin/nvcc    # nvcc的位置
$(TARGETBIN):$(TEST_SOURCE)
        # 一步到位 把源文件转化成目标文件-o    $(CC)是编译命令
        $(CC)  $(TEST_SOURCE) -o $(TARGETBIN)
.PHONY:clean
clean:
        -rm -rf $(TARGETBIN)     # 删掉目标文件
        -rm -rf *.o              # 删掉所有的中间文件

然后输入make；然后我们就可以得到一个名为hello_cuda.exe的程序，我们开始执行一下，这是执行命令./hello_cuda

多个文件协同编译，修改Makefile文件：

1、编译hello_from_gpu.cu文件生成hello_from_gpu.o

2、编译hello_cuda02-test.cu和上一步生成的hello_from_gpu.o，生成./hello_cuda_multi_file

TEST_SOURCE = hello_cuda02-test.cuTARGETBIN := ./hello_cuda_multi_fileCC = /usr/local/cuda/bin/nvcc# 目标文件的生成依赖于hello_cuda02-test.cu和hello_from_gpu.o$(TARGETBIN):hello_cuda02-test.cu hello_from_gpu.o $(CC) $(TEST_SOURCE) hello_from_gpu.o -o $(TARGETBIN)# hello_from_gpu.o依赖于hello_from_gpu.cu生成hello_from_gpu.o:hello_from_gpu.cu $(CC) --device-c hello_from_gpu.cu -o hello_from_gpu.o.PHONY:cleanclean: -rm -rf $(TARGETBIN) -rm -rf *.o

然后输入指令make -f Makefile_Multi_file，./hello_cuda_multi_file以及make -f Makefile_Multi_file clean

$ make -f rules.txt
 或者
$ make --file=rules.txt
上面代码指定make命令依据rules.txt文件中的规则，进行构建。

这时已经完成了第一个CUDA程序，接下来修改<<<...>>>里面的信息，查看显示效果，可以把<<<1,1>>>改成<<<2,4>>>查看效果

3.9.4.利用nvprof进行查看程序性能

sudo /usr/local/cuda/bin/nvprof ./hello_cuda

3.10 实验--CUDA编程模型（线程组织）

3.10.1 使用多个线程的和函数

1.当我们在讨论GPU和CUDA时，我们一定会考虑如何调用每一个线程，如何定为每一个线程。其实，在CUDA编程模型中，每一个线程都有一个唯一的标识符或者序号，而我们可以通过threadIdx来得到当前的线程在线程块中的序号,通过blockIdx来得到该线程所在的线程块在grid当中的序号，即：

threadIdx.x 是执行当前kernel函数的线程在block中的x方向的序号 blockIdx.x 是执行当前kernel函数的线程所在block，在grid中的x方向的序号

#include <stdio.h>

__global__ void hello_from_gpu()
{
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;
    printf("Hello World from block %d and thread %d!\n", bid, tid);
}

int main()
{
    hello_from_gpu<<<5, 5>>>();
    cudaDeviceSynchronize();
    return 0;
}

创建好了之后，我们开始编译make，执行./Index_of_thread

3.10.2 使用线程索引

那我们如何能够得到一个线程在所有的线程中的索引值？比如：我们申请了4个线程块，每个线程块有8个线程，那么我们就申请了32个线程，那么我需要找到第3个线程块（编号为2的block）里面的第6个线程（编号为5的thread）在所有线程中的索引值怎么办？

gridDim表示一个grid中包含多少个block，gridDim.x表示横向block数量，gridDim.y表示纵向block数量

blockDim表示一个block中包含多少个线程，blockDim.x表示横向block数量，blockDim.y表示纵向block数量

在上面的那个例子中，gridDim.x=4, blockDim.x=8，那么，我们要找的第22个线程（编号为21）的唯一索引就应该是，index = blockIdx.x * blockDim.x + threadIdx.x

接下来，我们通过完成一个向量加法的实例来实践一下，我们来实现的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;
}

为了完成这个程序，我们先要将数据传输给GPU，并在GPU完成计算的时候，将数据从GPU中传输给CPU内存。这时我们就需要考虑如何申请GPU存储单元，以及内存和显存之间的数据传输。

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

void __global__ add(const double *x, const double *y, double *z, int count)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;    //n是线程索引
        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 = 1000;
    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;
    }

    double *d_x, *d_y, *d_z;                            //设备端
    cudaMalloc((void **)&d_x, M);                        //申请GPU的存储单元
    cudaMalloc((void **)&d_y, M);//返回arraytype类型，不能返回直接分配好的指针，要返回一个指向那个指针的指针，要修改那个参量
    cudaMalloc((void **)&d_z, M);
    cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);    //完成数据从CPU到GPU的传输
    cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
   
    const int block_size = 128;                         //这是推荐设置
    const int grid_size = (N + block_size - 1) / block_size;
    add<<<grid_size, block_size>>>(d_x, d_y, d_z, N);    //调用add函数

    cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);     //从Device（GPU）端传输到Host（CPU）
    check(h_z, N);                                       //检查

    free(h_x);                                           //释放cpu的数据存储
    free(h_y);
    free(h_z);
    cudaFree(d_x);                                       //释放GPU的数据存储
    cudaFree(d_y);
    cudaFree(d_z);
    return 0;
}

然后输入make -f Makefile_vectorAdd，然后再输入./vectorAdd利用nvprof查看程序性能sudo /usr/local/cuda/bin/nvprof --print-api-trace ./vectorAdd