CUDA

CUDA-在Device上初始化

在GPU上初始化

绝大多数CUDA教材，都会提出一般CUDA程序的标准步骤，其中一步是将Host的数据拷贝到Device。也就是说Host中数据要先被初始化，后才能拷贝到Device。其实，可以直接在Device中初始化数据，如此既可以避免Host到Device的数据传输，又可以加快数据初始化的速度。如下是一个小例：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <time.h>

#define N 10
//#define N 16384 // N*N = 268435456   // 1GB 4Bytes 
//#define N 16384*2 // N*N = 268435456*4   // 4GB 4Bytes 
//#define N 16384*4 // N*N = 268435456*16   // 16GB 4Bytes 

__global__ void initOnGPU(int* dev_a){
	int ix = blockIdx.x * blockDim.x + threadIdx.x;
	int iy = blockIdx.y * blockDim.y + threadIdx.y;
	int tid = ix * N + iy;

	if (ix < N && iy < N){
        // initialize on Device
		dev_a[tid] = 321;
	}
}

int main(){

    // 1) 
	int* h_a;
	h_a = (int*)malloc(N*N*sizeof(int));

	int* dev_a;
	cudaMalloc((int**)&dev_a, N*N*sizeof(int));


	// 2) init on gpu
	dim3 block(2, 2);
	dim3 grid((N + block.x - 1), (N + block.y - 1));

	clock_t start = clock();
	initOnGPU <<< grid, block >>>(dev_a);
	clock_t end = clock();
	printf("Time init on GPU: %.10f \n", (double)(end - start) / CLOCKS_PER_SEC);
	
    // 3) 
	cudaMemcpy(h_a, dev_a, N*N*sizeof(int), cudaMemcpyDeviceToHost);

	// 4) show h_a
	for (int i = 0; i < N*N; i++){
		printf("%d ", h_a[i]);
		if ((i + 1) % N == 0)
			printf("\n");
	}
	printf("\n");

    // 5) 
	free(h_a);
	cudaFree(dev_a);

	return 0;
}

解释一下：
1）分别在Host和Device上开辟空间h_a和dev_a。
2）调用kernel，这个kernel的作用是初始化dev_a。注意并不是从Host中拷贝过去的。以及其他需要在GPU上执行的操作。
3）GPU上的操作完成后，得到的最终数据保存在dev_a中。拷贝到Hosth_a中。
4）显示最终计算结果。
5）释放资源。

其中kernel函数initOnGPU的作用是初始化dec_a中元素为321。

结果如下：

1
2
3
4
5
6
7
8
9
10
11
Time init on GPU: 0.0000160000 
321 321 321 321 321 321 321 321 321 321 
321 321 321 321 321 321 321 321 321 321 
321 321 321 321 321 321 321 321 321 321 
321 321 321 321 321 321 321 321 321 321 
321 321 321 321 321 321 321 321 321 321 
321 321 321 321 321 321 321 321 321 321 
321 321 321 321 321 321 321 321 321 321 
321 321 321 321 321 321 321 321 321 321 
321 321 321 321 321 321 321 321 321 321 
321 321 321 321 321 321 321 321 321 321

所以说，直接在Device上初始化，相对先在Host上初始化后拷贝到Device，是更高效的。
CPU与GPU的通信是通过PCIe 实现的。PCIe第三代 的极限速度是 16GB每秒。相比较，费米架构的GPU中，GPU芯片与GPU存储之间的数据交换速度高达144GB 每秒。所以说Host和Device间的数据传输是CUDA应用的 性能瓶颈。

敲黑板 CUDA程序的一个基本规则是，尽可能减少host 与device间的数据交换。