深度学习模型部署(十二)CUDA编程-绪

CUDA 运行时 API 与 CUDA 驱动 API 速度没有差别，实际中使用运行时 API 较多，运行时 API 是在驱动 API 上的一层封装。​

CUDA 是什么？​

CUDA(Compute Unified Device Architecture) 是 nvidia 推出的一个通用并行技术架构，用它来进行 GPU 编程。CUDA 本身并不是一门语言，而是一个 GPU 编程模型，是对 C++,Python 这种常见 CPU 编程语言的一个补充。​

为什么用 GPU 编程要用 CUDA？​

因为 GPU 的控制硬件少，所以编程模型要求非常严格，最早期 GPU 的唯一交互方式是通过 OpenGL 和 DirectX 这些图形 API，​

基本上所有的编程语言都是在 CPU 上运行的，所以催生了 CUDA 编程框架，为了更方便的与 GPU 交互和调用 GPU 的资源.​

CUDA编译运行流程

1、CUDA代码文件后缀.cu,使用nvcc进行编译
nvcc是nvidia基于LLVM开发的专门用于编译cuda代码的编译器，cuda代码有一套完整的工具链，称为cuda-toolkit，包括nvcc编译器，cuda-gdb调试工具等。
2、CUDA中的代码执行设备有两种，一种是device,一种是host,CPU被称为host,普通代码都是在host上执行，GPU被称为device,在device上执行的代码需要添加__global__ 或者__device__前缀
具体CUDA不同前缀意义见
3、CUDA中device使用的是显存，所以在device上执行的函数只能传入device上定义的变量，具体方法为定义好变量，然后使用cudaMalloc函数在显存中给变量分配空间，再使用cudaMemcpy将变量拷贝到显存中（cudaMemcpy是内存拷贝函数，可以根据给的参数将device中的memory拷贝到host中，也可以反过来拷贝），使用完后要使用cudaFree进行释放
4、通过__global__前缀定义的函数执行时需要设置执行的block数量和线程数
5、在host函数中进行完cuda调用后，需要使用cudaDeviceSynchronize()函数，因为启动内核是一个异步操作，只要发布了内核启动命令，不等内核执行完成，控制权就会立刻返回给调用内核的CPU线程。

CUDA kernel和线程管理

在CUDA中，函数称为kernel，每个kernel都有一个前缀，不同的前缀代表了kernel不同的运行要求。

#include<stdio.h>
//在GPU上执行的kernel中不允许使用C++的标准库 iostream
// warning #20096-D: address of a host variable "std::cout" cannot be directly taken in a device function
__global__ void hello_from_gpu()
{
    printf("hello world from gpu\n");
    //compute capacity 2.0以后才支持printf，也就是GeForce 830M以后的GPU
}

