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

仅作个人记录

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

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

上次的代码已经比上上次的快了很多，但是还是不够。

显卡上的内存是 DRAM，因此最有效率的存取方式，是以连续的方式存取。而上次的代码其实并不是连续存取的。

看核函数部分：

//__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;
}

加入线程 i 读取内存，等待结果的时候，我们知道GPU会切换到下一个线程 i+1，于是线程 i+1 也开始进行读取数据运算，接着又是 i+2，i+3 … 所以，实际上我们的线程并没有连续地存取（只是看起来像连续的，实际是跳跃的）。

所以我们可以改进存取模式，来实现连续存取，修改以下循环就可以：

for (int i = thread_id*size; i < (thread_id+1)*size; ++i) sum += num[i] * num[i] * num[i];//修改前

for (int i = thread_id; i < DATA_SIZE; i += THREAD_NUM)	sum += num[i] * num[i] * num[i];//修改后

总代码：

#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; i < DATA_SIZE; i += THREAD_NUM)	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，现在是2430166，提升了1.3倍，并没有原博客的7倍那么多，不过好歹也算是优化。