W

cuda学习笔记3

 1366  3 分钟

本文介绍如何在GPU上实现多线程块中开多个线程处理任务。

单线程块里开启多个线程

在代码实现上，和开启多个线程块，每个线程块只有一个线程的区别：

  • 将 add<<<N,1>>>(dev_a,dev_b,dev_c) 改为 add<<<1,N>>>(dev_a,dev_b,dev_c) 。

容易看出 add<<<N,1>>> 的第一个参数是开启的线程块的数目，第二个参数是一个线程块里的线程数。

  • 将 int tid = blockIdx.x; 改为 int tid = threadIdx.x;

代码：

 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
58
59
60
61

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <math.h>
#include <helper_cuda.h>
using namespace std;

#define N 100


__global__ void add_kernel(double *a, double *b, double *c) {
	int tid = threadIdx.x;
	if (tid < N)
	{
		c[tid] = a[tid] + b[tid];
	}

}

__global__ void value_init_kernel(double *a, double *b) {
	int tid = threadIdx.x;
	if (tid < N)
	{
		a[tid] = 1.0*tid;
		b[tid] = (1.0*tid*tid);
	}

}

int main(void)
{
	cudaError_t err1 = cudaSuccess, err2 = cudaSuccess, err3 = cudaSuccess;
	double a[N], b[N], c[N];
	double *dev_a, *dev_b, *dev_c;

	HANDLE_ERROR(cudaMalloc((void**)&dev_a, N * sizeof(double)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b, N * sizeof(double)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c, N * sizeof(double)));

	value_init_kernel << <1, N >> > (dev_a, dev_b);////在GPU上赋值操作
	add_kernel << <1, N >> > (dev_a, dev_b, dev_c);////在GPU上相加操作

	HANDLE_ERROR(cudaMemcpy(a, dev_a, N * sizeof(double), cudaMemcpyDeviceToHost));
	HANDLE_ERROR(cudaMemcpy(b, dev_b, N * sizeof(double), cudaMemcpyDeviceToHost));
	HANDLE_ERROR(cudaMemcpy(c, dev_c, N * sizeof(double), cudaMemcpyDeviceToHost));

	for (int i = 0; i < N; i++)
	{
		printf("%f + %f = %f\n", a[i], b[i], c[i]);
	}

	////释放GPU内存
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);


	return 0;

}

多个线程块开启多个（固定数量的）线程

和上文第一种情况相比，依然只需要改变核函数的索引计算方法和核函数的调用方式即可：

  • 核函数的索引计算方法：

将：

1

int tid = threadIdx.x;

改为：

1

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

注意的是， blockDim 其实是一个三维的线程数组，但在例子中只用到一维线程块，故只需要用到 blockIdx.x ，该变量是一个常数，保存的是线程块中每一维的线程数量。索引 tid 的值显然为：当前线程块中的线程索引 + 每一个线程块中线程数（这里只有一维） * 当前线程块索引（从0开始）。

  • 核函数的调用方式： 将：
1

add_kernel << <1, N >> >

改为：

1

add_kernel << <(N + 127) / 128, 128 >> >

由于 N 是一个整型，N+127/128 是为了向上取整（可多不可少）。

在GPU上实现任意长度的矢量求和

由于单线程格里的线程块有一个最大值（比如65535），当GPU处理的矢量数目超过，在情况2中当 (N + 127) / 128 超过65535时，函数就会调用失败，这时候可以这样处理：将调用的线程块数目和每个线程块中调用线程的数目都固定，比如都固定为128。然后利用这个线程格中的 128 * 128 个线程顺序处理GPU上的所有任务。比如：索引号为1的线程处理完id为 1 的任务后会跳转到id为 1+ 128 * 128 的任务进行处理，故我们只需要知道如何计算每一个并行线程的初始索引，以及如何确定递增量的大小，就可完成GPU上所有任务的处理。

和情况2相比：

  • 核函数的索引计算方法： 将：
1
2
3
4
5

int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
	{
		c[tid] = a[tid] + b[tid];
	}

改为：

1
2
3
4
5
6

int tid = threadIdx.x + blockIdx.x * blockDim.x;
	if (tid < N)
	{
		c[tid] = a[tid] + b[tid];
		tid += blockDim.x * gridDim.x; ////blockDim.x * gridDim.x 表示当前线程格上的所有在运行的线程，即128*128
 	}
  • 核函数的调用方式： 将：
1

add_kernel << <(N + 127) / 128, 128 >> >

改为：

1

add_kernel << < 128, 128 >> >

即可。

完整代码：

 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
58
59
60
61
62
63

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <math.h>
#include <helper_cuda.h>
using namespace std;

#define N 1280


__global__ void add_kernel(double *a, double *b, double *c) {
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	if (tid < N)
	{
		c[tid] = a[tid] + b[tid];
		tid += blockDim.x * gridDim.x;
 	}

}

__global__ void value_init_kernel(double *a, double *b) {
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	if (tid < N)
	{
		a[tid] = 1.0*tid;
		b[tid] = (1.0*tid*tid);
		tid += blockDim.x * gridDim.x;
	}

}

int main(void)
{
	cudaError_t err1 = cudaSuccess, err2 = cudaSuccess, err3 = cudaSuccess;
	double a[N], b[N], c[N];
	double *dev_a, *dev_b, *dev_c;

	HANDLE_ERROR(cudaMalloc((void**)&dev_a, N * sizeof(double)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b, N * sizeof(double)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c, N * sizeof(double)));

	value_init_kernel << <128, 128 >> > (dev_a, dev_b);////在GPU上赋值操作
	add_kernel << < 128, 128 >> > (dev_a, dev_b, dev_c);////在GPU上相加操作

	HANDLE_ERROR(cudaMemcpy(a, dev_a, N * sizeof(double), cudaMemcpyDeviceToHost));
	HANDLE_ERROR(cudaMemcpy(b, dev_b, N * sizeof(double), cudaMemcpyDeviceToHost));
	HANDLE_ERROR(cudaMemcpy(c, dev_c, N * sizeof(double), cudaMemcpyDeviceToHost));

	for (int i = 0; i < N; i++)
	{
		printf("%f + %f = %f\n", a[i], b[i], c[i]);
	}

	////释放GPU内存
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);


	return 0;

}
updatedupdated2019-12-282019-12-28

相关文章：