TBD，自用，图床To fix…

一般CUDA文件的后缀是.cu，当然用一般的.cpp也可以，只不过要在其头文件中加入CUDA调用API所在的库；

1
__global__ 核函数

<<<4, 4>>> 使用4 * 4 共16个线程；

cudaDeviceSynchronize()是CPU和GPU的同步代码；

1
2
3
4
5
6
7
8
9
10
11
12
13
14
#include <stdio.h>

__global__ void hello_from_gpu()
{
	printf("Hello World from the GPU\n");
}

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

上述代码用如下的指令编译：

1
2
nvcc test.cu -o test
./test

编写的kernel在一个4 * 4的线程块上运行，所以会打印出16个Hello world from GPU

一般GPU使用率高得时候，一般显存使用率也是很高的；

但是显存存满了，不一定GPU在高速工作；

即线程没有处于运算状态，所以GPU使用率不高；但是线程依旧需要被GPU显存存储；

核函数

GPU看成CPU的外设，需要CPU发送控制指令；

既要编写CPU上的控制代码，又要编写GPU上的调用代码；

  1. 核函数在GPU上进行并行执行：编写并行执行的函数，并行执行是GPU的特点；编写人员不需要自己考虑多线程；

CUDA和C++编写和函数是类似的

  1. 限定词
1
2
3
__global__ 核函数必须要有限定词global，这个global前后必须都是双下划线；
    
核函数返回值必须是void

global和void两个位置是可以互换的

核函数的编写规则：

（1）核函数只能访问GPU的内存；CPU和GPU异构，两者各自有各自的内存，两者的内存访问是通过PCIe进行；如果需要在GPU和CPU之间做内存交互，则需要使用运行时API做内存交互；

编写核函数不能访问主机内存！

（2）核函数不能使用变长参数；

（3）核函数不能使用静态变量；

（4）核函数不能使用函数指针；

（5）核函数具有异步性：和CPU、GPU的异构架构有关，CPU只是启动了GPU，但是CPU主机无法控制GPU的执行；CPU不会等待GPU执行完毕，所以我们需要显式地调用同步函数，同步主机CPU与设备GPU的工作进程；

有些线程的执行也是需要进行同步的；

带核函数的CUDA编写流程：

1
2
3
4
5
6
7
int main(void)
{
    主机代码：一般用来对GPU做配置、数据处理；
    核函数调用：并行调用核函数，对数据进行加速处理；
    主机代码：一般会讲GPU运算处理后的数据回传给主机，还会对GPU、CPU进行内存释放的工作；
    return 0; 
}

注意：核函数不支持c++的iostream类！

进行终端打印，要用printf来进行终端显示；

1
2
3
4
5
6
7
8
9
10
11
12
13
14
#include <stdio.h>

__global__ void hello_from_gpu()
{
    printf("Hello World from the GPU!\n");
}

int main(void)
{
    hello_from_gpu<<<1, 1>>>();
    cudaDeviceSynchronize(); //同步主机(CPU)与设备(GPU)，促使缓冲区刷新，打印hello world到终端
    return 0;
}

注意到，核函数调用是和C++调用有区别的，有个<<<1,1>>> 尖括号内的参数是用来设置执行核函数的线程的；

需要明确执行的线程模型；

尖括号内第一个参数是执行的线程块的数量、第二个参数是线程块内线程的个数；

问题：GPU没有执行完，但是CPU已经执行完了，所以需要同步。这样CPU就不会出现，在GPU执行完之前就顺序执行完控制流了；

主机与设备之间的同步是CPU和GPU的一个特点；

CUDA线程模型

一、线程模型结构

  1. 当一个核函数在主机host（CPU）中启动的时候，他所有线程构成了一个网格grid；

（1）grid 网格

一个网格grid中又包含了若干个线程块block

（2）block 线程块

一个线程块block中又包含了若干个线程；

线程是GPU编程中的最小单位；

  1. 线程分块是GPU逻辑上的划分，物理上线程并不分块；

  2. 核函数配置线程：<<<grid_size, block_size >>>

