cuda c++的一些记录（6）

仅作个人记录

参考：CUDA编程（七）共享内存与Thread的同步_MingChao_Sun-CSDN博客
顺便说一句，这位博主的cuda系列写的很清晰，关于环境配置，也建议参考这位博主，简单直接就行。

过程中碰到问题，还是建议查看官方API文档
CUDA Runtime API :: CUDA Toolkit Documentation

经过并行 + 连续存取 + 增加block的优化后，速度已经很快了

之前通过block，增大线程数量，虽然进行了提速，但是也产生了一个新的问题：在CPU端的加和压力变得很大，所以能不能在GPU上直接完成这个工作？

每个block内部的Thread之间是可以同步和通讯的，本篇我们将让每个block把它包好的thread的计算结果进行加和。

共享内存与Global内存

GlobalMemory

其实之前复制到显存的数据都是存放在gpu的global memory。CUDA 中，存储复制过来的数据的gpu内存，称为global memory。global memory没有 cache ，所以存取global memory 所需要的时间（ latency）是非常长的，通常是数百个 cycles(时钟周期)，也因此，我们之前才会想方设法隐藏或者减少这个latency。

之前采取了两个主要的措施分别取隐藏和减少latency：

• 一方面通过大量线程并行的方法去不断读取内存（当一个 thread 读取内存，开始等待结果的时候，GPU 就可以切换到下一个 thread，读取下一个内存位置）来尽可能的隐藏latency。
• 另一方面采取连续的内存存取模式，尽量减少latency

ShareMemory

接下来要使用的shared memory，是一个 block 中每个 thread 都共享的内存。它会使用在 GPU 上的内存，所以存取的速度相当快，不需要担心 latency 的问题。
声明一块ShareMemory十分简单，可以直接利用__shared__声明一个shared memory变量，例如：

__shared__ int sharedata[128];

Thread同步

CUDA中，想要完成block中的同步十分简单，就是使用一个CUDA 的内部函数：

__syncthreads()

表示block 中所有的 thread 都要同步到这个点才能继续执行。

总体代码：

#include <iostream>
#include <stdlib.h>
#include <time.h> //用于计时
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

using namespace std;

//32 个 block，每个 block 有 256个 thread，共有 32*256= 8192个thread
#define BLOCK_NUM 32
#define THREAD_NUM 256
#define DATA_SIZE 1048576
int data[DATA_SIZE];

//产生随机数
void generateNum(int *data, int size){
	for (int i = 0; i < size; ++i)	data[i] = rand() % 10;
}

void printDeviceProp(const cudaDeviceProp &prop){
	cout << "Device Name: " << prop.name << endl;
	cout << "totalGlobalMem: " << prop.totalGlobalMem << endl;
	cout << "sharedMemPerBlock: " << prop.sharedMemPerBlock << endl;
	cout << "regsPerBlock: " << prop.regsPerBlock << endl;
	cout << "warpSize: " << prop.warpSize << endl;
	cout << "memPitch: " << prop.memPitch << endl;
	cout << "maxThreadsPerBlock:" << prop.maxThreadsPerBlock << endl;
	cout << "maxThreadsDim[0 - 2]: " << prop.maxThreadsDim[0] << " " << prop.maxThreadsDim[1] << " " << prop.maxThreadsDim[2] << endl;
	cout << "maxGridSize[0 - 2]: " << prop.maxGridSize[0] << " " << prop.maxGridSize[1] << " " << prop.maxGridSize[2] << endl;
	cout << "totalConstMem:" << prop.totalConstMem << endl;
	cout << "major.minor:" << prop.major << " " << prop.minor << endl;
	cout << "clockRate:" << prop.clockRate << endl;
	cout << "textureAlignment:" << prop.textureAlignment << endl;
	cout << "deviceOverlap:" << prop.deviceOverlap << endl;
	cout << "multiProcessorCount:" << prop.multiProcessorCount << endl;
}

//cuda初始化
bool InitCuda(){
	int count;
	cudaGetDeviceCount(&count);//获取能够使用的gpu数量，编号从0开始
	if (count == 0)	return false;//没有支持cuda的gpu
	int device = 0;
	for (; device < count; ++device){
		cudaDeviceProp prop;
		if (cudaGetDeviceProperties(&prop, device) == cudaSuccess){
			printDeviceProp(prop);
			break;//寻找一个可用的gpu
		}
	}
	cudaSetDevice(device);//决定使用编号为device的gpu
	return true;
}

