2. 程式人生 > >CUDA之單thread單block&多thread單block&多thread多block

CUDA之單thread單block&多thread單block&多thread多block

阿新 • • 發佈：2018-11-15

用簡單的立方和歸約來舉例：

//單thread單block
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#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.\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 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 \n", sum);
int sum = 0;
for (int i = 0; i < DATA_SIZE; i++) {
sum += data[i] * data[i] * data[i];
}
printf("CPUsum: %d \n", sum);
getchar();
return 0;
}
 
//單block多thread

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
//CUDA RunTime API
#include <cuda_runtime.h>
#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);
	//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\n time:%d\n", 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 \n", final_sum);
	getchar();
	return 0;
}

 
//多block多thread

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
//CUDA RunTime API
#include <cuda_runtime.h>
#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 中執行函式語法：函式名稱<<<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);
	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\n", 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 \n", 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 <stdio.h>
#include <stdlib.h>
#include <time.h>
//CUDA RunTime API
#include <cuda_runtime.h>
#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 中執行函式語法：函式名稱<<<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);
	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\n", 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 \n", final_sum);
	getchar();
	return 0;
}

Block內完成部分加和工作，所以gputime增加了

//多block多thread

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
//CUDA RunTime API
#include <cuda_runtime.h>
#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 中執行函式語法：函式名稱<<<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);
	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\n", 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 \n", final_sum);
	getchar();
	return 0;
}

 

相關推薦

CUDA之單thread單block&amp;多thread單block&amp;多thread多block

用簡單的立方和歸約來舉例： //單thread單block #include <stdio.h> #include <stdlib.h> #include <cuda_runtime.h> #define DATA_SIZE 1048576 int data[

CUDA之單thread單block&多thread單block&多thread多block

用簡單的立方和歸約來舉例： //單thread單block #include <stdio.h> #include <stdlib.h> #include <cuda_runtime.h> #define DATA_SIZE 104857

系列之5-多入單出的一層神經網路能做什麼

全套教程請點選：微軟 AI 開發教程 知識點 多元線性迴歸 輸入層的矩陣運算設計 樣本資料歸一化 提出問題 假設有如下這樣一個問題：有1000個樣本，每個樣本有三個特徵值，一個標籤值。想讓你預測一下給定任意三個特徵值組合，其y值是多少？ 樣本序號 1

多程序單執行緒模型與單程序多執行緒模型之爭

伺服器，事件 多程序單執行緒模型典型代表：nginx 單程序多執行緒模型典型代表：memcached 另外redis, mongodb也可以說是走的“多程序單執行緒模”模型（叢集），只不過作為資料庫伺服器，需要進行防寫，只提供了讀同步。 原因很簡單，因為伺服器的發展大部分都

JS實例之滾動固定效果，實現頂部菜單固定效果。

rip bsp 實現 con tel body text 滾動 title 1 <head> 2 <meta http-equiv="Content-Type" content="text/html; charset=utf-8" /> 3

selenium測試（Java）--多表單切換（十二）

generate ati nqa void char support 編寫 generated key 采用下面的例子來編寫用例 <!DOCTYPE html> <html> <head> <meta http-equiv="co

學習筆記之05表格嵌套2（表單）

html 登錄 res cells htm bgcolor www radi charset <!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xh

使用foreach一次性添加多個單選按鈕

gallon ech 方法 單選按鈕 pri function 調用 cti down <?php //這個函數創建一個單選按鈕 //函數接收一個論據：值 //函數也會讓按鈕有“黏性”-->記住你是如何填寫它的。 functi

python並發編程之多進程(二)：互斥鎖（同步鎖）&amp;進程其他屬性&amp;進程間通信（queue）&amp;生產者消費者模型

互斥 數據 socket pan copy src too 如果 搶票 一，互斥鎖，同步鎖 進程之間數據不共享,但是共享同一套文件系統,所以訪問同一個文件,或同一個打印終端,是沒有問題的, 競爭帶來的結果就是錯亂，如何控制，就是加鎖處理 part1：多個進程共享同

bootstrap布局兩列或者多列表單

復制 AD BE 結束 body too ots 時間 IT 1, 代碼如下: <div class="form-group"> <labe

IIS https綁定主機頭方法（可實現禁止直接通過IP訪問，好像還可以實現IIS中綁定多個https<是使用多個單域名證書，非多域名證書>）

多個 nbsp app for iis 直接 fig protoc eight 步驟： 打開：%systemroot%\SYSTEM32\inetsrv\config\ 打開：applicationHost.config 查找：<binding protocol="

20180518VSTO多簿單表匯總

cte datetime tar ribbon nco book ont lap names using System; using System.Collections.Generic; using System.Linq; using System.Text; usi

20180518VSTO多簿單表匯總外接程序按鈕

dialog long ear picker clas sin PE forms tools using System; using System.Collections.Generic; using System.Linq; using System.Text; usi

4.9 多表單切換

b- 百度 class 可用 一個 fir switch AS oot 在 web 應用中經常會遇到 frame 嵌套頁面的應用，頁 WebDriver 每次只能在一個頁面上識別元素，對於 frame 嵌套內的頁面上的元素，直接定位是定位是定位不到的。這個時候就需要通過 s

Selenium 多表單(frame/iframe)切換

AR mes click com 行操作 多層 key 元素定位 info frame標簽有frameset、frame、iframe三種,frameset跟其他普通標簽沒有區別，不會影響到正常的定位，而frame與iframe需要切換進去才能定位到其中的元素 比如下面這個

繼承,多繼承,單繼承

self per int 深度優先 div 都是 sel nim ima 繼承子類以及子類實例化的對象可以父類的任何方法或變量類名可以訪問父類所有內容 class Animal: def __init__(self,name,sex,age): s

webpack入坑之旅（五）加載vue單文件組件

file loader clu html中 type 修改 mode hid extension 需要什麽？ 在經過前面的四個練習，相信已經對於webapck有了一定的了解，現在我們就來一個綜合案例，進一步加深對於webpack的理解。 首先我們應該思考要解析.vue類型的

html-4, form 表單 輸入、傳文件、單選、多選、下拉菜單、文本描述、重置、submit、按鈕限制輸入

part row man head 密碼 文本 inpu set 跳轉 <!-- form HTTP協議 action:提交的服務器網址 method:get(默認)| post(應用：

67 orm13查詢語句，外鍵查詢，多對多查詢 單表查詢雙下劃線方法

ever tle 基礎 lte ron -s 執行流程 first 什麽 主要內容：https://www.cnblogs.com/maple-shaw/articles/9403501.html 註意: 如果想要文件中運行整個django項目: 　　需要在文件中寫入:

selenium多表單操作與多窗口，以及警告框處理

itl 很快 hand phone from close 註冊頁面 需要 http 知識是需要經常溫習的，不然是很容易遺忘的。 以前自己操作IFRAME，多窗口的時候，覺得很簡單。半年沒有操作自動化了，知識又還了回去。 寫博客有一個好處，可以把自己記住的知識點記錄下來，這樣