grid_size用来配置，该核函数所包含的线程块Block的数目；block_size用来配置，一个线程块所包含的线程的数目和结构；

线程数目是允许远远高于计算核心数目的；

  1. 最大允许的线程块大小：1024；最大允许的网格大小：2^31 - 1（针对一维网格的情况，也即最大允许2^31 - 1个线程块）

总的线程数至少等于计算核心数才能充分利用计算资源；实际上，只有大于计算核心数时才能更充分利用好计算资源；

这样能使的GPU内存访问和计算同时进行，以节省计算时间；

使用CUDA的核心：让GPU的计算核心一直运作；

GPU和CPU访问之间的耗时很长，甚至大于GPU的运行时间；

一维线程模型：

  1. 每个线程在核函数中都有唯一的身份标识；

  2. 每一个线程的唯一标识由<<<grid_size, block_size >>>确定；两者保存在内建变量 build-in variable中；目前考虑一维的情况：

（1）gridDim.x：该变量的数值等于执行配置中变量grid_size的值；

（2）blockDim.x：该变量的数值等于执行配置中变量block_size的值；

  1. 线程索引保存成内建变量：build-in variable

（1）blockIdx.x：该变量指定一个线程在一个网格中的线程块索引值，范围为0~gridDim.x - 1

（2）threadIdx.x：该变量指定一个线程在一个线程块中的线程索引值，范围为0~blockDim.x - 1

举例：<<<2, 4>>>

我们的grid中有2个线程块，一个线程块中有4个线程；

线程块取值为0和1；

线程取值为0,1,2,3；

通过threadIdx，可以在核函数中找到其身份标识；

1

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
#include <stdio.h>

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

int main(void)
{
    hello_from_gpu<<<2, 4>>>();
    cudaDeviceSynchronize();
    
    return 0;
}

由上面的例子，我们利用blockIdx.x（线程块编号）；和threadIdx.x（线程块内部该线程的编号）；和我们计算的线程对应唯一的id，然后接着输出；

推广到多维线程：

2

3

在核函数以外，不能用剩余的变量；

gridDim和blockDim没有指定的维度，默认为1；

4

注意：

在三维及更高维的线程块处理的时候，优先计数的线程块是按照，先进行x维度的计数；再进行y维度的计数；最后进行z维度的计数；

（所以，高维计算id的时候是：（以3维线程块(x,y,z)为例）

tid = threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx;

5

7

线程全局索引计算方式

考虑一维网格、一维线程块；二维网格、二维线程块；三维网格、三维线程块三种不同的情况，如何计算线程全局的索引计算方式。

一维网格 一维线程块 计算对应的id:

1
2
3
4
5
int blockid = blockId.x;

int threadid = threadId.x;

int id = blockid * blockDim.x + threadid;

二维网格 二维线程块 计算对应的id:

1
2
3
4
5
int blockid = blockId.x + blockId.y * gridDim.x; //表示该线程块的y坐标乘以网格的宽度，然后再加上该线程块的x坐标

int threadid = threadId.y * blockDim.x + threadId.x; //表示该线程的y坐标乘以线程块的宽度（也即线程数）；然后再加上线程id

int id = blockid * (blockDim.x * blockDim.y) + threadid; //两部分组成，先计算前面的线程块所带来的偏移量，然后再计算在线程块内部的偏移量；

三维网格 三维线程块

1
2
3
4
5
6
7
//这里计算blockid的时候，注意符合C++数组的一般顺序，即坐标为(x,y,z)的block，其id对应为x * dim(y) * dim(z) + y * dim(z) + z
int blockid = blockId.x * gridDim.y * gridDim.z + blockId.y * gridDim.z + blockId.z;

//这里计算threadid的时候，注意和C++数组的一般顺序相反。相当于blockid计算(z,y,x)的顺序，所以结果如下所示：
int threadid = threadId.z * (blockDim.x * blockDim.y) + threadId.y * blockDim.x + threadId.x;

int id = blockid * (blockDim.x * blockDim.y * blockDim.z) + threadid;