int main(void)
{
    hello_from_gpu<<<4,4>>>();
    cudaDeviceSynchronize();

    return 0;
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16

<<<>>>是cuda调用kernel时的语法，<<<numBlocks, threadsPerBlock>>>意思为调用numBlocks个block，每个block中threadsPerBlock个线程。

函数前缀

1. device
   使用 device 限定符声明的函数具有以下特征：

• 在设备上执行；
• 仅可通过设备调用。

2. global
   使用 global 限定符可将函数声明为内核。此类函数：

• 在设备上执行；
• 仅可通过主机调用。

3. host
   使用 host 限定符声明的函数具有以下特征：

• 在主机上执行；
• 仅可通过主机调用。
  仅使用 host 限定符声明函数等同于不使用限定符声明函数，这两种情况下，函数都将仅为主机进行编译。
  函数前缀的一些限制
  device 和 global 函数不支持递归。
  device 和 global 函数的函数体内无法声明静态变量。
  device 和 global 函数不得有数量可变的参数。
  device 函数的地址无法获取，但支持 global 函数的函数指针。
  global 和 host 限定符无法一起使用。
  global 函数的返回类型必须为空。
  对 global 函数的任何调用都必须按规定指定其执行配置。
  global 函数的调用是异步的，也就是说它会在设备执行完成之前返回。
  global 函数参数将同时通过共享存储器传递给设备，且限制为 256 字节。
  线程管理
  CUDA有一套专门用于线程管理的机制，一个kernel调用时的配置<<<numBlocks, threadsPerBlock>>>中的numBlocks和threadPerBlock可以是int变量，也可以是线程管理的结构体，结构体更为常见。

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

__global__ void build_in_variables(void)
{
    // build-in variables
    // blockDim：等同于threadsPerBlock
    // gridDim：等同于numBlocks
    // blockIdx:一个block在grid中的id
    // threadIdx:一个thread在block中的id

    const int blockId = blockIdx.x + blockIdx.y * gridDim.x;
    const int threadId = threadIdx.x + blockDim.x * threadIdx.y;

    printf("blockIdx=(%d,%d) \n",blockIdx.x,blockIdx.y);
    printf("threadIdx=(%d,%d) \n",threadIdx.x,threadIdx.y);
    printf("blockid=:%d,threadId=%d \n",blockId,threadId);

}

int main(void)
{
    printf("*****device message*******\n");

    int dev=0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp,dev);
    printf("Using Device %d:%s\n",dev,deviceProp.name);
    printf("Device %d has compute capability %d.%d.\n",dev,deviceProp.major,deviceProp.minor);
    printf("Device %d has %d multi-processors.\n",dev,deviceProp.multiProcessorCount);
    printf("Device %d has %zu byte total global memory.\n",dev,deviceProp.totalGlobalMem);
    printf("Device %d has %zu byte total constant memory.\n",dev,deviceProp.totalConstMem);
    printf("Device %d has %zu byte shared memory per block.\n",dev,deviceProp.sharedMemPerBlock);
    printf("Device %d has %d total registers per block.\n",dev,deviceProp.regsPerBlock);
    printf("Device %d has %d max threads per block.\n",dev,deviceProp.maxThreadsPerBlock);
    printf("Device %d has %d max threads dimensions.\n",dev,deviceProp.maxThreadsDim[0]);
    printf("Device %d has %u max grid size.\n",dev,deviceProp.maxGridSize[0]);
    printf("Device %d has %d warp size.\n",dev,deviceProp.warpSize);
    printf("Device %d has %d clock rate.\n",dev,deviceProp.clockRate);
    printf("Device %d has %d max threads per multi-processor.\n",dev,deviceProp.maxThreadsPerMultiProcessor);
    

    dim3 numBlocks(2,2);
    // 2*2个block per grid
    // dim3,是一个包含xyz三个无符号整型数的结构体，默认值为1
    //三个维度，x变化最快，然后是y，最后是z
    dim3 threadsPerBlock(2,2);
    // 2*2个thread per block
    build_in_variables<<<numBlocks, threadsPerBlock>>>();
    cudaDeviceReset();
    return 0;
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52

在计算能力9.0以前的架构，thread的Hierarchy是二维的，只有两个层次，一个grid，一个block。在9.0以后的架构，thread的Hierarchy是三维的，新引入了一个可选层次：Cluster集群，每个Cluster中的block可以确保是在同一个GPC（GPU Processing Cluster）GPU集群中运行的。

#include<stdio.h>

__global__ void __cluster_dims__(2,1,1) hello_from_gpu()
{
    printf("Hello World from GPU!\n");
}

int main()
{
    hello_from_gpu<<<1,1>>>();
    cudaDeviceSynchronize();
    return 0;
}
1
2
3
4
5
6
7
8
9
10
11
12
13

内存管理

内存分为device memory和host memory，二者之间通过cudaMemcpy来进行管理。

GPU每个线程有自己单独的寄存器和内存，同一个block中有所有thread都能访问的shared memory，
在有cluster的架构中，同一个cluster中的block的shared memory组成了distributed shared memory，可以相互访问。除此之外还有专门的只读内存，用于存放texture(纹理)

#include <stdio.h>

__global__ void sharedMemoryExample(int* input)
{
    // Define shared memory array
    __shared__ int sharedArray[256];

    // Get the thread index
    int tid = threadIdx.x;

    // Load data from global memory to shared memory
    sharedArray[tid] = input[tid];

    // Synchronize threads to ensure all data is loaded
    __syncthreads();

    // Perform some computation using shared memory data
    sharedArray[tid] = sharedArray[tid] * 2;

    // Synchronize threads again before writing back to global memory
    __syncthreads();

    // Write the result back to global memory
    input[tid] = sharedArray[tid];
}

int main()
{
    // Define input data
    int input[256];

    // Initialize input data
    for (int i = 0; i < 256; i++)
    {
        input[i] = i;
    }

    // Allocate memory on the GPU
    int* d_input;
    cudaMalloc((void**)&d_input, sizeof(int) * 256);

    // Copy input data from host to device
    cudaMemcpy(d_input, input, sizeof(int) * 256, cudaMemcpyHostToDevice);

    // Launch the kernel
    sharedMemoryExample<<<1, 256>>>(d_input);

    // Copy the result back from device to host
    cudaMemcpy(input, d_input, sizeof(int) * 256, cudaMemcpyDeviceToHost);

    // Print the result
    for (int i = 0; i < 256; i++)
    {
        printf("%d ", input[i]);
    }

    // Free memory on the GPU
    cudaFree(d_input);

    return 0;
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61

变量前缀：

1.device
device 限定符声明位于设备上的变量。
在接下来的三节中介绍的其他类型限定符中，最多只能有一种可与 device 限定符一起使用，以更具体地指定变量属于哪个存储器空间。如果未出现其他任何限定符，则变量具有以下特征：

• 位于全局存储器空间中；
• 与应用程序具有相同的生命周期；
  可通过网格内的所有线程访问，也可通过运行时库从主机访问。
  2.constant
  constant 限定符可选择与 device 限定符一起使用，所声明的变量具有以下特征：
• 位于固定存储器空间中；
• 与应用程序具有相同的生命周期；
  可通过网格内的所有线程访问，也可通过运行时库从主机访问。
  3.shared
  shared 限定符可选择与 device 限定符一起使用，所声明的变量具有以下特征：
• 位于线程块的共享存储器空间中；
• 与块具有相同的生命周期；
• 尽可通过块内的所有线程访问。
  只有在 syncthreads()（参见第 4.4.2 节）的执行写入之后，才能保证共享变量对其他线程可见。除非变量被声明为瞬时变量，否则只要之前的语句完成，编译器即可随意优化共享存储器的读写操作。
  限制：
  不允许为在主机上执行的函数内的 struct 和 union 成员、形参和局部变量使用这些限定符。
  shared 和 constant 变量具有隐含的静态存储。
  device、shared 和 constant 变量无法使用 extern 关键字定义为外部变量。
  device 和 constant 变量仅允许在文件作用域内使用。
  不可为设备或从设备指派 constant 变量，仅可通过主机运行时函数从主机指派（参见第 4.5.2.3 节和第 4.5.3.6 节）。
  shared 变量的声明中不可包含初始化。
  访问速度顺序为：register>shared>constant>local>device
  下面是具体的一个应用：
  将共享存储器中的变量声明为外部数组时，例如：

extern __shared__ float shared[];
1

数组的大小将在启动时确定（参见第 4.2.3 节）。所有变量均以这种形式声明，在存储器中的同一地址开始，因此数组中的变量布局必须通过偏移显式管理。
例如:

// 如果一名用户希望在动态分配的共享存储器内获得与以下代码对应的内容：
// short array0[128];
// float array1[64];
// int array2[256];
// 则应通过以下方法声明和初始化数组：
extern __shared__ char array[];
__device__ void func() // device or global function
{
    short* array0 = (short*)array;
     float* array1 = (float*)&array0[128];
     int* array2 = (int*)&array1[64];
}
1
2
3
4
5
6
7
8
9
10
11
12

在设备代码中声明、不带任何限定符的自动变量通常位于寄存器中。但在某些情况下，编译器可能选择将其置于本地存储器中。
只要编译器能够确定在设备上执行的代码中的指针指向的是共享存储器空间还是全局存储器空间，此类指针即受支持，否则将仅限于指向在全局存储器空间中分配或声明的存储器。
通过获取 device、shared 或 constant 变量的地址而获得的地址仅可在设备代码中使用。通过 cudaGetSymbolAddress() 获取的 device 或 constant 变量的地址仅可在主机代码中使用。

对 global 函数进行配置

对 global 函数的任何调用都必须指定该调用的执行配置。
执行配置定义将用于在该设备上执行函数的网格和块的维度，以及相关的流。可通过在函数名称和括号参数列表之间插入 <<<Dg, Db, Ns, s>>> 形式的表达式来指定，其中：
Dg 的类型为 dim3，指定网格的维度和大小，Dg.x * Dg.y 等于所启动的块数量，Dg.z 无用；
Db 的类型为 dim3，指定各块的维度和大小，Db.x * Db.y * Db.z 等于各块的线程数量；
Ns 的类型为 size_t，指定各块为此调用动态分配的共享存储器（除静态分配的存储器之外），这些动态分配的存储器可供声明为外部数组的其他任何变量使用，Ns 是一个可选参数，默认值为 0；
S 的类型为 cudaStream_t，指定相关流；S 是一个可选参数，默认值为 0。
举例来说，

//一个函数的声明如下：
__global__ void Func(float* parameter);
//必须通过如下方法来调用此函数：
Func<<<Dg, Db, Ns>>>(parameter);
1
2
3
4

执行配置的参数将在实际函数参数之前被评估，与函数参数相同，通过共享存储器同时传递给设备。

如果 Dg 或 Db 大于设备允许的最大大小，或 Ns 大于设备上可用的共享存储器最大值，或者小于静态分配、函数参数和执行配置所需的共享存储器数量，则函数将失败。

​