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

前言

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

多执行绪的设定

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

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

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

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

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

程式撰写

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

定义阵列含 1024 个元素。

#define N1024

使用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];}

设定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,...,Nh_b[i] = i * 2; // 2,...,2xN}

主程式呼叫 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);

测试

如果更改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;

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

结语

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

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

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

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