Staskaer

Hi, this is Staskaer

0%

CUDA高性能并行计算(5)--CUDA二维计算网格

Posted on Edited on In Valine:

CUDA二维计算网格

二维计算网格

核函数创建二维计算网格

之前我们的核函数是按照如下方式来创建计算网格的。

xxx_kernel << <2, 2 >> > ();

前几篇也介绍过，这里会自动转换成xxx_kernel << <dim3(2,1,1), dim3(2,1,1) >> > ();，后面两个维度的参数默认是1。

需要注意的是，在开普勒、图灵架构中，线程块内线程总数不超过1024

所以我们需要创建一个二维网格只需要按照下面的方式调用核函数即可。

1
2
3
4
5
6
7
8

dim3 gridSize(2,2,1);
dim3 blockSize(2,2,1);

...

xxx_kernel<<<gridSize, blockSize>>>();

核函数内部确定线程位置

在一维计算网格中，我们使用blockIdx.x和threadIdx.x来获取当前线程处的索引。在二维计算网格中，也是类似的，我们给核函数传递的网格维度和线程块维度存储在gridDim，blockDim中，这两个变量有x,y,z三个成员存储传入dim3的三个维度。

与之类似，blockIdx和threadIdx也有x,y,z三个成员分别存储当前线程块索引和线程索引。

下面给出一个简单的二维计算网格的程序。

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
64

#include <device_launch_parameters.h>
#include <cuda_runtime.h>
#include <stdlib.h>
#include <stdio.h>
#include <math.h>

#define CHECK(call)                                   \
do                                                    \
{                                                     \
    const cudaError_t error_code = call;              \
    if (error_code != cudaSuccess)                    \
    {                                                 \
        printf("CUDA Error:\n");                      \
        printf("    File:       %s\n", __FILE__);     \
        printf("    Line:       %d\n", __LINE__);     \
        printf("    Error code: %d\n", error_code);   \
        printf("    Error text: %s\n",                \
            cudaGetErrorString(error_code));          \
        exit(1);                                      \
    }                                                 \
} while (0)



#define W 500
#define H 500
#define TX 32
#define TY 32

// 核函数，这里的float2是一个二维向量
__global__ void distanceKernel(float * d_out, int w, int h, float2 pos)
{
	const int c = blockIdx.x * blockDim.x + threadIdx.x; // 行索引
	const int r = blockIdx.y*blockDim.y + threadIdx.y; // 列索引
	const int i = r * w + c; // 行优先存储，获取到当前线程的实际位置

	if ((c >= w) || (r >= h))
		return;

	d_out[i] = sqrtf((c - pos.x)*(c - pos.x) + (r - pos.y)*(r - pos.y)); // 计算结果
	printf("thread (%d, %d) in grid (%d, %d) computes the result is %f\n", threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y,d_out[i]);
}

int main()
{
	float *out = (float *)calloc(W*H, sizeof(float)); // 主机内存
	float * d_out; // 设备内存
	cudaMalloc(&d_out, W*H* sizeof(float)); // 为设备分配内存
	const float2 pos = { 0.0f,0.0f }; // 初始化计算过程中传递的常量
	const dim3 blockSize(TX, TY); // 指定线程块大小
	const int bx = (W + TX - 1) / TX; // 向上舍入计算出x方向的计算网格数目
	const int by = (H + TY - 1) / TY; // 向上摄入计算处y方向的计算网格数目
	const dim3 gridSize = dim3(bx, by); // 指定计算网格大小

	distanceKernel << <gridSize, blockSize >> > (d_out, W, H, pos); // 启动核函数

	cudaDeviceSynchronize(); // 同步设备与主机
	cudaMemcpy(out, d_out, W*H * sizeof(float), cudaMemcpyDeviceToHost); // 拷贝计算结果

	cudaFree(d_out); // 释放内存
	free(out);
}

注意，执行这个程序的时候，可以观测到明显的输出卡顿，这就是前面提到的由于GPU高速缓存能力比较弱，所以当核函数执行简单的计算时内存IO将成为瓶颈。

总结

  • 线程块内线程总数不超过1024。
  • blockIdx和threadIdx存在三个成员来保存维度