CUDA极简入门

CUDA极简入门

CUDA，Compute Unified Device Architecture

典型流程

  1. 分配host内存，并进行数据初始化
  2. 分配device内存，并从host将数据拷贝到device上
  3. 调用CUDA的核函数在device上完成指定的运算
  4. 将device上的运算结果拷贝到host上
  5. 释放device和host上分配的内存

CUDA数组、对象和资源的使用

一般使用cudaArray_t、cudaTextureObj_t和cudaSurfaceObj_t三种，并且常将同一个变量的三者（or less，但至少使用）形式进行绑定。cudaChannelFormatDesc和cudaCreateChannelDesc的使用为cudaArray_t、cudaTextureObj_t和cudaSurfaceObj_t进行格式创建描述，使用MallocCudaArray对相关Resource进行显存分配。同时，如果需要将GPU中的数据与CPU中的数据进行同步或交换，采用cudaMemcpy2DToArray和cudaMemcpy2DFromArray，完成GPU和CPU的数据同步。

Error定位

由于CUDA是跑在GPU上的程序，所有的变量分配也是存放在GPU中，因此不能够和CPU调试一样有详细的输出或堆栈调用。但我们在编写时总可能会遇到在某一些代码段出错的情况，但CUDA只会抛出异常，并不会定位到代码段，因此需要进行GPU代码模块的Error定位。

一般而言，我们在完成GPU相应的代码模块功能后，可以使用cudaGetLastError和cudaStreamSynchtonize来定位可能出现的CUDA Error。这两个指令可以追踪CUDA中遇到的Error，在程序中可以使用多个追踪。在CUDA模块运行出错后，程序会根据Error追踪指令提示到出错CUDA模块后的第一个cudaGetLastError，用此方式可以快速定位到编写的哪一个模块在最终运行时出现错误。

实际上，所有的CUDA API都会返回一个CUDA状态值，但我们一般不需要将所有的CUDA命令都进行出错判断。

From simple example

示例代码功能为利用CUDA实现灰度图片对应像素相加。（当然我还没跑过，虽然是我写的）

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
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
#include <stdio.h>
#include <opencv2/core.hpp>
#include <opencv2/imgcodecs.hpp>
#include <cuda_runtime.h>

// __global__ declare
__global__ void add(cudaTextureObject_t* img1,cudaTextureObject_t* img2,cudaSurfaceObject_t* img3, int imgH, int imhW)
{
		int pidx = blockIdx.x * blockDim.x + threadIdx.x;
		int pidy = blockIdx.y * blockDim.y + threadIdx.y;
		if((pidx<imgW)&&(pidy<imgH)){
				uchar v1 = tex2D<uchar>(img1, pidx, pidy);
				uchar v2 = tex2D<uchar>(img2, pidx, pidy);
				uchar v3 = v1 + v2;
				surf2Dwrite(v3, img3, sizeof(uchar)*pidx, pidy);
		}
}

void Upload(cv::Mat &_img_host, cudaArray_t &_cuda_array)
	{
		cudaSafeCall(cudaMemcpy2DToArray(
			_cuda_array, 0, 0,
			_img_host.data,
			_img_host.cols * _img_host.elemSize(),
			_img_host.cols * _img_host.elemSize(),
			_img_host.rows,
			cudaMemcpyHostToDevice));
	}

void Download(cudaArray_t &_cuda_array, cv::Mat &_img_host)
	{
		cudaSafeCall(cudaMemcpy2DFromArray(
			_img_host.data,
			_img_host.cols * _img_host.elemSize(),
			_cuda_array,
			0, 0,
			_img_host.cols * _img_host.elemSize(),
			_img_host.rows,
			cudaMemcpyDeviceToHost));
	}

auto CreateCudaTextureAndSurfaceObject = [](
		cudaTextureObject_t &_tex,
		cudaSurfaceObject_t &_surf,
		cudaArray_t &_array,
		cudaResourceDesc &_res_desc,
		cudaTextureDesc &_tex_desc)
	{
		_res_desc.res.array.array = _array;
		cudaSafeCall(cudaCreateTextureObject(&_tex, &_res_desc, &_tex_desc, 0));
		cudaSafeCall(cudaCreateSurfaceObject(&_surf, &_res_desc));
	};

