借一栗子讲解基于C的CUDA并行计算

Love The Way You Lie 2022-11-14 11:49 49阅读 0赞

## 一、C语言接口回顾 ##

代码1

// ConsoleApplication6.cpp : 定义控制台应用程序的入口点。
    //
    
    #include "stdafx.h"
    #include <iostream>
    
    using namespace std;
    
    //自定义数据类型 数据对齐
    typedef struct student
    {
        
    	char label[1]; //1个字节
    	int number;    //4个字节
    	float score;   //4个字节
    }stu;
    
    int _tmain(int argc, _TCHAR* argv[])
    {
        
    	//在cpu-memory模式下面 数据是需要对齐的 因为这样cpu和memory的内存交互速度是最快的 
    	//在硬件设计的时候 就已经把这种属性 注入到硬件的平台 
    	//但是在gpu平台上面 对于数据对齐的处理要求会更严格
    	//因为并行计算 需要数据的提供延时更小 所以在不是4字节的类型里面 会自动进行填充 
    	//也就是说在gpu平台上面 这个student结构体12字节
    	std::cout << "sizeof(stu):" << sizeof(stu) << endl;
    	// printf(sizeof(stu));
    	//void*类型作为一个万能类型 专门给强制类型转换预留的接口 内核驱动里面和回掉函数里面 会有大量应用
    	// malloc free memset memcpy
    	char *p = (char *)malloc(100 * sizeof(char));
    
    	//申请的空间里面的值全部赋值成0 
    	memset(p, 0, 100);
    
    	char p_str[50] = "I love C++ program!!!";
    	//char *pstr = "I love C program!!!";
    
    	//把内存里面的值进行拷贝
    	memcpy(p, p_str, strlen(p_str));
    
    	//格式化输出
    	printf("%s", p);
    
    	//申请的内存空间需要释放 告诉编译器现在这段空间我不需要了 编译器进而通知系统 这段空间我现在不需要了 最后系统把这段空间重新利用
    	free(p);
    
    	system("pause");
    	return 0;
    }

**总结**：

1.  c/c++写一段代码流程：  
    1.申请内存 malloc / new  
    2.内存初始化 memset  
    3.处理数据(此处是复制数据，memcpy)  
    4.输出 （printf）  
    5.释放内存 free/ delete
2.  自定义数据，数据类型对齐  
    在cpu-memory模式下面，数据是需要对齐的，因为这样cpu和memory的内存交互速度是最快的。  
    在硬件设计的时候，就已经把这种属性，注入到硬件的平台。  
    但是在gpu平台上面，对于数据对齐的处理要求会更严格。  
    因为并行计算，需要数据的提供延时更小，所以在不是4字节的类型里面，会自动进行填充 。  
    void\*类型作为一个万能类型，专门给强制类型转换预留的接口，内核驱动里面和回掉函数里面会有大量应用。

## 二、 cuda编程基础概念 ##

基础概念  
主机：cpu和内存  
设备：gpu和显存  
API：  
warp： thread  
访问速度不同  
变量类型限定符：**device** **shared** **constant**  
函数类型限定符号 **global** 在cpu上定义，在gpu上显存上执行  
**device** gpu上执行  
**host** 主机上执行 cpu

thread：线程 cpu架构基本一致 同步  
多少个流就可以创建多少个线程  
block： 多个线程  
grid: 多个block  
SIMT：**单指令多线程**  
内置变量：threadid blockid blockdim gridid griddim  
显存带宽犹如高速公路的车道，优化的方法是尽可能接近理论的带宽。

代码2

