cuda《學(xué)習(xí)筆記二》——基本用法

2019-11-14 11:37:43

字體：大中小

供稿：網(wǎng)友

一、前言

本文對(duì)使用cuda加速一維數(shù)組運(yùn)算、二維圖像處理運(yùn)算的方法作總結(jié)，參考自《CUDA By Example》

二、一維數(shù)組并行運(yùn)算

經(jīng)過查詢，本人的老顯卡GT550M，可得其3維線程格，每維包含線程塊數(shù)量為（65536,65536,65536），相應(yīng)的每維包含線程數(shù)為（1024,1024,64），故可得知線程格的每一維可開啟的線程塊最大數(shù)均為65536，相應(yīng)線程的最大數(shù)為（1024,1024,64）。

1）以下為對(duì)一維數(shù)組并行運(yùn)算的例子，使用了3種基本方法

//main.cu#include <cuda_runtime.h>  #include <iostream>  #include "book.h"    //該頭文件定義了HANDLE_ERROR函數(shù)using namespace std;int mode = 0 ;#define N 1024void select(){	cout << "select the calculation mode:/n"		<< "1 -> only by blockIdx/n"		<< "2 -> only by threadIdx/n"		<< "3 -> by blockIdx and threadIdx /n";	cout << "mode =  ";	cin >> mode;	cout << endl;}__global__ void add_only_blockIdx(int *a, int *b, int *c)       //__global表示該函數(shù)可在主機(jī)調(diào)用，在器件執(zhí)行；{                                                 // 另外__device__表示該函數(shù)在器件調(diào)用，在期間執(zhí)行	int idx = blockIdx.x;     //計(jì)算位于線程塊索引處的數(shù)據(jù)	if (idx < N){		c[idx] = a[idx] + b[idx];	}}__global__ void add_only_threadIdx(int *a, int *b, int *c)       {                                                 	int idx = threadIdx.x;     //計(jì)算位于線程索引處的數(shù)據(jù)	if (idx < N){		c[idx] = a[idx] + b[idx];	}}__global__ void add_blockIdx_threadIdx(int *a, int *b, int *c)      {                                               	int idx = threadIdx.x +  blockIdx.x * blockDim.x;    //將線程塊索引與線程索引轉(zhuǎn)為線性	if (idx < N){		c[idx] = a[idx] + b[idx];	}}int main(){	while ( mode< 1 || mode >3 ){		select();	}	int a[N], b[N], c[N];	int *dev_a, *dev_b, *dev_c;	//gpu上分配內(nèi)存	HANDLE_ERROR(cudaMalloc((void**)&dev_a, N*sizeof(int)));	HANDLE_ERROR(cudaMalloc((void**)&dev_b, N*sizeof(int)));	HANDLE_ERROR(cudaMalloc((void**)&dev_c, N*sizeof(int)));	//為數(shù)組a,b初始化	for (int i = 0; i < N; ++i){		a[i] = i;		b[i] = i;	}	//講數(shù)組a,b數(shù)據(jù)復(fù)制至gpu	(cudaMemcpy(dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice));	(cudaMemcpy(dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice));	if (mode == 1)		add_only_blockIdx << < N, 1 >> >(dev_a, dev_b, dev_c);  //使用N個(gè)線程塊，其中每個(gè)線程塊使用1個(gè)線程,這里N小于65536，故可行	else if (mode == 2)		add_only_threadIdx << < 1, N >> >(dev_a, dev_b, dev_c);  //使用1個(gè)線程塊，其中該線程塊使用1024個(gè)線程,這里N小于1024，故可行	else		add_blockIdx_threadIdx << < (N+127)/128, 128 >> >(dev_a, dev_b, dev_c);  // 這里使用(N+127)/128個(gè)線程塊，每個(gè)線程塊128線程，為了	                                                                                 // 使所 (N+127)/128 * 128 ＞= N         	 	//將數(shù)組dev_c復(fù)制至cpu	HANDLE_ERROR(cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost));	//顯示結(jié)果	for (int i = 0; i < N; ++i){		PRintf("%d + %d = %d/n", a[i], b[i], c[i]);	}	//釋放在gpu分配的內(nèi)存	cudaFree(dev_a);	cudaFree(dev_b);	cudaFree(dev_c);	return 0;}
2）對(duì)任意長(zhǎng)度的一維數(shù)組并行求和運(yùn)算
       當(dāng)一維數(shù)組長(zhǎng)度大于65536*1024= 67108864 時(shí)，使用(1)的方法則不可行，因?yàn)榫€程格的每維最大線程塊不大于65536，且每維的線程不大于1024（第3維除外），故可使用如下方法。
