CUDA编程（五）关注内存的存取模式

CUDA编程（五）

关注内存的存取模式

上一篇博客我们使用Thread完成了简单的并行加速，虽然我们的程序运行速度有了50甚至上百倍的提升，但是根据内存带宽来评估的话我们的程序还远远不够， 
除了通过Block继续提高线程数量来优化性能，这次想给大家先介绍一个访存方面非常重要的优化，同样可以大幅提高程序的性能~

什么样的存取模式是高效的？

大家知道一般显卡上的内存是 DRAM，因此最有效率的存取方式，是以连续的方式存取，单纯说连续存取可能比较抽象，我们还是通过例子来看这个问题。

之前的程序，大家可以看到我们非常重要的核函数部分：

// __global__ 函数 (GPU上执行) 计算立方和__global__ static void sumOfSquares(int *num, int* result, clock_t* time){    //表示目前的 thread 是第几个 thread（由 0 开始计算）    const int tid = threadIdx.x;    //计算每个线程需要完成的量    const int size = DATA_SIZE / THREAD_NUM;    int sum = 0;    int i;    //记录运算开始的时间    clock_t start;    //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录    if (tid == 0) start = clock();    for (i = tid * size; i < (tid + 1) * size; i++) {        sum += num[i] * num[i] * num[i];    }    result[tid] = sum;    //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行    if (tid == 0) *time = clock() - start;}
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
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

在计算立方和的部分，虽然看起来是连续存取内存位置（每个 thread 对一块连续的数字计算平方和），但是实际上并不是这样的，我们要考虑到实际上 thread 的执行方式。

前面提过，当一个 thread 在等待内存的数据时，GPU 会切换到下一个 thread。也就是说，实际上线程执行的顺序是类似

thread 0 -> thread 1 -> thread 2 -> thread 3 -> thread 4 ->...
1
1

因此，在同一个 thread 中连续存取内存，在实际执行时反而不是连续了，下图很明显的反应了这个问题，我们的存取是跳跃式的。

要让实际执行结果是连续的存取，我们应该要让 thread 0 读取第一个数字，thread 1 读取第二个数字…依此类推，很容易可以想象，通过这种存储方式，我们取数字的时候就变成了连续存取。

改进存取模式

根据我们上面的分析，我们原本的核函数并不是连续存取的，读取数字完全是跳跃式的读取，这会非常影响内存的存取效率，因此我们下一步要将取数字的过程变成：

thread 0 读取第一个数字，thread 1 读取第二个数字…

这点通过对核函数的for循环进行一个小修改就可以达到了~

// __global__ 函数 (GPU上执行) 计算立方和__global__ static void sumOfSquares(int *num, int* result, clock_t* time){    //表示目前的 thread 是第几个 thread（由 0 开始计算）    const int tid = threadIdx.x;    int sum = 0;    int i;    //记录运算开始的时间    clock_t start;    //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录    if (tid == 0) start = clock();    //改为连续存取（thread 0 读取第一个数字，thread 1 读取第二个数字 …）    for (i = tid; i < DATA_SIZE; i += THREAD_NUM) {        sum += num[i] * num[i] * num[i];    }    result[tid] = sum;    //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行    if (tid == 0) *time = clock() - start;}
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
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

通过上面对for循环的一个小修改就可以达到目的了，那么这么一个微小的修改到底有多大作用呢？

完整代码 :