/***************************************************************************
    * first_cuda.cu
    *1.将数据从主机内存数据复制到设备显存
    *2.写好核函数
    *3.CUDA编译器执行核函数 在GPU上完成计算操作
    *4.把显存数据复制到主机内存
    *5.释放显存空间
    /***************************************************************************/
    
    #include <stdio.h>
    #include <stdlib.h>
    
    //CUDA RunTime API
    #include <cuda_runtime.h>
    #define DATA_SIZE 1048576
    int data[DATA_SIZE];
    
    //产生大量0-9之间的随机数
    // 指针不带内存大小，所以需要传入一个size，指明大小
    void GenerateNumbers(int *number, int size)
    {
        
    	for (int i = 0; i < size; i++) {
        
    		number[i] = rand() % 10;
    	}
    }
    
    //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;
    		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)
    {
        
    	int sum = 0;
    	int i;
    	for (i = 0; i < DATA_SIZE; i++) {
        
    		sum += num[i] * num[i] * num[i];
    	}
    	*result = sum;
    }
    
    int main()
    {
        
    	//CUDA 初始化
    	if (!InitCUDA()) {
        
    		return 0;
    	}
    	//生成随机数
    	GenerateNumbers(data, DATA_SIZE);
    	/*把数据复制到显卡内存中*/
    	int* gpudata, *result;
    	//cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果 )
    	cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
    	cudaMalloc((void**)&result, sizeof(int));
    	//cudaMemcpy 将产生的随机数复制到显卡内存中 
    	//cudaMemcpyHostToDevice - 从内存复制到显卡内存
    	//cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    	cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
    	// 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
    	sumOfSquares << <1, 1, 0 >> >(gpudata, result);
    	/*把结果从显示芯片复制回主内存*/
    	int sum;
    	//cudaMemcpy 将结果从显存中复制回内存
    	cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);
    	//Free
    	cudaFree(gpudata);
    	cudaFree(result);
    	printf("GPUsum: %d \n", sum);
    	sum = 0;
    	for (int i = 0; i < DATA_SIZE; i++) {
        
    		sum += data[i] * data[i] * data[i];
    	}
    	printf("CPUsum: %d \n", sum);
    	return 0;
    }

**总结cuda编程的流程**：

1.  申请显存空间 cudaMalloc
2.  将数据从主机内存数据复制到设备显存
3.  写好核函数
4.  CUDA编译器执行核函数 在GPU上完成计算操作
5.  把显存数据复制到主机内存
6.  释放显存空间

**API**：

1.  获取 CUDA 设备数：  
    函数原型：

cudaError_t cudaGetDeviceCount( int* count )

可以通过 cudaGetDeviceCount 函数获取 CUDA 的设备数，实例如上InitCUDA()函数通过引用传递 count 值，获取当前支持的 CUDA 设备数，即返回具有计算能力的设备的数量。  
以 \*count 形式返回可用于执行的计算能力大于等于 1.0 的设备数量。如果不存在此类设备，将返回 1  
**cudaSuccess**，注意，如果之前是异步启动，该函数可能返回错误码。

1.  获取 CUDA 设备属性  
    函数原型：

cudaError_t cudaGetDeviceProperties( struct cudaDeviceProp* prop,int dev )

可以通过 cudaGetDeviceProperties 函数获取 CUDA 设备的属性，具体用法  
函数通过引用传递 prop 关于属性的结构体，并且列出主设备号大于 1 的设备属性，其中设备属性通过函数 printDeviceProp 打印。  
以\*prop形式返回设备dev的属性。  
返回值：cudaSuccess、cudaErrorInvalidDevice，注，如果之前是异步启动，该函数可能返回错误码。  
3. 设置 CUDA 设备  
通过函数 cudaSetDevice 就可以设置 CUDA 设备了，具体用法

// set cuda device
    cudaSetDevice(i);

原型：

cudaError_t cudaSetDevice(int dev)

将dev记录为活动主线程将执行设备码的设备。  
cudaSuccess、cudaErrorInvalidDevice，注，如果之前是异步启动，该函数可能返回错误码。  
cudaDeviceProp 结构定义如下：

struct cudaDeviceProp {
        