这里算的id，就是blockid个线程块（线程块本身本部包含的线程数目为blockDim.x * blockDim.y * blockDim.z；

nvcc编译流程与GPU计算能力

可能GPU架构不同，nvcc编译出来的代码存在移植问题。

nvcc编译流程：

8

在指定计算能力的时候，GPU的真实计算能力一定是要大于虚拟计算能力的；

PTX: Parallel Thread Execution 是CUDA平台为基于GPU的通用计算而定义的虚拟机和指令集，PTX的作用是为GPU的演化提供稳定的指令集架构支持，同时上层C/C++解耦，提供稳定支持；

（也就是为了移植性设计的一个IR）

nvcc编译命令总是使用两个体系结构：一个是虚拟的中间体系结构，另一个是实际的GPU体系结构；
虚拟架构更像是对应用所需的GPU功能的声明；

虚拟架构应该尽可能选择低，以尽量适配更多实际GPU；

真实架构应该尽可能选择高，以充分发挥GPU性能；

PTX文档：https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

不同GPU之间的二进制兼容性不同，不能相互使用；

每一种GPU的指令集、指令编码与其他GPU是不同的；

9

次版本号高的GPU程序可以运行在次版本号较低的GPU程序；

注意，GPU版本和GPU架构是两个东西。如伏特Volta和图灵Turing两代GPU架构，其主版本号都是7。但是后者从X.Y=7.5以后就不认为是Volta架构的了；

GPU的计算能力和性能并不是正比的；

表明其计算能力的衡量单元：

浮点数运算峰值（还包含双精度浮点数、单精度浮点数）

而GPU性能还与显存容量、显存带宽等相关；

CUDA程序兼容性问题

CUDA程序可移植问题

1
2
nvcc helloworld.cu -o helloworld -arch=compute_61

10

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
#include <stdio.h>

__global__ void hello_from_gpu()
{
        int blockid = blockIdx.x;
        int threadid = threadIdx.x;
        int id = blockid * blockDim.x + threadid;
        printf("Hello from GPU! block id: %d, thread id: %d, and the global thread index: %d\n", blockid, threadid, id);

}

int main()
{
        printf("Hello from CPU!\n");
        hello_from_gpu<<<2, 2>>>();
        cudaDeviceSynchronize();
        return 0;
}


指定高于设备实际计算能力的GPU虚拟架构时候，GPU无法被正确调用！

虚拟架构尽量指定低一些，以兼容更多的GPU架构和版本；

PTX指令转化为二进制cubin代码与具体的GPU架构有关；

11

指定真实架构能力如果大于虚拟架构能力，则可能调用不了对应GPU；

指定多个版本的架构，生成指定多个GPU版本的编译文件（跨GPU版本的二进制文件）

12

称为胖二进制文件 fatbinary

（1）执行上述指令必须要求CUDA版本支持7.0计算能力，否则会报错；

（2）过多指定计算能力，会增加编译时间和可执行文件的大小；

13

14

CUDA程序例子

15

运行时API：

  1. 获取GPU设备数量：
1
2
int iDeviceCount = 0;
cudaGetDeviceCount(&iDeviceCount);

会返回计算能力>=2.0的设备数量；

注意，运行时API文档中有host，也有device；则说明该API既可以在hostCPU中调用，又可以在Device设备GPU中调用；

  1. 设置GPU执行时使用的设备
1
2
int iDev = 0;
cudaSetDevice(iDev);

只能在host主机上设置，来说明主机应该选择哪一个编号的GPU来执行设备代码；

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
#include <stdio.h>

int main(void)
{
    int iDeviceCount = 0;
    cudaError_t error = cudaGetDeviceCount(&iDeviceCount);
    if(error != cudaSuccess || iDeviceCount == 0)
    {
        printf("No CUDA compatable GPU found!\n");
        exit(-1);
    }
    else
    {
        printf("The count of GPUs is %d. \n", iDeviceCount);
    }
    
    int iDev = 0;
    error = cudaSetDevice(iDev);
    if (error != cudaSuccess)
    {
        printf("faile to set GPU 0 for computing!");
        exit(-1);
    }
    else
    {
        printf("set GPU 0 for computing.\n");
    }
    return 0;
}

