CUDA 编程

SATAN 先生

已于 2024-09-25 15:15:41 修改

阅读量745

点赞数 18

分类专栏： python 文章标签： python c++

于 2024-09-25 15:09:25 首次发布

本文链接：https://blog.csdn.net/qq_32527569/article/details/142466358

版权

python 专栏收录该内容

20 篇文章 0 订阅

订阅专栏

0. Abstract

前几天浅尝了使用 pybind11 将 C++ 代码打包成 Python 包, 假装我已经能通过 C++ 代码加速程序了. 现在, 搞一搞 CUDA 编程, 进而结合 pybind11 把 CUDA 代码打包成 Python 包. 主要目的是了解一下整体流程与框架. 主要参考CUDA编程基础入门系列.

1. CUDA 简介

可以看到 GPU 的架构和 CPU 是很相似的, 都有内存、缓存器、控制器和算术逻辑单元, 不同的是 GPU 有很多很多的控制器和算数逻辑单元, 提供了很高的并行计算能力.

CPU 和 GPU 之间的数据交换
打开京东查看显卡商品, 会发现有个属性叫接口类型: PCIE-4.0, 这是显卡连接到主板的协议. 所以, 内存和显存之间的数据交换是通过主板总线进行的, 那么其速度不会快, 故而像 .numpy(), .to("cuda:x") 这类操作还是尽量少为好.

在这里插入图片描述
多种语言都可以进行 CUDA 编程, 其中 C/C++ 语言是最为推荐的, 较为完善. 像 Python 的 Numba 进行 CUDA 编程, 性能相比 C++ 就略有不足.

CUDA 编程的架构如上图所示, 类似于 C++ 和汇编, 你可以写一个程序文件, CUDA 驱动可以直接运行, 也可以像写 C++ 程序一样写好 CUDA 程序, 再通过 nvcc 编译成驱动能执行的中间代码.

小结: 简介了 GPU 的大致架构以及和 CPU 的关系、数据交换方式, GPU 编程的大致逻辑架构.

2. CUDA Demo

创建一个后缀名为 .cu 的文件 demo.cu, 编写代码:

#include <stdio.h>

__global__ void hello_from_gpu()  // GPU 核函数
{
	printf("Hello World from the GPU!\n");
}

int main()
{
	hello_from_gpu<<<2, 2>>>();  // 2*2=4 个并行线程
	cudaDeviceSynchronize();     // 同步函数, 等待和函数执行完毕再继续往下.
	return 0;
}

与 C/C++ 文件不同的是, CUDA 编程的后缀名为 .cu, 这可以使 CUDA 编程用到的一些头文件自动导入, 从而关键字 __global__ 等就可以使用了, __global__ 标志着 hello_from_gpu() 是一个在 GPU 上执行的核函数. <<<2, 2>>> 指示并行线程数(后面有详细介绍).

使用 nvcc 编译:

nvcc demo.cu -o demo

执行:

./demo
# >>> output >>>
Hello World from the GPU!
Hello World from the GPU!
Hello World from the GPU!
Hello World from the GPU!

这个 nvcc 就是安装的 CUDA Toolkit, 在其 bin 目录下, 可以理解为支持 CUDA 程序编译的 GCC:

3. CUDA 核函数（Kernel Function）与并行计算

上面的 Demo 已经展示了 CUDA 核函数了, 它是由关键字 __global__ 标志的, 下面是一些特点:

可以利用核函数实现高度并行计算, 上面例子中输出的四个 Hello World from the GPU! 就是启动了 4 给我线程分别执行 hello_from_gpu(), 它们执行了相同的任务. 要想实现多线程协同处理同一个任务, 需要理解 <<<2, 2>>> 的意义, 它是线程模型.

3.1 线程模型

大概意思是多线程的组织管理器, 刚才的 <<<2, 2>>> 是指创建了 2 个 block, 每个 block 中 2 个线程. 这里比较容易误解: 以为 grid_size 是指 grid 的数量, block_size 是 block 的数量. 其实仔细看看图片, 发现:

主机调用一个核函数则启动一个 grid, 理解为一个容器, grid_size 指示这个容器的大小, 也即能装多少 block;
而 block_size 则指示了 grid 中每个 block 的大小, 也即能装多少线程;
本次调用核函数启动的线程总数为 grid_size * block_size, 也即 num_blocks * num_threads_perblock.