    char name [256];
    size_t totalGlobalMem;
    size_t sharedMemPerBlock;
    int regsPerBlock;
    int warpSize;
    size_t memPitch;
    int maxThreadsPerBlock;
    int maxThreadsDim [3];
    int maxGridSize [3];
    size_t totalConstMem;
    int major;
    int minor;
    int clockRate;
    size_t textureAlignment;
    int deviceOverlap;
    int multiProcessorCount;
    }

CUDA 初始化完整代码函数为InitCUDA()  
**cudaDeviceProp 结构中的各个变量意义如下**：  
**name**：用于标识设备的ASCII字符串;  
**totalGlobalMem**：设备上可用的全局存储器的总量,以字节为单位;  
**sharedMemPerBlock**：线程块可以使用的共享存储器的最大值,以字节为单位;多处理器上的所有线程块可以同时共享这些存储器;  
**regsPerBlock**：线程块可以使用的32位寄存器的最大值;多处理器上的所有线程块可以同时共享这些寄存器;  
**warpSize**：按线程计算的warp块大小;  
**memPitch**：允许通过cudaMallocPitch()为包含存储器区域的存储器复制函数分配的最大间距(pitch),以字节为单位;  
**maxThreadsPerBlock**：每个块中的最大线程数  
**maxThreadsDim\[3\]**：块各个维度的最大值:  
**maxGridSize\[3\]**：网格各个维度的最大值;  
**totalConstMem**：设备上可用的不变存储器总量,以字节为单位;  
**major,minor**：定义设备计算能力的主要修订号和次要修订号;  
**clockRate**：以千赫为单位的时钟频率;  
**textureAlignment**：对齐要求;与textureAlignment字节对齐的纹理基址无需对纹理取样应用偏移;  
**deviceOverlap**：如果设备可在主机和设备之间并发复制存储器,同时又能执行内核,则此值为 1;否则此值为 0;  
**multiProcessorCount**：设备上多处理器的数量。

## 三、 CUDA程序优化和评估的方式 ##

代码3：

#include <stdio.h>
    #include <stdlib.h>
    #include <time.h>
    
    //CUDA RunTime API
    #include <cuda_runtime.h>
    
    //1024 * 1024  1M
    #define DATA_SIZE 1048576
    
    int data[DATA_SIZE];
    float clockRate = 1.0;
    
    //产生大量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);
    		//打印gpu设备信息
    		printDeviceProp(prop);
    		//获得显卡的始终频率
    		clockRate = prop.clockRate;
    		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)
    {
        
    	int sum = 0;
    	int i;
    	clock_t start = clock();
    	for (i = 0; i < DATA_SIZE; i++) {
        
    		sum += num[i] * num[i] * num[i];
    	}
    	*result = sum;
    	*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));
    	cudaMalloc((void**)&time, sizeof(clock_t));
    	//cudaMemcpy 将产生的随机数复制到显卡内存中
    	//cudaMemcpyHostToDevice - 从内存复制到显卡内存
    	//cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    	cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
    	// 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
    	sumOfSquares << <1, 1, 0 >> >(gpudata, result, time);
    	/*把结果从显示芯片复制回主内存*/
    	int sum;
    	clock_t time_used;
    	//cudaMemcpy 将结果从显存中复制回内存，结果在cpu上显示出来
    	cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);
    	cudaMemcpy(&time_used, time, sizeof(clock_t), cudaMemcpyDeviceToHost);
    	//Free
    	cudaFree(gpudata);
    	cudaFree(result);
    	cudaFree(time);
    	printf("GPUsum: %d time_clock: %d time: %fs\n", sum, time_used, ((float)time_used / (clockRate * 1000)));
    	sum = 0;
    	for (int i = 0; i < DATA_SIZE; i++) {
          //循环1M的数据
    		sum += data[i] * data[i] * data[i];
    	}
    	printf("CPUsum: %d \n", sum);
    	return 0;
    }

