cuda實現向量相加

博客最后附上整體代碼

如果有說的不對的地方還請前輩指出, 因為cuda真的接觸沒幾天

一些總結(建議看)

1. cuda 並不純GPU在運行程序, 而是 cpu 與 gpu 一起在運行程序, cpu負責調度, gpu 負責運算, cpu稱為**HOST **, gpu 稱為 DEVICE
2. 記住三個東西 grid block thread ,關系分別是 grid 包含多個 block , block 包含多個 thread
3. 一個block中thread個數選取一般為32的整數倍, 原因和warp有關, 有興趣自行查閱
4. 一個grid中block的個數選取和你的kernel函數以及thread數量有關, 舉個例子, int a[1000] 加上 int b[1000] , 你的thread為64, 那么, block = 1000/64 = 16個合適
5. __global__函數一般表示一個內核函數，是一組由GPU執行的並行計算任務,由cpu調用
6. __host__一般是由CPU調用，由CPU執行的函數,
7. __device__一般表示由GPU中一個線程調用的函數

代碼實現

引入

#include <stdio.h>
#include <cuda_runtime.h>

kernel函數

__global__ void
vectorAdd(float *a, float *b, float *c, int num){
        int i = blockDim.x * blockIdx.x + threadIdx.x; //vector is 1-dim, blockDim means the number of thread in a block
        if(i < num){
                c[i] = a[i] + b[i];
        }
}

int i = blockDim.x * blockIdx.x + threadIdx.x;

這句代碼解釋一下:

blockDim.x 表示block的size行數(如果是一維的block的話,即一行有多少個thread)

blockIdx.x 表示當前運行到的第幾個block(一維grid的話,即該grid中第幾個block)

threadIdx.x 表示當前運行到的第幾個thread (一維的block的話.即該block中第幾個thread)

畫個圖解釋一下

比如上面這個圖的話, ABCDE各代表一個block, 總的為一個Grid, 每個block中有四個thread, 圖中我花了箭頭的也就是代表着第1個block中的第0個thread.

那么 i = blockDim.x * blockIdx.x + threadIdx.x 就是指 i = 4 * 1 + 0

申請內存空間與釋放

host中申請內存

float *a = (float *)malloc(size);
float *b = (float *)malloc(size);
float *c = (float *)malloc(size);

free(a);
free(b);
free(c);

device中申請內存

float *da = NULL;
float *db = NULL;
float *dc = NULL;

cudaMalloc((void **)&da, size);
cudaMalloc((void **)&db, size);
cudaMalloc((void **)&dc, size);

cudaFree(da);
cudaFree(db);
cudaFree(dc);

host中內存copy到device

cudaMemcpy(da,a,size,cudaMemcpyHostToDevice);
cudaMemcpy(db,b,size,cudaMemcpyHostToDevice);
cudaMemcpy(dc,c,size,cudaMemcpyHostToDevice);

上面的cudaMemcpyHostToDevice用於指定方向有四種關鍵詞

cudaMemcpyHostToDevice | cudaMemcpyHostToHost | cudaMemcpyDeviceToDevice | cudaMemcpyDeviceToHost

啟動 kernel函數

int threadPerBlock = 256;                        
int blockPerGrid = (num + threadPerBlock - 1)/threadPerBlock;
vectorAdd <<< blockPerGrid, threadPerBlock >>> (da,db,dc,num)

此處確定了block中的thread數量以及一個grid中block數量

利用kernel function <<< blockPerGrid, threadPerBlock>>> (paras,...) 來實現在cuda中運算

參考

https://zhuanlan.zhihu.com/p/345877391

https://docs.nvidia.com/cuda/cuda-c-programming-guide/

源碼展示

#include <stdio.h>

#include <cuda_runtime.h>

// vectorAdd run in device
__global__ void 
vectorAdd(float *a, float *b, float *c, int num){
	int i = blockDim.x * blockIdx.x + threadIdx.x; //vector is 1-dim, blockDim means the number of thread in a block
	if(i < num){
		c[i] = a[i] + b[i];
	}
}

// main run in host
int
main(void){
	int num = 10000; // size of vector
	size_t size = num * sizeof(float);

	// host memery
	float *a = (float *)malloc(size);
	float *b = (float *)malloc(size);
	float *c = (float *)malloc(size);

	// init the vector
	for(int i=1;i<num;++i){
		a[i] = rand()/(float)RAND_MAX;
		b[i] = rand()/(float)RAND_MAX;
	}

	// copy the host memery to device memery
	float *da = NULL;
	float *db = NULL;
	float *dc = NULL;

	cudaMalloc((void **)&da, size);
	cudaMalloc((void **)&db, size);
	cudaMalloc((void **)&dc, size);

	cudaMemcpy(da,a,size,cudaMemcpyHostToDevice);
	cudaMemcpy(db,b,size,cudaMemcpyHostToDevice);
	cudaMemcpy(dc,c,size,cudaMemcpyHostToDevice);

	// launch function add kernel
	int threadPerBlock = 256;
	int blockPerGrid = (num + threadPerBlock - 1)/threadPerBlock;
	printf("threadPerBlock: %d \nblockPerGrid: %d \n",threadPerBlock,blockPerGrid);

	vectorAdd <<< blockPerGrid, threadPerBlock >>> (da,db,dc,num);

	//copy the device result to host
	cudaMemcpy(c,dc,size,cudaMemcpyDeviceToHost);

	// Verify that the result vector is correct
	for (int i = 0; i < num; ++i){
		if (fabs(a[i] + b[i] - c[i]) > 1e-5){
			fprintf(stderr, "Result verification failed at element %d!\n", i);
			return 0;
		}
	}

	printf("Test PASSED\n");

	// Free device global memory
	cudaFree(da);
	cudaFree(db);
	cudaFree(dc);
	// Free host memory
	free(a);
	free(b);
	free(c);

	printf("free is ok\n");
	return 0;
}