注: 线程数量可以远远大于 GPU 的 CUDA 核心数, 也只有这样, 才能充分利用 GPU.

实质上, grid_size 和 block_size 都是类型为 dim3 的向量, 包含三个数值, 如示意图中:

dim3 grid_size(3, 2, 1);
dim3 block_size(5, 3, 1);

刚才的 <<<2, 2>>> 等价于:

dim3 grid_size(2, 1, 1);
dim3 block_size(2, 1, 1);
hello_from_gpu<<<grid_size, block_size>>>();

线程的身份标识

那么在完整的 3 维情况下, 内建变量总结如下:

gridDim: (gridDim.x, gridDim.y, gridDim.z), 就是所说的 gird_size;
blockDim: (blockDim.x, blockDim.y, blockDim.z), 就是所说的 block_size;
blockIdx: (blockIdx.x, blockIdx.y, blockIdx.z), 当前线程所处的 block 号;
blockIdx: (threadIdx.x, threadIdx.y, threadIdx.z), 当前线程局部号;

所以, 想获取线程的全局号, 需要:

int global_id = thread_id + block_id * block_size;

这里的 block_size, thread_id 和 block_id 是要经过换算的(类似二维数组算一维号):

int block_size = blockDim.x * blockDim.y * blockDim.z;
int thread_id = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;
int block_id = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y;

注意这里像是二维数组的列优先排列. 这适用于所有维度的组合, 如 2 维 grid, 1 维 block, 把省略的部分当作 1 就可以了.

其实多数时候不必考虑这么多关于线程全局 id 的事, 我估计线程这么组织和其常用于张量计算有关.

3.2 向量加法实现

这部分包括 GPU 的配置, CPU 内存和 GPU 显存的数据交换, 矩阵加法的并行计算等.

3.2.1 GPU 配置

在执行计算之前, 需要对 GPU 进行配置:

/**
 * 查询 GPU 数量并设置使用的设备
 * @param device 要使用的设备编号，从 0 开始
 */
bool setGPU(int device)
{
	int numDevices = 0;
	cudaError_t err = cudaGetDeviceCount(&numDevices);  // GPU 设备的数量

	if (err != cudaSuccess || numDevices == 0)         // 查询失败或者数量为 0
	{
		printf("No CUDA device found!\n");
		return false;
	}
	else
	{
		printf("Number of CUDA devices: %d\n", numDevices);
	}

	// set the device to use
	err = cudaSetDevice(device);  // 设置编号为 device 的 GPU 为默认计算设备

	if (err != cudaSuccess)
	{
		printf("Failed to set CUDA device %d!\n", device);
		return false;
	}
	else
	{
		printf("CUDA device %d set\n", device);
	}
	return true;
}

3.2.2 加法核函数

有了上面对线程组织方式的概念, 让我们来用 CUDA 试一试向量的加法计算:

__global__ void addFromGPU(float *a, float *b, float *c)
{
	const int block_id = blockIdx.x;    // 使用 1 维线程块
	const int thread_id = threadIdx.x;  // 使用 1 维网格
	const int id = thread_id + block_id * blockDim.x;
	c[id] = a[id] + b[id];
}

3.2.3 向量加法计算

// 随机初始化向量数据
void initialData(float *addr, int numElements)
{
	for (int i = 0; i < numElements; i++)
		addr[i] = (float)(rand() & 0xFF) / 10.0f;
}

