CUDA编程初试
第一个调用GPU计算的程序
在前面，我们已经写了一个串行的dist程序计算，现在我们尝试对其进行改进，从而得到一个使用GPU进行计算的程序。
为了调用GPU，所以我们需要写一个核函数，从而在主机端调用GPU。由于核函数是用于不断产生出线程，而线程才是真正用于计算的东西，所以我们在核函数内需要为每个线程分配计算任务，这个就是等价之前for循环内部的东西。
值得注意的是，每个被核函数调用出来的线程执行的计算都是相同的，所以在核函数内需要为每个线程分配数据，这样就可以执行计算。
与之类比，在C语言中fork出的新线程可以执行与原有线程不一样的指令。
每个线程分配计算数据
那如何为每个线程分配数据呢？
在核函数内，有以下的变量可以用于给不同的线程分配数据，这些变量是提供线程块核线程的维度数核索引变量
前两个的类型为dim3，这是一个1*3的向量，可以使用.x.y.z来访问。
后面两个类型为uint3，这也是一个1*3向量，遇上一个类似。
核函数调用时的<<<A,B>>>其实是<<<dim3(A,1,1),dim3(B,1,1)>>>，即目前声明的都是1维数组，当需要计算二维、三维的时候，就可以声明高维数组。
每个线程块的线程数不超过1024
主机端与GPU的数据通信
由于GPU不能直接访问主机端中的数据，所以，CUDA提供了一系列可以把数据从设备传回主机的API
位于头文件cuda_runtime.h
cudaMalloc()：分配设备端内存
cudaMemcpy()：将数据传入或传出设备
cudaFree()：释放设备内存
size_：内存大小的专用变量
cudaError_t：错误处理专用变量
编写程序
有了上面的内容，我们很容易就将之前的串行程序改成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
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
|
#include "device_launch_parameters.h"
#include "cuda_runtime.h"
#include <stdio.h>
#include <math.h>
#define N 64
#define TPB 32
__device__ float scale(int i, int n)
{
return ((float)i / (n - 1));
}
__device__ float distance(float x1, float x2)
{
return sqrt((x1 - x2)*(x1 - x2));
}
__global__ void distanceKernel(float *d_out, float ref, int len)
{
const int i = blockIdx.x*blockDim.x + threadIdx.x;
const float x = scale(i, len);
d_out[i] = distance(x, ref);
printf("i = %2d: dist from %f to %f is %f\n", i, ref, x, d_out[i]);
}
int main()
{
const float ref = 0.5f;
float *d_out = 0;
cudaMalloc(&d_out, N * sizeof(float));
distanceKernel << <N / TPB, TPB >> > (d_out, ref, N);
cudaFree(d_out);
return 0;
}
|
执行nvcc distcuda.cu来编译，运行后可以看到，程序的执行顺序是乱的，但是其结果都被打印出来了。
另一个更深入的CUDA程序
我们之前写过两个版本的dist，现在我们把第二个版本的dist并行化。
这是一个更加通用的结构
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
|
#if DIST == 2
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdlib.h>
#include <stdio.h>
#include <math.h>
#define N 64
#define TPB 32
float scale(int i, int n)
{
return ((float)i / (n - 1));
}
__device__ float distance(float x1, float x2)
{
return sqrt((x2 - x1)*(x2 - x1));
}
__global__ void distanceKernel(float *d_out, float * d_in, float ref)
{
const int i = blockIdx.x*blockDim.x + threadIdx.x;
const float x = d_in[i];
d_out[i] = distance(x, ref);
printf("i = %2d: dist from %f to %f is %f\n", i, ref, x, d_out[i]);
}
void distanceArray(float * out, float *in, float ref, int len)
{
float *d_out = 0;
float *d_in = 0;
cudaMalloc(&d_out, len * sizeof(float));
cudaMalloc(&d_in, len * sizeof(float));
cudaMemcpy(d_in, in, len * sizeof(float), cudaMemcpyHostToDevice);
distanceKernel << <N / TPB, TPB >> > (d_out, d_in, ref);
cudaMemcpy(out, d_out, len * sizeof(float), cudaMemcpyDeviceToHost);
cudaFree(d_in);
cudaFree(d_out);
}
int main()
{
const float ref = 0.5f;
float * in = (float *)calloc(N, sizeof(float));
float * out = (float *)calloc(N, sizeof(float));
for (int i = 0; i < N; i++)
{
in[i] = scale(i, N);
}
distanceArray(out, in, ref, N);
free(in);
free(out);
return 0;
}
|
程序执行过程
- 在设备上创建数组，这个数组大小和主机上的数组大小一致。
- 在内存上分配内存来存储输入的数据。
- 将主机端的数组复制到设备中(
cudaMemcpy函数)
- 启动核函数执行计算并将输出值存储在设备内存上的输出数组中
- 将设备上的结果赋值到主机端
- 释放内存
注意到，在GPU上的计算比在主机端的计算不一样的地方：
- 在设备上创建一个镜像数组，然后需要把数据拷贝过去
- 启动一个核函数并执行大量计算
- 把结果拷贝回主机
很明显，只有当核函数并行计算的收益超过拷贝数据时，使用GPU计算才是合理的。
统一内存与托管数组
在两个设备之间传递已有的很大数组这个步骤是必要且有点冗余的，所以出现了统一内存来简化开发流程。
统一内存就是一个主机和设备可以共同访问的内存，但实际上只是CUDA帮助我们完成了在主机与设备之间拷贝数组的活，这项工作本身没有消失，只是不需要显示地做了。另外，CUDA系统调度内存可能不如手工管理内存那样好。
其实现方式是使用cudaMallocManaged()即可。
我们对上一个程序使用统一内存来进行简化。
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
|
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdlib.h>
#include <stdio.h>
#include <math.h>
#define N 64
#define TPB 32
float scale(int i, int n)
{
return ((float)i / (n - 1));
}
__device__ float distance(float x1, float x2)
{
return sqrt((x2 - x1)*(x2 - x1));
}
__global__ void distanceKernel(float *d_out, float * d_in, float ref)
{
const int i = blockIdx.x*blockDim.x + threadIdx.x;
const float x = d_in[i];
d_out[i] = distance(x, ref);
printf("i = %2d: dist from %f to %f is %f\n", i, ref, x, d_out[i]);
}
void distanceArray(float * out, float *in, float ref, int len)
{
distanceKernel << <N / TPB, TPB >> > (out, in, ref);
}
int main()
{
const float ref = 0.5f;
float * in = 0;
float * out = 0;
cudaMallocManaged(&in, N * sizeof(float));
cudaMallocManaged(&out, N * sizeof(float));
for (int i = 0; i < N; i++)
{
in[i] = scale(i, N);
}
distanceArray(out, in, ref, N);
cudaFree(in);
cudaFree(out);
return 0;
}
|
注意到distanceArray函数中不再需要冗长的复制数据的过程了，这样就非常的方便了。
总结
- 核函数只能在设备上执行，不能向主机端返回变量，其可以访问设备内存，但是不能访问主机内存(这也是为什么需要用
cudaMalloc而不是malloc)。
- 可以使用
cudaMallocManaged()来自动传递内存，而避免使用cudaMemcpy()来显示传递数据。
- 一般每个线程块的线程数目是32的整倍数时，可以对应SM中的CUDA核心数目，这样性能更好。
- gridDim、blockDim、blockIdx、threadIdx这些变量可以用于定位线程的处理数据的索引。