上一篇我介绍了cuda的基本知识,本篇我将会介绍有关主存和显存的相关概念和二者的联系。
__host__,__device__与__global__修饰函数
cuda中引入了三个宏:__host__、__device__与__global__,用于修饰函数,使得函数被定位到不同的位置。
那修饰后的函数有什么作用呢?
__host__函数,其实就是我们平常写C/C++所定义的运行在CPU中的函数,这个修饰符通常可以不写,效果是等价的。而__device__函数和__global__函数则是必须运行在GPU的函数,因此必须要显式声明在函数前。
我们来看下例:
1 #include <cstring>
2 #include <cstdlib>
3 #include <cassert>
4
5 #include "cuda_runtime.h"
6 #include "device_launch_parameters.h"
7
8 __device__ double triple(double x) {
9 //返回x的三倍
10 return x * 3;
11 }
12
13 __global__ void kern_AddVector(double* c, double const* a, double const* b, size_t n) {
14 //求向量c = a + 3b
15 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x;
16 if(Idx >= n) return; //超过数组大小,直接返回
17
18 c[Idx] = a[Idx] + triple(b[Idx]); //实现向量相加
19 }
20
21 __host__ void addVector(double* c, double const* a, double const* b, size_t n) {
22 //申请显存内地址
23 double *device_c, *device_a, *device_b;
24 assert(cudaSuccess == cudaMalloc(&device_c, sizeof(double) * n));
25 assert(cudaSuccess == cudaMalloc(&device_a, sizeof(double) * n));
26 assert(cudaSuccess == cudaMalloc(&device_b, sizeof(double) * n));
27 //将数据拷贝到显存之中
28 assert(cudaSuccess == cudaMemcpy(device_a, a, sizeof(double) * n, cudaMemcpyHostToDevice));
29 assert(cudaSuccess == cudaMemcpy(device_b, b, sizeof(double) * n, cudaMemcpyHostToDevice));
30 //执行核函数
31 size_t thread_count = 1024;
32 size_t block_count = (n - 1) / thread_count + 1;
33 kern_AddVector<<<block_count, thread_count>>> (device_c, device_a, device_b, n);
34 cudaDeviceSynchronize();
35 cudaError_t ct = cudaGetLastError();
36 assert(cudaSuccess == ct);
37 //将显存中的数据拷贝到主存中
38 assert(cudaSuccess == cudaMemcpy(c, device_c, sizeof(double) * n, cudaMemcpyDeviceToHost));
39 //释放临时变量
40 assert(cudaSuccess == cudaFree(device_a));
41 assert(cudaSuccess == cudaFree(device_b));
42 assert(cudaSuccess == cudaFree(device_c));
43 }
44
45 #include <cstdio>
46
47 int main() {
48 const size_t N = 10;
49 double a[N] = {0.1, 0.2, -0.3, 0.1, 0.5, -0.2, 0.2, -0.3, 0.4, 0.1};
50 double b[N] = {0.2, -0.1, -0.1, 0.2, -0.2, 0.2, 0.1, 0.1, 0.1, 0.3};
51 double c[N];
52 addVector(c, a, b, N);
53 for(double& e: c) {
54 printf("%lf, ", e);
55 }
56
57 return 0;
58 }
例子中addVector函数就是__host__函数,当然main函数也是__host__函数。
__host__函数可以直接调用__host__函数,但不能直接调用__device__函数;__host__函数可以通过传递运行时参数来调用__global__函数,同样也不能像调用__host__函数那样直接调用。而能调用__device__函数的只有__global__函数或者__device__函数。
比如如果main函数这样写,就会报错:
1 int main() {
2 //__host__函数直接调用__device__函数
3 double e3 = triple(e); //error: calling a __device__ function("triple") from a __host__ function("main") is not allowed
4
5 //不传递运行时参数调用__global__函数
6 kern_AddVector(c, a, b, N); //error: a __global__ function call must be configured
7 }
事实上,我们现在大多数的显卡都已经支持了sm_50, compute_50及以上的计算能力(我们可以使用上一篇中提到的deviceQuery来获取自己显卡的计算能力),而这更加丰富了我们的调用关系——50之前cuda没有调用栈,所有__device__函数在编译的时候都是内联的;但50之后,__device__函数可以通过调用__device__函数实现直接或者间接的递归;而__device__和__global__函数也可以继续通过传递运行时参数调用__global__函数,实现二级甚至二级以上的并行。用拓扑结构图来表示则是:
cuda的函数调用关系
除了函数被宏修饰,变量也可以被修饰。
__device__,__shared__与__constant__修饰变量
__device__、__shared__与__constant__也是cuda的宏,用于修饰变量(别忘了__device__也可以修饰函数)。三种变量都不会被声明在CPU中,而是在GPU中。
__device__变量即设备端的全局变量,和C/C++的全局变量声明位置一样,只能在所有类和函数外声明。__host__函数无法直接访问__device__变量,但可以通过cuda运行库中的cudaMemcpyToSymbol()以及cudaMemcpyFromSymbol()函数传递或获取到它的值。__device__函数和__global__函数可以直接访问它们,只需要注意不要线程冲突就好。
__shared__变量即块内共享变量,只能在__device__函数或者__global__函数内被声明。__shared__变量不能跨过一个线程块,所以声明时其所在的__global__函数的运行时变量中的块数往往是1——当然也可以是更大的值,但某一个块中的__shared__变量就无法被其他块所访问到。变量声明时不能初始化,但可以对它进行赋值。
__constant__变量即设备端的常量,并不像它的名字那样一成不变——但至少它在__device__函数和__global__函数中的访问权限是只读的,这样它就可以被放在高速缓存中,极大地提升访问效率。声明方法又和C/C++不同:声明时赋初值是无效的,必须在__host__函数中通过cuda运行库中的cudaMemcpyToSymbol()函数传递给它;当然,__host__函数内部也可以用cudaMemcpyFromSymbol()函数获取到它的值。
于是我们又可以丰富上图:
cuda的函数调用关系,以及设备端变量的访问权限
图片是1920×1080的,所以想拿去做壁纸也是没问题的(狗头)。
而如果变量前面没有修饰,那就是寄存器变量(就像C/C++里的寄存器变量),如果是在__device__函数或者__global__函数内,那么每个线程分别持有一个该变量,不会共享,对其读取和修改也只会发生在该线程内。
不过要注意,__device__和__constant__只能声明在全局变量区域,__shared__变量只能声明在核函数内部,类的成员变量和其他函数内的局部变量是无法被上述关键词修饰的。
下面代码是一个例子:
1 #include <cstdio>
2 #include <cassert>
3
4 #include "cuda_runtime.h"
5 #include "device_launch_parameters.h"
6
7 #define N 10
8 __device__ int arr[N];
9
10 __global__ void print() {
11 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x;
12 if(Idx >= N) return;
13
14 printf("%d\n", arr[Idx]);
15 }
16
17 int main() {
18 int a[N] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
19 assert(cudaSuccess == cudaMemcpyToSymbol(arr, a, sizeof(a)));
20 print<<<10, 1>>> ();
21 cudaDeviceSynchronize();
22 assert(cudaSuccess == cudaGetLastError());
23 }
这里没有使用前文例子中的thread_count和block_count,因为我们明确知道线程数是远小于1024的,甚至核函数内也不需要写大于N则返回的逻辑。
当然,由于是多线程,所以输出是乱序的。但如果交换核数和线程数,因为一个核内线程是轮转调度的,所以输出是顺序的。
cudaMalloc、cudaFree、cudaMemset与cudaMemcpy
cuda_runtime.h库中包含了一些和标准C语言库中的函数非常相近的__host__函数——注意,他们只能在__host__函数中被调用,__global__函数和__device__函数要调用函数原型。
这些函数原型是:
1 cudaError_t cudaMalloc(void **devPtr, size_t size);
2 cudaError_t cudaFree(void *devPtr);
3 cudaError_t cudaMemset(void *devPtr, int value, size_t count);
4 cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);
5 cudaError_t cudaMemcpyToSymbol(const void *symbol, const void *src, size_t count, size_t offset = 0, enum cudaMemcpyKind kind = cudaMemcpyHostToDevice);
6 cudaError_t cudaMemcpyFromSymbol(void *dst, const void *symbol, size_t count, size_t offset = 0, enum cudaMemcpyKind kind = cudaMemcpyDeviceToHost);
前四个函数,我们可以通过名字找到它们在C/C++里的“近亲”:
1 void* malloc(size_t _Size);
2 void free(void* _Block);
3 void* memset(void* _Dst, int _Val, size_t _Size);
4 void* memcpy(void* _Dst, void const* _Src, size_t _Size);
但是这些函数又清一色地返回了cudaError_t这一枚举类型,所以,我们在申请显存空间时,写法应为:
1 double* p; const size_t N = 10;
2 cudaError_t ct = cudaMalloc(&p, sizeof(double) * N);
3 assert(cudaSuccess == ct);
4
5 //对比在内存中的malloc:
6 p = (double*)malloc(sizeof(double) * N);
相信大家也猜到了cudaFree和cudaMemset的用法,事实上我并不喜欢额外创建一个cudaError_t变量,而是直接放在assert中:
1 assert(cudaSuccess == cudaMemset(p, 0, sizeof(double) * N));
2 assert(cudaSuccess == cudaFree(p));
3
4 //对比在内存中的memset、free:
5 memset(p, 0, sizeof(double) * N);
6 free(p);
cudaMemcpy函数不同于memcpy,它有第四个参数,是cudaMemcpyKind枚举类型,其声明如下:
1 enum __device_builtin__ cudaMemcpyKind
2 {
3 cudaMemcpyHostToHost = 0, /**< Host -> Host */
4 cudaMemcpyHostToDevice = 1, /**< Host -> Device */
5 cudaMemcpyDeviceToHost = 2, /**< Device -> Host */
6 cudaMemcpyDeviceToDevice = 3, /**< Device -> Device */
7 cudaMemcpyDefault = 4 /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */
8 };
相信注释也写得非常清楚了:
- cudaMemcpyHostToHost就是从主机端拷贝到主机端,即此时cudaMemcpy等价于memcpy,不属于I/O,耗时最短;
- cudaMemcpyHostToDevice则是从主机端传送到设备端,即源数据在内存中,目标指针指向了一段显存范围,属于I/O,消耗时间较长;
- cudaMemcpyDeviceToHost则是从设备端传送到主机端,即源数据在显存中,目标指针指向了一段内存范围,同样属于I/O,消耗时间较长;
- cudaMemcpyDeviceToDevice则是从设备端拷贝到设备端,CPU只给显卡发送一个信号,不涉及数据交互,因此不属于I/O,不会消耗太多时间。大多数情况下可以异步执行。
下列代码则是一些例子:
1 /*
2 * host_a、host_b是经过malloc或new,或者全局、局部变量的数组,包含有N个int
3 * device_a、device_b是经过cudaMalloc的数组,同样包含有N个int
4 */
5
6 cudaMemcpy(host_a, host_b, sizeof(int) * N, cudaMemcpyHostToHost); //正确
7 cudaMemcpy(device_b, host_b, sizeof(int) * N, cudaMemcpyHostToDevice); //正确
8 cudaMemcpy(host_a, device_a, sizeof(int) * N, cudaMemcpyDeviceToHost); //正确
9 cudaMemcpy(device_b, device_a, sizeof(int) * N, cudaMemcpyDeviceToDevice); //正确
10
11 cudaMemcpy(host_a, device_a, sizeof(int) * N, cudaMemcpyHostToDevice); //错误,函数返回一个cudaErrorInvalidValue
12 cudaMemcpy(device_b, device_a, sizeof(int) * N, cudaMemcpyHostToHost); //错误,函数返回一个cudaErrorInvalidValue
13 cudaMemcpy(device_b, host_b, sizeof(int) * N, cudaMemcpyDeviceToDevice); //错误,函数返回一个cudaErrorInvalidValue
14 cudaMemcpy(host_a, host_b, sizeof(int) * N, cudaMemcpyDeviceToHost); //错误,函数返回一个cudaErrorInvalidValue
而cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()两个函数,前文也提到了,是用来初始化__device__显存全局变量和__constant__显存常量的。虽然函数有五个变量,但后两个变量我们一般只用其初始值,所以写法通常为:
1 __constant__ int arr[N];
2
3 __host__ void init() {
4 int a[N] = {9, 8, 7, 6, 5, 4, 3, 2, 1, 0};
5 int b[N];
6 assert(cudaSuccess == cudaMemcpyToSymbol(arr, a, sizeof(a)));
7 assert(cudaSuccess == cudaMemcpyFromSymbol(b, arr, sizeof(b)));
8 }
可是正如前文所说的,这些函数都是__host__函数,只能在__host__函数中被调用。那么__global__和__device__函数该如何申请、复制、修改和释放显存数据呢?
设备端的malloc、free、memset和memcpy
没错,这就是答案——在__global__和__device__函数中使用函数的原型:
1 template<typename T>
2 __global__ void buildList(T** arrs, size_t size, size_t tot_list) {
3 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x;
4 if(Idx >= tot_list) return;
5
6 arrs[Idx] = (T*)malloc(sizeof(T) * size);
7 memset(arrs[Idx], 0, sizeof(T) * size);
8 }
9
10 template<typename T>
11 __global__ void copyList(T** dsts, const T* const* srcs, size_t size, size_t tot_list) {
12 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x;
13 if(Idx >= tot_list) return;
14
15 memcpy(dsts[Idx], srcs[Idx], sizeof(T) * size);
16 }
17
18 template<typename T>
19 __global__ void clearList(T** arrs, size_t tot_list) {
20 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x;
21 if(Idx >= tot_list) return;
22
23 free(arrs[Idx]);
24 }
这三个函数实现了长度为tot_list的指针数组的每一个元素并行申请大小为size的内存并初始化,深拷贝指针数组,以及并行释放指针数组中的每一个元素的功能。
编程中可能出现的异常
说了半天,也没有提到cudaError_t具体会返回什么异常。
首先我们看一下枚举类型cudaError_t的常见值:
- cudaSuccess = 0。这是几乎所有程序继续运行下去的基础,即未发生任何错误。
- cudaErrorInvalidValue =1。在初学者身上比较常见意为传入API函数的值不在合法区间范围内。通常是一些低级错误,比如在初始化常量时没有使用cudaMemcpyToSymbol而是错误使用了cudaMemcpy、或是在cudaMalloc、cudaMemcpy等函数中传递了空指针等。
- cudaErrorMemoryAllocation = 2。通常是需要申请内存的函数如cudaMalloc会返回这个错误,一般是申请的内存超过了可用显存大小。
- cudaErrorInitializationError = 3。任何runtime库中的函数都有可能返回这个异常,但只有可能在第一次调用时返回。因为cuda的初始化方法是lazy context initialization,即直到调用才会初始化,并不会在程序一开始就初始化。
- cudaErrorCudartUnloading = 4。出现这个异常大多都是误删了cuda驱动。如果出现这个异常,请自行忏悔。
- cudaErrorInvalidConfiguration = 9。通常是传递运行时参数时超过了显卡的负载,如线程数大于deviceQuery输出的每个核的最大线程数、核数大于网格中最大核数等等。
- cudaErrorInvalidPitchValue = 12。一般是在使用cudaMemcpy2D、cudaMemcpy3D等拷贝高维数组时,Pitch出现了问题——可能没有申请Pitch,或者Pitch的地址出错等等。
- cudaErrorInvalidSymbol = 13。即对显存全局变量和常量进行相关操作时,符号名称出错,或进行了多余的格式转换。如你想将数组a拷贝给显存常量arr时,传递的第一个参数可以是单纯的arr,也可以是加引号的"arr",如果写成是转化过的(void*)arr,就会返回这一错误。
- cudaErrorDuplicateVariableName = 43。意为你在定义全局变量时,出现了多个变量重名的情况,可能在同一文件中,也可能在链接前的不同文件中。
- cudaErrorNoDevice = 100。你需要检查你的显卡是否支持cuda。
- cudaErrorFileNotFound = 301。找不到指定文件。
- cudaErrorSymbolNotFound = 500。找不到符号名。通常是在通过字符串寻找设备符号时出现的,此时需要检查你的拼写。
- cudaErrorIllegalAddress = 700。你可能搞错了传入API的指针究竟指向了内存空间还是显存空间,或者在核函数访问时发生了数组越界等等,产生了非法地址。一旦出现了这个问题,程序就必须终止才能继续使用cuda。
- cudaErrorLaunchOutOfResources = 701。你可能使用了过多的线程数或寄存器数,可以deviceQuery一下,然后在项目设置中限制一下寄存器的使用。
- cudaErrorAssert = 710。即在__global__或__device__函数中的断言assert被触发,在触发的同时cuda往往也会将具体行数、核坐标、线程坐标的信息打印出来。一旦出现了这个问题,程序就必须终止才能继续使用cuda。
- cudaErrorHardwareStackError = 714。通常是栈溢出,可能是你在递归__global__或__device__函数的层数太多,或函数内局部变量数组开得太大。一旦出现了这个问题,程序就必须终止才能继续使用cuda。
- cudaErrorLaunchFailure = 719。在执行核函数时发生了内核异常,通常是设备共享内存越界、取消引用无效设备指针等等。一旦出现了这个问题,程序同样必须终止之后才能继续使用cuda。
有关错误代码我将在CUDA教程四中详细介绍。
有关主存、显存、runtime库的基本函数以及cudaError_t的介绍就是这些了。