int main()
{
	// 1. 设置设备
	bool ret = setGPU(0);
	if (!ret)
		exit(-1);

	// 2. 分配主机内存和设备内存, 并初始化数据
	int numElements = 512;						   // 设置元素数量
	size_t numBytes = numElements * sizeof(float); // 字节数量

	// 2.1 分配主机内存, 并初始化数据
	float *fpHostA = (float *)malloc(numBytes);
	float *fpHostB = (float *)malloc(numBytes);
	float *fpHostC = (float *)malloc(numBytes);
	if (fpHostA != NULL && fpHostB != NULL && fpHostC != NULL)
	{  // 分配成功
		memset(fpHostA, 0, numBytes); // 主机内存初始化为 0
		memset(fpHostB, 0, numBytes);
		memset(fpHostC, 0, numBytes);
	}
	else
	{
		printf("Memory allocation failed!\n");
		return -1;
	}

	// 2.2 分配设备内存, 并初始化
	float *fpDeviceA, *fpDeviceB, *fpDeviceC;
	// 不一样的是, cudaMalloc() 返回值是 void * 类型, 所以指针以参数的形式传入
	cudaMalloc((void **)&fpDeviceA, numBytes); // & 再取指针的指针
	cudaMalloc((void **)&fpDeviceB, numBytes); // 要转化为 void **
	cudaMalloc((void **)&fpDeviceC, numBytes);
	if (fpDeviceA != NULL && fpDeviceB != NULL && fpDeviceC != NULL)
	{
		cudaMemset(fpDeviceA, 0, numBytes); // 设备内存初始化为 0
		cudaMemset(fpDeviceB, 0, numBytes);
		cudaMemset(fpDeviceC, 0, numBytes);
	}
	else
	{
		printf("Memory allocation failed!\n");
		free(fpHostA);
		free(fpHostB);
		free(fpHostC);
		return -1;
	}

	// 3. 初始化主机中的数据
	srand(666); // 设置随机种子
	initialData(fpHostA, numElements);
	initialData(fpHostB, numElements);

	// 4. 将主机内存拷贝到设备内存
	cudaMemcpy(fpDeviceA, fpHostA, numBytes, cudaMemcpyHostToDevice);
	cudaMemcpy(fpDeviceB, fpHostB, numBytes, cudaMemcpyHostToDevice);
	cudaMemcpy(fpDeviceC, fpHostC, numBytes, cudaMemcpyHostToDevice);

	// 5. 调用核函数
	dim3 blockSize(32);						 // 块大小
	dim3 gridSize(numElements / blockSize.x); // 块数量
	addFromGPU<<<gridSize, blockSize>>>(fpDeviceA, fpDeviceB, fpDeviceC);
	cudaDeviceSynchronize(); // 等待设备完成

	// 6. 将结果从设备内存拷贝到主机内存
	cudaMemcpy(fpHostC, fpDeviceC, numBytes, cudaMemcpyDeviceToHost);

	// 7. 输出结果
	for (int i = 0; i < 10; i++)
	{
		printf(
			"idx=%2d\tmatrix_A:%.2f\tmatrix_B:%.2f\tresult:%.2f\n",
			i + 1, fpHostA[i], fpHostB[i], fpHostC[i]);
	}
	// 8. 释放资源
	free(fpHostA);
	free(fpHostB);
	free(fpHostC);
	cudaFree(fpDeviceA);
	cudaFree(fpDeviceB);
	cudaFree(fpDeviceC);

	cudaDeviceReset(); // 释放设备资源
	return 0;
}

很简单, 说几个需要关注的地方:

CUDA 的内存管理和标准的 C 语言内存管理很像:

基本就是带前缀 cuda 的驼峰模式命名. cudaMalloc 似乎只支持二重指针 void**;
主机内存和 GPU 显存之间的数据交换由几种模式:
最后别忘了 cudaDeviceReset(); // 释放设备资源.

4. 编译时指定架构(CUDA程序兼容性问题)

在遇到含有 CUDA 代码的 Python 项目时, 说明书会让你先编译. 但当你编译后运行时, 却发现报与 CUDA 架构有关的错误.

4.1 CUDA 架构

注意, 安装的 CUDA 版本也有限制, CUDA10 最高支持 sm_75.

需要注意的是, -arch 是向后兼容的, 而 -code 只在大版本 X 内部向后兼容, 这是由于不同大版本之间 GPU 的实际架构不同.

5. 用 Pybind11 导出 CUDA 代码的 Python 接口

下面将用 CUDA 实现向量(Python中list,C++中vector)的加法运算, vector_add.cu:

#include <pybind11/pybind11.h>
#include <pybind11/stl.h>
#include <cstdio>

using namespace std;