**以上代码分析以及优化方式**：  
1M的数据（1024\*1024） int类型是4个字节，即4byte数据 ，一共是的数据量就是4M  
带宽：4MB/1.026178 = 3.89MB/s （4MB的数据，用时是1.026178s，计算得出3.89MB/s）  
我显卡的带宽是14Gb/s左右 也就是说 完全没有实现并行计算的威力  
**优化**：

1.  一定要先从显存带宽开始
2.  确定任务中并行和串行的算法
3.  需要两层线程并行的内核函数，每个SM上面至少有6个warp和2个block
4.  共享内存 shared memory

**nvcc 编译代码**  
nvcc 是 CUDA 的编译工具，它可以将 .cu 文件解析出在 GPU 和 host 上执行的部分，也就是说，它会帮忙把 GPU 上执行和主机上执行的代码区分开来，不许要我们手动去做了。在 GPU 执行的部分会通过 NVIDIA 提供的 编译器编译成中介码，主机执行的部分则调用 gcc 编译。

nvcc -o first_cuda first_cuda.cu

通过上述编译，生成可执行文件 first\_cuda

## 四、 thread多线程概念的引入 ##

cuda中GPU 的架构。它是由 grid 组成，每个 grid 又可以由 block 组成，而每个 block 又可以细分为 thread。所以，线程是我们处理的最小的单元。  
接下来的例子通过修改前一个例子，把数组分割成若干个组（每个组由一个线程实现），每个组计算出一个和，然后在 CPU 中将分组的这几个和加在一起，得到最终的结果。这种思想叫做**归约** 。其实和分治思想差不多，就是先将大规模问题分解为小规模的问题，最后这些小规模问题整合得到最终解。  
由于我的 GPU 支持的块内最大的线程数是 256个，即 **cudaGetDeviceProperties** 中的 **maxThreadsPerBlock** 属性。如何获取这个属性。  
我们使用 256 个线程来实现并行加速。

代码4：thread多线程概念的引入:用时0.0244065s

#include <stdio.h>
    #include <stdlib.h>
    #include <time.h>
    
    //CUDA RunTime API
    #include <cuda_runtime.h>
    
    //1M
    #define DATA_SIZE 1048576
    #define THREAD_NUM 256  //256线程
    
    int data[DATA_SIZE];
    int clockRate = 1.0;
    
    //产生大量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);
    		//获得显卡的时钟频率
    		clockRate = prop.clockRate;
    		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; //tid相当于一个向量，大小为256
    	//计算每个线程需要完成的量
    	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;
    }
    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 time_clock: %d time: %fs\n", final_sum, time_use, ((float)time_use / (clockRate * 1000)));
    	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;
    }
    //创建了256个线程帮我们并行计算 我们需要提前安排好每个线程计算的数据 防止线程同步的问题

重点代码分析

__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
    {
        
     //表示目前的 thread 是第几个 thread（由 0 开始计算）
     const int tid = threadIdx.x; //tid相当于一个向量，大小为256
     //计算每个线程需要完成的量
     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;
    }
    int main()
    {
        
     //C

代码5：内存连续存储:用时0.0053375s  
线程：0,1,2…,255, 0,1,2…255

#include <stdio.h>
    #include <stdlib.h>
    #include <time.h>
    
    //CUDA RunTime API
    #include <cuda_runtime.h>
    
    //1M
    #define DATA_SIZE 1048576
    #define THREAD_NUM 256
    
    int data[DATA_SIZE];
    int clockRate = 1.0;
    
    //产生大量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);
    		//获得显卡的时钟频率
    		clockRate = prop.clockRate;
    		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();
    	//改为连续存取（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;
    }
    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 time_clock: %d time: %f\n", final_sum, time_use, ((float)time_use / (clockRate * 1000)));
    	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;
    }
    //我们的程序需要尽可能连续操作内存，减少内存存取方面的时间浪费

## 五、block线程块概念 ##

