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

仅作个人记录

参考：CUDA编程（四）并行化我们的程序_MingChao_Sun-CSDN博客
顺便说一句，这位博主的cuda系列写的很清晰，关于环境配置，也建议参考这位博主，简单直接就行。

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

上一次的代码（单线程）速度很拉闸，原因是（来自：CUDA编程（四）并行化我们的程序_MingChao_Sun-CSDN博客）：

在 CUDA 中，gpu内存用于存储复制过来的数据的部分，称为global memory。global memory是没有 cache 的，而且，存取global memory 所需要的时间（即 latency，延迟）是非常长的，通常是数百个 cycles(时钟周期)。由于我们的程序只有一个 thread，所以每次它读取 global memory 的内容，都要等到读取数据、累加到 sum 之后，才能进行下一步。这就是为什么代码表现会这么的差，所使用的内存带宽这么的小的原因。

由于 global memory 没有 cache，所以避开巨大的 latency 的方法，是利用大量的threads。假设现在有大量的 threads 同时执行，那么当一个 thread 读取内存，开始等待结果的时候，GPU 就可以切换到下一个 thread，并读取下一个内存位置。因此，当thread 的数目够多的时候，就可以完全把 global memory 的巨大 latency 隐藏起来了，而此时就可以有效利用GPU的内存带宽。

所以我们需要并行化我们的程序

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

using namespace std;

#define THREAD_NUM 1024
#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){
	const int thread_id = threadIdx.x;//当前的线程编号（0开始）
	const int size = DATA_SIZE / THREAD_NUM;//分配给每个线程的量
	clock_t start;
	if (thread_id == 0)	start = clock();//计算时间，只在 threadid ==0 时进行
	int sum = 0;
	for (int i = thread_id*size; i < (thread_id+1)*size; ++i)	sum += num[i] * num[i] * num[i];
	result[thread_id] = sum;
	if (thread_id == 0) *time = clock() - start;
}

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)*THREAD_NUM);//线程数增多了，结果应当存储在数组里面
	cudaMalloc((void **)&time, sizeof(clock_t));

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

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

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

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

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

	cout << "GPU_sum: " << all_sum << " time cost: " << time_cost << 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;
}

运行结果：

耗时：
$$3189869/(1620000\ast 1000)=0.001969S$$

内存带宽：

DATA_SIZE 为 1048576 = 1024*1024 也就是 1M，1M 个 int(32bits) 数字的数据量是 1M * (32/8) byte= 4MB，内存带宽约为：

$$4MB/0.001969S=2031.488MB/S=2.03GB/S$$

已经是快了很多了，是之前的600多倍，不过1050ti的内存带宽在 112 GB/S，还是不够。

关于threadIdx，可以看CUDA（10）之深入理解threadIdx_林微的博客-CSDN博客

或者看Programming Guide :: CUDA Toolkit Documentation