GPU程序设计(2) -- 多执行绪

前言

GPU可以利用平行处理的方式，缩短执行时间，因此，这一次就来介绍多执行绪的程序设计方法及应用。

多执行绪的设定

上一篇 介绍GPU函数的设定是透过<<<...>>>指定三个参数：

1. 区块(Block)数量。
2. 执行绪(Thread)数量：一个区块要执行的执行绪数量。
3. 使用的共享记忆体(Shared memory)大小，此参数可不设定。

每一种GPU的数量均不相同，因此，我们先要侦测相关的数量限制，可以执行上一篇的deviceQuery.exe，预设建置的目录在【v11.2\bin\win64\Debug】。

1. deviceProp.maxThreadsPerMultiProcessor：一个处理器(Multiprocessor)的最大执行绪数量，笔者的GPU卡是【GTX1050 Ti】，数量是2048。
2. deviceProp.maxThreadsPerBlock：笔者的GPU卡区块的最大执行绪数量是1024。
3. deviceProp.sharedMemPerMultiprocessor：一个处理器(Multiprocessor)的最大共享记忆体大小，笔者的GPU卡是【98304 bytes】。
4. deviceProp.sharedMemPerBlock：一个区块的最大共享记忆体大小，笔者的GPU卡是【49152 bytes】。

所以，如果要在不同的卡都能执行，必须在程序中计算可设定的执行绪数量，总数量超过程序就会发生异常。共享记忆体也是一样的道理。【GTX1050 Ti】含6个处理器，我在以下的程序设定7*2048个执行绪也不会发生异常，我猜处理器有类似queue的功能，一旦执行绪超过容量限制，工作会排队等待，直到处理器有空时才被处理。

程序撰写

接下来，我们使用多执行绪进行向量的相加，一维阵列中的每个元素都独立计算，每个元素计算交由一个执行绪处理。

1. 定义阵列含 1024 个元素。

#define N	1024

2. 使用GPU进行向量相加：threadIdx.x 表执行绪序号。

__global__ void vectorAdd(int* d_a, int* d_b, int* d_c) {
	// 使用第几个执行绪计算 
	int tid = threadIdx.x;
	//printf("threadIdx = %d\n", tid);
	d_c[tid] = d_a[tid] + d_b[tid];
}

3. 设定CPU阵列(h_a、h_b)值。

    int h_a[N], h_b[N], h_c[N], h2_c[N];

	// 设定 a、b 的值
	for (int i = 0; i < N; i++) {
		h_a[i] = i; // 1,...,N
		h_b[i] = i * 2; // 2,...,2xN
	}

4. 主程序呼叫 vectorAdd 函数，采用 1 block x N threads，需先分配记忆体给GPU阵列(d_a、d_b)，再复制CPU阵列(h_a、h_b)值至GPU阵列(d_a、d_b)。

	// 分配记忆体
	cudaMalloc((void**)&d_a, N * sizeof(int));
	// 自 CPU变数 复制到 GPU变数
	cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
	cudaMalloc((void**)&d_b, N * sizeof(int));
	cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);

	// 分配记忆体
	cudaMalloc((void**) &d_c, N * sizeof(int));

    // 使用GPU进行向量相加，1 block x N threads
    vectorAdd <<<1, N >>> (d_a, d_b, d_c);
    
    // 等待所有执行绪处理完毕
	cudaDeviceSynchronize();

• cudaMemcpy 参数：第一个为目标变数，第二个为来源变数，cudaMemcpyHostToDevice为CPU变数复制到GPU变数。CPU/GPU 变数不可混合运算，需先将变数全部转移至GPU，才能使用GPU运算，CPU也是一样。其他注意事项如下：
• 要列印变数，必须将GPU变数复制到CPU变数，才可正确显示。
• vectorAdd 会被呼叫 1024 次，各由1个执行绪负责。
• 记得呼叫 cudaDeviceSynchronize()，等待所有执行绪处理完毕。
• GPU变数需在程序结束前释放记忆体。

	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);

测试

1. 如果更改N为1025，则GPU计算结果均为0，因为，笔者的GPU卡是【GTX1050 Ti】，区块的最大执行绪数量是1024。可以改成2个区块(2x1024)：

vectorAdd <<<2, 1024 >>> (d_a, d_b, d_c);

vectorAdd 函数的阵列索引值要改为：

// blockIdx.x：第几个block，blockDim.x：block内的执行绪数量
int tid = threadIdx.x + blockIdx.x * blockDim.x;

2. 使用clock()计算执行时间，CPU计算时间为0秒，GPU计算时间为0.188秒，含分配记忆体及CPU变数复制到GPU变数，若只考虑计算，则是0秒。

结语

在 GPU 撰写多执行绪真是方便，只要呼叫时，指定个数即可，但是，要以程序检查是否超出硬体限制条件，这部分可自deviceQuery截取相关程序码，计算应使用的区块数及执行绪数量。

平行处理要选择多区块或多执行绪呢? 使用多执行绪有两个好处：

1. 资料共享：可透过共享记忆体，多执行绪可读写同一块记忆体。
2. 同步：可以等非同步的多执行绪全部结束在进行下一段处理。
   我们会在下一篇实际运用这两个特性，可提升运算效能。

完整程序放在 『GitHub』的VectorAdd目录。使用多个区块的改良程序放在VectorAdd_Improved目录。