int main(void)
{
		// Host array init
    cv::Mat host_a, host_b, host_c;
		int width = 512, height = 512;
		host_a.create(width, height, CV_8CU1);
		host_b.create(width, height, CV_8CU1);
		host_c.create(width, height, CV_8CU1);

		// Device array, surface, texture init
		cudaArray_t dev_a_cuarray;
		cudaArray_t dev_b_cuarray;
		cudaArray_t dev_c_cuarray;

		cudaTextureObject_t dev_a_cutex;
		cudaTextureObject_t dev_b_cutex;
		cudaTextureObject_t dev_c_cutex;

		cudaSurfaceObject_t dev_a_cusurf;
		cudaSurfaceObject_t dev_b_cusurf;
		cudaSurfaceObject_t dev_c_cusurf;

		cudaChannelFormatDesc dev_a_desc = cudaCreateChannelDesc<uchar>();
		cudaChannelFormatDesc dev_b_desc = cudaCreateChannelDesc<uchar>();
		cudaChannelFormatDesc dev_c_desc = cudaCreateChannelDesc<uchar>();

		cudaMallocArray(&dev_a_cuarray, &dev_a_desc, width, height, cudaArraySurfaceLoadStore);
		cudaMallocArray(&dev_b_cuarray, &dev_b_desc, width, height, cudaArraySurfaceLoadStore);
		cudaMallocArray(&dev_c_cuarray, &dev_c_desc, width, height, cudaArraySurfaceLoadStore);

		cudaTextureDesc tex_desc;
		memset(&tex_desc, 0, sizeof(cudaTextureDesc));
		tex_desc.addressMode[0] = cudaAddressModeBorder;
		tex_desc.addressMode[1] = cudaAddressModeBorder;
		tex_desc.filterMode = cudaFilterModePoint;
		tex_desc.readMode = cudaReadModeElementType;
		tex_desc.normalizedCoords = 0;

		cudaResourceDesc res_desc;
		memset(&res_desc, 0, sizeof(cudaResourceDesc));
		res_desc.resType = cudaResourceTypeArray;

		CreateCudaTextureAndSurfaceObject(dev_a_cutex, dev_a_cusurf, dev_a_cuarray, res_desc, tex_desc);
		CreateCudaTextureAndSurfaceObject(dev_b_cutex, dev_b_cusurf, dev_b_cuarray, res_desc, tex_desc);
		CreateCudaTextureAndSurfaceObject(dev_c_cutex, dev_c_cusurf, dev_c_cuarray, res_desc, tex_desc);

		// Upload data from CPU to GPU
		Upload(host_a, dev_a_cuarray);
		Upload(host_b, dev_b_cuarray);
		Upload(host_c, dev_c_cuarray);

		// Block and Grid init
		dim3 block(16, 16);
		dim3 grid(width/block.x, height/block.y);
    add<<<grid, block>>>(dev_a_cutex, dev_b_cutex, dev_c_cusurf, height, width);

		cudaSafeCall(cudaGetLastError());
		cudaSafeCall(cudaDeviceSynchronize());

		// Download data from GPU to CPU
		Download(dev_c_cuarray, host_c);

		// Destroy CUDA resource
		cudaDestroyTextureObject(dev_a_cutex);
		cudaDestroyTextureObject(dev_b_cutex);
		cudaDestroyTextureObject(dev_c_cutex);

		cudaDestroySurfaceObject(dev_a_cusurf);
		cudaDestroySurfaceObject(dev_b_cusurf);
		cudaDestroySurfaceObject(dev_c_cusurf);

		cudaFreeArray(dev_a_cuarray);
		cudaFreeArray(dev_b_cuarray);
		cudaFreeArray(dev_c_cuarray);

		// Release cv::Mat
		host_a.release();
		host_b.release();
		host_c.release();
    return 0;
}

Grid and Block

Grid和Block是CUDA编程中最重要的概念之一，以下图[Reference 1]为例，代表一个CUDA核函数kernel调用的Grid和Block图解。

blockDim.x是每一个Block内部的x方向的维度，图中5，即每行5个线程。blockDim.y是Block内部的y方向的维度，这里是3，即每列3个线程。blockIdx.x是Block在grid中x方向的位置，图中放大的Block是Grid中的2，即为Grid中x方向的第2个。blockIdx.y是Block在grid中y向的位置，图中放大的Block是Grid中的2，即为Grid中y方向的第2个。blockIdx中的Idx是表示index的缩写，而不是表示x方向的ID。

Grid And Block

在CUDA kernel的使用时，需要传入<<<grid, block>>>参数。

