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

最新推荐文章于 2024-04-17 11:20:13 发布
最新推荐文章于 2024-04-17 11:20:13 发布
阅读量761 收藏 8
点赞数 1
分类专栏： 嵌入式AI C/C++
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/SMF0504/article/details/115293105
版权

一、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
在这里插入图片描述

#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;
}
荪荪
关注
  • 1
    点赞
  • 踩
  • 8
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
c语言并行结构程序,C语言并行程序编程（一）
weixin_35878962的博客
05-18 1098
这里先简单书写一个并行版的hello world 作为入门，后续的文章再更新相关的概念。#include #include #include void Hello(void);int main(int argc,char* argv[]){int thread_count = strtol(argv[1],NULL,10);//获取线程数# pragma omp parallel num_th...
c语言并行程序利用了openmp实现并行计算
05-30
openmp加c语言程序OpenMP是一个支持共享存储并行设计的库，特别适宜多核CPU上的并行程序设计。
C语言实现并行求和算法
流星的博客
11-03 2172
[b]1、问题描述[/b] 将数组A均匀划分成m个片段，每个数组片段最多有(n+m-1)/m 个元素。每个数组片段分别由一个线程负责局部求和，最后这些部分和加起来就得到数组中所有元素的总和。 [b]2、相关代码[/b] 此代码在gcc4.3下编译通过 [code="java"] #include #include #include #define NUM_THREADS ...
并行计算基础以及相关C语言API介绍
最新发布
qq_53324833的博客
04-17 791
功能描述：此函数用于初始化MPI环境。它接受两个参数，&argc和&argv，分别表示命令行参数的数量和值。这是开始任何MPI程序的第一步。参数描述：argc：指向命令行参数数量的指针。在MPI_Init调用后，这个值可能会被修改，因此最好传入实际使用的参数数量变量的地址。argv：指向命令行参数数组的指针的指针。这个数组同样可能在MPI_Init调用后被修改。功能描述：此函数表示MPI程序的结束。它是MPI程序的最后一条可执行语句，用于清理和关闭MPI环境。
【并行计算入门C语言篇】
qq_41666483的博客
01-04 1377
并行计算入门学习任务C语言部分
C语言实现并行计算
m0_65591847的博客
10-13 603
包含C语言static、inline、external关键字，pthread多线程库、OpenMP多线程库、错误处理等
行业资料-交通装置-一种新型板栗炒制销售车.zip
08-22
行业资料-交通装置-一种新型板栗炒制销售车.zip
行业文档-设计装置-一种节能炒栗子锅灶.zip
08-29
行业文档-设计装置-一种节能炒栗子锅灶.zip
一个栗子-序列化工具protobuf
07-08
一个栗子-序列化工具protobuf
CUDA By Example(三)——CUDA C并行编程
hhhcbw的博客，欢迎各位来访
11-09 1190
CUDA C并行编程，主要通过矢量求和与Julia集的绘制进行说明。以及对编译运行代码过程中遇到的问题及解决方案进行说明
使用cuda c进行并行编程
weixin_59999837的博客
11-18 282
使用cuda c进行并行编程
并行计算三种方式计算π
03-22
并行计算三种方式计算π：概率法，积分法，级数法，代码中包括计算量和线程个数设置。 使用时进行编译后输入相应提示的数值即可，例如N=100000 t=8
使用GPU对C语言代码进行加速_c语言代码大全和详细解释
12-07
使用GPU对C语言代码进行加速，其中以规则匹配为例，优化后的方法为：特征规则使用acbm算法编译成一个ptree、ptree匹配使用cuda编程在GPU运行、多个packet并发匹配，处理速度提升数倍。
c/c++并行计算
浩瀚的海洋
09-04 8756
 现在要跟老师做个关于遗传算法的程序，想通过几台机器运行，找到了几篇好的文章，但还没有仔细看和验证。网址一：http://www.cnblogs.com/cy163/archive/2007/01/19/624281.html网址二：http://blog.chinaunix.net/u1/37798/showart_320560.html关于在Windows环境下，利用Visual
多核并行计算：C语言Pthreads线程库小记
小那猿同学的博客
10-14 943
多核并行计算：C语言Pthreads线程库小记 在多核并行计算中，采用Pthreads进行共享内存编程，这篇笔记文章是对Pthreads库的小记。 对象 pthread_t pthread_t对象是一个不透明对象，对象存储的数据由系统绑定，用户级代码无法直接访问，pthread_t对象可以对他所从属的对象进行唯一的标识。 方法 pthread_create pthread_create用来生成线程，它的语法为： int pthread_create( pthread_t* thread_p
CUDA C编程4 - 动态并行理解
DU_YULIN的博客
03-01 1414
系列文章目录 文章目录系列文章目录前言一、动态并行概念二、嵌套执行三、动态并行的限制条件四、嵌套规约4.1 基本方法（线程同步及线程块同步）4.2 加速方法（去除线程同步以及线程块同步）4.3 进一步加速方法（待完善。。。）五、运行结果总结参考资料附录：完整代码 前言 之前学习了CUDA C编程中动态并行技术， 这里进行总结，以加深理解，并希望能够对大家有所帮助。 一、动态并行概念 之前接触的CUDA C设备函数的调用都是在CPU端进行的，即设备函数的执行在CPU的控制下。自然而然，我们就会
mpi大规模矩阵乘法C语言,C语言基于MPI并行计算矩阵的乘法
weixin_39753260的博客
05-24 396
浏览量：1,030(在苹果系统下，如果文章中的图片不能正常显示，请升级Safari浏览器到最新版本，或者使用Chrome、Firefox浏览器打开。)矩阵的计算是并行计算里的一个很重要的问题。矩阵是一种网格化的数据，是一组同类型数值的集合，矩阵的出现，使得代数系统更完善，对各种实际问题的求解产生了巨大的作用。但是其庞大的计算量往往令人生畏，稍微大一点的矩阵，计算就变得非常繁琐，不仅仅是人，而且串行...
c语言多核编程,多核并行计算——并发与并行编程问题
weixin_34991792的博客
05-18 2240
并发是用来描述并行执行的方式(看上去一起发生的事件,例如目前操作系统的多任务调度程序,看上去桌面上有许多程序在同时运行.),并发是形容词.并行是指通过并发将一个操作分解成一组粒度更细的工作单元,并且这些工作单元可以在不同的处理器内核上运行.并行是动词,指必须有两个以上的事件发生.工作单元,以粒度可分为最小到单一CPU指令,大到函数或系统任务.并行有数据并行性与任务并行性数据并行性指对许多数据执行相...
c语言怎么编程随机采板栗
03-24
您可以使用C语言的rand()函数来生成随机数，然后使用if语句来判断是否采摘板栗。以下是一个简单的示例代码： #include <stdio.h> #include <stdlib.h> #include <time.h> int main() { int rand_num; srand(time(NULL)); // 初始化随机数种子 rand_num = rand() % 2; // 生成0或1的随机数 if (rand_num == 0) { printf("不采摘板栗\n"); } else { printf("采摘板栗\n"); } return 0; }

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值