代码6：block线程块概念引入：用时：0.002125s  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1NNRjA1MDQ_size_16_color_FFFFFF_t_70]

#include <stdio.h>
    #include <stdlib.h>
    #include <time.h>
    
    //CUDA RunTime API
    #include <cuda_runtime.h>
    
    //1M
    #define DATA_SIZE 1048576  //1024*1024
    #define THREAD_NUM 256
    #define BLOCK_NUM 32  // block
    int clockRate = 1.0;
    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);
    		//获得显卡的时钟频率
    		clockRate = prop.clockRate;
    		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__ 函数 (GPU上执行) 计算立方和
    __global__ static void sumOfSquares(int *num, int* result, clock_t* time)
    {
        
    	//表示目前的 thread 是第几个 thread（由 0 开始计算）0~255
    	const int tid = threadIdx.x;
    	//表示目前的 thread 属于第几个 block（由 0 开始计算）0~31
    	const int bid = blockIdx.x;
    	int sum = 0;
    	int i;
    	//记录运算开始的时间
    	clock_t start;
    	//只在 thread 0（即 threadIdx.x = 0 的时候）进行记录，每个 block 都会记录开始时间及结束时间
    	if (tid == 0) time[bid] = clock();
    	//thread需要同时通过tid和bid来确定，同时不要忘记保证内存连续性
    	for (i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
        
    		sum += num[i] * num[i] * num[i];
    	}
    	//Result的数量随之增加
    	result[bid * THREAD_NUM + tid] = sum;
    	//计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行，每个 block 都会记录开始时间及结束时间
    	if (tid == 0) time[bid + BLOCK_NUM] = clock();
    }
    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* BLOCK_NUM);
    	cudaMalloc((void**)&time, sizeof(clock_t)* BLOCK_NUM * 2);
    	//cudaMemcpy 将产生的随机数复制到显卡内存中
    	//cudaMemcpyHostToDevice - 从内存复制到显卡内存
    	//cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    	cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
    	// 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
    	sumOfSquares << < BLOCK_NUM, THREAD_NUM, 0 >> >(gpudata, result, time);
    	/*把结果从显示芯片复制回主内存*/
    	int sum[THREAD_NUM*BLOCK_NUM];
    	clock_t time_use[BLOCK_NUM * 2];
    	//cudaMemcpy 将结果从显存中复制回内存
    	cudaMemcpy(&sum, result, sizeof(int)* THREAD_NUM*BLOCK_NUM, cudaMemcpyDeviceToHost);
    	cudaMemcpy(&time_use, time, sizeof(clock_t)* BLOCK_NUM * 2, cudaMemcpyDeviceToHost);
    	//Free
    	cudaFree(gpudata);
    	cudaFree(result);
    	cudaFree(time);
    	int final_sum = 0;
    	for (int i = 0; i < THREAD_NUM*BLOCK_NUM; i++) {
        
    		final_sum += sum[i];
    	}
    	//采取新的计时策略 把每个 block 最早的开始时间，和最晚的结束时间相减，取得总运行时间
    	clock_t min_start, max_end;
    	min_start = time_use[0];
    	max_end = time_use[BLOCK_NUM];
    	for (int i = 1; i < BLOCK_NUM; i++) {
        
    		if (min_start > time_use[i])
    			min_start = time_use[i];
    		if (max_end < time_use[i + BLOCK_NUM])
    			max_end = time_use[i + BLOCK_NUM];
    	}
    	printf("GPUsum: %d time_clock: %d time: %fs\n", final_sum, max_end - min_start, ((float)(max_end - min_start) / (clockRate * 1000)));
    	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;
    }

## 六、共享内存与同步 ##

代码7:共享内存机制：用时：0.003568s  
看似用时比上一次多了，其实主要节约了cpu中for循环的计算用时，同时也减少了数据copy的用时。总体来说，cpu和gpu的总时间减少了。  
线程同步问题。即一个block中所有的线程都操作完毕了。  
一个block中的所有线程可以共享内存  
通常，数据的量是大于线程的总和的