CUDA C++ Basic

  • cudaCreateChannelDesc，创建cuda的通道描述符，在使用cudaMallocArray的时候需要将描述符一同传入进行cudaArray的内存申请。使用cudaChannelFormatDesc dev_a_desc = cudaCreateChannelDesc<uchar>();的使用方式生成通道描述符
  • cudaArray_t, cudaSurfaceObject_t, cudaTextureObject_t为cuda中常用的三种数据类型格式，分别为数组类型、表面对象和纹理对象。表面对象和纹理对象通常会绑定到cuda数组
    1. 若将cudaArray_t, cudaSurfaceObject_t, cudaTextureObject_t利用CreateCudaTextureAndSurfaceObject的方式进行绑定，实际上cuda数组、表面对象和纹理对象都绑定在同一个显存区间中，仅是对象读写方式不同
    2. cudaArray_t为cuda数组对象，是最重要的一个对象，在完成cuda数组对象的申请和分配后，才可以将相应的纹理对象、表面对象等绑定到数组对象中，并且使用cuda数组对象能更方便地与Host对象进行数据同步
    3. cudaTextureObject_t可读不可写，其读取速度可能比cudaSurfaceObject_t的速度快（未证实）。纹理对象会自动实现插值，即若取值为uv中没有恰好落在某个像素的时候，纹理对象会自动插值生成对应值
    4. cudaSurfaceObject_t可写，在对cudaArray_t需要进行更改的情况下一般都会将数组对象和表面对象进行绑定，并通过cuda kernel中surf2Dwrite函数将更新后的数值写入表面对象中，由于表面对象和数组对象是同一个绑定，因此cuda数组完成了更新
  • cudaMallocArray是将数组对象在GPU中分配对应显存空间的函数，类似于CPU中分配内存的函数Malloc。但在对cuda对象进行显存分配的过程中，除了cuda数组对象、长宽和类型外，使用通道描述符来完成cuda数组对象类型的传入，如uchar, float4等
  • cudaCreateTextureObject, cudaCreateSurfaceObject是将cuda数组对象与表面对象和纹理对象进行绑定的函数，将表面对象和纹理对象与cuda数组对象绑定后就能够使用相应对象完成后续操作
  • cudaDestroyTextureObject, cudaDestroySurfaceObject是将表面对象和纹理对象进行销毁
  • cudaFreeArray是将cuda数组对象申请的GPU显存进行释放
  • cudaMemcpy2DToArray, cudaMemcpy2DFromArray分别是将Host对象拷贝到Device和Device对象拷贝到Host
  • surf2Dwrite是将对应2D点的值写入到cudaSurfaceObject_t的表面对象中，cuda数组对象与表面对象绑定后并创建格式为cudaArraySurfaceLoadStore后可以使用此方式进行更改，同样的表面对象读取可以使用surf2Dread的方式完成
  • cudaSafeCall，检查cuda函数返回值是否正确的判定宏

CUDA核函数

在CUDA C++的.cu文件中，函数前缀分为__host__, __global__, __device__三种

  • __host__为Host调用的函数，此类型函数无法被GPU调用，无法被__global__和__device__调用。若函数未表明前缀，默认为此前缀类型
  • __global__前缀代表此函数为核函数，可以在Host中调用，被Device执行，并且此函数可以调用__device__前缀的函数，函数使用严格按照<<<grid, block>>>的方式进行，函数返回类型必须是void，不支持可变参数参数，不能成为类成员函数。用__global__定义的kernel是异步的，这意味着Host不会等待Kernel执行完就执行下一步
  • __device__前缀表明此函数是在GPU中运行的，无法直接被CPU函数调用
  • __host__和__device__可以同时使用

由于CUDA核函数中有不同的线程，因此需要传入grid和block参数给CUDA核函数进行使用。在CUDA核函数编写的时候，需要先计算出当前所在的线程idx，然后在参与计算。一个线程块上的线程是放在同一个流式多处理器（Streaming Multiprocessor，SM）上的，但是单个SM的资源有限，这导致线程块中的线程数是有限制的，现在的GPU线程块可支持的线程数可达1024个。当我们要知道一个线程在blcok中的全局ID，就必须还要知道block的组织结构，这是通过线程的内置变量blockDim来获得。它获取线程块各个维度的大小。对于一个2-dim的block$(D_x,D_y)$，线程$(x,y)$的ID值为$x+y\cdot D_x$，如果是3-dim的block$(D_x,D_y,D_z)$，线程$(x,y,z)$的ID值为$x+y\cdot D_x+z\cdot D_x\cdot D_y$。另外线程还有内置变量gridDim，用于获得网格块各个维度的大小。

如在2-dim的核函数中，一般有

$$
int \ \ pidx = blockIdx.x * blockDim.x + threadIdx.x \ int \ \ pidy = blockIdx.y * blockDim.y + threadIdx.y
$$

CUDA的内存模型，如下图所示。可以看到，每个线程有自己的私有本地内存（Local Memory），而每个线程块有包含共享内存（Shared Memory）,可以被线程块中所有线程共享，其生命周期与线程块一致。此外，所有的线程都可以访问全局内存（Global Memory）。还可以访问一些只读内存块：常量内存（Constant Memory）和纹理内存（Texture Memory）。

CUDA Memory Model

一个SM的基本执行单元是包含32个线程的线程束，因此block大小一般设置为32的倍数。

Reference

CUDA编程入门极简教程

Programming Guide :: CUDA Toolkit Documentation

CUDA
#CUDA
CUDA极简入门
https://alschain.com/2022/06/25/CUDA极简入门/
作者
Alschain
发布于
2022年6月25日
许可协议