CUDA中的内存管理：

CUDA通过内存分配、数据传输、内存初始化、内存释放来进行内存管理；

内存分配：在GPU中开辟一块地址，分配一段内存；

数据传输：host和device之间传输数据；

内存初始化：内存开辟好了后对齐内容处死花花；

内存释放：类比c的free

例如：标准C语言内存管理函数对应cudaMalloc、cudaMemcpy、cudaMemset、cudaFree

  1. cudaMalloc内存分配函数
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
//一般的主机分配内存方式为：
extern void *malloc(unsigned int num_bytes);
float *fp = (float *)malloc(bytes_to_malloc);

//对于CUDA的设备分配内存
函数定义为：
__host____device__ cudaError_t cudaMalloc(void **devPtr, size_t size);

注，这里因为cudaMalloc的API返回值是cudaError_t，所以这里要传入一个双重指针，来将分配好的空间的地址绑定到我们上层的指针中；
    
    
理解：其实就是一个传址调用，只不过传址的对象是指针了。执行完上述API后，直接将我们分配好的地址设置到指针存储内容中去了。
而朴素的malloc是返回地址，在上层界面赋值给指针；
    
    
用法：
float *fpDevice_A;
cudaMalloc((float **)(&fpDevice_A), numbytes);
  1. cudaMemcpy 数据拷贝函数
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
//主机数据拷贝：void *memcpy(void *dst, const void *src, size_t n)
code: memcpy((void *)d, (void *)s, nBytes);

//设备数据拷贝：cudaMemcpy(Device_A, Host_A, nBytes, cudaMemcpyHostToHost);
API定义为:
__host__ cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind);

数据拷贝的几个kind:
cudaMemcpyHostToHost 	主机->主机
cudaMemcpyHostToDevice  主机->设备
cudaMemcpyDeviceToHost	设备->主机
cudaMemcpyDeviceToDevice 设备->设备
cudaMemcpyDefault		默认
    
默认方式只允许在支持统一虚拟寻址的系统上使用；

该API应该只能在CPU上调用运行？因为无论如何数据拷贝都需要CPU调用PCIe总线来转发或运输；
  1. 内存初始化
1
2
3
4
5
6
7
8
9
//主机内存初始化：memset(host_pA, 0, size_t n)

//设备内存初始化：
API具体定义：
__host__ cudaError_t cudaMemset(void *devPtr, int value, size_t count)
//按照字节对内存进行初始化
devPtr:指向设备内存的指针
value:设置每字节内存内容的变量
count:设置的bytes数
  1. 内存释放
1
2
3
4
5
6
7
//释放主机内存
free(p_hostA);

//释放设备内存：
cudaFree(p_DeviceA)
__host____device__ cudaError_t cudaFree(void *devPtr);
devPtr是设备内存指针;

下面用一个cuda矩阵加法的实例来展示上述运行时API之间的简单交互：

tools文件夹下创建一个common.cuh头文件，其内容如下所示：

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
#pragma once
#include<stdlib.h>
#include <stdio.h>

void setGPU()
{
    // 检验GPU的数量
    int iDeviceCount = 0;
    cudaEror_t error = cudaGetDeviceCount(&iDeviceCount);
    
    if (error != cudaSuccess || iDeviceCount == 0)
    {
        printf("No CUDA campatable GPU found!\n");
        exit(-1);
    }
    else
    {
        printf("The count of GPU is %d\n", iDeviceCount);
    }
    
    int iDev = 0;
    error = cudaSetDevice(iDev);
    if(error != cudaSuccess)
    {
        printf("fail to set GPU 0 for computing!.\n");
        exit(-1);
    }
    else
    {
        printf("Set GPU %d, for Computing!\n", iDev);
    }
}


然后主代码文件为：

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
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
#include "./tools/common.cuh"
#include <stdio.h>

__global__ void addFromGPU(float *A, float *B, float *C, const int N)
{
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;
    const int id = bid * blockDim.x + tid;
    C[id] = A[id] + B[id];
}