//__global__函数（GPU上执行），计算立方和
__global__ void sum_Squares(int *num, int *result, clock_t *time){
	extern __shared__ int shared[];//声明一块共享内存
	const int thread_id = threadIdx.x;//当前的线程编号（0开始）
	const int block_id = blockIdx.x;//当前的 thread 属于第几个 block（0 开始）
	shared[thread_id] = 0;

	clock_t start;
	if (thread_id == 0)	time[block_id] = clock();//计算时间，只在 threadid ==0 时进行，每个 block 都会记录开始时间及结束时间

	for (int i =block_id*THREAD_NUM + thread_id; i < DATA_SIZE; i += BLOCK_NUM*THREAD_NUM)	shared[thread_id] += num[i] * num[i] * num[i];

	__syncthreads();//同步 保证每个 thread 都已经把结果写到 shared[tid] 里面

	//使用线程0完成加和
	if (thread_id == 0)
	{
		for (int i = 1; i < THREAD_NUM; i++) shared[0] += shared[i];//全部累加到shared[0]
		result[block_id] = shared[0];
	}

	if (thread_id == 0) time[block_id+BLOCK_NUM] = clock();
}

int main(){
	if (!InitCuda())	return 0;
	//生成随机数
	generateNum(data, DATA_SIZE);

	int *gpudata, *result;
	clock_t *time;
	//gpu上开内存空间存储数组以及计算结果
	cudaMalloc((void **)&gpudata, sizeof(int)*DATA_SIZE);//第一个参数是指针的指针
	cudaMalloc((void **)&result, sizeof(int)*BLOCK_NUM);//thread,block增多
	cudaMalloc((void **)&time, sizeof(clock_t)*BLOCK_NUM*2);

	//数据从cpu搬运到gpu
	cudaMemcpy(gpudata, data, sizeof(int)*DATA_SIZE, cudaMemcpyHostToDevice);

	//CUDA 中执行函数 语法：函数名称<<<block数目, thread数目, shared memory大小>>>(args...)
	sum_Squares <<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>>(gpudata, result, time);//512个线程进行运算

	//运算结果又从gpu搬运回cpu
	int sum[BLOCK_NUM];//进行修改
	clock_t time_cost[BLOCK_NUM*2];
	cudaMemcpy(&sum, result, sizeof(int)*BLOCK_NUM, cudaMemcpyDeviceToHost);
	cudaMemcpy(&time_cost, time, sizeof(clock_t)*BLOCK_NUM * 2, cudaMemcpyDeviceToHost);

	//释放gpu上面开的内存
	cudaFree(gpudata);
	cudaFree(result);
	cudaFree(time);

	int all_sum = 0;//cpu端进行加和
	for (int i = 0; i < BLOCK_NUM; ++i)	all_sum += sum[i];


	//新的计时策略 把每个 block 最早开始时间和最晚结束时间之差，为总运行时间
	clock_t min_start = time_cost[0], max_end = time_cost[BLOCK_NUM];

	for (int i = 1; i < BLOCK_NUM; i++) {
		if (min_start > time_cost[i])
			min_start = time_cost[i];
		if (max_end < time_cost[i + BLOCK_NUM])
			max_end = time_cost[i + BLOCK_NUM];
	}
	cout << "GPU_sum: " << all_sum << " time cost: " << max_end - min_start << endl;

	all_sum = 0;//cpu上面也计算一次进行验证
	for (int i = 0; i < DATA_SIZE; ++i)	all_sum += data[i] * data[i] * data[i];
	cout << "CPU_sum: " << all_sum << endl;

	return 0;
}

运行结果：

运行时间显然有所变长，原因也很明显，因为在 GPU 上多做了一些动作，但是好的结果是这个时候CPU只需要加和32（BLOCK_NUM）个数字，同时还减少了需要拷贝到内存的数据量。

效率有所降低,一个原因是最后加和的操作只由每个block的Thread0完成，显然不是最有效率的方法。下篇博客中会通过并行加法的算法来解决这个问题。