#include <stdio.h>#include <stdlib.h>#include <time.h>//CUDA RunTime API#include <cuda_runtime.h>//1M#define DATA_SIZE 1048576#define THREAD_NUM 1024int data[DATA_SIZE];//产生大量0-9之间的随机数void GenerateNumbers(int *number, int size){    for (int i = 0; i < size; i++) {        number[i] = rand() % 10;    }}//打印设备信息void printDeviceProp(const cudaDeviceProp &prop){    printf("Device Name : %s.\n", prop.name);    printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);    printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);    printf("regsPerBlock : %d.\n", prop.regsPerBlock);    printf("warpSize : %d.\n", prop.warpSize);    printf("memPitch : %d.\n", prop.memPitch);    printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);    printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);    printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);    printf("totalConstMem : %d.\n", prop.totalConstMem);    printf("major.minor : %d.%d.\n", prop.major, prop.minor);    printf("clockRate : %d.\n", prop.clockRate);    printf("textureAlignment : %d.\n", prop.textureAlignment);    printf("deviceOverlap : %d.\n", prop.deviceOverlap);    printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);}//CUDA 初始化bool InitCUDA(){    int count;    //取得支持Cuda的装置的数目    cudaGetDeviceCount(&count);    if (count == 0) {        fprintf(stderr, "There is no device.\n");        return false;    }    int i;    for (i = 0; i < count; i++) {        cudaDeviceProp prop;        cudaGetDeviceProperties(&prop, i);        //打印设备信息        printDeviceProp(prop);        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {            if (prop.major >= 1) {                break;            }        }    }    if (i == count) {        fprintf(stderr, "There is no device supporting CUDA 1.x.\n");        return false;    }    cudaSetDevice(i);    return true;}// __global__ 函数 (GPU上执行) 计算立方和__global__ static void sumOfSquares(int *num, int* result, clock_t* time){    //表示目前的 thread 是第几个 thread（由 0 开始计算）    const int tid = threadIdx.x;    int sum = 0;    int i;    //记录运算开始的时间    clock_t start;    //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录    if (tid == 0) start = clock();    for (i = tid; i < DATA_SIZE; i += THREAD_NUM) {        sum += num[i] * num[i] * num[i];    }    result[tid] = sum;    //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行    if (tid == 0) *time = clock() - start;}int main(){    //CUDA 初始化    if (!InitCUDA()) {        return 0;    }    //生成随机数    GenerateNumbers(data, DATA_SIZE);    /*把数据复制到显卡内存中*/    int* gpudata, *result;    clock_t* time;    //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )    cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);    cudaMalloc((void**)&result, sizeof(int)*THREAD_NUM);    cudaMalloc((void**)&time, sizeof(clock_t));    //cudaMemcpy 将产生的随机数复制到显卡内存中    //cudaMemcpyHostToDevice - 从内存复制到显卡内存    //cudaMemcpyDeviceToHost - 从显卡内存复制到内存    cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);    // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);    sumOfSquares << < 1, THREAD_NUM, 0 >> >(gpudata, result, time);    /*把结果从显示芯片复制回主内存*/    int sum[THREAD_NUM];    clock_t time_use;    //cudaMemcpy 将结果从显存中复制回内存    cudaMemcpy(&sum, result, sizeof(int)* THREAD_NUM, cudaMemcpyDeviceToHost);    cudaMemcpy(&time_use, time, sizeof(clock_t), cudaMemcpyDeviceToHost);    //Free    cudaFree(gpudata);    cudaFree(result);    cudaFree(time);    int final_sum = 0;    for (int i = 0; i < THREAD_NUM; i++) {        final_sum += sum[i];    }    printf("GPUsum: %d  gputime: %d\n", final_sum, time_use);    final_sum = 0;    for (int i = 0; i < DATA_SIZE; i++) {        final_sum += data[i] * data[i] * data[i];    }    printf("CPUsum: %d \n", final_sum);    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
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
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
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
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183

运行结果 :

我们看到这次运行用了894297个时钟周期

不知道大家是否还记得上次我们用1024个线程运行的最终结果：6489302个时钟周期，现在我们只是很简单的改了一下存取模式，同样使用1024个线程最终只使用了894297个时钟周期

6489302/894297= 7.26

可以看到我们的速度居然提升了7.26倍，而我们只是单纯修改了一下存取模式罢了，所以我们可以看到连续存取这个存取优化还是十分重要的，在我们没法再单纯地从线程数量上继续优化的情况下，从存取模式上进行的这个优化是十分有效的。

我们还是从内存带宽的角度来进行一下评估：

首先计算一下使用的时间：

894297/ (797000 * 1000) =  0.0011221S

然后计算使用的带宽：

数据量仍然没有变 DATA_SIZE 1048576，也就是1024*1024 也就是 1M

1M 个 32 bits 数字的数据量是 4MB。

因此，这个程序实际上使用的内存带宽约为：

4MB / 0.0011221S = 3564.745MB/s = 3.48GB/s

注意我们没进行内存存取优化之前的内存带宽是491MB/s，可以看到，我们通过这个优化一下子就把内存带宽提升到了GB级别，不得不说这是一个非常令人满意的效果，我们在没有继续增加线程数量的情况下，通过把内存的存取模式变成连续的，取得了7倍左右的加速。

总结：

这篇博客主要讲解了通过如何尽可能的连续操作内存，减少内存存取方面的时间浪费。

通过最终的结果我们可以看到，看似不起眼的一个小改进（尽可能的去连续操作内存），竟然有这近7倍的性能提升，所以希望大家记住这个优化，在优化我们的CUDA程序的时候，一定不要忘记从内存存取角度去进行一些优化，这往往能取得出乎意料的结果。

希望我的博客能帮助到大家~

参考资料：《深入浅出谈CUDA》