void initialData(float *addr, int elemCount)
{
    for (int i = 0 ; i < elemCount ; i++)
    {
        addr[i] = (float)(rand() % 0xff) / 10.f;
    }
    return;
}

int main(void)
{
    //1. 设置GPU设备
    setGPU();
    //2. 分配主机内存和设备内存，并初始化
    int iElemCount = 512;//设置元素数量
    size_t stBytesCount = iElemCount * sizeof(float);//字节数
    
    // (1)分配主机内存，并初始化
    float *fpHost_A, *fpHost_B, *fpHost_C;
    fpHost_A = (float *)malloc(stBytesCount);
    fpHost_B = (float *)malloc(stBytesCount);
    fpHost_C = (float *)malloc(stBytesCount);
    
    if (fpHost_A != NULL && fpHost_B != NULL && fpHost_C != NULL)
    {
        memset(fpHost_A, 0, stBytesCount);
        memset(fpHost_B, 0, stBytesCount);
        memset(fpHost_C, 0, stBytesCount);
    }
    else
    {
        printf("Fail to allocate host memory!\n");
        exit(-1);
    }
    
    //(2)分配设备内存，并初始化
    float *fpDevice_A, *fpDevice_B, *fpDevice_C;
    cudaMalloc((float **)(&fpDevice_A), stBytesCount);
    cudaMalloc((float **)(&fpDevice_B), stBytesCount);
    cudaMalloc((float **)(&fpDevice_C), stBytesCount);
    if(fpDevice_A != NULL && fpDevice_B != NULL && fpDevice_C != NULL)
    {
        cudaMemset(fpDevice_A, 0, stBytesCount); //初始化设备内存为0
        cudaMemset(fpDevice_B, 0, stBytesCount);
        cudaMemset(fpDevice_C, 0, stBytesCount);
    }
    else
    {
        printf("Failed to allocate memory!\n");
        free(fpHost_A);
        free(fpHost_B);
        free(fpHost_C);
        exit(-1);
    }
    
    //3. 初始化主机中数据
    srand(666);//设置随机种子
    initData(fpHost_A, iElemCount);
    initData(fpHost_B, iElemCount);
    //4. 数据从主机复制到设备
    cudaMemcpy(fpDevice_A, fpHost_A, stBytesCount, cudaMemcpyHostToDevice);
    cudaMemcpy(fpDevice_B, fpHost_B, stBytesCount, cudaMemcpyHostToDevice);
    cudaMemcpy(fpDevice_C, fpHost_C, stBytesCount, cudaMemcpyHostToDevice);
    
    //5. 调用核函数在设备中进行计算
    dim3 block(32);			   //设置为一维线程块，线程块内部x维度含有32个线程；
    dim3 grid(iElemCount / 32);//设置的是一维网格，线程块数目显然就是1 * iElemCount矩阵的元素数除以线程块能容纳的线程（元素数目）
    addFromGPU<<<grid, block>>>(fpDevice_A, fpDevice_B, fpDevice_C, iElementCount); //调用核函数
    cudaDeviceSynchronize(); //同步GPU计算结果；
    //6. 将计算得到的数据块从设备传给主机
    cudaMemcpy(fpHost_C, fpDevice_C, stBytesCount, cudaMemcpyDeviceToHost); 
    
    for (int i = 0 ; i < 10 ; i++)
    {
        printf("idx=%2d\tmatrix_A:%.2f\tmatrix_B:%.2f\tresult=%.2f\n", i+1, fpHost_A[i], fpHost_B[i], fpHost_C[i]);
    }
    
    //7.释放主机内存与设备内存
    free(fpHost_A);
    free(fpHost_B);
    free(fpHost_C);
    cudaFree(fpDevice_A);
    cudaFree(fpDevice_B);
    cudaFree(fpDevice_C);
    
    cudaDeviceReset();
  	return 0;
}


上面是一个cuda编写GPU多线程的实例，本质是将一个一维数组的每一维都看作一个线程，对应元素的位置都同步加相应的内容；

