CUDA之单thread单block多thread单block多thread多block
用简单的立方和归约来举例：
//单thread单block
#include
#include
#include
#define DATA_SIZE 1048576
int data[DATA_SIZE];
//产生大量0-9之间的随机数
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.
");
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.
");
return false;
}
cudaSetDevice(i);
return true;
}
// __global__ 函数(GPU上执行) 计算立方和
__global__ static void sumOfcubes(int *num, int* result)
{
intsum = 0;
inti;
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((void**)&gpudata, sizeof(int)* DATA_SIZE);
cudaMalloc((void**)&result, sizeof(int));
//cudaMemcpy 将产生的随机数复制到显卡内存中
cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
sumOfcubes<< <1, 1, 0 >> > (gpudata, result);
cudaMemcpy(sum, result, sizeof(int), cudaMemcpyDeviceToHost);
cudaFree(gpudata);
cudaFree(result);
printf("GPUsum: %d
", sum);
int sum = 0;
for (int i = 0; i < DATA_SIZE; i++) {
sum += data[i] * data[i] * data[i];
}
printf("CPUsum: %d
", sum);
getchar();
return 0;
} //单block多thread#include
#include
#include
//CUDA RunTime API
#include
#include "device_launch_parameters.h"#define DATA_SIZE 1048576
#define THREAD_NUM 1024 //256--->1024
int data[DATA_SIZE];void GenerateNumbers(int *number, int size)
{for (int i = 0; i < size; i++) {number[i] = rand() % 10;}
}
// __global__ 函数(GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{ const int tid = threadIdx.x;
//计算每个线程需要完成的量
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; i < DATA_SIZE; i += THREAD_NUM)
//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 初始化//生成随机数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 将产生的随机数复制到显卡内存中cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);// 启动kernel函数cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);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);//FreecudaFree(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:%d
", final_sum,time_use);final_sum = 0;for (int i = 0; i < DATA_SIZE; i++) {final_sum += data[i] * data[i] * data[i];}printf("CPUsum: %d
", final_sum);getchar();return 0;
}
//多block多thread#include
#include
#include
//CUDA RunTime API
#include
#include "device_launch_parameters.h"#define DATA_SIZE 1048576
#define THREAD_NUM 256
#define BLOCK_NUM 32
int data[DATA_SIZE];void GenerateNumbers(int *number, int size)
{for (int i = 0; i < size; i++) {number[i] = rand() % 10;}
}
// __global__ 函数(GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{const int tid = threadIdx.x;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()
{GenerateNumbers(data, DATA_SIZE);int* gpudata, *result;clock_t* 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(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);// 在CUDA 中执行函数语法：函数名称<<>>(参数...);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);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 gputime: %d
", final_sum, max_end - min_start);final_sum = 0;for (int i = 0; i < DATA_SIZE; i++){final_sum += data[i] * data[i] * data[i];}printf("CPUsum: %d
", final_sum);getchar();return 0;
} ShareMemory
是一个block 中所有thread 都能使用的共享内存，存取的速度相当快，存取shared memory 的速度和存取寄存器相同，不需要担心latency 的问题。
可以直接利用__shared__声明一个shared memory变量
__shared__ float temp[THREAD_NUM * 3];
Shared memory 有时会出现存储体冲突（bank conflict）的问题：
例如：每个SM有16KB 的shared memory，分成16 个bank
•如果同时每个thread 是存取不同的bank，就不会有问题
•如果同时有两个（或更多）threads 存取同一个bank 的数据，就会发生bank conflict，这些threads 就必须照顺序去存取，而无法同时存取shared memory 了。
//多block多thread 使用sharememory#include
#include
#include
//CUDA RunTime API
#include
#include "device_launch_parameters.h"
#include "device_functions.h"#define DATA_SIZE 1048576
#define THREAD_NUM 256
#define BLOCK_NUM 32
int data[DATA_SIZE];void GenerateNumbers(int *number, int size)
{for (int i = 0; i < size; i++) {number[i] = rand() % 10;}
}
// __global__ 函数(GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{extern __shared__ int shared[];const int tid = threadIdx.x;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();//使用线程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()
{GenerateNumbers(data, DATA_SIZE);int* gpudata, *result;clock_t* 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(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);// 在CUDA 中执行函数语法：函数名称<<>>(参数...);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);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 gputime: %d
", final_sum, max_end - min_start);final_sum = 0;for (int i = 0; i < DATA_SIZE; i++){final_sum += data[i] * data[i] * data[i];}printf("CPUsum: %d
", final_sum);getchar();return 0;
}
Block内完成部分加和工作，所以gputime增加了
//多block多thread#include
#include
#include
//CUDA RunTime API
#include
#include "device_launch_parameters.h"
#include "device_functions.h"#define DATA_SIZE 1048576
#define THREAD_NUM 256
#define BLOCK_NUM 32
int data[DATA_SIZE];void GenerateNumbers(int *number, int size)
{for (int i = 0; i < size; i++) {number[i] = rand() % 10;}
}
// __global__ 函数(GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{extern __shared__ int shared[];const int tid = threadIdx.x;const int bid = blockIdx.x;shared[tid] = 0;int i;//记录运算开始的时间//只在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();}if (tid == 0){result[bid] = shared[0];time[bid + BLOCK_NUM] = clock();}
}int main()
{GenerateNumbers(data, DATA_SIZE);int* gpudata, *result;clock_t* 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(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);// 在CUDA 中执行函数语法：函数名称<<>>(参数...);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);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 gputime: %d
", final_sum, max_end - min_start);final_sum = 0;for (int i = 0; i < DATA_SIZE; i++){final_sum += data[i] * data[i] * data[i];}printf("CPUsum: %d
", final_sum);getchar();return 0;
}
更多相关：
-
Apache POI是一个开源的利用Java读写Excel，WORD等微软OLE2组件文档的项目。 我的需求是对Excel的数据进行导入或将数据以Excel的形式导出。先上简单的测试代码：package com.xing.studyTest.poi;import java.io.FileInputSt...
-
要取得[a,b)的随机整数，使用(rand() % (b-a))+ a; 要取得[a,b]的随机整数，使用(rand() % (b-a+1))+ a; 要取得(a,b]的随机整数，使用(rand() % (b-a))+ a + 1; 通用公式:a + rand() % n；其中的a是起始值，n是整数的范围。 要取得a到b之间的...
-
利用本征图像分解(Intrinsic Image Decomposition)算法，将图像分解为shading(illumination) image 和 reflectance(albedo) image，计算图像的reflectance image。 Reflectance Image 是指在变化的光照条件下能够维持不变的图像部分...
-
题目：面试题39. 数组中出现次数超过一半的数字 数组中有一个数字出现的次数超过数组长度的一半，请找出这个数字。 你可以假设数组是非空的，并且给定的数组总是存在多数元素。 示例 1: 输入: [1, 2, 3, 2, 2, 2, 5, 4, 2] 输出: 2 限制： 1 <= 数组长度 <= 50000 解题： cl...
-
题目：二叉搜索树的后序遍历序列 输入一个整数数组，判断该数组是不是某二叉搜索树的后序遍历结果。如果是则返回 true，否则返回 false。假设输入的数组的任意两个数字都互不相同。 参考以下这颗二叉搜索树： 5 / 2 6 / 1 3示例 1： 输入: [1,6,3,2,5] 输出...
-
给定一个以字符串表示的非负整数 num，移除这个数中的 k 位数字，使得剩下的数字最小。 注意: num 的长度小于 10002 且 ≥ k。 num 不会包含任何前导零。 示例 1 : 输入: num = “1432219”, k = 3 输出: “1219” 解释: 移除掉三个数字 4, 3, 和 2形成一个新的最小的数...
-
代码展示： http://paste.ubuntu.com/23693598/ #include
#include #include char * largeDiffer(char *a,char *b){ /* 使用说明 传入的a和b只能为整数 结果为a-b;返回... -
Description We all know that Bin-Laden is a notorious terrorist, and he has disappeared for a long time. But recently, it is reported that he hides in Hang Zhou of Ch...
-
/*Name: NYOJ--811--变态最大值Author: shen_渊 Date: 17/04/17 15:49Description: 看到博客上这道题浏览量最高，原来的代码就看不下去了 o(╯□╰)o */#include
#include #include u... -
生成唯一号：思路，根据yymmddhhmmss+自增长号+唯一服务器号( SystemNo)生成唯一码，总长度19，例如：1509281204550000101. public class UniqueNumber { private static long num = 0;//流水号 private sta...
-
如果你想从头学习Jmeter，可以看看这个系列的文章哦Charts 介绍包含了各种详细信息图表，比 GUI 模式的图表好看且易懂多了！做性能测试，如何发现是否有性能瓶颈？必须从结果图表中找到鸭！而 html 报告将性能测试可能需要用到的图表都加进去了，可谓是6666一共有三大模块Over TimeThroughputResponse...
-
控制（Controls） 1.PID控制简介 在工程实际中，应用最为广泛的调节器控制规律为比例、积分、微分控制，简称PID控制，又称PID调节。PID控制器问世至今已有近70年历史，它 以其结构简单、稳定性好、工作可靠、调整方便而成为工业控制的主要技术之一。当被控对象的结构和参数不能完全掌握，或得不到精确的数学模型时，控制理论的...
-
搞了很多年c/c++，有很多细小的东西，曾经不止一次遇到，可是一直都是放在零散的地方,要用的时候怎么也找不到，今天，我痛下决心，改掉不良习惯，把这些经验或是tips记录在这里，便于日后查找。 1.在统计网络下载信息时，如何表达文件大小？ 下面是输出结果 2.打印size_t类型数据的长度，使用%lu。 下面是一个使...
-
1,解决的问题。 2.如何实现。 面对大流量网站频繁访问数据库的一种优化，比如博客网站。不可能每个人查看都访问一次数据库。为了解决大量不必要访问的问题。 可以把第一次的内容保存为html页面。再以后定义的过期时间内都访问该静态页面。 以下是一个小的demo index.php来实现静态化的主要工作。 1
-
什么？有个 SQL 执行了 8 秒！ 哪里出了问题？臣妾不知道啊，得找 DBA 啊。 DBA 人呢？离职了！！擦！！！ 程序员在无处寻求帮助时，就得想办法自救，努力让自己变成 "伪 DBA"。 索引 按页编号查看数据表信息获取查询 SELECT 语句的执行次数排名看看哪些 Ad-hoc Query 在浪费资源查看当前处于等待状态的 T...