#include <stdio.h>
    #include <stdlib.h>
    #include <time.h>
    
    //CUDA RunTime API
    #include <cuda_runtime.h>
    
    //1M
    #define DATA_SIZE 1048576
    #define THREAD_NUM 256
    #define BLOCK_NUM 32
    int clockRate = 1.0;
    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);
    		//获得显卡的时钟频率
    		clockRate = prop.clockRate;
    		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)
    {
        
    	//声明一块共享内存
    	extern __shared__ int shared[];  
    	//表示目前的 thread 是第几个 thread（由 0 开始计算）
    	const int tid = threadIdx.x;
    	//表示目前的 thread 属于第几个 block（由 0 开始计算）
    	const int bid = blockIdx.x;
    	shared[tid] = 0;
    	int i;
    	//记录运算开始的时间
    	clock_t start;
    	//只在 thread 0（即 threadIdx.x = 0 的时候）进行记录，每个 block 都会记录开始时间及结束时间
    	if (tid == 0) time[bid] = clock();
    	//thread需要同时通过tid和bid来确定，同时不要忘记保证内存连续性
    	for (i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
        
    		shared[tid] += num[i] * num[i] * num[i];
    	}
    	//同步 保证每个 thread 都已经把结果写到 shared[tid] 里面
    	//cuda编程里，很大程度再解决如何并行和串行不出题
    	__syncthreads();//共享内存的基础上
    	//使用线程0完成加和
    	if (tid == 0)
    	{
        
    		for (i = 1; i < THREAD_NUM; i++)
    		{
        
    			shared[0] += shared[i];
    		}
    		result[bid] = shared[0];
    	}
    	//计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行，每个 block 都会记录开始时间及结束时间
    	if (tid == 0) time[bid + BLOCK_NUM] = clock();
    }
    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)* BLOCK_NUM);
    	cudaMalloc((void**)&time, sizeof(clock_t)* BLOCK_NUM * 2);
    	//cudaMemcpy 将产生的随机数复制到显卡内存中
    	//cudaMemcpyHostToDevice - 从内存复制到显卡内存
    	//cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    	cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
    	// 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
    	sumOfSquares << < BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int) >> >(gpudata, result, time);
    	/*把结果从显示芯片复制回主内存*/
    	int sum[BLOCK_NUM];
    	clock_t time_use[BLOCK_NUM * 2];
    	//cudaMemcpy 将结果从显存中复制回内存
    	cudaMemcpy(&sum, result, sizeof(int)*BLOCK_NUM, cudaMemcpyDeviceToHost);
    	cudaMemcpy(&time_use, time, sizeof(clock_t)* BLOCK_NUM * 2, cudaMemcpyDeviceToHost);
    	//Free
    	cudaFree(gpudata);
    	cudaFree(result);
    	cudaFree(time);
    	int final_sum = 0;
    	for (int i = 0; i < BLOCK_NUM; i++) {
        
    		final_sum += sum[i];
    	}
    	//采取新的计时策略 把每个 block 最早的开始时间，和最晚的结束时间相减，取得总运行时间
    	clock_t min_start, max_end;
    	min_start = time_use[0];
    	max_end = time_use[BLOCK_NUM];
    	for (int i = 1; i < BLOCK_NUM; i++) {
        
    		if (min_start > time_use[i])
    			min_start = time_use[i];
    		if (max_end < time_use[i + BLOCK_NUM])
    			max_end = time_use[i + BLOCK_NUM];
    	}
    	printf("GPUsum: %d time_clock: %d time: %fs\n", final_sum, max_end - min_start, ((float)(max_end - min_start) / (clockRate * 1000)));
    	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;
    }

## 七、树状求和算法加速 ##

代码8：树状求和算法简介: 用时：0.002941s  
算法优化