cudaMemcpy函数，在进行cudaMemcpyDeviceToHost，即将数据从设备传到主机的过程中，这个函数会自己进行一次同步；

16

自定义设备函数

  1. 设备函数 device function

（1）定义只能执行在GPU设备上的函数称为设备函数；

（2）设备函数只能被核函数或其他设备函数调用；

（3）设备函数用__device__修饰

  1. 核函数 kernel function

（1）用__global__修饰的函数称为核函数，一般由主机调用，在设备中执行；

（2）__global__修饰符既不能和__host__同时使用，也不可与__device__同时使用；

  1. 主机函数 host function

（1）主机端的普通C++函数可用__host__修饰

（2）对于主机端的函数，__host__修饰符可省略

（3）可以用__host__和__device__同时修饰一个函数减少冗余代码，编译器会针对主机和设备分别编译该函数；

注意到，我们上面计算线程块的数目（也即网格包含的线程块）时是直接用512除以32来完成的，如果不能整除，如总线程数为513，那么就需要用额外的线程块来确保能够整除。如修改为grid((iElemCount + block.x - 1) / 32)，那么

此时，注意我们需要修改核函数中的访存范 围。例如对于线程号大于514的，就不在数组范围内了，如果还是在核函数中做add的话就会有越界操作；所以需要仔细思考这里的细节，对于大于513编号的访问线程直接return

CUDA错误检查

运行时API错误代码

17

cudaSuccess = 0，运行时API调用成功；

剩余是一系列错误代码；

编写错误检查函数

1
2
3
4
5
6
7
8
9
cudaGetErrorName
__host____device__ const char *cudaGetErrorName(cudaError_t error)
    
返回一个字符串char *类型指针，返回错误代码名词；
    
__host____device__ const char *cudaGetErrorString(cudaError_t error)
    
返回一个含有错误代码描述信息的字符串；

整体封装一个错误检查函数，如下所示：

1
2
3
4
5
6
7
8
9
10
cudaError_t ErrorCheck(cudaError_t error_code, const char *filename, int lineNumber)
{
    if (error_code != cudaSuccess)
    {
        printf("CUDA error:\r\ncode=%d, name=%s, description=%s\r\nfile=%s, line=%d\r\n", error_code, cudaGetErrorName(error_code), cudaGetErrorString(error_code), filename, lineNumber);
    }
    return error_code;
    
    
}

参数filename和lineNumber可以分别用__FILE__或__LINE__来进行处理；

错误函数返回运行时API调用的错误代码；

一个纠错的例子代码：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
#include <stdio.h>
#include "./tools/common.cuh"

int main(void)
{
    float *fpHost_A;
    fpHost_A = (float *)malloc(4);
    memset(fpHost_A, 0, 4);
    
    float *fpDevice_A;
    cudaError_t error = ErrorCheck(cudaMalloc((float **)(&fpDevice_A), 4), __FILE__, __LINE__);
    cudaMemset(fpDevice_A, 0 ,4);
    
    ErrorCheck(cudaMemcpy(fpDevice_A, fpHost_A, 4, cudaMemcpyDeviceToHost), __FILE__, __LINE__); //这里，dst写的是Device_A，而src是Host，所以其实应该从主机向设备GPU传输数据，方向写错了，这里必定会报错的。
    free(fpHost_A);
    ErrorCheck(cudaFree(fpDevice_A), __FILE__, __LINE__);
 ErrorCheck(cudaDeviceReset(), __FILE__, __LINE__);
    return 0;
}

检查核函数，错误检测函数是不能捕捉调用核函数相关错误的，因为核函数的返回值必定是void；

我们捕捉调用核函数的方式是通过两个CUDA的API：

1
2
3
4
5
ErrorCheck(cudaGetLastError(), __FILE__, __LINE__);//检查同步函数主机最后一条之前的错误！也就是检查核函数有无错误；

ErrorCheck(cudaDeviceSynchronize(), __FILE__, __LINE__);


例子：

1
2
3
4
5
dim3 block(2048); //这里主动注入错误，一个线程块最大1024个线程！
dim3 grid((iElemCount + block.x - 1) / 2048);