bool setGPU(int device)
{
	int numDevices = 0;
	auto err = cudaGetDeviceCount(&numDevices);

	if (err != cudaSuccess || numDevices == 0)
	{
		printf("No CUDA device found!\n");
		return false;
	} else
	{
		printf("Number of CUDA devices: %d\n", numDevices);
	}

	// set the device to use
	err = cudaSetDevice(device);

	if (err != cudaSuccess)
	{
		printf("Failed to set CUDA device %d!\n", device);
		return false;
	} else
	{
		printf("CUDA device %d set\n", device);
	}
	return true;
}

__global__ void addFromGPU(float *a, float *b, float *c)
{
	const int block_id = blockIdx.x;
	const int thread_id = threadIdx.x;
	const int id = thread_id + block_id * blockDim.x;

	c[id] = a[id] + b[id];
}

std::vector<float> add(std::vector<float> a, std::vector<float> b)
{
	if (a.size() != b.size())
	{
		throw std::invalid_argument("Vectors must be the same size.");
	}

	// 1. 设置设备
	bool ret = setGPU(0);
	if (!ret)
		exit(-1);

	// 2. 分配主机内存和设备内存, 并初始化数据
	auto numElements = a.size();				   // 设置元素数量
	size_t numBytes = numElements * sizeof(float); // 字节数量
	std::vector<float> fpHostC(numElements, 0.0f);

	// 2.2 分配设备内存, 并初始化
	float *fpDeviceA, *fpDeviceB, *fpDeviceC;
	// 不一样的是, cudaMalloc() 返回值是 void * 类型, 所以指针以参数的形式传入
	cudaMalloc((void **)&fpDeviceA, numBytes); // & 再取指针的指针
	cudaMalloc((void **)&fpDeviceB, numBytes); // 要转化为 void **
	cudaMalloc((void **)&fpDeviceC, numBytes);
	if (fpDeviceA != NULL && fpDeviceB != NULL && fpDeviceC != NULL)
	{
		cudaMemset(fpDeviceA, 0, numBytes); // 设备内存初始化为 0
		cudaMemset(fpDeviceB, 0, numBytes);
		cudaMemset(fpDeviceC, 0, numBytes);
	}
	else
	{
		printf("Memory allocation failed!\n");
		exit(-1);
	}

	// 4. 将主机内存拷贝到设备内存
	cudaMemcpy(fpDeviceA, a.data(), numBytes, cudaMemcpyHostToDevice);
	cudaMemcpy(fpDeviceB, b.data(), numBytes, cudaMemcpyHostToDevice);

	// 5. 调用核函数
	dim3 blockSize(numElements);		  // 块大小
	dim3 gridSize(numElements / blockSize.x); // 块数量
	addFromGPU<<<gridSize, blockSize>>>(fpDeviceA, fpDeviceB, fpDeviceC);
	cudaDeviceSynchronize(); // 等待设备完成

	// 6. 将结果从设备内存拷贝到主机内存
	cudaMemcpy(fpHostC.data(), fpDeviceC, numBytes, cudaMemcpyDeviceToHost);

	cudaFree(fpDeviceA);
	cudaFree(fpDeviceB);
	cudaFree(fpDeviceC);

	cudaDeviceReset(); // 释放设备资源
	return fpHostC;
}

PYBIND11_MODULE(vector_module, m) { // 导出 Python 模块
	m.def("add", &add, "Add two vectors element-wise.");
}

CMakeLists.txt:

cmake_minimum_required(VERSION 3.18)
project(CuDemo LANGUAGES CUDA C CXX)

set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CUDA_ARCHITECTURES 75)
set(CMAKE_CUDA_COMPILER "/usr/local/cuda/bin/nvcc")

# set(PYTHON_BASE_PATH /root/Miniconda)
find_package(Python REQUIRED COMPONENTS Interpreter Development.Module)
find_package(pybind11 REQUIRED PATHS /root/Miniconda/share)

pybind11_add_module(vector_module ./vector_add.cu)

编译后就可以使用 vector_module 了:

>>> import vector_module
>>> vector_module.add([1, 2, 3, 4], [2, 3, 4, 5])
Number of CUDA devices: 2
CUDA device 0 set
[3.0, 5.0, 7.0, 9.0]

SATAN 先生

关注

18
点赞
踩
23

收藏

觉得还不错? 一键收藏
0
评论
复制链接

分享到 QQ

分享到新浪微博

扫一扫

专栏目录