//main2.cu#include <cuda_runtime.h>  #include <iostream>  #include "book.h"    //該頭文件定義了HANDLE_ERROR函數(shù)using namespace std;#define N 33*1024  //在gpu運(yùn)算的數(shù)組長(zhǎng)度受GPU內(nèi)存的限制，故這里使用33*1024表示長(zhǎng)度較大的數(shù)組__global__ void add(int *a, int *b, int *c)       {                                                	int idx = threadIdx.x + blockIdx.x * blockDim.x ;    //所計(jì)算的索引為 128個(gè)線程塊中的128個(gè)線程索引	                                                    // 由于128*128 < 33*1024,故某些線程需執(zhí)行多次運(yùn)算	while (idx < N){		c[idx] = a[idx] + b[idx]; 		idx += gridDim.x * blockDim.x;           //對(duì)索引進(jìn)行遞增，遞增步長(zhǎng)為gridDim.x * blockDim.x，gridDim.x為使用的線程塊總數(shù)，	}                                                //blockDim.x為使用的每個(gè)線程塊中的線程總數(shù)}int main(){	int a[N], b[N], c[N];	int *dev_a, *dev_b, *dev_c;	//gpu上分配內(nèi)存	HANDLE_ERROR(cudaMalloc((void**)&dev_a, N*sizeof(int)));	HANDLE_ERROR(cudaMalloc((void**)&dev_b, N*sizeof(int)));	HANDLE_ERROR(cudaMalloc((void**)&dev_c, N*sizeof(int)));	//為數(shù)組a,b初始化	for (int i = 0; i < N; ++i){		a[i] = i;		b[i] = i;	}	//講數(shù)組a,b數(shù)據(jù)復(fù)制至gpu	(cudaMemcpy(dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice));	(cudaMemcpy(dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice));	add <<< 128, 128 >>>(dev_a, dev_b, dev_c);  //使用128個(gè)線程塊，每個(gè)線程塊使用128個(gè)線程,128*128 < 33*1024 	    	//將數(shù)組dev_c復(fù)制至cpu	HANDLE_ERROR(cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost));	//顯示結(jié)果	for (int i = 0; i < N; ++i){		printf("%d + %d = %d/n", a[i], b[i], c[i]);	}	//釋放在gpu分配的內(nèi)存	cudaFree(dev_a);	cudaFree(dev_b);	cudaFree(dev_c);	return 0;}三、二維圖像處理并行運(yùn)算       對(duì)二維圖像數(shù)據(jù)進(jìn)行并行運(yùn)算，這里使用經(jīng)典的Reduce Color例子，其對(duì)圖像中的每個(gè)像素點(diǎn)進(jìn)行量化，如常見的RGB24圖像有256×256×256中顏色，通過Reduce Color將每個(gè)通道的像素減少8倍至256/8=32種，則圖像只有32×32×32種顏色。假設(shè)量化減少的倍數(shù)是N，則代碼實(shí)現(xiàn)時(shí)就是簡(jiǎn)單的value/N*N，通常我們會(huì)再加上N/2以得到相鄰的N的倍數(shù)的中間值。
//main3.cu#include <cuda_runtime.h>  #include <iostream>  #include "book.h"    //該頭文件定義了HANDLE_ERROR函數(shù)#include <opencv2/opencv.hpp>using namespace std;using namespace cv;__global__ void quantify(uchar *_src_dev, uchar *_dst_dev , int div){	int x_idx = threadIdx.x + blockIdx.x*blockDim.x;           //計(jì)算x坐標(biāo)索引	int y_idx = threadIdx.y + blockIdx.y*blockDim.y;           //計(jì)算y坐標(biāo)索引	int idx = x_idx + y_idx * gridDim.x * blockDim.x;        //將(x,y)坐標(biāo)轉(zhuǎn)換為線性	_dst_dev[idx ] =  _src_dev[idx] / div*div + div / 2;}int main(){	Mat src = imread("lena.bmp" );   //讀取圖像	Mat dst (src.size(), CV_8UC3);  		uchar *src_data = src.data;          // src_data指針指向圖像src的數(shù)據(jù)	if (!src.isContinuous())  return -1;   //判斷圖像數(shù)據(jù)字節(jié)是否填充	uchar *src_dev, *dst_dev;	int length = src.rows * src.cols * src.channels();	//在gpu上分配內(nèi)存空間	HANDLE_ERROR(cudaMalloc((void**)&src_dev, length*sizeof(uchar)));	HANDLE_ERROR(cudaMalloc((void**)&dst_dev, length*sizeof(uchar)));	HANDLE_ERROR(cudaMemcpy(src_dev, src_data, length*sizeof(uchar), cudaMemcpyHostToDevice));	dim3 blocks(src.cols*src.channels()/ 32, src.rows/ 32);  //寬和高均為512,所啟動(dòng)線程塊與線程數(shù)量差距不可過大	dim3 threads(32, 32);	quantify << < blocks, threads >> >(src_dev, dst_dev , 64);	//將數(shù)組dst_dev復(fù)制至cpu	HANDLE_ERROR(cudaMemcpy(dst.data, dst_dev, length*sizeof(uchar), cudaMemcpyDeviceToHost));	cudaFree(src_dev);	cudaFree(dst_dev);	imshow("src", src);	imshow("Dst", dst);	waitKey(0);	return 0;}運(yùn)行結(jié)果