addFromGPU<<<grid, block>>>(fpDevice_A, fpDevice_B, fpDevice_C, iElemCount); //调用核函数
ErrorCheck(cudaGetLastError(), __FILE__, __LINE__); //这里会检出上一条核函数执行出错了；

CUDA计时

分析性能等使用

18

CUDA可以使用事件计时来对程序执行性能进行度量，使用CUDA event事件计时方式。CUDA事件可以为主机代码、设备代码计时。

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
cudaEvent_t start, stop;
//创建两个用于计时的变量
ErrorCheck(cudaEventCreate(&start), __FILE__, __LINE__);
ErrorCheck(cudaEventCreate(&stop), __FILE__, __LINE__);
ErrorCheck(cudaEventRecord(start), __FILE__, __LINE__);

cudaEventQuery(start); //此处不可用错误检测函数，因为会返回一个错误，但是并不会造成影响；


/*

这部分是需要进行计时的代码块

*/


ErrorCheck(cudaEventRecord(stop), __FILE__, __LINE__);//停止stop的计时
ErrorCheck(cudaEventSynronize(stop), __FILE__, __LINE__);
float elapsed_time;
ErrorCheck(cudaEventElapsedTime(&elapsed_time, start, stop), __FILE__, __LINE__);
printf("Time = %g ms.\n", elapsed_time);

ErrorCheck(cudaEventDestroy(start), __FILE__, __LINE__);
ErrorCheck(cudaEventDestroy(stop), __FILE__, __LINE__);






计时的一些细节：往往计算效率的时候，需要进行多次采样计算，往往可以抛弃第一次调用核函数的时间。因为第一次调用核函数会耗费不少时间！

nvprof性能解析：

  1. nvprof是一个命令行程序，可以分析nv各个API使用的效率和耗时等
  2. 使用方式：nvprof ./ex3 这里的ex3为可执行程序;

它会显示GPU各个行为（具体的，调用不同API所用的最短时间、最大时间和平均时间等操作；

运行时API查询GPU信息

1
2
3
4
//调用:
cudaDeviceProp prop;
ErrorCheck(cudaGetDeviceProperties(&prop, device_id), __FILE__, __LINE__);//device_id为查询GPU的索引号
//只能在主机中调用

具体的内容查看手册即可；

组织线程模型

处理多维数组，建立多维数组与线程之间组织关系。

  1. 构建二维网格和二维线程块

想要发挥GPU中多线程的优点，就要让每个线程处理不同的计算；

要分配每个线程处理不同的数据，避免多个线程访问或处理同一个数据，也要避免多个线程胡乱访问内存；

多线程不能按照规定的顺序访问内存，就可能造成数据的混乱，程序break；

20

19

所以当前线程

1
2
ix = threadIdx.x + blockIdx.x * blockDim.x
iy = threadIdx.y + blockIdx.y * blockDim.y

线程与二维矩阵的映射关系：

1
idx = iy * nx + ix (nx为二维矩阵x方向上的维度)

用二维网格和二维线程来计算二维矩阵的加法：

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
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
#include <stdio.h>
#include "./tools/common.cuh"

cudaError_t ErrorCheck(cudaError_t error_code, const char *filename, int lineNumber)
{
    if (error_code != cudaSuccess)
    {
        printf("CUDA error:\r\ncode=%d, name=%s, description=%s\r\nfile=%s, line=%d\r\n", error_code, cudaGetErrorName(error_code), cudaGetErrorString(error_code), filename, lineNumber);
    }
    return error_code;
}


__global__ void addMatrix(int *A, int *B, int *C, const int nx, const int ny)
{
    int ix = threadIdx.x + blockIdx.x * blockDim.x;
    int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int id = iy * nx + ix;
    if((ix < nx) && (iy < ny))
    {
        C[id] = A[id] + B[id]; 
    }
}

