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

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

阿新 • • 發佈：2018-12-20

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

//單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&多thread單block&多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

CUDA之Thread、Wrap執行詳解

從硬體角度分析，支援CUDA的NVIDIA 顯示卡，都是由多個multiprocessors 組成。每個 multiprocessor 裡包含了8個stream processors，其組成是四個四個一組，也就是兩組4D的處理器。每個 multiprocessor 還具有 很多個（比如8192個）暫

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

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

Python實戰之多執行緒程式設計threading Thread

                在Python中可以使用繼承threading.Thread類來實現多執行緒程式設計，其中子類可以重寫父類的__init__和run方法來實現使用者執行緒的邏輯，如下是一個簡單的多執行緒類實現[python] view plain copy print?import threa

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

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

每天一個JS 小demo之通過事件委托實現菜單展開及選中特效。主要知識點：事件

實現 tween efi blog dem int position 16px eas <!DOCTYPE html><html lang="en"><head><meta charset="UTF-8"><titl

Java多線程系列二——Thread類的方法

index 表示 stop unit throws catch atomic while val Thread實現Runnable接口並實現了大量實用的方法 public static native void yield(); 此方法釋放CPU，但並不釋放已獲得的鎖，其它就

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

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

Unity3D編輯器之重寫Hierarchy的右鍵菜單

col art n) 條件 any rect obj lan dga 1 using UnityEngine; 2 using UnityEditor; 3 using System.Collections; 4 5 public class MyHiera

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

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

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="

Java多線程 5：Thread中的實例方法

守護 nds 屬性 exc pre 在線 結束 正在 ini Thread類中的方法調用方式：快速到底 學習 Thread 類中的方法是學習多線程的第一步。在學習多線程之前特別提出一點，調用 Thread 中的方法的時候，在線程類中（千萬別忘記了這個前提條件），有兩種方式

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需要切換進去才能定位到其中的元素 比如下面這個

用單進程、多線程並發、多線程分別實現爬一個或多個網站的所有鏈接，用瀏覽器打開所有鏈接並保存截圖 python

app imp mat 並發執行 cut h+ chrome 鏈接 目錄 #coding=utf-8import requestsimport re,os,time,ConfigParserfrom selenium import webdriverfrom multipr

繼承,多繼承,單繼承

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