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;
}在计算立方和的部分，虽然看起来是连续存取内存位置（每个 thread 对一块连续的数字计算平方和），但是实际上并不是这样的，我们要考虑到实际上 thread 的执行方式。
前面提过，当一个 thread 在等待内存的数据时，GPU 会切换到下一个 thread。也就是说，实际上线程执行的顺序是类似
thread 0 -> thread 1 -> thread 2 -> thread 3 -> thread 4 ->...因此，在同一个 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;
}通过上面对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 1024
int 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;
}运行结果 :
我们看到这次运行用了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》