#include <stdio.h>
    #include <stdlib.h>
    #include <time.h>
    
    //CUDA RunTime API
    #include <cuda_runtime.h>
    
    //1M
    #define DATA_SIZE 1048576
    #define THREAD_NUM 256
    #define BLOCK_NUM 32
    int clockRate = 1.0;
    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);
    		//获得显卡的时钟频率
    		clockRate = prop.clockRate;
    		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)
    {
        
    	//声明一块共享内存
    	extern __shared__ int shared[];
    	//表示目前的 thread 是第几个 thread（由 0 开始计算）
    	const int tid = threadIdx.x;
    	//表示目前的 thread 属于第几个 block（由 0 开始计算）
    	const int bid = blockIdx.x;
    	shared[tid] = 0;
    	int i;
    	//记录运算开始的时间
    	clock_t start;
    	//只在 thread 0（即 threadIdx.x = 0 的时候）进行记录，每个 block 都会记录开始时间及结束时间
    	if (tid == 0) time[bid] = clock();
    	//thread需要同时通过tid和bid来确定，同时不要忘记保证内存连续性
    	for (i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
        
    		shared[tid] += num[i] * num[i] * num[i];
    	}
    	//同步 保证每个 thread 都已经把结果写到 shared[tid] 里面
    	__syncthreads();
    	//树状加法
    	int offset = 1, mask = 1;
    	while (offset < THREAD_NUM)
    	{
        
    		if ((tid & mask) == 0)
    		{
        
    			shared[tid] += shared[tid + offset];
    		}
    		offset += offset;
    		mask = offset + mask;
    		__syncthreads();
    	}
    	//计算时间,记录结果，只在 thread 0（即 threadIdx.x = 0 的时候）进行，每个 block 都会记录开始时间及结束时间
    	if (tid == 0)
    	{
        
    		result[bid] = shared[0];
    		time[bid + BLOCK_NUM] = clock();
    	}
    }
    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)* BLOCK_NUM);
    	cudaMalloc((void**)&time, sizeof(clock_t)* BLOCK_NUM * 2);
    	//cudaMemcpy 将产生的随机数复制到显卡内存中
    	//cudaMemcpyHostToDevice - 从内存复制到显卡内存
    	//cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    	cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
    	// 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
    	sumOfSquares << < BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int) >> >(gpudata, result, time);
    	/*把结果从显示芯片复制回主内存*/
    	int sum[BLOCK_NUM];
    	clock_t time_use[BLOCK_NUM * 2];
    	//cudaMemcpy 将结果从显存中复制回内存
    	cudaMemcpy(&sum, result, sizeof(int)*BLOCK_NUM, cudaMemcpyDeviceToHost);
    	cudaMemcpy(&time_use, time, sizeof(clock_t)* BLOCK_NUM * 2, cudaMemcpyDeviceToHost);
    	//Free
    	cudaFree(gpudata);
    	cudaFree(result);
    	cudaFree(time);
    	int final_sum = 0;
    	for (int i = 0; i < BLOCK_NUM; i++) {
        
    		final_sum += sum[i];
    	}
    	//采取新的计时策略 把每个 block 最早的开始时间，和最晚的结束时间相减，取得总运行时间
    	clock_t min_start, max_end;
    	min_start = time_use[0];
    	max_end = time_use[BLOCK_NUM];
    	for (int i = 1; i < BLOCK_NUM; i++) {
        
    		if (min_start > time_use[i])
    			min_start = time_use[i];
    		if (max_end < time_use[i + BLOCK_NUM])
    			max_end = time_use[i + BLOCK_NUM];
    	}
    	printf("GPUsum: %d time_clock: %d time: %f\n", final_sum, max_end - min_start, ((float)(max_end - min_start) / (clockRate * 1000)));
    	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;
    }

[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1NNRjA1MDQ_size_16_color_FFFFFF_t_70]: /images/20221022/1e58b327ee2b422095e3d6c054843496.png