int main(void)
{
    // 1、设置GPU设备
    setGPU();

    // 2、分配主机内存和设备内存，并初始化
    int nx = 16;
    int ny = 8;
    int nxy = nx * ny;
    size_t stBytesCount = nxy * sizeof(int);
     
     // （1）分配主机内存，并初始化
    int *ipHost_A, *ipHost_B, *ipHost_C;
    ipHost_A = (int *)malloc(stBytesCount);
    ipHost_B = (int *)malloc(stBytesCount);
    ipHost_C = (int *)malloc(stBytesCount);
    if (ipHost_A != NULL && ipHost_B != NULL && ipHost_C != NULL)
    {
        for (int i = 0; i < nxy; i++)
            {
                ipHost_A[i] = i;
                ipHost_B[i] = i + 1;
            }
        memset(ipHost_C, 0, stBytesCount); 
    }
    else
    {
        printf("Fail to allocate host memory!\n");
        exit(-1);
    }
    

    // （2）分配设备内存，并初始化
    int *ipDevice_A, *ipDevice_B, *ipDevice_C;
    ErrorCheck(cudaMalloc((int**)&ipDevice_A, stBytesCount), __FILE__, __LINE__); 
    ErrorCheck(cudaMalloc((int**)&ipDevice_B, stBytesCount), __FILE__, __LINE__); 
    ErrorCheck(cudaMalloc((int**)&ipDevice_C, stBytesCount), __FILE__, __LINE__); 
    if (ipDevice_A != NULL && ipDevice_B != NULL && ipDevice_C != NULL)
    {
        ErrorCheck(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); 
        ErrorCheck(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); 
        ErrorCheck(cudaMemcpy(ipDevice_C, ipHost_C, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); 
    }   
    else
    {
        printf("Fail to allocate memory\n");
        free(ipHost_A);
        free(ipHost_B);
        free(ipHost_C);
        exit(1);
    }

    // calculate on GPU
    dim3 block(4, 4);
    dim3 grid((nx + block.x -1) / block.x, (ny + block.y - 1) / block.y);
    printf("Thread config:grid:<%d, %d>, block:<%d, %d>\n", grid.x, grid.y, block.x, block.y);
    
    addMatrix<<<grid, block>>>(ipDevice_A, ipDevice_B, ipDevice_C, nx, ny);  // 调用内核函数
    
    ErrorCheck(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount, cudaMemcpyDeviceToHost), __FILE__, __LINE__); 
    for (int i = 0; i < 10; i++)
    {
        printf("id=%d, matrix_A=%d, matrix_B=%d, result=%d\n", i + 1,ipHost_A[i], ipHost_B[i], ipHost_C[i]);
    }

    //
    for (int j = 0 ; j < ny ; j++) //i is the column
    {
        int *aj = ipHost_A + j * nx;
        int *bj = ipHost_B + j * nx;
        int *cj = ipHost_C + j * nx;
        for (int i = 0 ; i < nx ; i++)
        {
            int aji = *(aj + i);
            int bji = *(bj + i);
            int cji = *(cj + i);
            printf("a[%d][%d](%d) + b[%d][%d](%d) = c[%d][%d](%d)\n", j,i,aji, j,i, bji, j,i, cji);
        }
    }
    
    free(ipHost_A);
    free(ipHost_B);
    free(ipHost_C);

    ErrorCheck(cudaFree(ipDevice_A), __FILE__, __LINE__); 
    ErrorCheck(cudaFree(ipDevice_B), __FILE__, __LINE__); 
    ErrorCheck(cudaFree(ipDevice_C), __FILE__, __LINE__); 

    ErrorCheck(cudaDeviceReset(), __FILE__, __LINE__); 
    return 0;
}


二维网格一维线程块：

21

22

一维网格一维线程块：

每一个线程块对应的是二维矩阵一列的计算，需要设计循环，依次计算每一列的情况。这种因为效率低，略；

CUDA编写的算子和Pytorch编译

https://zhuanlan.zhihu.com/p/645330027

https://godweiyang.com/2021/03/18/torch-cpp-cuda/#toc-heading-5

https://godweiyang.com/2021/03/21/torch-cpp-cuda-2/

https://godweiyang.com/2021/03/24/torch-cpp-cuda-3/